search

clij / clij2

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

ClearCLException "problem while setting argument 'src'" on MacOS #25

Open haesleinhuepf opened 4 years ago

haesleinhuepf commented 4 years ago

Dear future self,

there is a not-reproducible bug, potentially only on Macs with specific Hardware (suspicion).

Sytem config:

ImageJ: Version: 2.1.0/1.53c Build: 5f23140693

macOS Catalina. Version 10.15.5 MacBook Pro (16-inch, 2019) Processor 2.3 GHz 8-Core Intel Core i9 Memory: 32 GB 2667 MHz DDR4 Graphics: AMD Radeon Pro 5500M 8 GB Intel UHD Graphics 630 1536 M

Executed code

// top hat
run("T1 Head (2.4M, 16-bits)");
run("CLIJ2 Macro Extensions", "cl_device=[Intel(R) UHDGraphics 620]");
image1 = "t1-head.tif";
Ext.CLIJ2_push(image1);
image2 = "top_hat-427502308";
radius_x = 10.0;
radius_y = 10.0;
radius_z = 10.0;
// study time tracing of the Top Hat filter
Ext.CLIJ2_startTimeTracing();
Ext.CLIJ2_topHatBox(image1, image2, radius_x, radius_y,radius_z);
Ext.CLIJ2_stopTimeTracing();
Ext.CLIJ2_pull(image2);// determine and print time 
Ext.CLIJ2_getTimeTracing(time_traces);
print(time_traces);

Error message:

Overwriting image in cache.
net.haesleinhuepf.clij.clearcl.exceptions.ClearCLException: problem while setting argument 'dst' at index 0 

    at net.haesleinhuepf.clij.clearcl.ClearCLKernel.setArgumentsInternal(ClearCLKernel.java:426)
    at net.haesleinhuepf.clij.clearcl.ClearCLKernel.lambda$run$0(ClearCLKernel.java:481)
    at net.haesleinhuepf.clij.clearcl.util.ElapsedTime.measure(ElapsedTime.java:97)
    at net.haesleinhuepf.clij.clearcl.util.ElapsedTime.measure(ElapsedTime.java:64)
    at net.haesleinhuepf.clij.clearcl.ClearCLKernel.run(ClearCLKernel.java:477)
    at net.haesleinhuepf.clij.clearcl.ClearCLKernel.run(ClearCLKernel.java:459)
    at net.haesleinhuepf.clij.utilities.CLKernelExecutor.lambda$enqueue$0(CLKernelExecutor.java:333)
    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.utilities.CLKernelExecutor.enqueue(CLKernelExecutor.java:331)
    at net.haesleinhuepf.clij.CLIJ.lambda$execute$0(CLIJ.java:278)
    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.CLIJ.execute(CLIJ.java:255)
    at net.haesleinhuepf.clij.CLIJ.execute(CLIJ.java:235)
    at net.haesleinhuepf.clij.kernels.Kernels.copyInternal(Kernels.java:835)
    at net.haesleinhuepf.clij.kernels.Kernels.copy(Kernels.java:847)
    at net.haesleinhuepf.clij.macro.CLIJHandler.pushInternal(CLIJHandler.java:316)
    at net.haesleinhuepf.clij.macro.CLIJHandler.pushToGPU(CLIJHandler.java:268)
    at net.haesleinhuepf.clij2.plugins.Push.executeCL(Push.java:30)
    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:4967)
    at ij.macro.Functions.getStringFunction(Functions.java:278)
    at ij.macro.Interpreter.getStringTerm(Interpreter.java:1475)
    at ij.macro.Interpreter.getString(Interpreter.java:1453)
    at ij.macro.Interpreter.doStatement(Interpreter.java:333)
    at ij.macro.Interpreter.doStatements(Interpreter.java:264)
    at ij.macro.Interpreter.run(Interpreter.java:160)
    at ij.macro.Interpreter.run(Interpreter.java:93)
    at ij.macro.Interpreter.run(Interpreter.java:104)
    at ij.plugin.Macro_Runner.runMacro(Macro_Runner.java:161)
    at ij.IJ.runMacro(IJ.java:153)
    at ij.IJ.runMacro(IJ.java:142)
    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:165)
    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)
