Mergesort CUDA

Question

Ich habe versucht, mein eigenes Mergesort zu implementieren, basierend auf von unten nach oben/iterativer mergesort Algorithmus. Dieser Algorithmus teilt die Daten nach 2 Elementen und sortiert sie. Dann um 4 Elemente und sortiert und so weiter, bis alle Daten sortiert sind. Also, mein Plan ist jedem Thread durch 2 Elemente zuzuordnen. So mache ich das:

__global__ void mergeBU(int *d_a, int *d_aux, int sz, int N) 
{ 
    int idk = blockIdx.x*blockDim.x+threadIdx.x; 
    int lo = 2 * sz * idk; 
    int mid = lo + sz - 1; 
    float hi = fminf(lo + sz + sz - 1, N - 1); 
    merge(d_a, d_aux, lo, mid, hi); 
} 

__device__ void merge(int *d_a, int *d_aux, int lo, int mid, float hi) 
{ 
int i = lo; 
int j = mid + 1; 

    for (int k = lo; k <= hi; k++) 
    { 
     d_aux[k] = d_a[k]; 
    } 

    for (int k = lo; k <= hi; k++) 
    { 
     if (i > mid)     { d_a[k] = d_aux[j]; j++; } 
     else if (j > hi)    { d_a[k] = d_aux[i]; i++; } 
     else if (d_aux[j] < d_aux[i]) { d_a[k] = d_aux[j]; j++; } 
     else        { d_a[k] = d_aux[i]; i++; } 
    } 
} 

Sagen wir, ich meine Kernel aufrufen < < < 2,4 >>> (die 8 Threads ist), so kann ich nur Art 16 Elemente max. Wenn ich 32 Elemente eingabe, bleiben die restlichen Datenindexe unberührt (16-31). Wie kann der Thread-Index weiterhin den Rest des Datenindex verarbeiten? Mit fortfahren meine ich, dass der Thread-Index (0,1,2,3,4,5,6,7) den Rest des Datenindex weiter verarbeitet, er sollte wie threadindex (dataindex, dataindex) sein -> 0 (16, 17); 1 (18,19); 2 (20,21); und so weiter. Jeder Kommentar ist willkommen.

Answer 1

Ohne an Ihrem eigentlichen Code zu suchen: Sortierung Merge ist ein Multi-Pass-Algorithmus. Da verschiedene Blöcke beim Ausführen eines Kernels normalerweise nicht synchronisiert werden (es sei denn, Sie verwenden atomare Einheiten), sollten Sie mehrere aufeinanderfolgende Kernel-Starts in Erwägung ziehen, einen für jeden Durchgang. Zum Beispiel sortiert beim ersten Start jeder Threadblock n_1 Elemente; Beim zweiten Start verbindet jedes Blockpaar 2 * n_1 Elemente und so weiter. Das ist natürlich nicht so einfach, wie es sich anhört: Wie kann man sagen, welcher Block was genau tun soll?

Auch möchten Sie vielleicht am approach used in the ModernGPU library für andere Ideen einen Blick haben.

Answer 2

Ihr Ansatz besteht darin, ein Array der Größe n in n/2 Sub-Arrays aufzuteilen, Paare dieser Sub-Arrays zu kombinieren, um schließlich mit n/4 Sub-Arrays zu enden und so weiter. Dieser Ansatz würde jedoch wahrscheinlich Speicherbandbreite begrenzt sein.

Angenommen, Sie verwenden 8 "Gewinde". Teilen Sie das Array in 8 Sub-Arrays der Größe n/8 auf (das letzte Sub-Array kann eine andere Größe haben), verwenden Sie dann 8 Threads, um das Sub-Array zu sortieren, und 4 Threads, um die 4 Paare des sortierten Subs zusammenzuführen - Arrays, dann 2 Threads zum Zusammenführen der 2 Paare von zusammengefügten Sub-Arrays, dann 1 Thread zum Zusammenführen der letzten 2 Paare.

Basierend auf meiner Erfahrung mit einer Multi-Thread-Sortierung, erreichen Sie die Speicherbandbreite bei 8 Threads für eine CPU, aber wenn der GPU-Speicher verwendet werden kann, um große Teile des Arrays zu halten, dann mehr als 8 Threads nützlich sein. Ich weiß nicht, welche Operationen (vergleichen, bewegen, ...) in der GPU und ihrem Speicher möglich sind.