二维数组的CUDA动态共享内存分配答案

【问题标题】：CUDA dynamic shared memory allocation of 2d array二维数组的CUDA动态共享内存分配
【发布时间】：2016-04-26 18:19:53
【问题描述】：

我想在 CUDA 的共享内存中分配一个二维数组。我知道要分配一维共享内存数组，您必须将每个块的大小作为参数传递给内核。

我也知道不可能 2 在共享内存中动态创建一个实际的二维数组。

但是我想知道如果其中一个维度是已知的，是否可以这样做。

extern __shared__ int array[COMPILE_TIME_SIZE][];

这可以吗？如果是这样，我如何传递第二维的大小？

【问题讨论】：

标签： c++ arrays cuda

【解决方案1】：

这样做是不可能的，因为编译器可能无法正确实现寻址，因此首先将已知维度（最高顺序 - 第一个方括号条目）放在首位。

但是，可以只在编译时设置第二个参数。这是一个示例代码：

extern __shared__ int shared2Darray[][17] ;

__global__ void kernel(int* output)
{
    shared2Darray[threadIdx.y][threadIdx.x] = threadIdx.x + 2*threadIdx.y ;
    __syncthreads();
    output [threadIdx.y * blockDim.x + threadIdx.x] = shared2Darray[threadIdx.y][threadIdx.x] ;
    __syncthreads();
}

int main()
{
    int* h_output, *d_output ;

    cudaMalloc(&d_output, 16*16*sizeof(int));

    kernel<<<1, dim3(16,16,1), 16*17*sizeof(int)>>> (d_output) ;

    h_output = new int[16*16] ;
    cudaMemcpy (h_output, d_output, 16*16*sizeof(int), cudaMemcpyDeviceToHost) ;

    cudaDeviceReset();

    for (int x = 0 ; x < 16 ; ++x)
    {
        for (int y = 0 ; y < 16 ; ++y)
        {
            if (h_output[y*16+x] != x+2*y)
                printf ("ERROR\n");
        }
    }

    printf ("DONE\n");

    delete[] h_output ;

    return 0 ;
}

数组的大小由三尖括号表示法中的共享内存参数定义。因此，第二维的大小是通过将共享内存大小（以字节为单位）除以单个条目的大小（以字节为单位）得出的。

【讨论】：