search

clij / clij2

GPU-accelerated image processing for everyone
https://clij.github.io/clij2
Other
48 stars 14 forks source link

Error when computing of Standard Deviation in Clij2 #60

Closed schmiedc closed 2 years ago

schmiedc commented 2 years ago

Standard Deviation of all on GPU in Clij2 throws and Error. Interestingly the masked standard deviation works.

(Fiji Is Just) ImageJ 2.3.0/1.53s; Java 1.8.0_322 [64-bit]; Linux 4.18.0-348.20.1.el8_5.x86_64; 443MB of 19319MB (2%)

java.lang.NullPointerException
    at net.haesleinhuepf.clij2.CLIJ2.execute(CLIJ2.java:537)
    at net.haesleinhuepf.clij2.plugins.VarianceOfAllPixels.varianceOfAllPixels(VarianceOfAllPixels.java:62)
    at net.haesleinhuepf.clij2.plugins.StandardDeviationOfAllPixels.standardDeviationOfAllPixels(StandardDeviationOfAllPixels.java:46)
    at net.haesleinhuepf.clij2.plugins.StandardDeviationOfAllPixels.standardDeviationOfAllPixels(StandardDeviationOfAllPixels.java:42)
    at net.haesleinhuepf.clij2.CLIJ2Ops.standardDeviationOfAllPixels(CLIJ2Ops.java:1769)
    at net.haesleinhuepf.clij2.plugins.StandardDeviationOfAllPixels.executeCL(StandardDeviationOfAllPixels.java:31)
    at net.haesleinhuepf.clij.macro.AbstractCLIJPlugin.run(AbstractCLIJPlugin.java:478)
    at ij.plugin.filter.PlugInFilterRunner.processOneImage(PlugInFilterRunner.java:266)
    at ij.plugin.filter.PlugInFilterRunner.<init>(PlugInFilterRunner.java:114)
    at ij.IJ.runUserPlugIn(IJ.java:239)
    at ij.IJ.runPlugIn(IJ.java:203)
    at ij.Executer.runCommand(Executer.java:152)
    at ij.Executer.run(Executer.java:67)
    at java.lang.Thread.run(Thread.java:750)
package net.haesleinhuepf.clij2.plugins, clij2_, version 0.0
sumPixels is deprecated. Check the documentation for a replacement. https://clij.github.io/clij2-doccs/reference
Error when trying to create kernel squared_sum_project
net.haesleinhuepf.clij.clearcl.exceptions.OpenCLException: OpenCL error: -45 -> CL_INVALID_PROGRAM_EXECUTABLE
    at net.haesleinhuepf.clij.clearcl.backend.BackendUtils.checkOpenCLErrorCode(BackendUtils.java:352)
    at net.haesleinhuepf.clij.clearcl.backend.jocl.ClearCLBackendJOCL.lambda$getKernelPeerPointer$19(ClearCLBackendJOCL.java:601)
    at net.haesleinhuepf.clij.clearcl.backend.BackendUtils.checkExceptions(BackendUtils.java:156)
    at net.haesleinhuepf.clij.clearcl.backend.jocl.ClearCLBackendJOCL.getKernelPeerPointer(ClearCLBackendJOCL.java:593)
    at net.haesleinhuepf.clij.clearcl.ClearCLCompiledProgram.createKernel(ClearCLCompiledProgram.java:137)
    at net.haesleinhuepf.clij.clearcl.ClearCLProgram.createKernel(ClearCLProgram.java:685)
    at net.haesleinhuepf.clij.clearcl.util.CLKernelExecutor.getKernel(CLKernelExecutor.java:382)
    at net.haesleinhuepf.clij.clearcl.util.CLKernelExecutor.enqueue(CLKernelExecutor.java:258)
    at net.haesleinhuepf.clij2.CLIJ2.lambda$executeSubsequently$1(CLIJ2.java:579)
    at net.haesleinhuepf.clij.clearcl.util.ElapsedTime.measure(ElapsedTime.java:97)
    at net.haesleinhuepf.clij.clearcl.util.ElapsedTime.measure(ElapsedTime.java:28)
    at net.haesleinhuepf.clij2.CLIJ2.executeSubsequently(CLIJ2.java:569)
    at net.haesleinhuepf.clij2.CLIJ2.executeSubsequently(CLIJ2.java:556)
    at net.haesleinhuepf.clij2.CLIJ2.executeSubsequently(CLIJ2.java:551)
    at net.haesleinhuepf.clij2.CLIJ2.execute(CLIJ2.java:536)
    at net.haesleinhuepf.clij2.plugins.VarianceOfAllPixels.varianceOfAllPixels(VarianceOfAllPixels.java:62)
    at net.haesleinhuepf.clij2.plugins.StandardDeviationOfAllPixels.standardDeviationOfAllPixels(StandardDeviationOfAllPixels.java:46)
    at net.haesleinhuepf.clij2.plugins.StandardDeviationOfAllPixels.standardDeviationOfAllPixels(StandardDeviationOfAllPixels.java:42)
    at net.haesleinhuepf.clij2.CLIJ2Ops.standardDeviationOfAllPixels(CLIJ2Ops.java:1769)
    at net.haesleinhuepf.clij2.plugins.StandardDeviationOfAllPixels.executeCL(StandardDeviationOfAllPixels.java:31)
    at net.haesleinhuepf.clij.macro.AbstractCLIJPlugin.run(AbstractCLIJPlugin.java:478)
    at ij.plugin.filter.PlugInFilterRunner.processOneImage(PlugInFilterRunner.java:266)
    at ij.plugin.filter.PlugInFilterRunner.<init>(PlugInFilterRunner.java:114)
    at ij.IJ.runUserPlugIn(IJ.java:239)
    at ij.IJ.runPlugIn(IJ.java:203)
    at ij.Executer.runCommand(Executer.java:152)
    at ij.Executer.run(Executer.java:67)
    at java.lang.Thread.run(Thread.java:750)
haesleinhuepf commented 2 years ago

Hi Christopher @schmiedc ,

thanks for the report! I'm afraid I need more information. I've tested the function in this scenario and it works:

run("Blobs (25K)");

run("CLIJ2 Macro Extensions", "cl_device=");

image = getTitle();
Ext.CLIJ2_push(image);

Ext.CLIJ2_standardDeviationOfAllPixels(image);

image

Can you please tell me

Thanks again!

Best, Robert

schmiedc commented 2 years ago

That is interesting. If I use one of the sample images in Fiji it also works for me. I tested it with a 685x834x64 stack single channel 16-bit 70MB in size.

I have a Quadro RTX 3000 Mobile / Max-Q with Driver Version: 510.47.03 CUDA Version: 11.6

The code that reproduces that: 

