GPU 2D shared memory dynamic allocation

Question

I am aware of the dynamic allocation when 1D arrays are used, but how can it be done when 2D arrays are used?

myKernel<<<blocks, threads,sizeofSharedMemoryinBytes>>>();
         ....

__global__ void myKernerl(){
 __shared__ float sData[][];
     .....
}

Say I want to allocate a 2D shared memory array:

__shared__ float sData[32][32];

How can it be done dynamically? would be:

myKernel<<< blocks, threads, sizeof(float)*32*32 >>>();

Your statically declared "2D shared memory array" isn't two dimensional, it is just linear memory and the compiler generates row-major order access to it. Based on your endless number of questions about multidimensional arrays, perhaps it is time to sit down with some reference material and learn about how arrays work in C++.. — talonmies

stuhlo stuhlo · Accepted Answer · 2012-11-26T19:49:19

As you have correctly written you have to specify size of dynamically allocated shared memory before each kernel calling in configuration of execution (in <<<blocks, threads, sizeofSharedMemoryinBytes>>>). This specifies the number of bytes in shared memory that is dynamically allocated per block for this call in addition to the statically allocated memory. IMHO there is no way to access such memory as 2D array, you have to use 1D array and use it like 2D. Last think, don't forget qualifier extern. So your code should look like this:

   sizeofSharedMemoryinBytes = dimX * dimY * sizeof(float);

   myKernel<<<blocks, threads,sizeofSharedMemoryinBytes>>>();
     ....

   __global__ void myKernerl() {

       extern __shared__ float sData[];
       .....
       sData[dimX * y + x] = ...
   }

GPU 2D shared memory dynamic allocation

1 Answers