0

假设我们有一个内核，它需要一个唯一的块索引，并且应该尽可能地扩展，以便它应该在 3D 网格中扩展。

计算看起来很复杂，让一个线程执行它并将其存储在共享内存中。这是个好主意吗？
在所有文献中它总是存储在寄存器中，但是共享内存的缺点是什么？

我不确定，但共享内存读写访问是 4 个周期，而寄存器是 1 个周期？

代替：

__global__ kernel()
{
    //get unique 3D block index
    const unsigned long long int blockId = blockIdx.x //1D
        + blockIdx.y * gridDim.x //2D
        + gridDim.x * gridDim.y * blockIdx.z; //3D
}

也许使用：（假设只使用块的 x 维度）

__global__ kernel()
{
    __shared__ unsigned long long int blockId_s;

    if(threadIdx.x == 0)
        blockId_s = blockIdx.x //1D
            + blockIdx.y * gridDim.x //2D
            + gridDim.x * gridDim.y * blockIdx.z; //3D
    __syncthreads();
}

这将为每个线程节省一个寄存器，这在计算能力 1.x 中是昂贵的。

我没有测试，也不知道它对性能是好是坏。__syncthreads()cc 1.x 上的另一个可用寄存器是一个参数，但使用语句时性能应该会慢一些。

4

2 回答 2

2

在这种情况下，答案是肯定的，但主要是因为生成的代码使用了更少的寄存器，这允许更高的整体占用率和一些速度。尽管代码在第一个 warp 中有一些分支分歧，加上一个同步原语。

但是，这不应被视为通用规则，唯一确定的方法是编写代码并在目标 GPU 上对其进行基准测试。

此答案是作为社区 wiki 条目从评论中添加的，以将此问题从未回答的队列中删除。

于 2014-05-26T05:55:42.743 回答
-2

在共享人的情况下，每次访问 blockId_s 时都会发生 n 路银行冲突 (n=WarpSize)。[CUDA C 编程指南，5.3.2.3]

并且 + shared mem 比 register 慢。

于 2012-04-28T19:43:33.977 回答