run("CLIJ2 Macro Extensions", "cl_device=[Quadro RTX 3000]");

// standard deviation of all pixels
image1 = "C4-Crop-1-20220201-14-1-002-0.tif";
Ext.CLIJ2_push(image1);
Ext.CLIJ2_standardDeviationOfAllPixels(image1);

I can send a link to the image.

schmiedc commented 2 years ago

BTW when using a mask that is the entire size of the image the masked standard deviation works for me.

haesleinhuepf commented 2 years ago

Ok, I can reproduce the error using this macro:

newImage("Untitled", "16-bit ramp", 685, 834, 64);
run("CLIJ2 Macro Extensions", "cl_device=RTX");

// standard deviation of all pixels
image1 = "Untitled";
Ext.CLIJ2_push(image1);
Ext.CLIJ2_standardDeviationOfAllPixels(image1);

And I have a more detailed error message:

sumPixels is deprecated. Check the documentation for a replacement. https://clij.github.io/clij2-doccs/reference

 //###########################################################################
// Preamble:
#pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable

#pragma OPENCL EXTENSION cl_amd_printf : enable

#pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable

#ifndef M_PI
    #define   M_PI 3.14159265358979323846f /* pi */
#endif

#ifndef M_LOG2E
    #define   M_LOG2E   1.4426950408889634074f /* log_2 e */
#endif

#ifndef M_LOG10E
    #define   M_LOG10E   0.43429448190325182765f /* log_10 e */
#endif

#ifndef M_LN2
    #define   M_LN2   0.69314718055994530942f  /* log_e 2 */
#endif

#ifndef M_LN10
    #define   M_LN10   2.30258509299404568402f /* log_e 10 */
#endif

#ifndef BUFFER_READ_WRITE
    #define BUFFER_READ_WRITE 1

#define MINMAX_TYPE int

inline char2 read_buffer3dc(int read_buffer_width, int read_buffer_height, int read_buffer_depth, __global char * buffer_var, sampler_t sampler, int4 position )
{
    int4 pos = (int4){position.x, position.y, position.z, 0};
    if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
        pos.x = max((MINMAX_TYPE)pos.x, (MINMAX_TYPE)0);
        pos.y = max((MINMAX_TYPE)pos.y, (MINMAX_TYPE)0);
        pos.z = max((MINMAX_TYPE)pos.z, (MINMAX_TYPE)0);
        pos.x = min((MINMAX_TYPE)pos.x, (MINMAX_TYPE)read_buffer_width - 1);
        pos.y = min((MINMAX_TYPE)pos.y, (MINMAX_TYPE)read_buffer_height - 1);
        pos.z = min((MINMAX_TYPE)pos.z, (MINMAX_TYPE)read_buffer_depth - 1);
    }
    int pos_in_buffer = pos.x + pos.y * read_buffer_width + pos.z * read_buffer_width * read_buffer_height;
    if (pos.x < 0 || pos.x >= read_buffer_width || pos.y < 0 || pos.y >= read_buffer_height || pos.z < 0 || pos.z >= read_buffer_depth) {
        return (char2){0, 0};
    }
    return (char2){buffer_var[pos_in_buffer],0};
}

inline uchar2 read_buffer3duc(int read_buffer_width, int read_buffer_height, int read_buffer_depth, __global uchar * buffer_var, sampler_t sampler, int4 position )
{
    int4 pos = (int4){position.x, position.y, position.z, 0};
    if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
        pos.x = max((MINMAX_TYPE)pos.x, (MINMAX_TYPE)0);
        pos.y = max((MINMAX_TYPE)pos.y, (MINMAX_TYPE)0);
        pos.z = max((MINMAX_TYPE)pos.z, (MINMAX_TYPE)0);
        pos.x = min((MINMAX_TYPE)pos.x, (MINMAX_TYPE)read_buffer_width - 1);
        pos.y = min((MINMAX_TYPE)pos.y, (MINMAX_TYPE)read_buffer_height - 1);
        pos.z = min((MINMAX_TYPE)pos.z, (MINMAX_TYPE)read_buffer_depth - 1);
    }
    int pos_in_buffer = pos.x + pos.y * read_buffer_width + pos.z * read_buffer_width * read_buffer_height;
    if (pos.x < 0 || pos.x >= read_buffer_width || pos.y < 0 || pos.y >= read_buffer_height || pos.z < 0 || pos.z >= read_buffer_depth) {
        return (uchar2){0, 0};
    }
    return (uchar2){buffer_var[pos_in_buffer],0};
}

inline short2 read_buffer3ds(int read_buffer_width, int read_buffer_height, int read_buffer_depth, __global short * buffer_var, sampler_t sampler, int4 position )
{
    int4 pos = (int4){position.x, position.y, position.z, 0};
    if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
        pos.x = max((MINMAX_TYPE)pos.x, (MINMAX_TYPE)0);
        pos.y = max((MINMAX_TYPE)pos.y, (MINMAX_TYPE)0);
        pos.z = max((MINMAX_TYPE)pos.z, (MINMAX_TYPE)0);
        pos.x = min((MINMAX_TYPE)pos.x, (MINMAX_TYPE)read_buffer_width - 1);
        pos.y = min((MINMAX_TYPE)pos.y, (MINMAX_TYPE)read_buffer_height - 1);
        pos.z = min((MINMAX_TYPE)pos.z, (MINMAX_TYPE)read_buffer_depth - 1);
    }
    int pos_in_buffer = pos.x + pos.y * read_buffer_width + pos.z * read_buffer_width * read_buffer_height;
    if (pos.x < 0 || pos.x >= read_buffer_width || pos.y < 0 || pos.y >= read_buffer_height || pos.z < 0 || pos.z >= read_buffer_depth) {
        return (short2){0, 0};
    }
    return (short2){buffer_var[pos_in_buffer],0};
}

inline ushort2 read_buffer3dus(int read_buffer_width, int read_buffer_height, int read_buffer_depth, __global ushort * buffer_var, sampler_t sampler, int4 position )
{
    int4 pos = (int4){position.x, position.y, position.z, 0};
    if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
        pos.x = max((MINMAX_TYPE)pos.x, (MINMAX_TYPE)0);
        pos.y = max((MINMAX_TYPE)pos.y, (MINMAX_TYPE)0);
        pos.z = max((MINMAX_TYPE)pos.z, (MINMAX_TYPE)0);
        pos.x = min((MINMAX_TYPE)pos.x, (MINMAX_TYPE)read_buffer_width - 1);
        pos.y = min((MINMAX_TYPE)pos.y, (MINMAX_TYPE)read_buffer_height - 1);
        pos.z = min((MINMAX_TYPE)pos.z, (MINMAX_TYPE)read_buffer_depth - 1);
    }
    int pos_in_buffer = pos.x + pos.y * read_buffer_width + pos.z * read_buffer_width * read_buffer_height;
    if (pos.x < 0 || pos.x >= read_buffer_width || pos.y < 0 || pos.y >= read_buffer_height || pos.z < 0 || pos.z >= read_buffer_depth) {
        return (ushort2){0, 0};
    }
    return (ushort2){buffer_var[pos_in_buffer],0};
}