Caused by: net.haesleinhuepf.clij.clearcl.exceptions.ClearCLException
    at net.haesleinhuepf.clij.clearcl.backend.BackendUtils.checkExceptions(BackendUtils.java:178)
    at net.haesleinhuepf.clij.clearcl.backend.jocl.ClearCLBackendJOCL.setKernelArgument(ClearCLBackendJOCL.java:614)
    at net.haesleinhuepf.clij.clearcl.ClearCLKernel.setArgumentsInternal(ClearCLKernel.java:417)
    ... 50 more
Caused by: java.lang.NullPointerException
    at net.haesleinhuepf.clij.clearcl.backend.jocl.ClearCLBackendJOCL.lambda$setKernelArgument$20(ClearCLBackendJOCL.java:707)
    at net.haesleinhuepf.clij.clearcl.backend.BackendUtils.checkExceptions(BackendUtils.java:171)
    ... 52 more

 //###########################################################################
// 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 long

inline char2 read_buffer4dc(long read_buffer_width, long read_buffer_height, long read_buffer_depth, __global char * buffer_var, sampler_t sampler, int4 position )
{
    int4 pos = (int4){position.x, position.y, position.z, position.w};
    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);
        // todo: correct pos.w
    }
    long pos_in_buffer = (long)pos.x + (long)pos.y * read_buffer_width + (long)pos.z * read_buffer_width * read_buffer_height + pos.w * read_buffer_width * read_buffer_height * read_buffer_depth ;
    if (pos.x < 0 || (long)pos.x >= read_buffer_width || pos.y < 0 || (long)pos.y >= read_buffer_height || pos.z < 0 || (long)pos.z >= read_buffer_depth) { // todo: check pos.w
        return (char2){0, 0};
    }
    return (char2){buffer_var[pos_in_buffer],0};
}

inline uchar2 read_buffer4duc(long read_buffer_width, long read_buffer_height, long read_buffer_depth, __global uchar * buffer_var, sampler_t sampler, int4 position )
{
    int4 pos = (int4){position.x, position.y, position.z, position.w};
    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);
        // todo: correct pos.w
    }
    long pos_in_buffer = (long)pos.x + (long)pos.y * read_buffer_width + (long)pos.z * read_buffer_width * read_buffer_height + pos.w * read_buffer_width * read_buffer_height * read_buffer_depth;
    if (pos.x < 0 || (long)pos.x >= read_buffer_width || pos.y < 0 || (long)pos.y >= read_buffer_height || pos.z < 0 || (long)pos.z >= read_buffer_depth) { // todo: check pos.w
        return (uchar2){0, 0};
    }
    return (uchar2){buffer_var[pos_in_buffer],0};
}

inline short2 read_buffer4di(long read_buffer_width, long read_buffer_height, long read_buffer_depth, __global short * buffer_var, sampler_t sampler, int4 position )
{
    int4 pos = (int4){position.x, position.y, position.z, position.w};
    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);
        // todo: correct pos.w
    }
    long pos_in_buffer = (long)pos.x + (long)pos.y * read_buffer_width + (long)pos.z * read_buffer_width * read_buffer_height + pos.w * read_buffer_width * read_buffer_height * read_buffer_depth;
    if (pos.x < 0 || (long)pos.x >= read_buffer_width || pos.y < 0 || (long)pos.y >= read_buffer_height || pos.z < 0 || (long)pos.z >= read_buffer_depth) { // todo: check pos.w
        return (short2){0, 0};
    }
    return (short2){buffer_var[pos_in_buffer],0};
}

