Rimozione __syncthreads() nella riduzione del livello ordito CUDA

Question

Il seguente somme codice ogni 32 elementi di una matrice al primo elemento di ciascun gruppo 32 elemento:

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(); 
} 

ho pensato posso eliminare tutto il __syncthreads() in il codice, dal momento che tutte le operazioni sono fatte nello stesso ordito. Ma se li elimini, ottengo risultati spazzatura. Non deve influire troppo sulle prestazioni, ma voglio sapere perché ho bisogno di __syncthreads() qui.

Answer 1

Forse dai un'occhiata a queste diapositive di Mark Harris. Perché reinventare la ruota.

www.uni-graz.at/~haasegu/Lectures/GPU_CUDA/Lit/reduction.pdf?page=35 passo

Ogni riduzione dipende dall'altra. Quindi si può solo tralasciare la sincronizzazione nell'ultimo ordito eccitato uguale a 32 fili attivi nella fase di riduzione. Un passo prima di aver bisogno di 64 thread e quindi bisogno di una sincronizzazione poiché l'esecuzione parallela non è garantita poiché si usano 2 warps.

Answer 2

Fornisco una risposta qui perché ritengo che i due precedenti non siano completamente soddisfacenti. La "proprietà intellettuale" di questa risposta appartiene a Mark Harris, che ha segnalato questo problema in questo presentation (diapositiva 22) e in @talonmies, che ha segnalato questo problema all'OP nei commenti sopra.

Vorrei prima provare a riprendere quello che l'OP chiedeva, filtrando i suoi errori.

L'OP sembra avere a che fare con l'ultimo passaggio di riduzione della riduzione della memoria condivisa, riduzione ordito per ciclo di srotolamento. Sta facendo qualcosa di simile

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]; 
} 

Come sottolineato da Mark Harris e talonmies, la memoria condivisa variabili sdata deve essere dichiarato come volatile, per evitare che le ottimizzazioni del compilatore. Quindi, il modo giusto per definire la funzione __device__ sopra è:

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]; 
} 

Vediamo ora i codici smontati corrispondenti ai due casi sopra esaminati, cioè, sdata dichiarati come non volatile o volatile (codice compilato per l'architettura Fermi).

Non 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 */ 

Linee /*0128*/-/*0148*/, /*0160*/ e /*0170*/ corrispondono ai carichi memoria condivisa tra i registri e la linea /*0190*/ al negozio memoria condivisa dal registro. Le linee intermedie corrispondono alle sommatorie, come eseguite nei registri. Quindi, i risultati intermedi vengono mantenuti nei registri (che sono privati ​​di ciascun thread) e non vengono scaricati ogni volta nella memoria condivisa, impedendo ai thread di avere una visibilità completa dei risultati intermedi.

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 */ 

Come si può notare dalle linee /*0130*/-/*01e8*/, ora ogni volta che viene eseguita una somma, il risultato intermedio è immediatamente lavata alla memoria condivisa per visibilità completa thread.