不连贯的商店答案

【问题标题】：Incoherent stores不连贯的商店
【发布时间】：2011-09-07 07:23:06
【问题描述】：

为什么这个内核会产生不连贯的存储

__global__ void reverseArrayBlock(int *d_out, int *d_in)
{
    int inOffset  = blockDim.x * blockIdx.x;
    int outOffset = blockDim.x * (gridDim.x - 1 - blockIdx.x);
    int in  = inOffset + threadIdx.x;
    int out = outOffset + (blockDim.x - 1 - threadIdx.x);
    d_out[out] = d_in[in];
}

这个没有

__global__ void reverseArrayBlock(int *d_out, int *d_in)
{
    extern __shared__ int s_data[];

    int inOffset  = blockDim.x * blockIdx.x;
    int in  = inOffset + threadIdx.x;

    // Load one element per thread from device memory and store it 
    // *in reversed order* into temporary shared memory
    s_data[blockDim.x - 1 - threadIdx.x] = d_in[in];

    // Block until all threads in the block have written their data to shared mem
    __syncthreads();

    // write the data from shared memory in forward order, 
    // but to the reversed block offset as before

    int outOffset = blockDim.x * (gridDim.x - 1 - blockIdx.x);

    int out = outOffset + threadIdx.x;
    d_out[out] = s_data[threadIdx.x];
}

我知道第二个正在使用共享内存。但是当我查看 d_out 的指标时，它们在两个内核中似乎是相同的。你能帮我理解一下吗？

【问题讨论】：

标签： cuda

【解决方案1】：

合并要求地址在 warp 中遵循“base + tid”模式，其中 tid 是线程索引的缩写。换句话说，随着 tid 的增加，地址也会增加。您的评论将此称为“转发订单”。在第一个内核中，地址的生成使得随着 tid 增加，地址减少，即访问是“倒序”的。

【讨论】：

请注意，示例代码应该只在 Compute Capability 1.1 和早期设备上产生不连贯的存储。从 1.2 设备开始，这种访问模式是完全合并/连贯的。

【解决方案2】：

在我们开始之前，您需要了解写入共享内存比写入全局内存要便宜得多。

考虑到这一点，假设我们正在反转一个数组 1->32

方法一这样做：写作时线程 1 从位置 x 读取，线程 2 从位置 (x + 1) 读取，线程 3 从位置 (x + 2) 读取......线程 32 从位置 (x + 31) 读取。

您可以在 2 次（如果对齐）或 3 次（如果未对齐）读取中读取整个内存块，因为这些操作是在半扭曲块（16 个线程）中完成的。

写作时线程 1 写入位置 (y + 31)，线程 2 写入位置 (y + 30)，线程 3 写入位置 (y + 29) ...线程 32 写入位置 (y)。

虽然它们正在写入连续的内存块，但它们的顺序是相反的。除非您使用一些最新的硬件（即使使用它我也很怀疑），否则执行该操作需要 32 次写入。

至于第二种情况，您正在对共享内存进行 32 次反向写入和从共享内存进行 32 次反向读取，这并不昂贵。

现在您已经以相反的顺序读取了数据，您可以按照正确的顺序写入全局内存。

线程 1 写入位置 y，线程 2 写入位置 y+1，依此类推。

最重要的是，您可以节省执行 32（方法 1）- 3（方法 2）= 29 次写入全局内存所花费的时间。

【讨论】：