inline ushort2 read_buffer4dui(long read_buffer_width, long read_buffer_height, long read_buffer_depth, __global ushort * buffer_var, sampler_t sampler, int4 position )
{
    int4 pos = (int4){position.x, position.y, position.z, position.w};
    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);
        // todo: correct pos.w
    }
    long pos_in_buffer = (long)pos.x + (long)pos.y * read_buffer_width + (long)pos.z * read_buffer_width * read_buffer_height + pos.w * read_buffer_width * read_buffer_height * read_buffer_depth;
    if (pos.x < 0 || (long)pos.x >= read_buffer_width || pos.y < 0 || (long)pos.y >= read_buffer_height || pos.z < 0 || (long)pos.z >= read_buffer_depth) { // todo: check pos.w
        return (ushort2){0, 0};
    }
    return (ushort2){buffer_var[pos_in_buffer],0};
}

inline float2 read_buffer4df(long read_buffer_width, long read_buffer_height, long read_buffer_depth, __global float* buffer_var, sampler_t sampler, int4 position )
{
    int4 pos = (int4){position.x, position.y, position.z, position.w};
    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);
        // todo: correct pos.w
    }
    long pos_in_buffer = (long)pos.x + (long)pos.y * read_buffer_width + (long)pos.z * read_buffer_width * read_buffer_height + pos.w * read_buffer_width * read_buffer_height * read_buffer_depth;
    if (pos.x < 0 || (long)pos.x >= read_buffer_width || pos.y < 0 || (long)pos.y >= read_buffer_height || pos.z < 0 || (long)pos.z >= read_buffer_depth) { // todo: check pos.w
        return (float2){0, 0};
    }
    return (float2){buffer_var[pos_in_buffer],0};
}

inline void write_buffer4dc(long write_buffer_width, long write_buffer_height, long write_buffer_depth, __global char * buffer_var, int4 pos, char value )
{
    long pos_in_buffer = (long)pos.x + (long)pos.y * write_buffer_width + (long)pos.z * write_buffer_width * write_buffer_height + pos.w * write_buffer_width * write_buffer_height * write_buffer_depth;
    if (pos.x < 0 || (long)pos.x >= write_buffer_width || pos.y < 0 || (long)pos.y >= write_buffer_height || pos.z < 0 || (long)pos.z >= write_buffer_depth) { // todo: check pos.w
        return;
    }
    buffer_var[pos_in_buffer] = value;
}

inline void write_buffer4duc(long write_buffer_width, long write_buffer_height, long write_buffer_depth, __global uchar * buffer_var, int4 pos, uchar value )
{
    long pos_in_buffer = (long)pos.x + (long)pos.y * write_buffer_width + (long)pos.z * write_buffer_width * write_buffer_height + pos.w * write_buffer_width * write_buffer_height * write_buffer_depth;
    if (pos.x < 0 || (long)pos.x >= write_buffer_width || pos.y < 0 || (long)pos.y >= write_buffer_height || pos.z < 0 || (long)pos.z >= write_buffer_depth) { // todo: check pos.w
        return;
    }
    buffer_var[pos_in_buffer] = value;
}

inline void write_buffer4di(long write_buffer_width, long write_buffer_height, long write_buffer_depth, __global short * buffer_var, int4 pos, short value )
{
    long pos_in_buffer = (long)pos.x + (long)pos.y * write_buffer_width + (long)pos.z * write_buffer_width * write_buffer_height + pos.w * write_buffer_width * write_buffer_height * write_buffer_depth;
    if (pos.x < 0 || (long)pos.x >= write_buffer_width || pos.y < 0 || (long)pos.y >= write_buffer_height || pos.z < 0 || (long)pos.z >= write_buffer_depth) { // todo: check pos.w
        return;
    }
    buffer_var[pos_in_buffer] = value;
}

