Cuda effiziente Einfügung von Daten in unsortiertem Array

Question

Ich habe zwei Arrays innerhalb Cuda;

int *main; // unsorted 
int *source; // sorted 

Teil meines Algorithmus erfordert, dass ich neue Daten in das Hauptarray aus dem Quell-Array einfügen regulary. Wenn eine Position innerhalb des Hauptarrays Null ist, nimmt sie an, dass sie leer ist. Daher kann sie mit einem Wert aus dem Quellenarray gefüllt werden.

Ich frage mich nur, was die effizienteste Methode ist, dies zu tun, habe ich ein paar Ansätze versucht, aber immer noch denke, es gibt noch mehr Leistungsgewinne hier gemacht werden.

Derzeit verwende ich eine modifizierte Version einer Radix-Sortierung, um den Inhalt des Haupt-Arrays bis zum Ende des Haupt-Arrays zu "mischen", wobei alle Null-Werte am Anfang des Arrays liegen bleiben aus der Quelle trivial. Die Sortierung wurde so geändert, dass sie über ein einzelnes Bit anstatt über 32 Bits iteriert wird. Dies funktioniert mit einem einfachen Schalter am Eingang.

input[i] = source[i] > 1 ? 1 : 0 

Ich frage mich, ob dies schon recht ein effizienter Weg, dies zu tun? Ich frage mich, ob ich nicht etwas mit einem taktisch eingesetzten atomicAdd wie z.

Ich füge im Moment nicht so viele Elemente über das Quell-Array ein, aber das könnte sich in der Zukunft ändern.

Das fühlt sich an, als sollte es ein allgemeines Problem sein, das vorher gelöst wurde, ich frage mich, ob die Schubbibliothek helfen kann, aber nach geeigneten Funktionen suchend, fühlt es sich nicht richtig für das, was ich versuche (nicht sehr sauber passend mit dem Code, den ich bereits habe)

Gedanken von erfahrenen Cuda-Entwicklern geschätzt!

Answer 1

Sie können Ihren Suchalgorithmus, der als Stream-Komprimierungsverfahren kategorisiert ist, und Ihre Einfügung, die als Streuungsvorgang kategorisiert wird, entkoppeln. Sie können jedoch die Funktionalität der beiden zusammenführen.

Angenommen, srcPtr ist ein Zeiger, der seinen Inhalt im globalen Speicher befindet und bereits vor dem Start des Kernels auf Null gesetzt ist.

__global__ void find_and_insert(int* destination, int const* source, int const N, int* srcPtr) { // Assuming N is the length of the destination buffer and also the length of the source buffer is less than N. 

int const idx = blockIdx.x * blockDim.x + threadIdx.x; 

// Get the assigned element. 
int const dstElem = destination[ idx ]; 
bool const pred = (dstElem == 0); 

// Intra-warp binary reduction to count the total number of lanes with empty elements. 
int const predBallot = __ballot(pred); 
int const intraWarpRed = __popc(predBallot); 

// Warp-aggregated atomics to reduce the contention over the srcPtr content. 
unsigned int laneID; asm("mov.u32 %0, %laneid;" : "=r"(laneID)); //const uint laneID = tidWithinCTA & (WARP_SIZE - 1); 
int posW; 
if(laneID == 0) 
    posW = atomicAdd(srcPtr, intraWarpRed); 
posW = __shfl(posW, 0); 

// Threads that have found empty elements can fill out their assigned positions from the src. Intra-warp binary prefix sum is used here. 
uint laneMask; asm("mov.u32 %0, %lanemask_lt;" : "=r"(laneMask)); //const uint laneMask = 0xFFFFFFFF >> (WARP_SIZE - laneID) ; 
int const positionToRead = posW + __popc(predBallot & laneMask); 
if(pred) 
    destination[ idx ] = source[ positionToRead ]; 

} 

Ein paar Dinge:

1. Dieser Kernel ist nur ein Vorschlag, wie Sie es tun können. Hier arbeiten Threads innerhalb der Warps bei der Aufgabe zusammen. Sie können die Binärreduktion und die Präfixsumme über den Thread-Block erweitern.
2. Ich schrieb diesen Kernel in den Browser und habe es nicht getestet. Also sei vorsichtig.
3. Das ganze Design ist nicht etwas Neues. Ähnliche Ansätze wurden implementiert (zum Beispiel this paper) und basieren größtenteils auf the work done by Mark Harris and Michael Garland.