CUDA中指令重放开销的原因

Question

我在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使用全局内存中的数据执行操作后，它们将写入输出缓冲区。

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

Answer 1

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

在需要对小型只读数据集进行非均匀访问的情况下，将数据绑定到纹理可能比将其存储为恒定内存要好。