是什么导致 CUDA 中的指令重放开销答案

【问题标题】：What Causes Instruction Replay Overhead in CUDA是什么导致 CUDA 中的指令重放开销
【发布时间】：2012-07-20 08:39:57
【问题描述】：

我在我的一个 CUDA 应用程序上运行了可视化分析器。如果数据太大，应用程序会多次调用单个内核。这个内核没有分支。

分析器报告 高指令重播开销 83.6% 和 高全局内存指令重播开销 83.5%。

这是内核的一般外观：

// Decryption kernel
__global__ void dev_decrypt(uint8_t *in_blk, uint8_t *out_blk){

    __shared__ volatile word sdata[256];
    register uint32_t data;

    // Thread ID
#define xID (threadIdx.x + blockIdx.x * blockDim.x)
#define yID (threadIdx.y + blockIdx.y * blockDim.y)
    uint32_t tid = xID + yID * blockDim.x * gridDim.x;
#undef xID
#undef yID

    register uint32_t pos4 = tid%4;
    register uint32_t pos256 = tid%256;
    uint32_t blk = pos256&0xFC;

    // Indices
    register uint32_t index0 = blk + (pos4+3)%4;
    register uint32_t index1 = blk + (pos4+2)%4;

    // Read From Global Memory
    b0[pos256] = ((word*)in_blk)[tid+4] ^ dev_key[pos4];

    data  = tab(0,sdata[index0]);
    data ^= tab(1,sdata[index1]);
    sdata[pos256] = data ^ tab2[pos4];

    data  = tab(0,sdata[index0]);
    data ^= tab(1,sdata[index1]);
    sdata[pos256] = data ^ tab2[2*pos4];

    data  = tab(0,sdata[index0]);
    data ^= tab(1,sdata[index1]);
    data ^= tab2[3*pos4];

    ((uint32_t*)out_blk)[tid] = data + ((uint32_t*)in_blk)[tid];
}

如您所见，没有分支。线程最初将根据线程 ID + 16 字节从全局内存中读取。然后，它们将根据线程 ID 对全局内存中的数据执行操作后写入输出缓冲区。

知道为什么这个内核会有这么多开销吗？

【问题讨论】：

@talonmies 这是我用来访问常量内存的宏。 tab2 也是常量内存。这不是我的内核实际的样子。然而，这就是它通常的行为方式。
好的，那么序列化可能是因为不断的内存访问。如果一个 warp 中的所有线程都不会访问常量内存中的同一个单词，则可以进行序列化。
@talonmies 所以即使是恒定的内存访问也必须以某种方式访问才能合并？ CUDA C 编程指南没有提到这一点。
@talonmies 我更改了代码以创建表的 32 个实例（仍然有足够的常量内存）。探查器正在报告同样的事情。我是否应该重新排列表格，以便相邻的线程访问相邻的单词？
除非 warp 中的每个线程都访问常量内存中的相同单词，否则您将获得序列化。如果您想要随机或半随机访问，请使用全局内存或纹理，而不是常量内存。

标签： memory cuda gpu overhead

【解决方案1】：

在这种情况下，指令重放的来源是扭曲内的非均匀常量内存访问。在您的代码中，tab 存储在常量内存中，并根据线程索引和数据存储共享内存的某种组合进行索引。结果似乎是同一扭曲中的非统一访问线程。常量内存实际上适用于warp中的所有线程访问同一个字的情况，然后可以在单个操作中从常量内存缓存中广播该值，否则会发生warp序列化。

在需要对小型只读数据集进行非统一访问的情况下，将数据绑定到纹理可能比将其存储在常量内存中更好。

【讨论】：