CUDA共享内存占用的空间是所需空间的两倍

Question

我只是注意到我的CUDA内核使用的空间是"理论"计算的空间的两倍.例如

__global__ void foo( )
{
    __shared__ double t;
    t = 1;
}

PTX信息显示:
ptxas info :_Z3foov的函数属性,0字节堆栈帧,0字节溢出存储,0字节溢出加载
ptxas info:使用4个寄存器,16字节smem,32字节cmem [0]

但双人的大小只有8.

更多例子:

__global__ void foo( )
{
    __shared__ int t[1024];
    t[0] = 1;
}

ptxas info:使用3个寄存器,8192个字节的smem,32个字节的cmem [0]

有人能解释为什么吗？

Answer 1

看来这个问题在当前的 CUDA 编译器中已经消失了。

__shared__ int a[1024];

用命令编译

nvcc -m64 -Xptxas -v -ccbin /opt/gcc-4.6.3/bin/g++-4.6.3 shmem.cu

给出

ptxas info    : Used 1 registers, 4112 bytes smem, 4 bytes cmem[1]

这种情况下会有一些共享内存开销，但使用量并没有增加一倍。