1

我有 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 方法以包括减少方法的任何帮助或建议，我将不胜感激。

4

1 回答 1

1

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

检查您的初始内核，我们注意到对于 的每个值a，您在计算四个增量时最多使用 12 个值相加（可能少于 12 个，我没有准确计算）。这一切都非常适合您的寄存器文件 - 即使对于双精度值： 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]添加之后都改变 global 中的值，而且你当然不关心在你要写的时候读取那里的值。就您的内核而言，单个线程计算一个新值- 因此它可以将寄存器中的内容相加，并在完成时写入（甚至无需查看其地址）。resN[idx]resN[idx]

如果您按照我在第 1 点中的建议更改内存访问模式，那么实施第 2 点中的建议会变得更加棘手，因为您需要在同一个扭曲中将多个通道的值相加，并且可能要确保您不要'不要通过与单个计算相关的读取来跨越扭曲边界。要了解如何做到这一点，我建议您查看有关基于 shuffle 的缩减的演示文稿。

于 2017-12-20T00:05:28.033 回答