Cuda atomare Operationen

Question

Ich habe folgende kernel:

__global__ 
void collect_boundary(const int64_t* cvert, const csr_node* neighb, const bool* affected, int64_t* parent, const uint64_t* dist, uint64_t* ndist, bool* mask, int64_t numvertices){ 
    int64_t tid = blockIdx.x*blockDim.x + threadIdx.x; 
    if(tid >= numvertices || affected[tid] || dist[tid]==MY_INFINITY) 
     return; 
    for(int64_t index = cvert[tid]; index<cvert[tid+1]; index++){ 
     auto vtex = neighb[index]; 
     if(affected[vtex.head]){ 
      int64_t ndistent = dist[tid] + vtex.weight; 
      atomicMin((unsigned long long int*)(ndist + vtex.head),(unsigned long long int)ndistent); 
      /*if(ndist[vtex.head] == ndistent){ 
       parent[vtex.head] = tid; 
      }*/ 
     } 
    } 
} 

Grundsätzlich i jeder Thread als gegeben berechnen ndistent wollte und i NDIST [vtex.head] als das Minimum aller ndistents aktualisieren.

ich erreicht dies mit:

atomicMin((unsigned long long int*)(ndist + vtex.head),(unsigned long long int)ndistent); 

//That is each thread will update ndist[vtex.head] if and only if 
//it's own value of ndistent is less than the ndist[vtex.head] 
//which was initialized to INFINITY before the kernel launch 

Aber jetzt wollte ich die tid speichern, die die minimale ndistent gibt.

Ich habe versucht, so etwas wie dieses

if(ndist[vtex.head] == ndistent){ // prob_condition 1 
    parent[vtex.head] = tid;  // prob_statment 1 
} 

//That is each thread will check wether the value in 
//ndist[vtex.head] is equal to it's own ndistent 
// and then store the tid if it is. 

Das obige Snippet wird nicht funktionieren, weil einige Thread X kann, dass prob_condition 1 um wahr zu sein finden, aber bevor es prob_statement führt 1 lassen Sie uns den Faden sagen die gibt den minimalen Wert sagen Thread Y führt prob_statement 1 und speichert es ist tid. Jetzt wird Thread X fortgesetzt und speichert es ist tid, so dass die min tid verloren ist.

Also ich prob_condition 1 und prob_statement 1 atomically ausgeführt werden soll.

Oder alternativ muss ich die folgenden 3 opertaions atomar tun:

1. Überprüfen Sie, ob ndistent < NDIST [vtex.head]

2. Update NDIST [vtex.head]

3. store tid in übergeordnetes Element [vtex.head]

Hat jemand irgendwelche Vorschläge wie ich das machen kann?

EDIT: Bitte beachten Sie, dass ich diesen Kernel mit variabler Anzahl von Blöcken und variable Anzahl von Threads laufen muß.

Answer 1

Es könnte Ihr Nebenläufigkeitsproblem nicht so lösen, wie Sie es vorhaben, aber Sie könnten einen zweiphasigen Ansatz haben: Berechnen Sie zuerst das Min und suchen Sie dann die Typen, die dieses Min haben.

Wenn mehrere tid den gleichen ndistenten Wert haben, kann die Ausgabe von einer Ausführung zur anderen variieren, in der Tat, wie Taro darauf hingewiesen hat, gehorcht die Ausführungsreihenfolge der Warps nicht vorhersagbaren Regeln. Dieser Ansatz mit zwei Phasen kann Ihnen dabei helfen, ein vorhersagbares Muster für die Minima-Liste zu erstellen.

In einer Hacky Ansatz, wenn ndistent Wert und tid kann passen sowohl in 64 Bit, können Sie mit ndistent und niedriger Ordnung höherwertigen Bits eines 64-Bit-Wert zugeführt versuchen, die Bits, die den tid zu halten und zu tun das atomicMin in einer Anweisung.