Удаление __syncthreads() в CUDA сокращении уровня варпа
Следующий код суммирует каждый 32 элементы в массиве до самого первого элемента каждого 32 группа элементов:
int i = threadIdx.x;
int warpid = i&31;
if(warpid < 16){
s_buf[i] += s_buf[i+16];__syncthreads();
s_buf[i] += s_buf[i+8];__syncthreads();
s_buf[i] += s_buf[i+4];__syncthreads();
s_buf[i] += s_buf[i+2];__syncthreads();
s_buf[i] += s_buf[i+1];__syncthreads();
}
Я думал, что могу устранить все __syncthreads() в коде, так как все операции выполняются в одной и той же деформации. Но если я устраню их, я получу результаты мусора обратно. Это не должно сильно влиять на производительность, но я хочу знать, зачем мне это нужно __syncthreads() Вот.
2 ответа
Я даю ответ здесь, потому что я думаю, что вышеупомянутые два не полностью удовлетворительны. "Интеллектуальная собственность" этого ответа принадлежит Марку Харрису, который указал на эту проблему в этой презентации (слайд 22), и @talonmies, который указал на эту проблему ФП в комментариях выше.
Позвольте мне сначала попытаться возобновить то, что спрашивал ОП, фильтруя его ошибки.
Похоже, OP имеет дело с последним этапом сокращения общего объема памяти, уменьшения деформации при развертывании цикла. Он делает что-то вроде
template <class T>
__device__ void warpReduce(T *sdata, int tid) {
sdata[tid] += sdata[tid + 32];
sdata[tid] += sdata[tid + 16];
sdata[tid] += sdata[tid + 8];
sdata[tid] += sdata[tid + 4];
sdata[tid] += sdata[tid + 2];
sdata[tid] += sdata[tid + 1];
}
template <class T>
__global__ void reduce4_no_synchthreads(T *g_idata, T *g_odata, unsigned int N)
{
extern __shared__ T sdata[];
unsigned int tid = threadIdx.x; // Local thread index
unsigned int i = blockIdx.x*(blockDim.x*2) + threadIdx.x; // Global thread index - Fictitiously double the block dimension
// --- Performs the first level of reduction in registers when reading from global memory.
T mySum = (i < N) ? g_idata[i] : 0;
if (i + blockDim.x < N) mySum += g_idata[i+blockDim.x];
sdata[tid] = mySum;
// --- Before going further, we have to make sure that all the shared memory loads have been completed
__syncthreads();
// --- Reduction in shared memory. Only half of the threads contribute to reduction.
for (unsigned int s=blockDim.x/2; s>32; s>>=1)
{
if (tid < s) { sdata[tid] = mySum = mySum + sdata[tid + s]; }
// --- At the end of each iteration loop, we have to make sure that all memory operations have been completed
__syncthreads();
}
// --- Single warp reduction by loop unrolling. Assuming blockDim.x >64
if (tid < 32) warpReduce(sdata, tid);
// --- Write result for this block to global memory. At the end of the kernel, global memory will contain the results for the summations of
// individual blocks
if (tid == 0) g_odata[blockIdx.x] = sdata[0];
}
Как отметил Марк Харрис и talonmies, переменная общей памяти sdata должен быть объявлен как volatile, чтобы предотвратить оптимизацию компилятора. Итак, правильный способ определить __device__ Функция выше это:
template <class T>
__device__ void warpReduce(volatile T *sdata, int tid) {
sdata[tid] += sdata[tid + 32];
sdata[tid] += sdata[tid + 16];
sdata[tid] += sdata[tid + 8];
sdata[tid] += sdata[tid + 4];
sdata[tid] += sdata[tid + 2];
sdata[tid] += sdata[tid + 1];
}
Давайте теперь посмотрим разобранные коды, соответствующие двум рассмотренным выше случаям, т.е. sdata объявлен как не volatile или же volatile (код скомпилирован для архитектуры Fermi).
Не volatile
/*0000*/ MOV R1, c[0x1][0x100]; /* 0x2800440400005de4 */
/*0008*/ S2R R0, SR_CTAID.X; /* 0x2c00000094001c04 */
/*0010*/ SHL R3, R0, 0x1; /* 0x6000c0000400dc03 */
/*0018*/ S2R R2, SR_TID.X; /* 0x2c00000084009c04 */
/*0020*/ IMAD R3, R3, c[0x0][0x8], R2; /* 0x200440002030dca3 */
/*0028*/ IADD R4, R3, c[0x0][0x8]; /* 0x4800400020311c03 */
/*0030*/ ISETP.LT.U32.AND P0, PT, R3, c[0x0][0x28], PT; /* 0x188e4000a031dc03 */
/*0038*/ ISETP.GE.U32.AND P1, PT, R4, c[0x0][0x28], PT; /* 0x1b0e4000a043dc03 */
/*0040*/ @P0 ISCADD R3, R3, c[0x0][0x20], 0x2; /* 0x400040008030c043 */
/*0048*/ @!P1 ISCADD R4, R4, c[0x0][0x20], 0x2; /* 0x4000400080412443 */
/*0050*/ @!P0 MOV R5, RZ; /* 0x28000000fc0161e4 */
/*0058*/ @!P1 LD R4, [R4]; /* 0x8000000000412485 */
/*0060*/ @P0 LD R5, [R3]; /* 0x8000000000314085 */
/*0068*/ SHL R3, R2, 0x2; /* 0x6000c0000820dc03 */
/*0070*/ NOP; /* 0x4000000000001de4 */
/*0078*/ @!P1 IADD R5, R4, R5; /* 0x4800000014416403 */
/*0080*/ MOV R4, c[0x0][0x8]; /* 0x2800400020011de4 */
/*0088*/ STS [R3], R5; /* 0xc900000000315c85 */
/*0090*/ BAR.RED.POPC RZ, RZ, RZ, PT; /* 0x50ee0000ffffdc04 */
/*0098*/ MOV R6, c[0x0][0x8]; /* 0x2800400020019de4 */
/*00a0*/ ISETP.LT.U32.AND P0, PT, R6, 0x42, PT; /* 0x188ec0010861dc03 */
/*00a8*/ @P0 BRA 0x118; /* 0x40000001a00001e7 */
/*00b0*/ NOP; /* 0x4000000000001de4 */
/*00b8*/ NOP; /* 0x4000000000001de4 */
/*00c0*/ MOV R6, R4; /* 0x2800000010019de4 */
/*00c8*/ SHR.U32 R4, R4, 0x1; /* 0x5800c00004411c03 */
/*00d0*/ ISETP.GE.U32.AND P0, PT, R2, R4, PT; /* 0x1b0e00001021dc03 */
/*00d8*/ @!P0 IADD R7, R4, R2; /* 0x480000000841e003 */
/*00e0*/ @!P0 SHL R7, R7, 0x2; /* 0x6000c0000871e003 */
/*00e8*/ @!P0 LDS R7, [R7]; /* 0xc10000000071e085 */
/*00f0*/ @!P0 IADD R5, R7, R5; /* 0x4800000014716003 */
/*00f8*/ @!P0 STS [R3], R5; /* 0xc900000000316085 */
/*0100*/ BAR.RED.POPC RZ, RZ, RZ, PT; /* 0x50ee0000ffffdc04 */
/*0108*/ ISETP.GT.U32.AND P0, PT, R6, 0x83, PT; /* 0x1a0ec0020c61dc03 */
/*0110*/ @P0 BRA 0xc0; /* 0x4003fffea00001e7 */
/*0118*/ ISETP.GT.U32.AND P0, PT, R2, 0x1f, PT; /* 0x1a0ec0007c21dc03 */
/*0120*/ @P0 BRA.U 0x198; /* 0x40000001c00081e7 */
/*0128*/ @!P0 LDS R8, [R3]; /* 0xc100000000322085 */
/*0130*/ @!P0 LDS R5, [R3+0x80]; /* 0xc100000200316085 */
/*0138*/ @!P0 LDS R4, [R3+0x40]; /* 0xc100000100312085 */
/*0140*/ @!P0 LDS R7, [R3+0x20]; /* 0xc10000008031e085 */
/*0148*/ @!P0 LDS R6, [R3+0x10]; /* 0xc10000004031a085 */
/*0150*/ @!P0 IADD R8, R8, R5; /* 0x4800000014822003 */
/*0158*/ @!P0 IADD R8, R8, R4; /* 0x4800000010822003 */
/*0160*/ @!P0 LDS R5, [R3+0x8]; /* 0xc100000020316085 */
/*0168*/ @!P0 IADD R7, R8, R7; /* 0x480000001c81e003 */
/*0170*/ @!P0 LDS R4, [R3+0x4]; /* 0xc100000010312085 */
/*0178*/ @!P0 IADD R6, R7, R6; /* 0x480000001871a003 */
/*0180*/ @!P0 IADD R5, R6, R5; /* 0x4800000014616003 */
/*0188*/ @!P0 IADD R4, R5, R4; /* 0x4800000010512003 */
/*0190*/ @!P0 STS [R3], R4; /* 0xc900000000312085 */
/*0198*/ ISETP.NE.AND P0, PT, R2, RZ, PT; /* 0x1a8e0000fc21dc23 */
/*01a0*/ @P0 BRA.U 0x1c0; /* 0x40000000600081e7 */
/*01a8*/ @!P0 ISCADD R0, R0, c[0x0][0x24], 0x2; /* 0x4000400090002043 */
/*01b0*/ @!P0 LDS R2, [RZ]; /* 0xc100000003f0a085 */
/*01b8*/ @!P0 ST [R0], R2; /* 0x900000000000a085 */
/*01c0*/ EXIT; /* 0x8000000000001de7 */
линии /*0128*/-/*0148*/, /*0160*/ а также /*0170*/ соответствуют разделяемой памяти загружается в регистры и строки /*0190*/ в хранилище общей памяти из реестра. Промежуточные строки соответствуют суммам, которые выполняются в регистрах. Таким образом, промежуточные результаты хранятся в регистрах (которые являются частными для каждого потока) и не сбрасываются каждый раз в общую память, что не позволяет потокам иметь полную видимость промежуточных результатов.
volatile
/*0000*/ MOV R1, c[0x1][0x100]; /* 0x2800440400005de4 */
/*0008*/ S2R R0, SR_CTAID.X; /* 0x2c00000094001c04 */
/*0010*/ SHL R3, R0, 0x1; /* 0x6000c0000400dc03 */
/*0018*/ S2R R2, SR_TID.X; /* 0x2c00000084009c04 */
/*0020*/ IMAD R3, R3, c[0x0][0x8], R2; /* 0x200440002030dca3 */
/*0028*/ IADD R4, R3, c[0x0][0x8]; /* 0x4800400020311c03 */
/*0030*/ ISETP.LT.U32.AND P0, PT, R3, c[0x0][0x28], PT; /* 0x188e4000a031dc03 */
/*0038*/ ISETP.GE.U32.AND P1, PT, R4, c[0x0][0x28], PT; /* 0x1b0e4000a043dc03 */
/*0040*/ @P0 ISCADD R3, R3, c[0x0][0x20], 0x2; /* 0x400040008030c043 */
/*0048*/ @!P1 ISCADD R4, R4, c[0x0][0x20], 0x2; /* 0x4000400080412443 */
/*0050*/ @!P0 MOV R5, RZ; /* 0x28000000fc0161e4 */
/*0058*/ @!P1 LD R4, [R4]; /* 0x8000000000412485 */
/*0060*/ @P0 LD R5, [R3]; /* 0x8000000000314085 */
/*0068*/ SHL R3, R2, 0x2; /* 0x6000c0000820dc03 */
/*0070*/ NOP; /* 0x4000000000001de4 */
/*0078*/ @!P1 IADD R5, R4, R5; /* 0x4800000014416403 */
/*0080*/ MOV R4, c[0x0][0x8]; /* 0x2800400020011de4 */
/*0088*/ STS [R3], R5; /* 0xc900000000315c85 */
/*0090*/ BAR.RED.POPC RZ, RZ, RZ, PT; /* 0x50ee0000ffffdc04 */
/*0098*/ MOV R6, c[0x0][0x8]; /* 0x2800400020019de4 */
/*00a0*/ ISETP.LT.U32.AND P0, PT, R6, 0x42, PT; /* 0x188ec0010861dc03 */
/*00a8*/ @P0 BRA 0x118; /* 0x40000001a00001e7 */
/*00b0*/ NOP; /* 0x4000000000001de4 */
/*00b8*/ NOP; /* 0x4000000000001de4 */
/*00c0*/ MOV R6, R4; /* 0x2800000010019de4 */
/*00c8*/ SHR.U32 R4, R4, 0x1; /* 0x5800c00004411c03 */
/*00d0*/ ISETP.GE.U32.AND P0, PT, R2, R4, PT; /* 0x1b0e00001021dc03 */
/*00d8*/ @!P0 IADD R7, R4, R2; /* 0x480000000841e003 */
/*00e0*/ @!P0 SHL R7, R7, 0x2; /* 0x6000c0000871e003 */
/*00e8*/ @!P0 LDS R7, [R7]; /* 0xc10000000071e085 */
/*00f0*/ @!P0 IADD R5, R7, R5; /* 0x4800000014716003 */
/*00f8*/ @!P0 STS [R3], R5; /* 0xc900000000316085 */
/*0100*/ BAR.RED.POPC RZ, RZ, RZ, PT; /* 0x50ee0000ffffdc04 */
/*0108*/ ISETP.GT.U32.AND P0, PT, R6, 0x83, PT; /* 0x1a0ec0020c61dc03 */
/*0110*/ @P0 BRA 0xc0; /* 0x4003fffea00001e7 */
/*0118*/ ISETP.GT.U32.AND P0, PT, R2, 0x1f, PT; /* 0x1a0ec0007c21dc03 */
/*0120*/ SSY 0x1f0; /* 0x6000000320000007 */
/*0128*/ @P0 NOP.S; /* 0x40000000000001f4 */
/*0130*/ LDS R5, [R3]; /* 0xc100000000315c85 */
/*0138*/ LDS R4, [R3+0x80]; /* 0xc100000200311c85 */
/*0140*/ IADD R6, R5, R4; /* 0x4800000010519c03 */
/*0148*/ STS [R3], R6; /* 0xc900000000319c85 */
/*0150*/ LDS R5, [R3]; /* 0xc100000000315c85 */
/*0158*/ LDS R4, [R3+0x40]; /* 0xc100000100311c85 */
/*0160*/ IADD R6, R5, R4; /* 0x4800000010519c03 */
/*0168*/ STS [R3], R6; /* 0xc900000000319c85 */
/*0170*/ LDS R5, [R3]; /* 0xc100000000315c85 */
/*0178*/ LDS R4, [R3+0x20]; /* 0xc100000080311c85 */
/*0180*/ IADD R6, R5, R4; /* 0x4800000010519c03 */
/*0188*/ STS [R3], R6; /* 0xc900000000319c85 */
/*0190*/ LDS R5, [R3]; /* 0xc100000000315c85 */
/*0198*/ LDS R4, [R3+0x10]; /* 0xc100000040311c85 */
/*01a0*/ IADD R6, R5, R4; /* 0x4800000010519c03 */
/*01a8*/ STS [R3], R6; /* 0xc900000000319c85 */
/*01b0*/ LDS R5, [R3]; /* 0xc100000000315c85 */
/*01b8*/ LDS R4, [R3+0x8]; /* 0xc100000020311c85 */
/*01c0*/ IADD R6, R5, R4; /* 0x4800000010519c03 */
/*01c8*/ STS [R3], R6; /* 0xc900000000319c85 */
/*01d0*/ LDS R5, [R3]; /* 0xc100000000315c85 */
/*01d8*/ LDS R4, [R3+0x4]; /* 0xc100000010311c85 */
/*01e0*/ IADD R4, R5, R4; /* 0x4800000010511c03 */
/*01e8*/ STS.S [R3], R4; /* 0xc900000000311c95 */
/*01f0*/ ISETP.NE.AND P0, PT, R2, RZ, PT; /* 0x1a8e0000fc21dc23 */
/*01f8*/ @P0 BRA.U 0x218; /* 0x40000000600081e7 */
/*0200*/ @!P0 ISCADD R0, R0, c[0x0][0x24], 0x2; /* 0x4000400090002043 */
/*0208*/ @!P0 LDS R2, [RZ]; /* 0xc100000003f0a085 */
/*0210*/ @!P0 ST [R0], R2; /* 0x900000000000a085 */
/*0218*/ EXIT; /* 0x8000000000001de7 */
Как видно из строк /*0130*/-/*01e8*/ теперь каждый раз, когда выполняется суммирование, промежуточный результат немедленно сбрасывается в общую память для полной видимости потока.
Может быть, посмотрите на эти слайды от Марка Харриса. Зачем изобретать велосипед.
www.uni-graz.at/~haasegu/Lectures/GPU_CUDA/Lit/reduction.pdf?page=35
Каждый шаг сокращения зависит от другого. Таким образом, вы можете пропустить синхронизацию только в последней исключенной деформации, равной 32 активным потокам в фазе сокращения. За один шаг до того, как вам понадобится 64 потока и, следовательно, потребуется синхронизация, поскольку параллельное выполнение не гарантируется, поскольку вы используете 2 перекоса.