优化带有不规则内存访问的CUDA核函数

Question

我有以下的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] 数组时出现了不规则的内存访问模式。我的想法是使用共享内存。但对于这种情况，使用共享内存来合并全局内存访问似乎非常困难。

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

- user3089852

---

也许您可以通过以错开的方式启动处理重新映射缓冲区的内核，与变换内核同时进行，从而隐藏一些延迟。第一次迭代仅是重新映射内核。第二次迭代同时启动重新映射内核和处理第一个重新映射结果的内核等。也可以直接将此功能纳入下一个内核中（使内核在分步位置提取其值）。 - Roger Dahl

Answer 1

Trove库是一个支持AoS的CUDA/C++库，可能会为随机AoS访问提供接近最优的性能。从GitHub页面上看，对于16字节结构，Trove似乎比朴素方法快2倍。

https://github.com/BryanCatanzaro/trove

- Maddy Scientist

---

只需提到Trove适用于计算能力为3.0及以上的设备。 - Vitality

Answer 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 架构白皮书。

- Vitality