CUDA 对全局内存的低效访问模式

我编写了一个 CUDA 内核，它执行两个数组arr1和arr2. 将哪些索引arr1与哪些索引相加的信息arr2存储在一个数组中idx。

这是一个代码示例：

__global__ add(float* arr1, float* arr2, int* idx, int length)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;

    // each thread performs (length) additions,
    // arr2 is (lenght) times larger than arr1
    for (int j=threadIdx.x; j<length*blockDim.x; j+=blockDim.x)
    {
        arr1[i] += arr2[ idx[blockIdx.x*blockDim.x + j] ]; // edited here
    }
}

该代码产生正确的输出，但几乎不比具有 8 个线程的 CPU 上的 openmp-parallel 代码快。我为不同的块大小尝试了这个。

我怀疑访问模式arr2效率低下，因为arr2它在全局内存中并且是准随机访问的——数组idx包含唯一的、排序的、但不连续的索引（可能是 2、3、57、103……） . 因此，没有利用 L1 缓存。此外，数组非常大，不能完全适合共享内存。

有没有办法绕过这个障碍？您对如何优化访问模式有想法arr2吗？