Entfernen __syncthreads) in CUDA warp-Pegelreduzierung

Die folgenden Code Summen alle 32 Elemente in einem Array auf das erste Element jeder 32 Elementgruppe:Entfernen __syncthreads) in CUDA warp-Pegelreduzierung

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

I dachte ich beseitigen kann die __syncthreads() in der Code, da alle Operationen in demselben Warp ausgeführt werden. Aber wenn ich sie beseitige, bekomme ich Müll zurück. Es wird die Leistung nicht zu sehr beeinträchtigen, aber ich möchte wissen, warum ich hier __syncthreads() brauche.

Quelle

2012-05-23 small_potato

Verwenden Sie eine Fermi-GPU? – talonmies

Ja, es ist eine Quadro 6000, und ich verwende CUDA4.0. In der Tat habe ich ähnliche Technik auf einer GTX 580 verwendet. Ich war überrascht, dass dies nicht funktioniert ohne __synctreads() –

Sie erkennen, dass 'threadIdx.x & 31' ist nicht die Warp-Nummer und' (threadIdx.x & 31) <16' wählt keine Threads innerhalb desselben Warps aus? – talonmies

Vielleicht werfen Sie einen Blick auf diese Folien von Mark Harris. Warum erfinden Sie das Rad neu?

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

Jeder Reduktionsschritt auf dem anderen abhängig. Sie können also nur die Synchronisation in der letzten Ausführung von Warp auf 32 aktive Threads in der Reduktionsphase auslassen. Ein Schritt bevor Sie 64 Threads benötigen und daher eine Synchronisation benötigen, da die parallele Ausführung nicht garantiert ist, da Sie 2 Warps verwenden.

Quelle

2012-05-24 02:23:11 djmj

Das ist ziemlich viel, was ich tun möchte. Das Problem ist wirklich, wenn ich __synctreads() verlasse, fangen die Dinge an zu brechen. Und der Code arbeitet tatsächlich im Debug-Modus, während er im Freigabemodus bricht. –

Möchten Sie eine Warp-basierte Reduktion implementieren? Inneres Warp reduzieren, um Daten um den Faktor 32 zu reduzieren? also bei 1024 threads/elements sind nur 2 synchreads nötig? Dies könnte die Leistung im Vergleich zur herkömmlichen Implementierung möglicherweise erheblich verbessern. Wird diese Idee später überprüfen. – djmj

Das Problem, mit dem ich konfrontiert bin, besteht darin, 128 Zahlen im gemeinsamen Speicher zu summieren. Ich stehe nicht vor einem weltweiten Reduzierungsproblem, aber was Sie sagen, könnte auch funktionieren. –

Ich gebe eine Antwort hier, weil ich denke, dass die oben genannten zwei nicht voll befriedigend sind. Das "geistige Eigentum" dieser Antwort gehört Mark Harris, der auf dieses Problem in dieser presentation (Folie 22) hingewiesen hat, und auf @talonmies, der dieses Problem auf das OP in den obigen Kommentaren hingewiesen hat.

Lassen Sie mich zuerst versuchen, das wieder aufzunehmen, was der OP verlangte, seine Fehler filternd.

Das OP scheint sich mit dem letzten Schritt der Reduzierung der gemeinsamen Speicherreduktion zu befassen, Warpreduktion durch Schleife entrolling. Er tut so etwas wie

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

Wie von Mark Harris und talonmies wies darauf hin, das gemeinsame Speichervariable sdata müssen volatile deklariert werden, um Compiler-Optimierungen zu verhindern. Also der richtige Weg, um die __device__ Funktion oben zu definieren ist:

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

wir nun die zerlegten Codes entsprechend den beiden Fällen siehe oben untersucht, dh sdata erklärt nicht volatile oder volatile (Code für Fermi-Architektur kompiliert).

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

Linien /*0128*/-/*0148*/, /*0160*/ und /*0170*/ entsprechen die gemeinsam genutzten Speicherlasten zu den Registern und Leitungs /*0190*/ an den gemeinsamen Speicher Speicher von Registern. Die Zwischenzeilen entsprechen den in Registern ausgeführten Summierungen. Daher werden die Zwischenergebnisse in Registern (die für jeden Thread privat sind) gespeichert und nicht jedes Mal in den gemeinsamen Speicher geleert, wodurch verhindert wird, dass die Threads volle Sichtbarkeit der Zwischenergebnisse haben.

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

Wie aus Linien /*0130*/-/*01e8*/, jetzt eine Summierung jedes Mal durchgeführt wird, das Zwischenergebnis ist für Vollgewinde Sichtbarkeit zu gemeinsam genutzten Speicher sofort gespült gesehen werden kann.

Quelle

2014-08-30 16:32:16 JackOLantern

Entfernen __syncthreads) in CUDA warp-Pegelreduzierung

Antwort

Verwandte Themen