inline int2 read_buffer3di(int read_buffer_width, int read_buffer_height, int read_buffer_depth, __global int * buffer_var, sampler_t sampler, int4 position )
{
    int4 pos = (int4){position.x, position.y, position.z, 0};
    if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
        pos.x = max((MINMAX_TYPE)pos.x, (MINMAX_TYPE)0);
        pos.y = max((MINMAX_TYPE)pos.y, (MINMAX_TYPE)0);
        pos.z = max((MINMAX_TYPE)pos.z, (MINMAX_TYPE)0);
        pos.x = min((MINMAX_TYPE)pos.x, (MINMAX_TYPE)read_buffer_width - 1);
        pos.y = min((MINMAX_TYPE)pos.y, (MINMAX_TYPE)read_buffer_height - 1);
        pos.z = min((MINMAX_TYPE)pos.z, (MINMAX_TYPE)read_buffer_depth - 1);
    }
    int pos_in_buffer = pos.x + pos.y * read_buffer_width + pos.z * read_buffer_width * read_buffer_height;
    if (pos.x < 0 || pos.x >= read_buffer_width || pos.y < 0 || pos.y >= read_buffer_height || pos.z < 0 || pos.z >= read_buffer_depth) {
        return (int2){0, 0};
    }
    return (int2){buffer_var[pos_in_buffer],0};
}

inline uint2 read_buffer3dui(int read_buffer_width, int read_buffer_height, int read_buffer_depth, __global uint * buffer_var, sampler_t sampler, int4 position )
{
    int4 pos = (int4){position.x, position.y, position.z, 0};
    if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
        pos.x = max((MINMAX_TYPE)pos.x, (MINMAX_TYPE)0);
        pos.y = max((MINMAX_TYPE)pos.y, (MINMAX_TYPE)0);
        pos.z = max((MINMAX_TYPE)pos.z, (MINMAX_TYPE)0);
        pos.x = min((MINMAX_TYPE)pos.x, (MINMAX_TYPE)read_buffer_width - 1);
        pos.y = min((MINMAX_TYPE)pos.y, (MINMAX_TYPE)read_buffer_height - 1);
        pos.z = min((MINMAX_TYPE)pos.z, (MINMAX_TYPE)read_buffer_depth - 1);
    }
    int pos_in_buffer = pos.x + pos.y * read_buffer_width + pos.z * read_buffer_width * read_buffer_height;
    if (pos.x < 0 || pos.x >= read_buffer_width || pos.y < 0 || pos.y >= read_buffer_height || pos.z < 0 || pos.z >= read_buffer_depth) {
        return (uint2){0, 0};
    }
    return (uint2){buffer_var[pos_in_buffer],0};
}

inline long2 read_buffer3dl(int read_buffer_width, int read_buffer_height, int read_buffer_depth, __global long * buffer_var, sampler_t sampler, int4 position )
{
    int4 pos = (int4){position.x, position.y, position.z, 0};
    if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
        pos.x = max((MINMAX_TYPE)pos.x, (MINMAX_TYPE)0);
        pos.y = max((MINMAX_TYPE)pos.y, (MINMAX_TYPE)0);
        pos.z = max((MINMAX_TYPE)pos.z, (MINMAX_TYPE)0);
        pos.x = min((MINMAX_TYPE)pos.x, (MINMAX_TYPE)read_buffer_width - 1);
        pos.y = min((MINMAX_TYPE)pos.y, (MINMAX_TYPE)read_buffer_height - 1);
        pos.z = min((MINMAX_TYPE)pos.z, (MINMAX_TYPE)read_buffer_depth - 1);
    }
    int pos_in_buffer = pos.x + pos.y * read_buffer_width + pos.z * read_buffer_width * read_buffer_height;
    if (pos.x < 0 || pos.x >= read_buffer_width || pos.y < 0 || pos.y >= read_buffer_height || pos.z < 0 || pos.z >= read_buffer_depth) {
        return (long2){0, 0};
    }
    return (long2){buffer_var[pos_in_buffer],0};
}

inline ulong2 read_buffer3dul(int read_buffer_width, int read_buffer_height, int read_buffer_depth, __global ulong * buffer_var, sampler_t sampler, int4 position )
{
    int4 pos = (int4){position.x, position.y, position.z, 0};
    if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
        pos.x = max((MINMAX_TYPE)pos.x, (MINMAX_TYPE)0);
        pos.y = max((MINMAX_TYPE)pos.y, (MINMAX_TYPE)0);
        pos.z = max((MINMAX_TYPE)pos.z, (MINMAX_TYPE)0);
        pos.x = min((MINMAX_TYPE)pos.x, (MINMAX_TYPE)read_buffer_width - 1);
        pos.y = min((MINMAX_TYPE)pos.y, (MINMAX_TYPE)read_buffer_height - 1);
        pos.z = min((MINMAX_TYPE)pos.z, (MINMAX_TYPE)read_buffer_depth - 1);
    }
    int pos_in_buffer = pos.x + pos.y * read_buffer_width + pos.z * read_buffer_width * read_buffer_height;
    if (pos.x < 0 || pos.x >= read_buffer_width || pos.y < 0 || pos.y >= read_buffer_height || pos.z < 0 || pos.z >= read_buffer_depth) {
        return (ulong2){0, 0};
    }
    return (ulong2){buffer_var[pos_in_buffer],0};
}

inline float2 read_buffer3df(int read_buffer_width, int read_buffer_height, int read_buffer_depth, __global float* buffer_var, sampler_t sampler, int4 position )
{
    int4 pos = (int4){position.x, position.y, position.z, 0};
    if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
        pos.x = max((MINMAX_TYPE)pos.x, (MINMAX_TYPE)0);
        pos.y = max((MINMAX_TYPE)pos.y, (MINMAX_TYPE)0);
        pos.z = max((MINMAX_TYPE)pos.z, (MINMAX_TYPE)0);
        pos.x = min((MINMAX_TYPE)pos.x, (MINMAX_TYPE)read_buffer_width - 1);
        pos.y = min((MINMAX_TYPE)pos.y, (MINMAX_TYPE)read_buffer_height - 1);
        pos.z = min((MINMAX_TYPE)pos.z, (MINMAX_TYPE)read_buffer_depth - 1);
    }
    int pos_in_buffer = pos.x + pos.y * read_buffer_width + pos.z * read_buffer_width * read_buffer_height;
    if (pos.x < 0 || pos.x >= read_buffer_width || pos.y < 0 || pos.y >= read_buffer_height || pos.z < 0 || pos.z >= read_buffer_depth) {
        return (float2){0, 0};
    }
    return (float2){buffer_var[pos_in_buffer],0};
}