inline void write_buffer4dui(long write_buffer_width, long write_buffer_height, long write_buffer_depth, __global ushort * buffer_var, int4 pos, ushort value )
{
    long pos_in_buffer = (long)pos.x + (long)pos.y * write_buffer_width + (long)pos.z * write_buffer_width * write_buffer_height + pos.w * write_buffer_width * write_buffer_height * write_buffer_depth;
    if (pos.x < 0 || (long)pos.x >= write_buffer_width || pos.y < 0 || (long)pos.y >= write_buffer_height || pos.z < 0 || (long)pos.z >= write_buffer_depth) { // todo: check pos.w
        return;
    }
    buffer_var[pos_in_buffer] = value;
}

inline void write_buffer4df(long write_buffer_width, long write_buffer_height, long write_buffer_depth, __global float* buffer_var, int4 pos, float value )
{
    long pos_in_buffer = (long)pos.x + (long)pos.y * write_buffer_width + (long)pos.z * write_buffer_width * write_buffer_height + pos.w * write_buffer_width * write_buffer_height * write_buffer_depth;
    if (pos.x < 0 || (long)pos.x >= write_buffer_width || pos.y < 0 || (long)pos.y >= write_buffer_height || pos.z < 0 || (long)pos.z >= write_buffer_depth) { // todo: check pos.w
        return;
    }
    buffer_var[pos_in_buffer] = value;
}

inline char2 read_buffer3dc(long read_buffer_width, long read_buffer_height, long 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);
    }
    long pos_in_buffer = (long)pos.x + (long)pos.y * read_buffer_width + (long)pos.z * read_buffer_width * read_buffer_height;
    if (pos.x < 0 || (long)pos.x >= read_buffer_width || pos.y < 0 || (long)pos.y >= read_buffer_height || pos.z < 0 || (long)pos.z >= read_buffer_depth) {
        return (char2){0, 0};
    }
    return (char2){buffer_var[pos_in_buffer],0};
}

inline uchar2 read_buffer3duc(long read_buffer_width, long read_buffer_height, long 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);
    }
    long pos_in_buffer = (long)pos.x + (long)pos.y * read_buffer_width + (long)pos.z * read_buffer_width * read_buffer_height;
    if (pos.x < 0 || (long)pos.x >= read_buffer_width || pos.y < 0 || (long)pos.y >= read_buffer_height || pos.z < 0 || (long)pos.z >= read_buffer_depth) {
        return (uchar2){0, 0};
    }
    return (uchar2){buffer_var[pos_in_buffer],0};
}

inline short2 read_buffer3di(long read_buffer_width, long read_buffer_height, long 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);
    }
    long pos_in_buffer = (long)pos.x + (long)pos.y * read_buffer_width + (long)pos.z * read_buffer_width * read_buffer_height;
    if (pos.x < 0 || (long)pos.x >= read_buffer_width || pos.y < 0 || (long)pos.y >= read_buffer_height || pos.z < 0 || (long)pos.z >= read_buffer_depth) {
        return (short2){0, 0};
    }
    return (short2){buffer_var[pos_in_buffer],0};
}

inline ushort2 read_buffer3dui(long read_buffer_width, long read_buffer_height, long 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);
    }
    long pos_in_buffer = (long)pos.x + (long)pos.y * read_buffer_width + (long)pos.z * read_buffer_width * read_buffer_height;
    if (pos.x < 0 || (long)pos.x >= read_buffer_width || pos.y < 0 || (long)pos.y >= read_buffer_height || pos.z < 0 || (long)pos.z >= read_buffer_depth) {
        return (ushort2){0, 0};
    }
    return (ushort2){buffer_var[pos_in_buffer],0};
}

inline float2 read_buffer3df(long read_buffer_width, long read_buffer_height, long 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);
    }
    long pos_in_buffer = (long)pos.x + (long)pos.y * read_buffer_width + (long)pos.z * read_buffer_width * read_buffer_height;
    if (pos.x < 0 || (long)pos.x >= read_buffer_width || pos.y < 0 || (long)pos.y >= read_buffer_height || pos.z < 0 || (long)pos.z >= read_buffer_depth) {
        return (float2){0, 0};
    }
    return (float2){buffer_var[pos_in_buffer],0};
}

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

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

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

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

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

