__syncthreads Extracción() en la reducción de nivel de urdimbre CUDA

el siguiente código de sumas cada 32 elementos de una matriz a la muy primer elemento de cada grupo 32 elemento:__syncthreads Extracción() en la reducción de nivel de urdimbre CUDA

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

pensé que puede eliminar todo el __syncthreads() en el código, ya que todas las operaciones se realizan en la misma urdimbre. Pero si los elimino, recupero los resultados de la basura. No afectará demasiado el rendimiento, pero quiero saber por qué necesito __syncthreads() aquí.

Fuente

2012-05-23 small_potato

¿Está utilizando una GPU de Fermi? – talonmies

sí, es un Quadro 6000, y estoy usando CUDA4.0. De hecho, he usado una técnica similar en un GTX 580. Me sorprendió que esto no funcione sin __syncthreads() –

. Te das cuenta de que 'threadIdx.x & 31' no es el número de warp y' (threadIdx.x & 31) <16 'no selecciona hilos dentro de la misma urdimbre? – talonmies

Quizás eche un vistazo a estas diapositivas de Mark Harris. Por qué reinventar la rueda.

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

Cada reducción es dependiente de la otra. Así que solo puede omitir la sincronización en el último warpe ejecutado igual a 32 hilos activos en la fase de reducción. Un paso antes de necesitar 64 subprocesos y, por lo tanto, necesita una sincronización, ya que la ejecución en paralelo no está garantizada ya que utiliza 2 distorsiones.

Fuente

2012-05-24 02:23:11 djmj

Eso es más o menos lo que quiero hacer. El problema es realmente, cuando dejo __syncthreads() fuera, las cosas comienzan a romperse. Y el código realmente funciona en modo de depuración mientras se rompe en modo de lanzamiento. –

¿Pretende implementar la reducción basada en warp? ¿Reducir el interior de la urdimbre para reducir los datos en el factor 32? entonces, con 1024 hilos/elementos solo son necesarios 2 sincrones? Esto podría mejorar el rendimiento mucho más en comparación con la implementación convencional. Verificará esta idea más tarde. – djmj

El problema que estoy enfrentando es solo sumar 128 números que residen en la memoria compartida. No estoy enfrentando un problema de reducción global, pero lo que dices podría funcionar también. –

Estoy dando una respuesta aquí porque creo que los dos anteriores no son totalmente satisfactorios. La "propiedad intelectual" de esta respuesta pertenece a Mark Harris, que ha señalado este problema en este presentation (diapositiva 22), y a @talonmies, que ha señalado este problema al OP en los comentarios anteriores.

Permítanme primero intentar resumir lo que el OP estaba pidiendo, filtrando sus errores.

El OP parece tratar con el último paso de reducción en la reducción de memoria compartida, reducción de deformación mediante desenrollado de bucle. Él está haciendo algo así como

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

Como se ha señalado por Mark Harris y talonmies, la memoria compartida variables sdata deben ser declarados como volatile, para evitar que las optimizaciones del compilador. Por lo tanto, de la manera correcta para definir la función __device__ anterior es:

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

Veamos ahora los códigos desmontados correspondientes a los dos casos anteriormente examinados, es decir, sdata declarados como no volatile o volatile (código compilado para la arquitectura Fermi)

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

Líneas /*0128*/-/*0148*/, /*0160*/ y /*0170*/ corresponden a las cargas de memoria compartida a los registros y la línea /*0190*/ a la tienda de la memoria compartida de registro. Las líneas intermedias corresponden a las sumas, tal como se realizan en los registros. Por lo tanto, los resultados intermedios se guardan en registros (que son privados para cada hilo) y no se enjuagan cada vez a la memoria compartida, lo que impide que los hilos tengan una visibilidad completa de los resultados intermedios.

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

Como se puede observar a partir de líneas /*0130*/-/*01e8*/, ahora cada vez que se realiza una suma, el resultado intermedio es inmediatamente enrojecida a la memoria compartida para la visibilidad rosca completa.

Fuente

2014-08-30 16:32:16 JackOLantern

__syncthreads Extracción() en la reducción de nivel de urdimbre CUDA

Respuesta

Cuestiones relacionadas