Совместно используемая память CUDA медленнее, чем глобальная, даже когда нет конфликтов с банками и многократное повторное использование данных

Я попытался использовать общую память для ускорения работы моего ядра. Оригинальная версия с использованием глобальной памяти выглядит так:

__global__ void my_kernel(float* inout, float* in, float* const_array)
{
int y = blockIdx.y * blockDim.y + threadIdx.y;
int x = blockIdx.x * blockDim.x + threadIdx.x;

if (y < gpu_ny && x < gpu_nx) {
for (int z = 0; z < gpu_nz; ++z) {
int index = x + y * gpu_nx + gpu_nxy * z;
inout[index] = const_array[index] * (
C0 * in[index] +
C1 * (in[index - gpu_nx] + in[index + gpu_nx]) +
C2 * (in[index - 2 * gpu_nx] + in[index + 2 * gpu_nx]) +
C3 * (in[index - 3 * gpu_nx] + in[index + 3 * gpu_nx]) +
C4 * (in[index - 4 * gpu_nx] + in[index + 4 * gpu_nx]) +
C5 * (in[index - 1] + in[index + 1]) +
C6 * (in[index - 2] + in[index + 2]) +
C7 * (in[index - 3] + in[index + 3]) +
C8 * (in[index - 4] + in[index + 4]) +
C9 * (in[index - 1 * gpu_nxy] + in[index + 1 * gpu_nxy]) +
C10 * (in[index - 2 * gpu_nxy] + in[index + 2 * gpu_nxy]) +
C11 * (in[index - 3 * gpu_nxy] + in[index + 3 * gpu_nxy]) +
C12 * (in[index - 4 * gpu_nxy] + in[index + 4 * gpu_nxy])
) + in[index] + in[index] - inout[index];
}
}
}

Типичный подход в таких задачах — загрузить глобальные значения в общую память и повторно использовать их между потоками.

__global__ void my_kernel(float* inout, float* in, float* const_array)
{
__shared__ float in_shared[BLOCK_NY + 2 * RADIUS][BLOCK_NX + 2 * RADIUS];

int y = blockIdx.y * blockDim.y + threadIdx.y;
int x = blockIdx.x * blockDim.x + threadIdx.x;

int sx = threadIdx.x + RADIUS;
int sy = threadIdx.y + RADIUS;

for (int z = 0; z < gpu_nz; ++z) {
bool in_bounds = y < gpu_ny && x < gpu_nx;
int index = x + y * gpu_nx + z * gpu_nxy;
if (in_bounds) {
in_shared[sy][sx] = in[index];
/*if (threadIdx.x < RADIUS) {
in_shared[sy][threadIdx.x] = in[index - RADIUS];
in_shared[sy][sx + BLOCK_SIZE] = in[index + BLOCKX_SIZE];
}
if (threadIdx.y < RADIUS) {
in_shared[threadIdx.y][sx] = in[index - gpu_nx * RADIUS];
in_shared[sy + BLOCK_SIZE][sx] = in[index + gpu_nx * BLOCKY_SIZE];
}*/
}

// __syncthreads();

if (in_bounds) {
inout[index] = const_array[index] * (
C0 * in_shared[sy][sx] +
C1 * (in_shared[sy-1][sx] + in_shared[sy+1][sx]) +
C2 * (in_shared[sy-2][sx] + in_shared[sy+2][sx]) +
C3 * (in_shared[sy-3][sx] + in_shared[sy+3][sx]) +
C4 * (in_shared[sy-4][sx] + in_shared[sy+4][sx]) +
C5 * (in_shared[sy][sx-1] + in_shared[sy][sx+1]) +
C6 * (in_shared[sy][sx-2] + in_shared[sy][sx+2]) +
C7 * (in_shared[sy][sx-3] + in_shared[sy][sx+3]) +
C8 * (in_shared[sy][sx-4] + in_shared[sy][sx+4]) +
C9 * (in[index - 1 * gpu_nxy] + in[index + 1 * gpu_nxy]) +
C10 * (in[index - 2 * gpu_nxy] + in[index + 2 * gpu_nxy]) +
C11 * (in[index - 3 * gpu_nxy] + in[index + 3 * gpu_nxy]) +
C12 * (in[index - 4 * gpu_nxy] + in[index + 4 * gpu_nxy])
) + in_shared[sy][sx] + in_shared[sy][sx] - inout[index];
}
}
}

Вопреки ожиданиям, это дало мне худшую производительность, чем ядро, использующее только глобальную память (примерно в 4 раза медленнее). Я пытался выяснить причину.

Прежде всего, я закомментировал код для чтения граничных значений и __syncthreads(), Я знаю, что это сделало мой код неправильным, но я хотел оценить влияние на производительность. Тем не менее, производительность была намного хуже, когда я использовал общую память.

Я знаю, что здесь нет банковских конфликтов. Прежде всего, все 32 потока в варпе обращаются к последовательным числам с плавающей точкой, что означает, что каждый поток будет обрабатываться отдельным банком памяти. Более того, я побежал nvprof искать shared_ld_bank_conflict а также shared_st_bank_conflict события, и они не были обнаружены.

Повторное использование данных является высоким, и неправильная версия должна потреблять в 9 раз меньше глобальных чтений. Я даже пробовал разные подходы, например использование общей памяти для повторного использования данных в направлении z. Тем не менее ядро ​​работает значительно медленнее. Что я сделал не так?