Cuda::Entry 函数共享数据

Question

有人知道为什么下面的函数使用 16 432 B 的共享数据吗？ 在我看来应该是：32x32x8x2 = 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 设备。

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

尝试为 cc 2.x 设备编译：

nvcc -arch=sm_20 ...