inline char2 read_buffer2dc(long read_buffer_width, long read_buffer_height, long 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);
    }
    long pos_in_buffer = (long)pos.x + (long)pos.y * read_buffer_width;
    if (pos.x < 0 || (long)pos.x >= read_buffer_width || pos.y < 0 || (long)pos.y >= read_buffer_height) {
        return (char2){0, 0};
    }
    return (char2){buffer_var[pos_in_buffer],0};
}

inline uchar2 read_buffer2duc(long read_buffer_width, long read_buffer_height, long 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);
    }
    long pos_in_buffer = (long)pos.x + (long)pos.y * read_buffer_width;
    if (pos.x < 0 || (long)pos.x >= read_buffer_width || pos.y < 0 || (long)pos.y >= read_buffer_height) {
        return (uchar2){0, 0};
    }
    return (uchar2){buffer_var[pos_in_buffer],0};
}

inline short2 read_buffer2di(long read_buffer_width, long read_buffer_height, long 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);
    }
    long pos_in_buffer = (long)pos.x + (long)pos.y * read_buffer_width;
    if (pos.x < 0 || (long)pos.x >= read_buffer_width || pos.y < 0 || (long)pos.y >= read_buffer_height) {
        return (short2){0, 0};
    }
    return (short2){buffer_var[pos_in_buffer],0};
}

inline ushort2 read_buffer2dui(long read_buffer_width, long read_buffer_height, long 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);
    }
    long pos_in_buffer = (long)pos.x + (long)pos.y * read_buffer_width;
    if (pos.x < 0 || (long)pos.x >= read_buffer_width || pos.y < 0 || (long)pos.y >= read_buffer_height) {
        return (ushort2){0, 0};
    }
    return (ushort2){buffer_var[pos_in_buffer],0};
}

inline float2 read_buffer2df(long read_buffer_width, long read_buffer_height, long 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);
    }
    long pos_in_buffer = (long)pos.x + (long)pos.y * read_buffer_width;
    if (pos.x < 0 || (long)pos.x >= read_buffer_width || pos.y < 0 || (long)pos.y >= read_buffer_height) {
        return (float2){0, 0};
    }
    return (float2){buffer_var[pos_in_buffer],0};
}

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

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

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

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

