passing a jagged array to a kernel in opencl

Question

So I know how to do this in CUDA, but basically, i want to pass a small number (0-5, variable) of _global ptrs to a function, then load those pointers into either a local or a private (because of the small number and the fact that I already have a local memory fence in the kernel, i'm not sure which is fastest, I'm going to determine this experimentally after I get it working). So i wrote the kernel like so:

__kernel foo(
  __global int* img,
  __global int** img_history,
  __private int** private_history,
  uint history_length)//could be local
{
    for (int i = 0; i < history_length; i++)
       private_history[i] = img_history[i];
}

To clarify, in cuda I do this like this

__global__ foo(int* img, int** img_history, uint history_length)
{
   int* private_history[10];//max values 10
   for (int i = 0; i < history_length; i++)
      private_history[i] = img_history[i];
 }

and to load it

int** host_array = new int*[history_length];
for (int i = 0; i < history_length; i++)
    cudaMalloc(host_array+i,size);
int** device_array;
cudaMalloc(&device_array,sizeof(int*)*history_length);
cudaMemcpy(device_array, host_array,sizeof(int*)*history_length,cudaMemcpyHostToDevice)

However, I get the error error: invalid address space for pointee of pointer argument to __kernel function. What is the correct way to do this?

Answer 1

I dunno how you are doing in CUDA. But this is completely forbiden as an argument to an OpenCL kernel.

You can't copy to the device a pointer value and then use it directly because the memory adresses are different.

In order to do it you need to:

1. Copy only the indexes of img_history (not the pointer) referencing an image table.
2. Operate with thouse indexes as you like (integer operation).
3. Use these indexes to access the image table or do whatever you want. If you need to acess the img with these indexes, then it has to be a parameter to the kernel. And you have to copy all of it. (the full length img array)

Example:

__kernel foo(
  __global int* img,
  __global int* img_history,
  __private int* private_history,
  uint history_length)//could be local
{
    for (int i = 0; i < history_length; i++)
       private_history[i] = img_history[i];

    /* img[private_history[i]] */ //Use it as you wish
}