GPU 2D shared memory dynamic allocation

Question

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

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

__global__ void myKernel(){
    __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 it 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, Nov 02 '12 at 13:08

score 6 · Answer 1 · answered Nov 26 '12 at 19:49

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 Answers1