inline void write_buffer2df(long write_buffer_width, long write_buffer_height, long write_buffer_depth, __global float* buffer_var, int2 pos, float value )
{
    long pos_in_buffer = (long)pos.x + (long)pos.y * write_buffer_width;
    if (pos.x < 0 || (long)pos.x >= write_buffer_width || pos.y < 0 || (long)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 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_IMAGE_2D(a,b,c)   write_buffer2dui(GET_IMAGE_WIDTH(a),GET_IMAGE_HEIGHT(a),GET_IMAGE_DEPTH(a),a,b,c)
#define DTYPE_IMAGE_IN_3D   __global ushort*
#define GET_IMAGE_OUT_HEIGHT(image_key)     IMAGE_SIZE_ ## image_key ## _HEIGHT
#define DTYPE_IN    ushort
#define DTYPE_IMAGE_OUT_3D  __global ushort*
#define GET_IMAGE_OUT_DEPTH(image_key)  IMAGE_SIZE_ ## image_key ## _DEPTH
#define GET_IMAGE_HEIGHT(image_key)     IMAGE_SIZE_ ## image_key ## _HEIGHT
#define IMAGE_SIZE_dst_HEIGHT   256
#define GET_IMAGE_IN_HEIGHT(image_key)  IMAGE_SIZE_ ## image_key ## _HEIGHT
#define READ_IMAGE_3D(a,b,c)    read_buffer3dui(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 READ_IMAGE_2D(a,b,c)    read_buffer2dui(GET_IMAGE_WIDTH(a),GET_IMAGE_HEIGHT(a),GET_IMAGE_DEPTH(a),a,b,c)
#define IMAGE_SIZE_dst_WIDTH    256
#define GET_IMAGE_OUT_WIDTH(image_key)  IMAGE_SIZE_ ## image_key ## _WIDTH
#define GET_IMAGE_IN_WIDTH(image_key)   IMAGE_SIZE_ ## image_key ## _WIDTH
#define WRITE_IMAGE_3D(a,b,c)   write_buffer3dui(GET_IMAGE_WIDTH(a),GET_IMAGE_HEIGHT(a),GET_IMAGE_DEPTH(a),a,b,c)
#define DTYPE_IMAGE_IN_2D   __global ushort*
#define GET_IMAGE_IN_DEPTH(image_key)   IMAGE_SIZE_ ## image_key ## _DEPTH
#define DTYPE_IMAGE_OUT_2D  __global ushort*
#define IMAGE_SIZE_src_WIDTH    256
#define CONVERT_DTYPE_OUT   clij_convert_ushort_sat
#define MAX_ARRAY_SIZE  1000
#define GET_IMAGE_WIDTH(image_key)  IMAGE_SIZE_ ## image_key ## _WIDTH
#define IMAGE_SIZE_dst_DEPTH    129
#define DTYPE_OUT   ushort
#define IMAGE_SIZE_src_HEIGHT   256
#define IMAGE_SIZE_src_DEPTH    129

 //###########################################################################
// Source: 'duplication.cl' relative to Kernels
__kernel void copy_3d (DTYPE_IMAGE_OUT_3D dst, DTYPE_IMAGE_IN_3D src) {
  const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;

  const int dx = get_global_id(0);
  const int dy = get_global_id(1);
  const int dz = get_global_id(2);

  const int4 pos = (int4){dx,dy,dz,0};

  const DTYPE_IN out = READ_IMAGE_3D(src,sampler,pos).x;
  WRITE_IMAGE_3D(dst, pos, CONVERT_DTYPE_OUT(out));
}

__kernel void copy_2d(DTYPE_IMAGE_OUT_2D dst, DTYPE_IMAGE_IN_2D src) {
  const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;

  const int dx = get_global_id(0);
  const int dy = get_global_id(1);

  const int2 pos = (int2){dx,dy};

  const DTYPE_IN out = READ_IMAGE_2D(src,sampler,pos).x;
  WRITE_IMAGE_2D(dst,pos, CONVERT_DTYPE_OUT(out));
}

__kernel void copySlice(DTYPE_IMAGE_OUT_2D dst, DTYPE_IMAGE_IN_3D src, int slice) {
  const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;

  const int dx = get_global_id(0);
  const int dy = get_global_id(1);

  const int4 pos4 = (int4){dx,dy,slice,0};
  const int2 pos2 = (int2){dx,dy};

  const DTYPE_IN out = READ_IMAGE_3D(src,sampler,pos4).x;
  WRITE_IMAGE_2D(dst,pos2, CONVERT_DTYPE_OUT(out));
}

__kernel void putSliceInStack(DTYPE_IMAGE_OUT_3D dst, DTYPE_IMAGE_IN_2D src, int slice) {
  const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;

  const int dx = get_global_id(0);
  const int dy = get_global_id(1);

  const int2 pos2 = (int2){dx,dy};
  const int4 pos4 = (int4){dx,dy,slice,0};

  const DTYPE_IN out = READ_IMAGE_2D(src,sampler,pos2).x;
  WRITE_IMAGE_3D(dst,pos4, CONVERT_DTYPE_OUT(out));
}

__kernel void crop_3d(DTYPE_IMAGE_OUT_3D dst, DTYPE_IMAGE_IN_3D src, int start_x, int start_y, int start_z) {
  const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;

  const int dx = get_global_id(0);
  const int dy = get_global_id(1);
  const int dz = get_global_id(2);

  const int sx = start_x + dx;
  const int sy = start_y + dy;
  const int sz = start_z + dz;

  const int4 dpos = (int4){dx,dy,dz,0};
  const int4 spos = (int4){sx,sy,sz,0};

  const DTYPE_IN out = READ_IMAGE_3D(src,sampler,spos).x;
  WRITE_IMAGE_3D(dst,dpos, CONVERT_DTYPE_OUT(out));
}

__kernel void crop_2d(DTYPE_IMAGE_OUT_2D dst, DTYPE_IMAGE_IN_2D src, int start_x, int start_y) {
  const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;

  const int dx = get_global_id(0);
  const int dy = get_global_id(1);

  const int sx = start_x + dx;
  const int sy = start_y + dy;

  const int2 dpos = (int2){dx,dy};
  const int2 spos = (int2){sx,sy};

  const DTYPE_IN out = READ_IMAGE_2D(src,sampler,spos).x;
  WRITE_IMAGE_2D(dst,dpos, CONVERT_DTYPE_OUT(out));
}

CLIJ Error: For support please contact the CLIJ2 developers via the forum on https://image.sc or create an issue on https://github.com/clij/clij2/issues .
Therefore, please report the complete error message, the code snippet or workflow you were running, an example image if possible and details about your graphics hardware.

net.haesleinhuepf.clij.clearcl.exceptions.ClearCLException: problem while setting argument 'src' at index 7 

    at net.haesleinhuepf.clij.clearcl.ClearCLKernel.setArgumentsInternal(ClearCLKernel.java:426)
    at net.haesleinhuepf.clij.clearcl.ClearCLKernel.lambda$run$0(ClearCLKernel.java:481)
    at net.haesleinhuepf.clij.clearcl.util.ElapsedTime.measure(ElapsedTime.java:97)
    at net.haesleinhuepf.clij.clearcl.util.ElapsedTime.measure(ElapsedTime.java:64)
    at net.haesleinhuepf.clij.clearcl.ClearCLKernel.run(ClearCLKernel.java:477)
    at net.haesleinhuepf.clij.clearcl.ClearCLKernel.run(ClearCLKernel.java:459)
    at net.haesleinhuepf.clij.clearcl.util.CLKernelExecutor.lambda$enqueue$1(CLKernelExecutor.java:267)
    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.clearcl.util.CLKernelExecutor.enqueue(CLKernelExecutor.java:266)
    at net.haesleinhuepf.clij2.CLIJ2.lambda$executeSubsequently$0(CLIJ2.java:466)
    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:456)
    at net.haesleinhuepf.clij2.CLIJ2.executeSubsequently(CLIJ2.java:443)
    at net.haesleinhuepf.clij2.CLIJ2.executeSubsequently(CLIJ2.java:438)
    at net.haesleinhuepf.clij2.CLIJ2.execute(CLIJ2.java:423)
    at net.haesleinhuepf.clij2.utilities.CLIJUtilities.executeSeparableKernel(CLIJUtilities.java:77)
    at net.haesleinhuepf.clij2.plugins.Minimum3DBox.minimumBox(Minimum3DBox.java:47)
    at net.haesleinhuepf.clij2.plugins.Minimum3DBox.minimum3DBox(Minimum3DBox.java:43)
    at net.haesleinhuepf.clij2.CLIJ2Ops.minimum3DBox(CLIJ2Ops.java:4250)
    at net.haesleinhuepf.clij2.plugins.TopHatBox.topHatBox(TopHatBox.java:44)
    at net.haesleinhuepf.clij2.CLIJ2Ops.topHatBox(CLIJ2Ops.java:688)
    at net.haesleinhuepf.clij2.plugins.TopHatBox.executeCL(TopHatBox.java:34)
    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:4967)
    at ij.macro.Functions.getStringFunction(Functions.java:278)
    at ij.macro.Interpreter.getStringTerm(Interpreter.java:1475)
    at ij.macro.Interpreter.getString(Interpreter.java:1453)
    at ij.macro.Interpreter.doStatement(Interpreter.java:333)
    at ij.macro.Interpreter.doStatements(Interpreter.java:264)
    at ij.macro.Interpreter.run(Interpreter.java:160)
    at ij.macro.Interpreter.run(Interpreter.java:93)
    at ij.macro.Interpreter.run(Interpreter.java:104)
    at ij.plugin.Macro_Runner.runMacro(Macro_Runner.java:161)
    at ij.IJ.runMacro(IJ.java:153)
    at ij.IJ.runMacro(IJ.java:142)
    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:165)
    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)
