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

Question

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

__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. Тем не менее ядро ​​работает значительно медленнее. Что я сделал не так?