我有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"循环，特斯拉K40卡对于特斯拉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方法以包括减少方法。
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。在全球内存中的位置过多
此问题更特定于您的情况。您会看到，您真的不需要更改之后的global中的 resN[idx]值，而每个 添加的添加物之一，当然，您肯定不在乎阅读每当您要去的价值写。当您的内核站立时，单个线程为resN[idx]计算一个新值 - 因此，它可以在寄存器中加起来，然后在完成后写入resN[idx]（甚至不查看其地址）。
如果您像我在第1点所建议的那样更改内存访问模式。当然，您不会与与单个计算相关的读取跨越边界。要了解如何做到这一点，我建议您查看有关基于洗牌的减少的演示文稿。