Parallel Reduction

Question

Ich habe den Artikel Optimierung der parallelen Reduktion in CUDA von Mark Harris gelesen, und ich fand es wirklich sehr nützlich, aber immer noch kann ich manchmal 1 oder 2 Konzepte nicht verstehen. Es steht geschrieben, auf S. 18:

//First add during load 

// each thread loads one element from global to shared mem 

unsigned int tid = threadIdx.x; 

unsigned int i = blockIdx.x*blockDim.x + threadIdx.x; 

sdata[tid] = g_idata[i]; 
__syncthreads(); 

optimierte Code: Mit 2 Lasten und ersten Zusatz der Reduktion:

// perform first level of reduction, 

// reading from global memory, writing to shared memory 
unsigned int tid = threadIdx.x;         ...1 

unsigned int i = blockIdx.x*(blockDim.x*2) + threadIdx.x;   ...2 

sdata[tid] = g_idata[i] + g_idata[i+blockDim.x];     ...3 

__syncthreads();             ...4 

vermögen mich nicht zu verstehen, Linie 2; Wenn ich 256 Elemente habe und 128 als meine Blockgröße wähle, warum multipliziere ich dann mit 2? Bitte erläutern, wie man die Blockgröße ermittelt?

Answer 1

Es Durchführung grundsätzlich der Betrieb in der Abbildung unten gezeigt:

Dieser Code ist im Grunde „sagen“, dass die Hälfte der Fäden werden die Leistung der Lesung aus dem globalen Speicher und in dem gemeinsamen Speicher zu schreiben, als zeigte auf dem Bild.

Sie führen eine Kernell, und jetzt möchten Sie einige Werte reduzieren, Sie beschränken den Zugriff auf den Code oben auf nur die Hälfte der gesamten laufenden Threads. Imagining Sie 4 Blöcke haben, die jeweils mit 512 Threads, begrenzen Sie den obigen Code in nur durch die beiden ersten Blöcke ausführen, und Sie haben eine g_idate [4 * 512]:

unsigned int i = blockIdx.x*(blockDim.x*2) + threadIdx.x; 

sdata[tid] = g_idata[i] + g_idata[i+blockDim.x]; 

So:

thread 0 of block = 0 will copy the position 0 and 512, 
thread 1 of block = 0 position 1 and 513; 
thread 511 of block = 0 position 511 and 1023; 
thread 0 of block 1 position 1024 and 1536 
thread 511 of block = 1 position 1535 and 2047 

Der BlockDim.x * 2 wird verwendet, da jeder Thread auf die Positionen i und i+blockDim.x zugreift, so dass Sie mehrfache 2 vergeben, um sicherzustellen, dass die Threads des nächsten ID-Blocks nicht dieselbe Position von g_idata bereits berechnet überlappen.

Answer 2

In optimiertem Code führen Sie den Kernel mit halb so großen Blöcken aus wie in der nicht optimierten Implementierung. Nennen wir die Größe des Blocks in nicht optimiertem Code work, lassen Sie die Hälfte dieser Größe unit heißen, und lassen Sie diese Größen denselben numerischen Wert für den optimierten Code haben.

In nicht optimiertem Code führen Sie den Kernel mit so vielen Threads wie der work ist, das ist blockDim = 2 * unit. Der Code in jedem Block kopiert nur einen Teil von g_idata in ein Array im Shared Memory der Größe 2 * unit.

In dem optimierten Code blockDim = unit, so gibt es nun 1/2 der Threads, und das Array im Shared Memory ist 2x kleiner. In Zeile 3 kommt der erste Summand aus geraden Einheiten, der zweite aus ungeraden Einheiten. Auf diese Weise werden alle für die Reduzierung erforderlichen Daten berücksichtigt.

Beispiel: Wenn Sie nicht optimierten Kernel mit blockDim=256=work (Einzelblock, unit=128) laufen, dann optimierten Code hat einen einzigen Block von blockDim=128=unit. Da dieser Block blockIdx=0 bekommt, spielt die *2 keine Rolle; der erste Thread tut g_idata[0] + g_idata[0 + 128].

Wenn Sie hatte 512 Elemente und führen nicht optimierten mit 2 Blöcke der Größe 256 (work=256, unit=128), so optimierten Code besteht aus 2 Blöcken, aber jetzt der Größe 128. Das erste Gewinde in dem zweiten Block (blockIdx=1) tut g_idata[2*128] + g_idata[2*128+128].