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

cl25kdpy  于 2023-05-20  发布在  其他
关注(0)|答案(1)|浏览(178)

我想在CUDA的共享内存中分配一个2D数组。我知道要分配一个1D共享内存数组,你必须将每个块的大小作为一个参数传递给内核。
我还知道在共享内存中动态创建实际的2D数组是不可能的。
然而,我想知道如果其中一个维度是已知的,是否可以做到这一点。

extern __shared__ int array[COMPILE_TIME_SIZE][];

这可以做到吗?如果是这样,我如何传递第二维的大小?

2eafrhcq

2eafrhcq1#

这样做是不可能的,因为编译器可能无法正确地实现寻址。
但是,可以只在编译时设置第二个参数。下面是一个示例代码:

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

数组的大小由三尖括号表示法中的共享内存参数定义。因此,通过将以字节为单位的共享存储器大小除以单个条目的以字节为单位的大小来推导第二维度的大小。

相关问题