inline void write_buffer3dc(int write_buffer_width, int write_buffer_height, int write_buffer_depth, __global char * buffer_var, int4 pos, char value )
{
    int pos_in_buffer = pos.x + pos.y * write_buffer_width + pos.z * write_buffer_width * write_buffer_height;
    if (pos.x < 0 || pos.x >= write_buffer_width || pos.y < 0 || pos.y >= write_buffer_height || pos.z < 0 || pos.z >= write_buffer_depth) {
        return;
    }
    buffer_var[pos_in_buffer] = value;
}

inline void write_buffer3duc(int write_buffer_width, int write_buffer_height, int write_buffer_depth, __global uchar * buffer_var, int4 pos, uchar value )
{
    int pos_in_buffer = pos.x + pos.y * write_buffer_width + pos.z * write_buffer_width * write_buffer_height;
    if (pos.x < 0 || pos.x >= write_buffer_width || pos.y < 0 || pos.y >= write_buffer_height || pos.z < 0 || pos.z >= write_buffer_depth) {
        return;
    }
    buffer_var[pos_in_buffer] = value;
}

inline void write_buffer3ds(int write_buffer_width, int write_buffer_height, int write_buffer_depth, __global short * buffer_var, int4 pos, short value )
{
    int pos_in_buffer = pos.x + pos.y * write_buffer_width + pos.z * write_buffer_width * write_buffer_height;
    if (pos.x < 0 || pos.x >= write_buffer_width || pos.y < 0 || pos.y >= write_buffer_height || pos.z < 0 || pos.z >= write_buffer_depth) {
        return;
    }
    buffer_var[pos_in_buffer] = value;
}

inline void write_buffer3dus(int write_buffer_width, int write_buffer_height, int write_buffer_depth, __global ushort * buffer_var, int4 pos, ushort value )
{
    int pos_in_buffer = pos.x + pos.y * write_buffer_width + pos.z * write_buffer_width * write_buffer_height;
    if (pos.x < 0 || pos.x >= write_buffer_width || pos.y < 0 || pos.y >= write_buffer_height || pos.z < 0 || pos.z >= write_buffer_depth) {
        return;
    }
    buffer_var[pos_in_buffer] = value;
}

inline void write_buffer3di(int write_buffer_width, int write_buffer_height, int write_buffer_depth, __global int * buffer_var, int4 pos, int value )
{
    int pos_in_buffer = pos.x + pos.y * write_buffer_width + pos.z * write_buffer_width * write_buffer_height;
    if (pos.x < 0 || pos.x >= write_buffer_width || pos.y < 0 || pos.y >= write_buffer_height || pos.z < 0 || pos.z >= write_buffer_depth) {
        return;
    }
    buffer_var[pos_in_buffer] = value;
}

inline void write_buffer3dui(int write_buffer_width, int write_buffer_height, int write_buffer_depth, __global uint * buffer_var, int4 pos, uint value )
{
    int pos_in_buffer = pos.x + pos.y * write_buffer_width + pos.z * write_buffer_width * write_buffer_height;
    if (pos.x < 0 || pos.x >= write_buffer_width || pos.y < 0 || pos.y >= write_buffer_height || pos.z < 0 || pos.z >= write_buffer_depth) {
        return;
    }
    buffer_var[pos_in_buffer] = value;
}

inline void write_buffer3dl(int write_buffer_width, int write_buffer_height, int write_buffer_depth, __global long * buffer_var, int4 pos, long value )
{
    int pos_in_buffer = pos.x + pos.y * write_buffer_width + pos.z * write_buffer_width * write_buffer_height;
    if (pos.x < 0 || pos.x >= write_buffer_width || pos.y < 0 || pos.y >= write_buffer_height || pos.z < 0 || pos.z >= write_buffer_depth) {
        return;
    }
    buffer_var[pos_in_buffer] = value;
}

inline void write_buffer3dul(int write_buffer_width, int write_buffer_height, int write_buffer_depth, __global ulong * buffer_var, int4 pos, ulong value )
{
    int pos_in_buffer = pos.x + pos.y * write_buffer_width + pos.z * write_buffer_width * write_buffer_height;
    if (pos.x < 0 || pos.x >= write_buffer_width || pos.y < 0 || pos.y >= write_buffer_height || pos.z < 0 || pos.z >= write_buffer_depth) {
        return;
    }
    buffer_var[pos_in_buffer] = value;
}

inline void write_buffer3df(int write_buffer_width, int write_buffer_height, int write_buffer_depth, __global float* buffer_var, int4 pos, float value )
{
    int pos_in_buffer = pos.x + pos.y * write_buffer_width + pos.z * write_buffer_width * write_buffer_height;
    if (pos.x < 0 || pos.x >= write_buffer_width || pos.y < 0 || pos.y >= write_buffer_height || pos.z < 0 || pos.z >= write_buffer_depth) {
        return;
    }
    buffer_var[pos_in_buffer] = value;
}

inline char2 read_buffer2dc(int read_buffer_width, int read_buffer_height, int read_buffer_depth, __global char * buffer_var, sampler_t sampler, int2 position )
{
    int2 pos = (int2){position.x, position.y};
    if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
        pos.x = max((MINMAX_TYPE)pos.x, (MINMAX_TYPE)0);
        pos.y = max((MINMAX_TYPE)pos.y, (MINMAX_TYPE)0);
        pos.x = min((MINMAX_TYPE)pos.x, (MINMAX_TYPE)read_buffer_width - 1);
        pos.y = min((MINMAX_TYPE)pos.y, (MINMAX_TYPE)read_buffer_height - 1);
    }
    int pos_in_buffer = pos.x + pos.y * read_buffer_width;
    if (pos.x < 0 || pos.x >= read_buffer_width || pos.y < 0 || pos.y >= read_buffer_height) {
        return (char2){0, 0};
    }
    return (char2){buffer_var[pos_in_buffer],0};
}

