Cuda ::输入函数共享数据

有谁知道为什么以下函数使用共享数据的16 432 B？在我看来，这应该是：32x32x8x2 = 16 384乙Cuda ::输入函数共享数据

__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)

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

来源

2014-02-20 b1es

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

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

尝试编译为CC 2.x版本的设备：

nvcc -arch=sm_20 ...

来源

2014-02-20 17:50:54

谢谢罗伯特。我已经编译了nvcc -arch = sm_30。我只是好奇它为什么这样工作。问候 – b1es

Cuda ::输入函数共享数据

回答

相关问题