Caused by: net.haesleinhuepf.clij.clearcl.exceptions.ClearCLException
    at net.haesleinhuepf.clij.clearcl.backend.BackendUtils.checkExceptions(BackendUtils.java:178)
    at net.haesleinhuepf.clij.clearcl.backend.jocl.ClearCLBackendJOCL.setKernelArgument(ClearCLBackendJOCL.java:614)
    at net.haesleinhuepf.clij.clearcl.ClearCLKernel.setArgumentsInternal(ClearCLKernel.java:417)
    ... 54 more
Caused by: java.lang.NullPointerException
    at net.haesleinhuepf.clij.clearcl.backend.jocl.ClearCLBackendJOCL.lambda$setKernelArgument$20(ClearCLBackendJOCL.java:707)
    at net.haesleinhuepf.clij.clearcl.backend.BackendUtils.checkExceptions(BackendUtils.java:171)
    ... 56 more
net.haesleinhuepf.clij.clearcl.exceptions.ClearCLException
    at net.haesleinhuepf.clij.clearcl.backend.BackendUtils.checkExceptions(BackendUtils.java:178)
    at net.haesleinhuepf.clij.clearcl.backend.jocl.ClearCLBackendJOCL.enqueueReadFromBuffer(ClearCLBackendJOCL.java:790)
    at net.haesleinhuepf.clij.clearcl.ClearCLBuffer.writeTo(ClearCLBuffer.java:400)
    at net.haesleinhuepf.clij.clearcl.ClearCLBuffer.writeTo(ClearCLBuffer.java:366)
    at net.haesleinhuepf.clij.converters.implementations.ClearCLBufferToImagePlusConverter.convert(ClearCLBufferToImagePlusConverter.java:61)
    at net.haesleinhuepf.clij.converters.implementations.ClearCLBufferToImagePlusConverter.convert(ClearCLBufferToImagePlusConverter.java:26)
    at net.haesleinhuepf.clij.CLIJ.convert(CLIJ.java:475)
    at net.haesleinhuepf.clij.CLIJ.show(CLIJ.java:358)
    at net.haesleinhuepf.clij.macro.CLIJHandler.pullFromGPU(CLIJHandler.java:253)
    at net.haesleinhuepf.clij2.plugins.Pull.executeCL(Pull.java:24)
    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:4967)
    at ij.macro.Functions.getStringFunction(Functions.java:278)
    at ij.macro.Interpreter.getStringTerm(Interpreter.java:1475)
    at ij.macro.Interpreter.getString(Interpreter.java:1453)
    at ij.macro.Interpreter.doStatement(Interpreter.java:333)
    at ij.macro.Interpreter.doStatements(Interpreter.java:264)
    at ij.macro.Interpreter.run(Interpreter.java:160)
    at ij.macro.Interpreter.run(Interpreter.java:93)
    at ij.macro.Interpreter.run(Interpreter.java:104)
    at ij.plugin.Macro_Runner.runMacro(Macro_Runner.java:161)
    at ij.IJ.runMacro(IJ.java:153)
    at ij.IJ.runMacro(IJ.java:142)
    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:165)
    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)
Caused by: java.lang.NullPointerException
    at net.haesleinhuepf.clij.clearcl.backend.jocl.ClearCLBackendJOCL.lambda$enqueueReadFromBuffer$24(ClearCLBackendJOCL.java:792)
    at net.haesleinhuepf.clij.clearcl.backend.BackendUtils.checkExceptions(BackendUtils.java:171)
    ... 40 more

The bug cannot be reproduced on our Windows/Linux/Mac test systems. We keep it here for later reference. Thanks to Dominic Waithe for reporting.