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方法，以包括减少方法。

CUDA：重载共享内存以实现具有多个数组的简化方法

1. What you need is not shared memory

2. Suboptimal memory access patterns

3. Excessive addition to locations in global memory