Manage multiarray in kernel

[KaiK]

Hi!

I was wondering to know if there’s any way of managing multiarrays in the kernel as arrays, and not as a pointers. Specially when the size of the multiarray is decided in host because it does not depend on a constant.

What I mean is that, as far as I know, the way to manage multiarrays in kernel is passing the pointer to first possition as argument, and then using the pointer as if it was a 1D array:

[i] global void myKernel ( int * my3dArrayIn, …)

{

//instead of my3dArayIn[threadIdx.x][threadIdx.y][threadIdx.z]

   //I have to do

my3dArrayIn[threadIdx.zblockDim.xblockDim.y + threadIdx.y*blockDim.x + threadIdx.x]

…

  }[/i]

Is there any option to use the array as I’ve wrote in the comment?

I cannot pass the array as argument like this:

global void myKernel ( int my3dArrayIn, …)

because size is declared in the host. I’ve seen the code about the “external shared variable array”, but it’s not the case because here the array just have the information, the problem is not to write in it but reading from it, so even if I declare it as an external variable array I’ll made the same access as I’ve wrote in the code above.

Any help?

Thanks in advanced!

Enrique

[YDD]

You could set it up that way, and move each block of pointers to the GPU, but you’d be in for a lot of pain. Pointer chasing on the GPU (and remember that C doesn’t have multi-dimensional arrays; your ‘3D’ array is actually an array of pointers, each of which points to another array of pointers, and each of those pointers points to a 1D array of numbers) is very, very bad news. Each pointer dereference is going to take a full 200 cycles (no cache), and after the first, you’ll probably stop coalescing your reads.