inline uchar2 read_buffer2duc(int read_buffer_width, int read_buffer_height, int read_buffer_depth, __global uchar * buffer_var, sampler_t sampler, int2 position )
{
    int2 pos = (int2){position.x, position.y};
    if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
        pos.x = max((MINMAX_TYPE)pos.x, (MINMAX_TYPE)0);
        pos.y = max((MINMAX_TYPE)pos.y, (MINMAX_TYPE)0);
        pos.x = min((MINMAX_TYPE)pos.x, (MINMAX_TYPE)read_buffer_width - 1);
        pos.y = min((MINMAX_TYPE)pos.y, (MINMAX_TYPE)read_buffer_height - 1);
    }
    int pos_in_buffer = pos.x + pos.y * read_buffer_width;
    if (pos.x < 0 || pos.x >= read_buffer_width || pos.y < 0 || pos.y >= read_buffer_height) {
        return (uchar2){0, 0};
    }
    return (uchar2){buffer_var[pos_in_buffer],0};
}

inline short2 read_buffer2ds(int read_buffer_width, int read_buffer_height, int read_buffer_depth, __global short * buffer_var, sampler_t sampler, int2 position )
{
    int2 pos = (int2){position.x, position.y};
    if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
        pos.x = max((MINMAX_TYPE)pos.x, (MINMAX_TYPE)0);
        pos.y = max((MINMAX_TYPE)pos.y, (MINMAX_TYPE)0);
        pos.x = min((MINMAX_TYPE)pos.x, (MINMAX_TYPE)read_buffer_width - 1);
        pos.y = min((MINMAX_TYPE)pos.y, (MINMAX_TYPE)read_buffer_height - 1);
    }
    int pos_in_buffer = pos.x + pos.y * read_buffer_width;
    if (pos.x < 0 || pos.x >= read_buffer_width || pos.y < 0 || pos.y >= read_buffer_height) {
        return (short2){0, 0};
    }
    return (short2){buffer_var[pos_in_buffer],0};
}

inline ushort2 read_buffer2dus(int read_buffer_width, int read_buffer_height, int read_buffer_depth, __global ushort * buffer_var, sampler_t sampler, int2 position )
{
    int2 pos = (int2){position.x, position.y};
    if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
        pos.x = max((MINMAX_TYPE)pos.x, (MINMAX_TYPE)0);
        pos.y = max((MINMAX_TYPE)pos.y, (MINMAX_TYPE)0);
        pos.x = min((MINMAX_TYPE)pos.x, (MINMAX_TYPE)read_buffer_width - 1);
        pos.y = min((MINMAX_TYPE)pos.y, (MINMAX_TYPE)read_buffer_height - 1);
    }
    int pos_in_buffer = pos.x + pos.y * read_buffer_width;
    if (pos.x < 0 || pos.x >= read_buffer_width || pos.y < 0 || pos.y >= read_buffer_height) {
        return (ushort2){0, 0};
    }
    return (ushort2){buffer_var[pos_in_buffer],0};
}

inline int2 read_buffer2di(int read_buffer_width, int read_buffer_height, int read_buffer_depth, __global int * buffer_var, sampler_t sampler, int2 position )
{
    int2 pos = (int2){position.x, position.y};
    if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
        pos.x = max((MINMAX_TYPE)pos.x, (MINMAX_TYPE)0);
        pos.y = max((MINMAX_TYPE)pos.y, (MINMAX_TYPE)0);
        pos.x = min((MINMAX_TYPE)pos.x, (MINMAX_TYPE)read_buffer_width - 1);
        pos.y = min((MINMAX_TYPE)pos.y, (MINMAX_TYPE)read_buffer_height - 1);
    }
    int pos_in_buffer = pos.x + pos.y * read_buffer_width;
    if (pos.x < 0 || pos.x >= read_buffer_width || pos.y < 0 || pos.y >= read_buffer_height) {
        return (int2){0, 0};
    }
    return (int2){buffer_var[pos_in_buffer],0};
}

inline uint2 read_buffer2dui(int read_buffer_width, int read_buffer_height, int read_buffer_depth, __global uint * buffer_var, sampler_t sampler, int2 position )
{
    int2 pos = (int2){position.x, position.y};
    if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
        pos.x = max((MINMAX_TYPE)pos.x, (MINMAX_TYPE)0);
        pos.y = max((MINMAX_TYPE)pos.y, (MINMAX_TYPE)0);
        pos.x = min((MINMAX_TYPE)pos.x, (MINMAX_TYPE)read_buffer_width - 1);
        pos.y = min((MINMAX_TYPE)pos.y, (MINMAX_TYPE)read_buffer_height - 1);
    }
    int pos_in_buffer = pos.x + pos.y * read_buffer_width;
    if (pos.x < 0 || pos.x >= read_buffer_width || pos.y < 0 || pos.y >= read_buffer_height) {
        return (uint2){0, 0};
    }
    return (uint2){buffer_var[pos_in_buffer],0};
}

inline long2 read_buffer2dl(int read_buffer_width, int read_buffer_height, int read_buffer_depth, __global long * buffer_var, sampler_t sampler, int2 position )
{
    int2 pos = (int2){position.x, position.y};
    if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
        pos.x = max((MINMAX_TYPE)pos.x, (MINMAX_TYPE)0);
        pos.y = max((MINMAX_TYPE)pos.y, (MINMAX_TYPE)0);
        pos.x = min((MINMAX_TYPE)pos.x, (MINMAX_TYPE)read_buffer_width - 1);
        pos.y = min((MINMAX_TYPE)pos.y, (MINMAX_TYPE)read_buffer_height - 1);
    }
    int pos_in_buffer = pos.x + pos.y * read_buffer_width;
    if (pos.x < 0 || pos.x >= read_buffer_width || pos.y < 0 || pos.y >= read_buffer_height) {
        return (long2){0, 0};
    }
    return (long2){buffer_var[pos_in_buffer],0};
}

inline ulong2 read_buffer2dul(int read_buffer_width, int read_buffer_height, int read_buffer_depth, __global ulong * buffer_var, sampler_t sampler, int2 position )
{
    int2 pos = (int2){position.x, position.y};
    if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
        pos.x = max((MINMAX_TYPE)pos.x, (MINMAX_TYPE)0);
        pos.y = max((MINMAX_TYPE)pos.y, (MINMAX_TYPE)0);
        pos.x = min((MINMAX_TYPE)pos.x, (MINMAX_TYPE)read_buffer_width - 1);
        pos.y = min((MINMAX_TYPE)pos.y, (MINMAX_TYPE)read_buffer_height - 1);
    }
    int pos_in_buffer = pos.x + pos.y * read_buffer_width;
    if (pos.x < 0 || pos.x >= read_buffer_width || pos.y < 0 || pos.y >= read_buffer_height) {
        return (ulong2){0, 0};
    }
    return (ulong2){buffer_var[pos_in_buffer],0};
}

