Эффективность общей памяти CUDA на 50%?

Question

У меня есть следующий код, который выполняет мозаичное транспонирование матрицы с использованием общей памяти для повышения производительности. Совместно используемая память дополняется 1 столбцом, чтобы избежать конфликта банков для блока потока 32x32.

__global__ void transpose_tiled_padded(float *A, float *B, int n)
{
    int i_in = blockDim.x*blockIdx.x + threadIdx.x;
    int j_in = blockDim.y*blockIdx.y + threadIdx.y;
    int i_out = blockDim.x*blockIdx.y + threadIdx.x;
    int j_out = blockDim.y*blockIdx.x + threadIdx.y;

    extern __shared__ float tile[];

    // coalesced read of A rows to (padded) shared tile column (transpose)
    tile[threadIdx.y + threadIdx.x*(blockDim.y+1)] = A[i_in + j_in*n];
    __syncthreads();

    // coalesced write from (padded) shared tile column to B rows
    B[i_out + j_out*n] = tile[threadIdx.x + threadIdx.y*(blockDim.x+1)];
}

Запустив этот код, я получаю 100% эффективность использования общей памяти в визуальном профилировщике NVIDIA, как и ожидалось. Но когда я запускаю его с блоком потоков 16x16, я получаю только 50% эффективности. Это почему? Насколько я могу судить, ни один поток в варпе не читает из того же банка с таким макетом. Или я ошибаюсь?

Answer 1

Да вы ошибаетесь

Учитывая этот доступ (чтение) для деформации 0 в блоке 16x16:

tile[threadIdx.x + threadIdx.y*(blockDim.x+1)];
     ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
                     "index"

Вот соответствующие вычисления для каждой нити в основе:

warp lane:    0  1  2  3  4  5  6  7  8  9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 23 25 26 27 28 29 30 31
threadIdx.x:  0  1  2  3  4  5  6  7  8  9 10 11 12 13 14 15  0  1  2  3  4  5  6  7  8  9 10 11 12 13 14 15
threadIdx.y:  0  0  0  0  0  0  0  0  0  0  0  0  0  0  0  0  1  1  1  1  1  1  1  1  1  1  1  1  1  1  1  1
"index":      0  1  2  3  4  5  6  7  8  9 10 11 12 13 14 15 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32
bank:         0  1  2  3  4  5  6  7  8  9 10 11 12 13 14 15 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31  0

Итак, мы видим, что для этой деформации первый и последний потоки оба читают из банка 0. Это приводит к двухстороннему конфликту банков, двухсторонней сериализации и 50% эффективности.