Банковский конфликт CUDA поделился памятью?

Question

Я сталкиваюсь (как я полагаю) с конфликтами банков с общей памятью в ядре CUDA. Сам код довольно сложный, но я воспроизвел его в простом примере, приведенном ниже.

В этом случае он упрощается до простого копирования из глобальной -> общей -> глобальной памяти двумерного массива размером 16x16 с использованием массива общей памяти, который может быть дополнен справа (переменная ng).

Если я скомпилирую код с ng=0 и изучить шаблон доступа к разделяемой памяти с NVVP, он говорит мне, что "нет проблем". Например, ng=2 Я получаю "Транзакции общего хранилища / Доступ = 2, Идеальные транзакции / Доступ = 1" в строках, помеченных как "Предупреждение NVVP". Я не понимаю, почему (или более конкретно: почему заполнение вызывает предупреждения).

РЕДАКТИРОВАТЬ, как упомянуто Грегом Смитом ниже, в Kepler есть 32 банка шириной 8 байт ( http://gpgpu.org/wp/wp-content/uploads/2013/09/08-opti-smem-instr.pdf, слайд 18). Но я не вижу, как это меняет проблему.

Если я правильно понимаю, с 32 банками (B1, B2, ..) 4 байта, двойники (D01, D02, ..) хранятся как:

B1   B2   B3   B4   B5    ..   B31
----------------------------------
D01       D02       D03   ..   D15
D16       D17       D18   ..   D31
D32       D33       D34   ..   D47

Без прокладки пол варпа пишите (as[ijs] = in[ij]) к разделяемой памяти D01 .. D15, D16 .. D31 и т. д. С отступом (размером 2) первая половина деформации пишет в D01 .. D15 второй после отступа D18 .. D33, который все еще не должен вызывать банковский конфликт?

Есть идеи, что здесь может пойти не так?

Упрощенный пример (проверено cuda 6.5.14):

// Compiled with nvcc -O3 -arch=sm_35 -lineinfo 

__global__ void copy(double * const __restrict__ out, const double * const __restrict__ in, const int ni, const int nj, const int ng)

{
    extern __shared__ double as[];
    const int ij=threadIdx.x + threadIdx.y*blockDim.x;
    const int ijs=threadIdx.x + threadIdx.y*(blockDim.x+ng);

    as[ijs] = in[ij]; // NVVP warning
    __syncthreads();
    out[ij] = as[ijs]; // NVVP warning
}

int main()
{
    const int itot = 16;
    const int jtot = 16;
    const int ng = 2;
    const int ncells = itot * jtot;

    double *in  = new double[ncells];
    double *out = new double[ncells];
    double *tmp = new double[ncells];
    for(int n=0; n<ncells; ++n)
        in[n]  = 0.001 * (std::rand() % 1000) - 0.5;

    double *ind, *outd;
    cudaMalloc((void **)&ind,  ncells*sizeof(double));
    cudaMalloc((void **)&outd, ncells*sizeof(double));
    cudaMemcpy(ind,  in,  ncells*sizeof(double), cudaMemcpyHostToDevice);
    cudaMemcpy(outd, out, ncells*sizeof(double), cudaMemcpyHostToDevice);

    dim3 gridGPU (1, 1 , 1);
    dim3 blockGPU(16, 16, 1);

    copy<<<gridGPU, blockGPU, (itot+ng)*jtot*sizeof(double)>>>(outd, ind, itot, jtot, ng);

    cudaMemcpy(tmp, outd, ncells*sizeof(double), cudaMemcpyDeviceToHost);

    return 0;
}

Answer 1

Оказывается, я не правильно понял архитектуру Кепплера. Как указано в одном из приведенных выше комментариев Грега Смита, Кепплер может быть настроен на 32 банка общей памяти по 8 байт. В таком случае, используя cudaDeviceSetSharedMemConfig( cudaSharedMemBankSizeEightByte )расположение общей памяти выглядит так:

bank:  B0   B1   B2   B3   B4    ..   B31
       ----------------------------------
index: D00  D01  D02  D03  D04   ..   D31
       D32  D33  D34  D35  D36   ..   D63   

Теперь для моего простого примера (используя itot=16) запись / чтение в / из общей памяти, например, в первых двух строках (threadIdx.y=0, threadIdx.y=1) обрабатывается в пределах одной основы. Это означает, что для threadIdx.y=0 ценности D00..D15 хранятся в B0..B15, то есть дополнение двух двойных, после чего в пределах одинаковых значений деформации D18..D33 хранятся в B18..B31+B00..B01, что вызывает конфликт банка на B00-B01, Без обивки (ng=0) первая строка написана D00..D15 в B00..B15, второй ряд в D16..D31 в B16..B31Таким образом, никакого банковского конфликта не происходит.

Для ниточного блока blockDim.x>=32 проблема не должна возникать. Например, для itot=32, blockDim.x=32, ng=2, первый ряд хранится в банках B00..B31затем две ячейки заполнения, второй ряд в B02..B31+B00..B01, так далее.