Cuda::入口函数共享数据

Question

有人知道为什么下面的函数使用 16,432 B 的共享数据吗？ 在我看来应该是:2 x 32 x 32 x 8 B = 16,384 B

__global__ void matrixMulKernel(double *c, const double *a, const double *b, unsigned int size)
{
    __shared__ double as[32][32];
    __shared__ double bs[32][32];
    unsigned int bx = blockIdx.x, by = blockIdx.y;
    unsigned int tx = threadIdx.x, ty = threadIdx.y;
    unsigned int row = bx * TILE_WIDTH + tx;
    unsigned int col = by * TILE_WIDTH + ty;
    double Pval = 0.0;
    for(unsigned int q = 0; q < size / TILE_WIDTH; q++)
    {
        as[tx][ty] = a[row * size + q * TILE_WIDTH + ty];
        bs[ty][tx] = b[(q * TILE_WIDTH + tx) * size + col];
        __syncthreads();
        
        for(unsigned int k = 0; k < TILE_WIDTH; k++)
            Pval += as[tx][k] * bs[k][ty];
        __syncthreads();
    }

    c[row * size + col] = Pval;
}

编译器给出以下错误:

Entry function '_Z15matrixMulKernelPdPKdS1_j' uses too much shared data (0x4030 bytes, 0x4000 max)

我感兴趣的是为什么会这样，而不是作为解决方法:)

Answer 1

您可能正在为 cc 1.x 设备进行编译。 documentation表示 cc 1.x 设备的全局内核参数通过共享内存传递。

因此，您的显式 __shared__ 声明有 16,384 字节。 其余部分来自显式内核参数所需的 28 字节(假设 64 位目标)以及通过共享内存通信的其他开销。

尝试针对 cc 2.x 设备进行编译:

nvcc -arch=sm_20 ...