Wie packt man Bits (effizient) in CUDA?

Ich habe ein Array von Bytes, wobei jedes Byte entweder 0 oder 1 ist. Jetzt möchte ich diese Werte in Bits packen, so dass 8 Original-Bytes 1 Ziel-Byte belegen, wobei Original-Byte 0 in Bit 0, Byte 1 in Bit geht 1 usw. Bisher habe ich folgendes im Kernel:Wie packt man Bits (effizient) in CUDA?

const uint16_t tid = threadIdx.x; 
__shared__ uint8_t packing[cBlockSize]; 

// ... Computation of the original bytes in packing[tid] 
__syncthreads(); 

if ((tid & 4) == 0) 
{ 
    packing[tid] |= packing[tid | 4] << 4; 
} 
if ((tid & 6) == 0) 
{ 
    packing[tid] |= packing[tid | 2] << 2; 
} 
if ((tid & 7) == 0) 
{ 
    pOutput[(tid + blockDim.x*blockIdx.x)>>3] = packing[tid] | (packing[tid | 1] << 1); 
}

Ist das richtig und effizient?

Quelle

2016-09-14 Serge Rogatch

Dies kann nicht funktionieren. Das ist ein Erinnerungsrennen. Es gibt keine Transaktionen mit parallelen Bitgrößen in CUDA – talonmies

@talonmies, ich dachte, es gäbe keine Konkurrenz, weil Threads, die dasselbe Byte verarbeiten, zum selben Warp gehören. –

Im selben Warp zu sein ist keine Garantie für Sicherheit. Keine zwei Threads können dasselbe Byte gleichzeitig ändern, ohne ein Rennen zu verursachen. – talonmies

Die Warp-Voting-Funktion __ballot() ist sehr praktisch. Unter der Annahme, dass Sie pOutput sein von uint32_t Typ neu definieren kann, und dass Ihre Blockgröße ist ein Vielfaches der Kettschlichte (32):

unsigned int target = __ballot(packing[tid]); 
if (tid % warpSize == 0) { 
    pOutput[(tid + blockDim.x*blockIdx.x)/warpSize] = target; 
}

Streng genommen ist die, wenn bedingt ist nicht einmal notwendig, da alle Threads des Warps schreiben die gleichen Daten an dieselbe Adresse. So eine hoch optimierte Version wäre nur

pOutput[(tid + blockDim.x*blockIdx.x)/warpSize] = __ballot(packing[tid]);

Quelle

2016-09-14 10:57:20 tera

Tolle Lösung, danke. Dann brauche ich keinen gemeinsamen Speicher und '__synctreads()'. –

Schöne Antwort auf eine gute Frage. – harrism

Wie ich verstehe, skaliert dies nicht zum Verpacken von 2-Bit-Werten? Auch wenn wir 2 '__ballot' Aufrufe verwenden können, um niedrigere und höhere Bits in 2 separaten 32-Bit Variablen zu erhalten, dann verschachteln Sie die Bits (zB mit http://stackoverflow.com/questions/39490345/interleave-bits-efficiently) ist auf CUDA teurer als der Algorithmus, den ich in der Frage gegeben habe. Für das Packen von 2-Bit-Werten müsste dieser Algorithmus die Ausgabe innerhalb von 'if ((tid & 6) == 0)' schreiben. –

Für zwei Bits pro Faden, uint2 *pOutput

int lane = tid % warpSize; 
uint2 target; 
target.x = __ballot(__shfl(packing[tid], lane/2)    & (lane & 1) + 1)); 
target.y = __ballot(__shfl(packing[tid], lane/2 + warpSize/2) & (lane & 1) + 1)); 
pOutput[(tid + blockDim.x*blockIdx.x)/warpSize] = target;

mit Sie Benchmark haben werden, ob dies noch schneller als die herkömmliche Lösung.

Quelle

2016-09-15 15:12:13 tera

Wie packt man Bits (effizient) in CUDA?

Antwort

Verwandte Themen