inline float2 read_buffer2df(int read_buffer_width, int read_buffer_height, int read_buffer_depth, __global float* buffer_var, sampler_t sampler, int2 position )
{
    int2 pos = (int2){position.x, position.y};
    if (true) { // if (CLK_ADDRESS_CLAMP_TO_EDGE & sampler) {
        pos.x = max((MINMAX_TYPE)pos.x, (MINMAX_TYPE)0);
        pos.y = max((MINMAX_TYPE)pos.y, (MINMAX_TYPE)0);
        pos.x = min((MINMAX_TYPE)pos.x, (MINMAX_TYPE)read_buffer_width - 1);
        pos.y = min((MINMAX_TYPE)pos.y, (MINMAX_TYPE)read_buffer_height - 1);
    }
    int pos_in_buffer = pos.x + pos.y * read_buffer_width;
    if (pos.x < 0 || pos.x >= read_buffer_width || pos.y < 0 || pos.y >= read_buffer_height) {
        return (float2){0, 0};
    }
    return (float2){buffer_var[pos_in_buffer],0};
}

inline void write_buffer2dc(int write_buffer_width, int write_buffer_height, int write_buffer_depth, __global char * buffer_var, int2 pos, char value )
{
    int pos_in_buffer = pos.x + pos.y * write_buffer_width;
    if (pos.x < 0 || pos.x >= write_buffer_width || pos.y < 0 || pos.y >= write_buffer_height) {
        return;
    }
    buffer_var[pos_in_buffer] = value;
}

inline void write_buffer2duc(int write_buffer_width, int write_buffer_height, int write_buffer_depth, __global uchar * buffer_var, int2 pos, uchar value )
{
    int pos_in_buffer = pos.x + pos.y * write_buffer_width;
    if (pos.x < 0 || pos.x >= write_buffer_width || pos.y < 0 || pos.y >= write_buffer_height) {
        return;
    }
    buffer_var[pos_in_buffer] = value;
}

inline void write_buffer2ds(int write_buffer_width, int write_buffer_height, int write_buffer_depth, __global short * buffer_var, int2 pos, short value )
{
    int pos_in_buffer = pos.x + pos.y * write_buffer_width;
    if (pos.x < 0 || pos.x >= write_buffer_width || pos.y < 0 || pos.y >= write_buffer_height) {
        return;
    }
    buffer_var[pos_in_buffer] = value;
}

inline void write_buffer2dus(int write_buffer_width, int write_buffer_height, int write_buffer_depth, __global ushort * buffer_var, int2 pos, ushort value )
{
    int pos_in_buffer = pos.x + pos.y * write_buffer_width;
    if (pos.x < 0 || pos.x >= write_buffer_width || pos.y < 0 || pos.y >= write_buffer_height) {
        return;
    }
    buffer_var[pos_in_buffer] = value;
}

inline void write_buffer2di(int write_buffer_width, int write_buffer_height, int write_buffer_depth, __global int * buffer_var, int2 pos, int value )
{
    int pos_in_buffer = pos.x + pos.y * write_buffer_width;
    if (pos.x < 0 || pos.x >= write_buffer_width || pos.y < 0 || pos.y >= write_buffer_height) {
        return;
    }
    buffer_var[pos_in_buffer] = value;
}

inline void write_buffer2dui(int write_buffer_width, int write_buffer_height, int write_buffer_depth, __global uint * buffer_var, int2 pos, uint value )
{
    int pos_in_buffer = pos.x + pos.y * write_buffer_width;
    if (pos.x < 0 || pos.x >= write_buffer_width || pos.y < 0 || pos.y >= write_buffer_height) {
        return;
    }
    buffer_var[pos_in_buffer] = value;
}

inline void write_buffer2dl(int write_buffer_width, int write_buffer_height, int write_buffer_depth, __global long * buffer_var, int2 pos, long value )
{
    int pos_in_buffer = pos.x + pos.y * write_buffer_width;
    if (pos.x < 0 || pos.x >= write_buffer_width || pos.y < 0 || pos.y >= write_buffer_height) {
        return;
    }
    buffer_var[pos_in_buffer] = value;
}

inline void write_buffer2dul(int write_buffer_width, int write_buffer_height, int write_buffer_depth, __global ulong * buffer_var, int2 pos, ulong value )
{
    int pos_in_buffer = pos.x + pos.y * write_buffer_width;
    if (pos.x < 0 || pos.x >= write_buffer_width || pos.y < 0 || pos.y >= write_buffer_height) {
        return;
    }
    buffer_var[pos_in_buffer] = value;
}

inline void write_buffer2df(int write_buffer_width, int write_buffer_height, int write_buffer_depth, __global float* buffer_var, int2 pos, float value )
{
    int pos_in_buffer = pos.x + pos.y * write_buffer_width;
    if (pos.x < 0 || pos.x >= write_buffer_width || pos.y < 0 || pos.y >= write_buffer_height) {
        return;
    }
    buffer_var[pos_in_buffer] = value;
}

inline uchar clij_convert_uchar_sat(float value) {
    if (value > 255) {
        return 255;
    }
    if (value < 0) {
        return 0;
    }
    return (uchar)value;
}

inline char clij_convert_char_sat(float value) {
    if (value > 127) {
        return 127;
    }
    if (value < -128) {
        return -128;
    }
    return (char)value;
}

inline ushort clij_convert_ushort_sat(float value) {
    if (value > 65535) {
        return 65535;
    }
    if (value < 0) {
        return 0;
    }
    return (ushort)value;
}

inline short clij_convert_short_sat(float value) {
    if (value > 32767) {
        return 32767;
    }
    if (value < -32768) {
        return -32768;
    }
    return (short)value;
}

inline uint clij_convert_uint_sat(float value) {
    if (value > 4294967295) {
        return 4294967295;
    }
    if (value < 0) {
        return 0;
    }
    return (uint)value;
}

inline int clij_convert_int_sat(float value) {
    if (value > 2147483647) {
        return 2147483647;
    }
    if (value < -2147483648) {
        return -2147483648;
    }
    return (int)value;
}

inline uint clij_convert_ulong_sat(float value) {
    if (value > 18446744073709551615) {
        return 18446744073709551615;
    }
    if (value < 0) {
        return 0;
    }
    return (ulong)value;
}

inline int clij_convert_long_sat(float value) {
    if (value > 9223372036854775807) {
        return 9223372036854775807;
    }
    if (value < -9223372036854775808 ) {
        return -9223372036854775808 ;
    }
    return (long)value;
}

inline float clij_convert_float_sat(float value) {
    return value;
}

#define READ_IMAGE(a,b,c) READ_ ## a ## _IMAGE(a,b,c)
#define WRITE_IMAGE(a,b,c) WRITE_ ## a ## _IMAGE(a,b,c)

#endif

 //###########################################################################
