1

我正在使用 GTX9800，它报告可用共享内存为 16384 字节

给定以下内核代码，使用 T = int (4 bytes) 运行

template <typename T>
__global__ void foo(unsigned n, T *x)
{
    unsigned idx    = blockIdx.x * blockDim.x + threadIdx.x;
    __shared__ T sx[4096];
    x[idx] = 0;
}

我得到了预期的结果，即最初非零的数组 x 将用零填充。

但是，添加一行不执行任何操作的代码：

template <typename T>
__global__ void foo(unsigned n, T *x)
{
    unsigned idx    = blockIdx.x * blockDim.x + threadIdx.x;
    __shared__ T sx[4096];
    sx[0] = 0;
    x[idx] = 0;
}

现在 x 在调用内核后根本不包含任何零！

但是，如果我将 sx 的大小更改为 <= 4088，我会再次得到预期的结果。

这是怎么回事？我比较困惑。

编辑：

更正错字：将 16384 "KB" 更改为 "bytes"

4

1 回答 1

3

计算能力 1.x 设备上的共享内存大小是每个 SM 16384字节，而不是千字节。

此外，每个块将消耗 16 个字节用于内部目的（存储块索引等），以及内核参数的额外存储空间。

所以不幸的是，你不能在一个块中使用完整的 16kb 共享内存。

对于更高的计算能力，这些数据将存储在其他地方（常量内存和特殊寄存器），因此整个共享内存都在那里可用。

于 2013-04-09T17:32:09.397 回答