CUDA：重载共享内存以实现多个数组的归约方法

Posted 2023-04-15

【中文标题】CUDA：重载共享内存以实现多个数组的归约方法

【英文标题】：CUDA: overloading of shared memory to implement reduction approach with multiple arrays

【发布时间】：2018-06-02 10:07:29

【问题描述】：

我有 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 文件夹中看到了缩减示例，但该示例 仅对共享的一个向量求和，而无需对共享内存中的多个数组进行任何复杂的算术运算。对于改进我现有的 kernel_shared 方法以包含减少方法的任何帮助或建议，我将不胜感激。

【问题讨论】：

第一段代码中的 for 循环对我来说没有意义。因为您正在使用赋值运算符 (=)，例如在 for 循环完成时，res1[idx] 将只包含在最后一次循环迭代中计算的值。应该这样做 += 而不是 = 吗？

嗨罗伯特，我做了那个更正。谢谢

我投票决定将此问题作为离题结束，因为它属于 codereview.stackexchange.com（实际上我的回答也是如此）。 @RobertCrovella：你怎么看？

您正在使用 AoS 存储和访问方案。如果我追求性能，我会做的第一件事是convert it to SoA。在考虑共享内存之前，我会这样做。我认为您已经得到的答案大致相同。

@einpoklum cmets 并不是真正用于侧边栏讨论（为此，您可以使用元数据或聊天）。做你认为合适的事。没事的。我的意见在这里无关紧要。

【参考方案1】：

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

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

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

2。次优的内存访问模式

这是您可以在任何 CUDA 编程演示文稿中阅读的内容；我只是简单地说一下，不是每个线程自己处理几个连续的数组元素，而是应该将其交错在扭曲的通道之间，或者更好的是块的线程。因此而不是线程全局索引 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 点中的建议会变得更加棘手，因为您需要将同一条线中多个通道的值相加，并且可能会使确保您不会通过与单个计算相关的读取来跨越扭曲边界。要了解如何做到这一点，我建议您查看 this presentation 了解基于 shuffle 的缩减。

【讨论】：

您好 einpoklum，感谢您抽出宝贵时间提供如此详细的回复。我想在您的回复中澄清方法＃1，而主题仍然是开放的。您是否建议创建一个局部变量 A_reg（默认情况下将存储在寄存器中），然后为每个线程对该局部变量执行所有算术运算？ A_reg[a]=A[idx*5+a]; B_reg[a]=B[idx*5+a] 对于每个 idx？

@user2415927 而不是 A_reg[a]=A[idx*5+a] ，使用单个寄存器（即A_reg = A[idx*5+a]）并在五次迭代中重复使用它

@user2415927：我建议您为每个多次访问的值（读取或写入）设置一个本地（可寄存器分配）变量；并且在您可能仍在使用它时不要摆脱该值。结果也一样 - 将它们放入寄存器中。