使用不规则内存访问优化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]数组时的不规则内存访问模式。我的想法是使用共享内存。但在这种情况下，似乎很难使用共享内存来合并全局内存访问。

有谁有关于如何优化这个内核的建议吗？

Answer 1

Trove库是一个支持AoS的CUDA/C++库，可能会为随机AoS访问提供接近最佳的性能。从GitHub页面上看，对于16字节的结构，与朴素的方法相比，trove将获得大约2倍的性能。

https://github.com/BryanCatanzaro/trove

​

​

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编程指南的B.2节。对于开普勒体系结构，const和__restrict的关键字可能会被const编译器添加标签，以便通过只读数据和缓存重新加载，请参阅Kepler architecture whitepaper。