// Defines:
#define WRITE_dst_IMAGE(a,b,c)  write_buffer2df(GET_IMAGE_WIDTH(a),GET_IMAGE_HEIGHT(a),GET_IMAGE_DEPTH(a),a,b,c)
#define CONVERT_dst_PIXEL_TYPE  clij_convert_float_sat
#define GET_IMAGE_HEIGHT(image_key)     IMAGE_SIZE_ ## image_key ## _HEIGHT
#define IMAGE_SIZE_dst_HEIGHT   834
#define READ_src_IMAGE(a,b,c)   read_buffer3dus(GET_IMAGE_WIDTH(a),GET_IMAGE_HEIGHT(a),GET_IMAGE_DEPTH(a),a,b,c)
#define GET_IMAGE_DEPTH(image_key)  IMAGE_SIZE_ ## image_key ## _DEPTH
#define CONVERT_src_PIXEL_TYPE  clij_convert_ushort_sat
#define IMAGE_SIZE_dst_WIDTH    685
#define IMAGE_dst_PIXEL_TYPE    float
#define IMAGE_src_TYPE  long image_size_src_width, long image_size_src_height, long image_size_src_depth, __global ushort*
#define READ_dst_IMAGE(a,b,c)   read_buffer2df(GET_IMAGE_WIDTH(a),GET_IMAGE_HEIGHT(a),GET_IMAGE_DEPTH(a),a,b,c)
#define POS_dst_TYPE    int2
#define IMAGE_src_PIXEL_TYPE    ushort
#define POS_src_INSTANCE(pos0, pos1, pos2, pos3)    ((int4)(pos0, pos1, pos2, pos3))
#define POS_dst_INSTANCE(pos0, pos1, pos2, pos3)    ((int2)(pos0, pos1))
#define POS_src_TYPE    int4
#define IMAGE_SIZE_src_WIDTH    685
#define IMAGE_dst_TYPE  long image_size_dst_width, long image_size_dst_height, long image_size_dst_depth, __global float*
#define MAX_ARRAY_SIZE  1000
#define WRITE_src_IMAGE(a,b,c)  write_buffer3dus(GET_IMAGE_WIDTH(a),GET_IMAGE_HEIGHT(a),GET_IMAGE_DEPTH(a),a,b,c)
#define GET_IMAGE_WIDTH(image_key)  IMAGE_SIZE_ ## image_key ## _WIDTH
#define IMAGE_SIZE_dst_DEPTH    1
#define IMAGE_SIZE_src_HEIGHT   834
#define IMAGE_SIZE_src_DEPTH    64

 //###########################################################################
// Source: 'variance_projection_x.cl' relative to VarianceOfMaskedPixels
__kernel void squared_sum_project(
    IMAGE_dst_TYPE dst,
    IMAGE_src_TYPE src,
    float mean_intensity
) {
  const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;

  const int x = get_global_id(0);
  const int y = get_global_id(1);
  float sum = 0;
  for(int z = 0; z < GET_IMAGE_DEPTH(src); z++)
  {
    float value = READ_src_IMAGE(src,sampler,POS_src_INSTANCE(x,y,z,0)).x;
    sum = sum + pow(value - mean_intensity, 2);
  }
  WRITE_dst_IMAGE(dst,POS_dst_INSTANCE(x,y,0,0), CONVERT_sum_PIXEL_TYPE(sum));
}

ptxas fatal   : Unresolved extern function 'CONVERT_sum_PIXEL_TYPE'
Error when trying to create kernel squared_sum_project
net.haesleinhuepf.clij.clearcl.exceptions.OpenCLException: OpenCL error: -45 -> CL_INVALID_PROGRAM_EXECUTABLE
    at net.haesleinhuepf.clij.clearcl.backend.BackendUtils.checkOpenCLErrorCode(BackendUtils.java:352)
    at net.haesleinhuepf.clij.clearcl.backend.jocl.ClearCLBackendJOCL.lambda$getKernelPeerPointer$19(ClearCLBackendJOCL.java:601)
    at net.haesleinhuepf.clij.clearcl.backend.BackendUtils.checkExceptions(BackendUtils.java:156)
    at net.haesleinhuepf.clij.clearcl.backend.jocl.ClearCLBackendJOCL.getKernelPeerPointer(ClearCLBackendJOCL.java:593)
    at net.haesleinhuepf.clij.clearcl.ClearCLCompiledProgram.createKernel(ClearCLCompiledProgram.java:137)
    at net.haesleinhuepf.clij.clearcl.ClearCLProgram.createKernel(ClearCLProgram.java:685)
    at net.haesleinhuepf.clij.clearcl.util.CLKernelExecutor.getKernel(CLKernelExecutor.java:382)
    at net.haesleinhuepf.clij.clearcl.util.CLKernelExecutor.enqueue(CLKernelExecutor.java:258)
    at net.haesleinhuepf.clij2.CLIJ2.lambda$executeSubsequently$1(CLIJ2.java:579)
    at net.haesleinhuepf.clij.clearcl.util.ElapsedTime.measure(ElapsedTime.java:97)
    at net.haesleinhuepf.clij.clearcl.util.ElapsedTime.measure(ElapsedTime.java:28)
    at net.haesleinhuepf.clij2.CLIJ2.executeSubsequently(CLIJ2.java:569)
    at net.haesleinhuepf.clij2.CLIJ2.executeSubsequently(CLIJ2.java:556)
    at net.haesleinhuepf.clij2.CLIJ2.executeSubsequently(CLIJ2.java:551)
    at net.haesleinhuepf.clij2.CLIJ2.execute(CLIJ2.java:536)
    at net.haesleinhuepf.clij2.plugins.VarianceOfAllPixels.varianceOfAllPixels(VarianceOfAllPixels.java:62)
    at net.haesleinhuepf.clij2.plugins.StandardDeviationOfAllPixels.standardDeviationOfAllPixels(StandardDeviationOfAllPixels.java:46)
    at net.haesleinhuepf.clij2.plugins.StandardDeviationOfAllPixels.standardDeviationOfAllPixels(StandardDeviationOfAllPixels.java:42)
    at net.haesleinhuepf.clij2.CLIJ2Ops.standardDeviationOfAllPixels(CLIJ2Ops.java:1769)
    at net.haesleinhuepf.clij2.plugins.StandardDeviationOfAllPixels.executeCL(StandardDeviationOfAllPixels.java:31)
    at net.haesleinhuepf.clij.macro.CLIJHandler.lambda$handleExtension$0(CLIJHandler.java:163)
    at net.haesleinhuepf.clij.clearcl.util.ElapsedTime.measure(ElapsedTime.java:97)
    at net.haesleinhuepf.clij.clearcl.util.ElapsedTime.measure(ElapsedTime.java:28)
    at net.haesleinhuepf.clij.macro.CLIJHandler.handleExtension(CLIJHandler.java:53)
    at ij.macro.ExtensionDescriptor.dispatch(ExtensionDescriptor.java:288)
    at ij.macro.Functions.doExt(Functions.java:5073)
    at ij.macro.Functions.getStringFunction(Functions.java:279)
    at ij.macro.Interpreter.getStringTerm(Interpreter.java:1520)
    at ij.macro.Interpreter.getString(Interpreter.java:1498)
    at ij.macro.Interpreter.doStatement(Interpreter.java:336)
    at ij.macro.Interpreter.doStatements(Interpreter.java:267)
    at ij.macro.Interpreter.run(Interpreter.java:163)
    at ij.macro.Interpreter.run(Interpreter.java:93)
    at ij.macro.Interpreter.run(Interpreter.java:107)
    at ij.plugin.Macro_Runner.runMacro(Macro_Runner.java:162)
    at ij.IJ.runMacro(IJ.java:159)
    at ij.IJ.runMacro(IJ.java:148)
    at net.imagej.legacy.IJ1Helper$3.call(IJ1Helper.java:1148)
    at net.imagej.legacy.IJ1Helper$3.call(IJ1Helper.java:1144)
    at net.imagej.legacy.IJ1Helper.runMacroFriendly(IJ1Helper.java:1095)
    at net.imagej.legacy.IJ1Helper.runMacro(IJ1Helper.java:1144)
    at net.imagej.legacy.plugin.IJ1MacroEngine.eval(IJ1MacroEngine.java:145)
    at org.scijava.script.ScriptModule.run(ScriptModule.java:157)
    at org.scijava.module.ModuleRunner.run(ModuleRunner.java:163)
    at org.scijava.module.ModuleRunner.call(ModuleRunner.java:124)
    at org.scijava.module.ModuleRunner.call(ModuleRunner.java:63)
    at org.scijava.thread.DefaultThreadService.lambda$wrap$2(DefaultThreadService.java:225)
    at java.util.concurrent.FutureTask.run(FutureTask.java:266)
    at java.util.concurrent.ThreadPoolExecutor.runWorker(ThreadPoolExecutor.java:1149)
    at java.util.concurrent.ThreadPoolExecutor$Worker.run(ThreadPoolExecutor.java:624)
    at java.lang.Thread.run(Thread.java:748)
