In python how do I pass a scalar argument to an OpenCL kernel?

Question

I am using the python bindings for OpenCL and I have written a kernel that expects a scalar (float) argument, but I can not figure out the proper way to pass it in.

If I simply invoke

prg.bloop(queue, [width,height], None, centers_g, 16/9, result_g)

I get this error:

pyopencl.cffi_cl.LogicError: when processing argument #2 (1-based): 'float' does not support the buffer interface

If I wrap it with numpy.float32(16/9) then the kernel behaves as if it were passed 0 instead of 1.7777777 .

Here is more of the source code in case that helps you understand what I'm doing.

def mission2(cells, width, height):
    ctx = cl.create_some_context()
    queue = cl.CommandQueue(ctx)

    centers = numpy.array(list(cells), dtype=numpy.float32)

    centers_g = cl.Buffer(ctx, cl.mem_flags.READ_ONLY|cl.mem_flags.COPY_HOST_PTR, hostbuf = centers)

    prg = cl.Program(ctx, """
__kernel void test1(__global char *out) {
    out[0] = 77;
}

float sphere(float r, float x, float y)
{
    float q = r*r-(x*x+y*y);
    if (q<0)
        return 0;
    return sqrt(q);
}

__kernel void bloop(__global const float centers[][6], float aspect, __global char * out)
{
    int u = get_global_id(0);
    int v = get_global_id(1);
    int width = get_global_size(0);
    int height = get_global_size(1);

    float x = u/(float)width * aspect;
    float y = v/(float)height;

    float max = sphere(0.3, x-centers[0][0], y-centers[0][1]);

    int idx = u+v*width;
    out[idx] = 255*max;

}

    """).build()

    result_g = cl.Buffer(ctx, cl.mem_flags.WRITE_ONLY, width*height)


    if False:
        prg.test1(queue, [width,height], None, result_g)
    else:
        prg.bloop(queue, [width,height], None, centers_g, numpy.float32(16/9), result_g)


    result = numpy.zeros([height,width], dtype=numpy.uint8)
    future = cl.enqueue_copy(queue, result, result_g)
    future.wait()

    print(result)

    imsave("/tmp/bloop.png", result, cmap = matplotlib.pyplot.get_cmap('gray'))

Answer 1

In order to pass scalars to an OpenCL kernel you can use the set_scalar_arg_dtypes function like this:

kernel = prg.bloop
kernel.set_scalar_arg_dtypes( [None, numpy.float32, None] )
kernel(queue, [width, height], None, centers_g, aspect, result_g)

The manual page clearly states that it will not work if you code it like this:

prg.bloop.set_scalar_arg_dtypes( [None, numpy.float32, None] )
# this will fail:
prg.bloop(queue, [width, height], None, centers_g, 16/9, result_g)

because "The information set by this rountine is attached to a single kernel instance. A new kernel instance is created every time you use program.kernel attribute access."