Beschleunigen CUDA Atomic Berechnung für viele Bins/einige Bins

Question

Ich versuche, meine Histogramm-Berechnungen in CUDA zu optimieren. Es gibt mir eine ausgezeichnete Beschleunigung gegenüber der entsprechenden OpenMP-CPU-Berechnung. Ich vermute jedoch (nach Intuition), dass die meisten Pixel in ein paar Eimer fallen. Nehmen wir an, dass wir 256 Pixel haben, die in zwei Buckets fallen.

Der einfachste Weg, es zu tun ist, es zu tun

1. laden die Variablen in den gemeinsamen Speicher
   • Sie vektorisiert Lasten für unsigned char usw. zu sein scheint, wenn nötig.
2. ein Atom Fügen Sie im gemeinsam genutzten Speicher
3. auf globale einen verschmolzenen schreiben Sie.

Etwas wie folgt:

__global__ void shmem_atomics_reducer(int *data, int *count){ 
    uint tid = blockIdx.x*blockDim.x + threadIdx.x; 

    __shared__ int block_reduced[NUM_THREADS_PER_BLOCK]; 
    block_reduced[threadIdx.x] = 0; 

    __syncthreads(); 

    atomicAdd(&block_reduced[data[tid]],1); 
    __syncthreads(); 

    for(int i=threadIdx.x; i<NUM_BINS; i+=NUM_BINS) 
    atomicAdd(&count[i],block_reduced[i]); 

} 

Die Leistung dieses kernel fällt (natürlich), wenn wir eine Verringerung der Anzahl der Bins, von etwa 45 GB/s bei 32 Behältern auf etwa 10 GB/s bei 1 Tonne. Konflikte und geteilte Speicherbankkonflikte werden als Gründe angegeben. Ich weiß nicht, ob es irgendeinen Weg gibt, um einen von diesen für diese Berechnung in bedeutender Weise zu entfernen.

Ich habe auch mit einer anderen (schönen) Idee aus dem Parallelforall-Blog experimentiert, die Warp-Level-Reduzierungen mit __ballot zum Erfassen von Warp-Ergebnissen und dann mit __popc() zur Reduzierung des Warp-Levels verwendet.

__global__ void ballot_popc_reducer(int *data, int *count){ 
    uint tid = blockIdx.x*blockDim.x + threadIdx.x; 
    uint warp_id = threadIdx.x >> 5; 

    //need lane_ids since we are going warp level 
    uint lane_id = threadIdx.x%32; 

    //for ballot 
    uint warp_set_bits=0; 

    //to store warp level sum 
    __shared__ uint warp_reduced_count[NUM_WARPS_PER_BLOCK]; 
    //shared data 
    __shared__ uint s_data[NUM_THREADS_PER_BLOCK]; 

//load shared data - could store to registers 
    s_data[threadIdx.x] = data[tid]; 

    __syncthreads(); 


//suspicious loop - I think we need more parallelism 
    for(int i=0; i<NUM_BINS; i++){ 
     warp_set_bits = __ballot(s_data[threadIdx.x]==i); 

     if(lane_id==0){ 
     warp_reduced_count[warp_id] = __popc(warp_set_bits); 
     } 

    __syncthreads(); 

     //do warp level reduce 
     //could use shfl, but it does not change the overall picture 
     if(warp_id==0){ 
     int t = threadIdx.x; 
     for(int j = NUM_WARPS_PER_BLOCK/2; j>0; j>>=1){ 
      if(t<j) warp_reduced_count[t] += warp_reduced_count[t+j]; 
      __syncthreads(); 
     } 
     }                                                                 



     __syncthreads(); 


     if(threadIdx.x==0){ 
     atomicAdd(&count[i],warp_reduced_count[0]); 
     } 

    }                                                            

    } 

Dies gibt anständige Zahlen (gut, dass strittig ist - Spitzengerät mem bw 133 GB/s, die Dinge scheinen auf Startkonfiguration abhängig) für die einzelnen Behälter Fall (35-40 GB/s für 1 bin, verglichen mit 10-15 GB/s mit Atomics), aber die Leistung sinkt drastisch, wenn wir die Anzahl der Bins erhöhen. Wenn wir mit 32 Bins arbeiten, sinkt die Leistung auf etwa 5 GB/s. Der Grund könnte vielleicht darin liegen, dass der einzelne Thread alle Bins durchläuft und nach der Parallelisierung der NUM_BINS-Schleife fragt.

Ich habe mehrere Möglichkeiten der Parallelisierung der NUM_BINS-Schleife versucht, von denen keine ordnungsgemäß zu funktionieren scheint. Zum Beispiel könnte man (sehr unelegant) den Kernel manipulieren, um einige Blöcke für jedes Bin zu erzeugen. Dies scheint sich auf die gleiche Weise zu verhalten, möglicherweise, weil wir wieder unter dem Konflikt mit mehreren Blöcken leiden würden, die versuchen, aus dem globalen Speicher zu lesen. Plus, die Programmierung ist klobig. Gleichermaßen ergibt die Parallelisierung in der y-Richtung für Behälter ähnlich wenig inspirierende Ergebnisse.