java.lang.NullPointerException
    at net.haesleinhuepf.clij2.CLIJ2.execute(CLIJ2.java:537)
    at net.haesleinhuepf.clij2.plugins.VarianceOfAllPixels.varianceOfAllPixels(VarianceOfAllPixels.java:62)
    at net.haesleinhuepf.clij2.plugins.StandardDeviationOfAllPixels.standardDeviationOfAllPixels(StandardDeviationOfAllPixels.java:46)
    at net.haesleinhuepf.clij2.plugins.StandardDeviationOfAllPixels.standardDeviationOfAllPixels(StandardDeviationOfAllPixels.java:42)
    at net.haesleinhuepf.clij2.CLIJ2Ops.standardDeviationOfAllPixels(CLIJ2Ops.java:1769)
    at net.haesleinhuepf.clij2.plugins.StandardDeviationOfAllPixels.executeCL(StandardDeviationOfAllPixels.java:31)
    at net.haesleinhuepf.clij.macro.CLIJHandler.lambda$handleExtension$0(CLIJHandler.java:163)
    at net.haesleinhuepf.clij.clearcl.util.ElapsedTime.measure(ElapsedTime.java:97)
    at net.haesleinhuepf.clij.clearcl.util.ElapsedTime.measure(ElapsedTime.java:28)
    at net.haesleinhuepf.clij.macro.CLIJHandler.handleExtension(CLIJHandler.java:53)
    at ij.macro.ExtensionDescriptor.dispatch(ExtensionDescriptor.java:288)
    at ij.macro.Functions.doExt(Functions.java:5073)
    at ij.macro.Functions.getStringFunction(Functions.java:279)
    at ij.macro.Interpreter.getStringTerm(Interpreter.java:1520)
    at ij.macro.Interpreter.getString(Interpreter.java:1498)
    at ij.macro.Interpreter.doStatement(Interpreter.java:336)
    at ij.macro.Interpreter.doStatements(Interpreter.java:267)
    at ij.macro.Interpreter.run(Interpreter.java:163)
    at ij.macro.Interpreter.run(Interpreter.java:93)
    at ij.macro.Interpreter.run(Interpreter.java:107)
    at ij.plugin.Macro_Runner.runMacro(Macro_Runner.java:162)
    at ij.IJ.runMacro(IJ.java:159)
    at ij.IJ.runMacro(IJ.java:148)
    at net.imagej.legacy.IJ1Helper$3.call(IJ1Helper.java:1148)
    at net.imagej.legacy.IJ1Helper$3.call(IJ1Helper.java:1144)
    at net.imagej.legacy.IJ1Helper.runMacroFriendly(IJ1Helper.java:1095)
    at net.imagej.legacy.IJ1Helper.runMacro(IJ1Helper.java:1144)
    at net.imagej.legacy.plugin.IJ1MacroEngine.eval(IJ1MacroEngine.java:145)
    at org.scijava.script.ScriptModule.run(ScriptModule.java:157)
    at org.scijava.module.ModuleRunner.run(ModuleRunner.java:163)
    at org.scijava.module.ModuleRunner.call(ModuleRunner.java:124)
    at org.scijava.module.ModuleRunner.call(ModuleRunner.java:63)
    at org.scijava.thread.DefaultThreadService.lambda$wrap$2(DefaultThreadService.java:225)
    at java.util.concurrent.FutureTask.run(FutureTask.java:266)
    at java.util.concurrent.ThreadPoolExecutor.runWorker(ThreadPoolExecutor.java:1149)
    at java.util.concurrent.ThreadPoolExecutor$Worker.run(ThreadPoolExecutor.java:624)
    at java.lang.Thread.run(Thread.java:748)
haesleinhuepf commented 2 years ago

I presume replacing CONVERT_sum_PIXEL_TYPE with CONVERT_dst_PIXEL_TYPE here might solve the issue. Do you want to give it a try @schmiedc and send a PR, or do you want to wait until I find time to fix this? It may take a week or two.

Thanks again for reporting!

schmiedc commented 2 years ago

Only one answer possible :D https://github.com/clij/clij2/pull/61

haesleinhuepf commented 2 years ago

Hey @schmiedc ,

I just uploaded clij 2.5.3.1 to the Fiji update site. It contains your bugfix. Thanks again and let me know if there are any further issues!

Best, Robert