2D Array as OpenCL kernel argument

Question

Pretty straight forward question: How would one use a 2D array as an OpenCL kernel argument?

Common sense suggests using

__kernel void main(__global <datatype> **<name>),

however the compiler doesn't seem to be quite amused by this idea:

kernel parameter cannot be declared as a pointer to a pointer.

Am I overseeing the obvious, or what exactly is it, I am doing wrong here?

Edit:

The hosts (c++) datastructure looks like this:

vector<vector<Element>>,

where Element is a struct containing the indexes of the child nodes inside the very same array. Basicly pointers.

Answer 1

You need to reduce the 2D array down into a 1D array.

Host:

int array[50][50];
int * ptr_to_array_data = array[0];
int width = 50, height = 50;
cl_mem device_array = clCreateBuffer(/*context*/, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, 50 * 50 * sizeof(int), ptr_to_array_data, /*&err*/);
clSetKernelArg(/*kernel*/, 0, sizeof(cl_mem), &device_array);
clSetKernelArg(/*kernel*/, 1, sizeof(cl_int), &width);
clSetKernelArg(/*kernel*/, 2, sizeof(cl_int), &height);

Device:

kernel function(global int * array, int width, int height) {
    int id = get_global_id(0);
    int our_value = array[id];
    int x = id % width; //This will depend on how the memory is laid out in the 2d array. 
    int y = id / width; //If it's not row-major, then you'll need to flip these two statements.
    /*...*/
}

If your 2D array is not stored contiguously in memory like my example implies, you'll need to roll your own function to make sure that the entire memory is stored contiguously in a single heap-allocated object.