Die andere Idee, die ich nur für Kicks versuchte, war dynamische Parallelität, Starten eines Kernels für jeden Behälter. Dies war katastrophal langsam, möglicherweise aufgrund der Tatsache, dass keine echte Rechenarbeit für die Kind-Kernel und den Start-Overhead geleistet wurde.

Der vielversprechendste Ansatz scheint zu sein - von Nicholas Wilts article

zur Verwendung der sogenannten privatisierten Histogramme enthalten Bins für jeden Thread im gemeinsam genutzten Speicher, die angeblich sehr schwer auf shmem Nutzung wäre (und wir nur haben 48 kB pro SM auf Maxwell).

Vielleicht könnte jemand Einblick in das Problem geben? Ich finde, dass man stattdessen den Algorithmus ändern sollte, um keine Histogramme zu verwenden, um etwas weniger Frequentist zu verwenden. Ansonsten nehme ich einfach die Atomic-Version.

Bearbeiten: Der Kontext für mein Problem besteht in der Berechnung von Wahrscheinlichkeitsdichtefunktionen, die für die Musterklassifizierung verwendet werden sollen. Wir können ungefähre Histogramme (genauer gesagt, PDFs) berechnen, indem wir nichtparametrische Methoden wie Parzen Windows oder Kernel Density Estimation verwenden. Dies überwindet jedoch nicht das Problem der Dimensionalität, da wir alle Datenpunkte für jeden Behälter summieren müssen, was teuer wird, wenn die Anzahl von Behältern groß wird. Siehe hier: Parzen

Answer 1

Ich war mit ähnlichen Chalanges konfrontiert, um mit Clustering zu arbeiten, aber am Ende war die beste Lösung, das Scanmuster zu verwenden, um die Verarbeitung zu gruppieren. Also, ich denke nicht, dass es für dich funktionieren würde. Da du nach etwas Erfahrung gefragt hast, werde ich meine mit dir teilen.

Die Themen

In Ihrem ersten Code, ich denke, dass das Geschäft mit der niedrigen Leistung mit der Anzahl der Bins Reduktion Stall verziehen wird verbunden, da sie alle ausgewerteten Daten sehr wenig Verarbeitung durchführen können. Wenn die Anzahl der Bins erhöht wird, erhöht sich auch die Beziehung zwischen der Verarbeitung und dem Laden des globalen Speichers (Dateninfo) für diesen Kernel. Das können Sie sehr einfach mit den "Issue Efficiency" Experiments bei der Performance Analyse von Nsight überprüfen. Wahrscheinlich erhalten Sie eine niedrige Rate von Zyklen mit mindestens einer eleganten Kette (Warp Issue Efficiency).

Da ich nicht in der Lage war in der Nähe von 95%, die Anzahl der wahlberechtigten Warps irgendwo zu verbessern, habe ich diesen Ansatz auf, da in einigen Fällen es noch schlimmer (der Stall Speicherabhängigkeit erhält 90% meiner Verarbeitungszyklen.

die Shuffle und Abstimmung Reduktion ist sehr nützlich, wenn die Anzahl der Behälter nicht zu groß ist. wenn es zu groß ist, eine kleine Menge von Threads ist, sollte für jeden bin Filter aktiv sein. So können Sie mit einem bis Ende Mai viel Code-Divergenz, und das ist sehr unerwünscht für die parallele Verarbeitung.Sie können versuchen, die Divergenz zu gruppieren, um Verzweigungen zu entfernen und einen guten Kontrollfluss zu haben, so dass der gesamte Warp/Block eine ähnliche Verarbeitung darstellt, aber a viel Chance über Blöcke.

Eine machbare Lösung

Ich weiß nicht, wo, aber es gibt sehr gute Lösungen für Ihr Problem um, dass ich gesehen habe. Hast du versucht this one?

Sie können auch eine vectorized load verwenden und so etwas versuchen, aber ich bin nicht sicher, wie viel würde es Ihre Leistung zu verbessern:

__global__ hist(int4 *data, int *count, int N, int rem, unsigned int init) { 

__shared__ unsigned int sBins[N_OF_BINS]; // you may want to declare this one dinamically 
int idx = blockIdx.x * blockDim.x + threadIdx.x; 
if (threadIdx.x < N_OF_BINS) sBins[threadIdx.x] = 0; 

for (int i = 0; i < N; i+= warpSize) { 
    atomicAdd(&sBins[data[i + init].w], 1); 
    atomicAdd(&sBins[data[i + init].x], 1); 
    atomicAdd(&sBins[data[i + init].y], 1); 
    atomicAdd(&sBins[data[i + init].z], 1); 
} 

//process remaining elements if the data is not multiple of 4 
// using recast and a additional control 
for (int i = 0; i < rem; i++) { 
    atomicAdd(&sBins[reinterpret_cast<int*>(data)[N * 4 + init + i]], 1); 
} 
//update your histogram data here 
}