Passing a matrix to openCL

Hello,

I have an OpenCL kernel where all the threads need to access the same 3x4 matrix. So, I am guessing constant memory would be the best place for such a matrix.

However, I am unable to find an example where this constant memory is used. In CUDA, one has this cudaMemcpyToSymbol where one can bind data from CPU to this constant memory.

How do I achieve something similar in openCL? Sorry for the n00b question.

Thanks,
xarg

Depending on your use case, the easiest way is to put it along your kernel sources in the __constant address space, like so:

const __constant float data[4] = { 1.2f, 3.4f, 5.6f, 7.8f };

__kernel void transfer(__global float *output)
{
    int idx = get_global_id(0); 
    output[idx] = data[idx];
}

Thank you for the reply. The problem is that the matrix is only known at run time and not compile time. So, I read an image and an associated transformation matrix and then I transform every pixel location in the image with this transformation matrix. Hence, every thread needs to access the same matrix but it is known only when I get the input from the user.

Thanks,

xarg

Well, OpenCL code is compiled at run-time, so you could prepend your kernel source with the user matrix and compile this on the fly.

If this is not an option, you can also create a buffer with the CL_MEM_READ_ONLY flag set

cl_mem mem = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, data_size, host_ptr, NULL);
clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &mem);

and pass it to your kernel

__kernel void transfer(__constant float *data, __global float *output)
{
    int idx = get_global_id(0);
    output[idx] = data[idx];
}

Ahhhhhh…yes, that is a great idea. I will do the read only flag bit as that is probably a bit easier to understand for the other developers.

Thank you for the reply. Much appreciated.

I am eally starting to like OpenCL. I will be a total convert as soon as they add support for some C++ specs like templates :slight_smile:

Thank again.

/xarg

ArrayFire is a matrix-based library for OpenCL that I’m working on. You might find that useful.

I agree that OpenCL is in dire need of template support.

Thanks! I will have a look at ArrayFire this week :slight_smile: