c++ - CUDA:重载共享内存以实现多个数组的缩减方法

Question

我有 5 个大型数组 A(N*5)、B(N*5)、C(N*5)、D(N*5)、E(N*2) 数字 5 和 2 表示这些变量在不同平面/轴中的分量。 这就是我以这种方式构建数组的原因，这样我就可以在编写代码时可视化数据。 N ~ 200^3 ~ 8e06 个节点

例如:这是我的内核最简单的形式，我在其中对全局内存进行所有计算。

#define N 200*200*200
__global__ void kernel(doube *A, double *B, double *C, 
            double *D, double *E, double *res1, double *res2, 
            double *res3, double *res4 )
    {
       int a, idx=threadIdx.x + blockIdx.x * blockDim.x;
        if(idx>=N) {return;}
        res1[idx]=0.; res2[idx]=0.; 
        res3[idx]=0.; res4[idx]=0.

        for (a=0; a<5; a++)
        {
            res1[idx] += A[idx*5+a]*B[idx*5+a]+C[idx*5+a] ;
            res2[idx] += D[idx*5+a]*C[idx*5+a]+E[idx*2+0] ;
            res3[idx] += E[idx*2+0]*D[idx*5+a]-C[idx*5+a] ;
            res4[idx] += C[idx*5+a]*E[idx*2+1]-D[idx*5+a] ;
        }

    }

我知道可以去掉“for”循环，但我把它留在这里，因为这样看代码很方便。 这行得通，但显然即使在删除“for”循环后，它对于 Tesla K40 卡来说也是极其低效和缓慢的。 “for”循环中显示的算法只是为了提供一个想法，实际的计算要长得多，并且与 res1、res2... 混杂在一起。

我已经实现了以下改进有限，但是 我想通过共享内存的过载进一步改进它。

    #define THREADS_PER_BLOCK 256
    __global__ void kernel_shared(doube *A, double *B, double *C, 
               double *D, double *E, double *res1, double *res2, 
               double *res3, double *res4  )
    {
       int a, idx=threadIdx.x + blockIdx.x * blockDim.x;
       int ix = threadIdx.x;
       __shared__ double A_sh[5*THREADS_PER_BLOCK];
       __shared__ double B_sh[5*THREADS_PER_BLOCK];
       __shared__ double C_sh[5*THREADS_PER_BLOCK];
       __shared__ double D_sh[5*THREADS_PER_BLOCK];
       __shared__ double E_sh[2*THREADS_PER_BLOCK];

       //Ofcourse this will not work for all arrays in shared memory; 
        so I am allowed  to put any 2 or 3 variables (As & Bs) of  
         my choice in shared and leave rest in the global memory. 

       for(int a=0; a<5; a++)
     {
        A_sh[ix*5 + a] = A[idx*5 + a] ;
        B_sh[ix*5 + a] = B[idx*5 + a] ;
     }
            __syncthreads();



    if(idx>=N) {return;}
        res1[idx]=0.; res2[idx]=0.; 
        res3[idx]=0.; res4[idx]=0.
    for (a=0; a<5; a++)
    {
        res1[idx] += A_sh[ix*5+a]*B_sh[ix*5+a]+C[idx*5+a];
        res2[idx] += B_sh[ix*5+a]*C[idx*5+a]+E[idx*2+0]  ;
        res3[idx] += E[idx*2+0]*D[idx*5+a]-C[idx*5+a]    ;
        res4[idx] += B_sh[ix*5+a]*E[idx*2+1]-D[idx*5+a]  ;
    }

}

这有点帮助，但我想实现其中一项减少 方法(没有银行冲突)来提高性能，我可以把所有 我共享的变量(可能是平铺方法)然后进行计算部分。 我在 CUDA_Sample 文件夹中看到了缩减示例，但是那个示例 仅对共享中的一个 vector 求和，而不涉及共享内存中多个数组的任何复杂算术。我将不胜感激任何帮助或建议来改进我现有的 kernel_shared 方法以包括减少方法。

Answer 1

1。你需要的不是共享内存

检查您的初始内核，我们注意到对于 a 的每个值，您在计算要加起来的四个增量时最多使用 12 个值(可能少于 12 个，我没有准确计数)。这一切都非常适合您的寄存器文件 - 即使是 double 值: 12 * sizeof(double) ，加上 4 * sizeof(double) 中间结果使每个线程有 32 个 4 字节寄存器。即使每个 block 有 1024 个线程，也远远超出了限制。

现在，你的内核运行缓慢的原因主要是

2。次优内存访问模式

这是您可以在任何 CUDA 编程演示中读到的内容；我只是简单地说，不是每个线程自己处理几个连续的数组元素，而是应该将其交错在 warp 的 channel 之间，或者更好的是在 block 的线程之间。因此，而不是线程全局索引 idx 处理

5 * idx
5 * idx + 1
...
5 * idx + 4

让它处理

5 * blockDim.x * blockIdx.x + threadIdx.x
5 * blockDim.x * blockIdx.x + threadIdx.x + blockDim.x
...
5 * blockDim.x * blockIdx.x + threadIdx.x + 4 * blockDim.x

这样，每当线程读取或写入时，它们的读取和写入合并。在您的情况下，这可能有点棘手，因为您的某些访问模式略有不同，但您明白了。

3。过度添加到全局内存中的位置

这个问题更具体到你的情况。你看，你真的不需要在 每一个 添加后更改全局的 resN[idx] 值，而且你当然不关心阅读每当你要写的时候，它就在那里。正如您的内核所代表的那样，单个线程为 resN[idx] 计算一个新值 - 因此它可以将寄存器中的内容相加，然后写入 resN[idx]当它完成时(甚至不看它的地址)。

---

如果您按照我在第 1 点中的建议更改内存访问模式，则实现第 2 点中的建议会变得更加棘手，因为您需要将同一 warp 中的多个 channel 的值相加，并且可能使确保您不会跨越与单个计算相关的读取的扭曲边界。要了解如何执行此操作，我建议您查看 this presentation关于基于洗牌的减少。