Usuwanie __syncthreads() w CUDA redukcji osnowy poziomu

Następujący kod sumy co 32 elementów w macierzy do pierwszego elementu w każdej grupie 32 Element:Usuwanie __syncthreads() w CUDA redukcji osnowy poziomu

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

myślałem można wyeliminować wszystkie __syncthreads() w kod, ponieważ wszystkie operacje są wykonywane w tej samej osnowie. Ale jeśli je wyeliminuję, otrzymam wyniki śmieci. Nie wpłynie to zbytnio na wydajność, ale chcę wiedzieć, dlaczego potrzebuję tutaj __syncthreads().

Źródło

2012-05-23 small_potato

Czy używasz procesora graficznego Fermi? – talonmies

tak, to jest Quadro 6000 i używam CUDA4.0. W rzeczywistości użyłem podobnej techniki na GTX 580. Byłem zaskoczony, że to nie działa bez __syncthreads() –

Zdajesz sobie sprawę, że 'threadIdx.x & 31' nie jest numerem osnowy i' (threadIdx.x i 31) <16' nie wybiera wątków w obrębie tej samej osnowy? – talonmies

Może spojrzeć na te slajdy od Marka Harrisa. Po co wyważać koło.

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

Każda redukcja jest zależny od drugiego. Można więc pominąć tylko synchronizację w ostatniej wypreparowanej osnowie, która wynosi 32 aktywne wątki w fazie redukcji. Jeden krok, zanim będziesz potrzebować 64 wątków, a więc potrzebujesz synchronizacji, ponieważ wykonanie równoległe nie jest gwarantowane, ponieważ używasz 2 osnów.

Źródło

2012-05-24 02:23:11 djmj

To prawie chcę zrobić. Problem polega na tym, że kiedy opuszczam __syncthreads(), rzeczy zaczynają się łamać. I kod faktycznie działa w trybie debugowania, gdy łamie się w trybie zwolnienia. –

Czy masz zamiar wdrożyć redukcję opartą na osnowie? Zredukuj wnętrze osnowy, aby zmniejszyć dane o czynnik 32? tak więc przy 1024 wątkach/elementach potrzebne są tylko 2 synchreadi? Może to poprawić znacznie wydajność w porównaniu do konwencjonalnej implementacji. Sprawdzę ten pomysł później. – djmj

Problem, przed którym stoję, polega jedynie na zsumowaniu 128 numerów przechowywanych w pamięci współdzielonej. Nie mam do czynienia z problemem globalnej redukcji, ale to, co mówisz, może również zadziałać. –

Podam tutaj odpowiedź, ponieważ uważam, że powyższe dwa elementy nie są w pełni satysfakcjonujące. "Własność intelektualna" tej odpowiedzi należy do Marka Harrisa, który zwrócił uwagę na tę kwestię w tym numerze: presentation (slajd 22) oraz na @talonmies, który wskazał ten problem na PO w komentarzach powyżej.

Pozwól mi najpierw spróbować wznowić to, co OP zadawał, filtrując jego błędy.

OP wydaje się radzić sobie z ostatnim krokiem redukcji redukcji pamięci współdzielonej, redukcja zniekształceń przez rozwijanie pętli. On robi coś jak

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

Jak zauważył Mark Harris i talonmies, pamięci współdzielonej zmiennej sdata musi być zadeklarowany jako volatile, aby zapobiec optymalizacje kompilatora. Tak więc, we właściwy sposób zdefiniować funkcję __device__ powyżej jest:

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

Zobaczmy teraz zdemontowane kody odpowiadające obu przypadkach powyżej badanych, tj sdata zadeklarowana jako nie volatile lub volatile (kod skompilowany dla Fermiego architektury).

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

Linie /*0128*/-/*0148*/, /*0160*/ i /*0170*/ odpowiadają współdzielonych obciążenia pamięci do rejestrów i linii /*0190*/ do współużytkowanej pamięci sklepu z rejestru. Linie pośrednie odpowiadają sumacjom, jak wykonywane w rejestrach. Tak więc wyniki pośrednie są przechowywane w rejestrach (które są prywatne dla każdego wątku) i nie są przepłukiwane za każdym razem do pamięci współdzielonej, uniemożliwiając wątkom pełną widoczność wyników pośrednich.

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

Jak widać z linii /*0130*/-/*01e8*/ teraz za każdym razem gdy przeprowadza się sumowanie, w wyniku pośredniego natychmiast przedmuchano do pamięci wspólnej, aby je widoczność nici.

Źródło

2014-08-30 16:32:16 JackOLantern

Usuwanie __syncthreads() w CUDA redukcji osnowy poziomu

Odpowiedz

Powiązane problemy