优化具有不规则内存访问的 CUDA 内核

Posted 2023-04-15

【中文标题】优化具有不规则内存访问的 CUDA 内核

【英文标题】：Optimizing a CUDA kernel with irregular memory accesses

【发布时间】：2013-12-29 00:28:39

【问题描述】：

我有以下 CUDA 内核，似乎很难优化：

__global__ void DataLayoutTransformKernel(cuDoubleComplex* d_origx, cuDoubleComplex* d_origx_remap, int n, int filter_size, int ai )

    for(int idx = blockIdx.x * blockDim.x + threadIdx.x; idx < filter_size; idx+=blockDim.x * gridDim.x)
    
        int index = (idx * ai) & (n-1);
        d_origx_remap[idx] = d_origx[index];
    


//Parameters were defined before
int permute[loops] = 29165143,3831769,17603771,9301169,32350975, ...
int n = 33554432;
int filter_size = 1783157;

for(int i=0; i<loops; i++)

    DataLayoutTransformKernel<<<dimGrid, dimBlock, 0, stream[i]>>>((cuDoubleComplex*) d_origx,(cuDoubleComplex*)d_origx_remap+i*filter_size, n, filter_size, permute[i]);

内核的目的是将d_origx[]的数据布局从不规则重新排序为规则（d_origx_remap）。内核以不同的访问步幅启动多次 (ai)。

这里的挑战是引用d_origx[index] 数组时的不规则内存访问模式。我的想法是使用共享内存。但是对于这种情况，似乎很难使用共享内存来合并全局内存访问。

有人对如何优化这个内核有建议吗？

【问题讨论】：

也许你可以通过启动内核来隐藏一些延迟，该内核以交错的方式同时处理重新映射的缓冲区和转换内核。第一次迭代只是重新映射内核。第二次迭代同时启动重映射内核和处理第一次重映射结果的内核等。也许还考虑将此功能直接滚动到下一个内核中（让内核在​​跨步位置提取其值）。

【参考方案1】：

Trove 库是一个支持 AoS 支持的 CUDA/C++ 库，并且可能为随机 AoS 访问提供接近最佳的性能。从 GitHub 页面来看，对于 16 字节结构的原始方法，trove 将获得大约 2 倍的收益。

https://github.com/BryanCatanzaro/trove

【讨论】：

只需提到 Trove 可用于计算能力3.0 及以上。

【参考方案2】：

我不确定您是否可以做很多事情来优化您的代码。

根本没有线程合作，所以我想说共享内存不是要走的路。

你可以尝试改变

__global__ void DataLayoutTransformKernel(cuDoubleComplex* d_origx, cuDoubleComplex* d_origx_remap, int n, int filter_size, int ai)

到

__global__ void DataLayoutTransformKernel(const cuDoubleComplex* __restrict__ d_origx, cuDoubleComplex* __restrict__ d_origx_remap, const int n, const int filter_size, const int ai)

即，使用 const 和 __restrict__ 关键字。特别是__restrict__ 将使nvcc 能够执行一些优化，请参阅CUDA C 编程指南的B.2 节。对于 Kepler 架构，const 和 __restrict 关键字可能会被编译器标记为通过只读数据缓存加载，请参阅Kepler architecture whitepaper。