CUDA: перегрузка разделяемой памяти для реализации подхода сокращения с несколькими массивами

У меня есть 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 можно исключить, но я оставил его здесь, так как на него удобно смотреть код.
Это работает, но, очевидно, это крайне неэффективно и медленно для карты Tesla K40 даже после удаления цикла for. Арифметика, показанная внутри цикла 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, чтобы включить подход сокращения.