Kooperative Gruppen in CUDA

Question

seit der Veröffentlichung von CUDA 9 ist es anscheinend möglich, verschiedene Threads und Blöcke in derselben Gruppe zu gruppieren, so dass Sie sie gemeinsam verwalten können. Das ist sehr nützlich für mich, weil ich einen Kernel mit mehreren Blöcken starten muss und warten bis alle synchronisiert sind (cudaThreadSynchronize() ist nicht wert für mich, denn nachdem die Threads synchronisiert sind, muss ich weiter in meinem Kernel arbeiten).

Was ich dachte ist, diese Blöcke von Threads in die gleiche Gruppe aufzunehmen und zu warten, bis alle synchronisiert sind, wie die Beispiele der Nvdia-Hauptseite vermuten lassen.

Sie etwas tun, wie folgt aus:

__device__ int reduce_sum(thread_group g, int *temp, int val) 
{ 
    int lane = g.thread_rank(); 

// Each iteration halves the number of active threads 
// Each thread adds its partial sum[i] to sum[lane+i] 
for (int i = g.size()/2; i > 0; i /= 2) 
{ 
    temp[lane] = val; 
    g.sync(); // wait for all threads to store 
    if(lane<i) val += temp[lane + i]; 
    g.sync(); // wait for all threads to load 
} 

Mein Problem ist, wie diese Blöcke zu einer Gruppe in die g Gruppe. Diese ist, wie ich meinen Kernel ursprünglich ins Leben gerufen:

asap << <5, 1000 >> > (cuda_E2, cuda_A2, cuda_temp, Nb, *binM, Nspb); 

Jedes Mal, dass ich versuche, thread_group die Compiler sagt zu verwenden, dass es undefied wird. Ich verwende die kooperative_groups.h Header.

Kann jemand damit umgehen? Danke im Voraus.

Answer 1

Zitat aus dem documentation:

Cooperative Gruppen erfordert CUDA 9.0 oder höher. So verwenden Cooperative Gruppen, umfassen die Header-Datei:

#include <cooperative_groups.h> 

und verwenden Sie die Cooperative Gruppen Namespace:

using namespace cooperative_groups; 

Dann Code enthält jede Intrablock Cooperative Gruppen-Funktionalität kann in der normalen kompiliert werden Weg mit nvcc.

Der Namensraum ist, was Sie vermissen.