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”循环后，对于特斯拉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文件夹中看到了缩减示例，但该示例仅适用于共享中的一个向量的总和，而没有任何复杂的算术涉及来自共享内存的多个数组。我将不胜感激任何帮助或建议，以改进我现有的kernel_shared方法，以包括减少方法。

Answer 1

1. What you need is not shared memory

检查你的初始内核，我们注意到对于a的每个值，你在计算四个增量时最多使用12个值（可能小于12，我没有完全计算）。这一切都非常适合您的寄存器文件 - 即使是双值：12 * sizeof（double），加上4 * sizeof（double）的中间结果，每个线程有32个4字节寄存器。即使每个块有1024个线程，也超出限制。

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

2. Suboptimal memory access patterns

在任何CUDA编程的演示中，您都可以阅读这些内容;我只是简单地说，不是每个线程自己处理几个连续的数组元素，而是应该在warp的通道之间交错，或者更好的是块的线程。从而代替线程全局索引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. Excessive addition to locations in global memory

此问题更适用于您的情况。你知道，在每次添加之后你真的不需要在全局中更改resN[idx]值，而且当你打算写时，你肯定不在乎读取那里的值。正如你的内核所代表的那样，单个线程为resN[idx]计算一个新值 - 所以它可以在寄存器中添加东西，并在完成后写入resN[idx]（甚至不查看其地址）。

---

如果您按照我在第1点中的建议更改了内存访问模式，那么在第2点中实现建议会变得更加棘手，因为您需要在同一个warp中添加来自多个通道的值，并且可能确保您没有与单个计算相关的读取交叉经线边界。要了解如何做到这一点，我建议你看看this presentation关于基于shuffle的减少。