Будет ли тот же поток доступ к одному и тому же блоку памяти дважды вызвать конфликты?

Question

Я работаю над ядром, которое выполняет векторное сокращение. Это в основном добавляет все позиции в векторе, и сохраняет результат в положении 0.

Я следую этой схеме, с блоками 512 поплавковых элементов:

Код:

//scratch[] is a vector located in shared memory with all 512 elements 
NUM_ELEMENTS = 512; 
for(stride=NUM_ELEMENTS/2; stride>=1; stride = stride/2) { 
    if (threadIdx.x < stride){ 
    scratch[threadIdx.x] += scratch[threadIdx.x + stride]; 
    } 
    __syncthreads(); 
} 

Странно, я ожидаю, что вы столкнетесь с конфликтами в общем банке, а я нет. На первой итерации нить 0 складывает положение 0 и позицию 256, которые находятся в одном банке. Резьба 1 складывает положение 1 и позицию 257 и так далее.

Все эти операции требуют, чтобы каждая нить в варп, чтобы получить 2 различные значения из того же банка, тем не менее, я не получаю никаких конфликтов вообще:

Что мне не хватает?

Answer 1

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