c++ - 二维数组的CUDA动态共享内存分配

Question

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

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

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

extern __shared__ int array[COMPILE_TIME_SIZE][];

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

Answer 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 ;
}

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