Cuda атомарные операции

Question

У меня есть следующие ядра:

__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; 
      }*/ 
     } 
    } 
} 

В основном я хотел каждая нить вычислить ndistent как данность и я обновить ndist [vtex.head] как минимум всех ndistents.

Я достиг этого с помощью:

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 

Но теперь я хотел сохранить TID, который дает минимальную ndistent.

Я пытался что-то вроде этого

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. 

Это выше фрагмент кода не будет работать, потому что некоторые нить X может обнаружить, что prob_condition 1, чтобы быть правдой, но, прежде чем он выполняет prob_statement 1 допустим нить, которая даст минимальное значение скажем, что нить Y выполняет prob_statement 1 и сохраняет его. Теперь поток X возобновится и сохранит его, поэтому min tid будет потерян.

Поэтому я хочу, чтобы prob_condition 1 и prob_statement 1 выполнялись атомарно.

Или же мне нужно сделать следующие 3 opertaions атомарно:

1. Проверьте, ndistent < ndist [vtex.head]

2. обновление ndist [vtex.head]

3. хранить tid в исходном состоянии [vtex.head]

У кого-нибудь есть предложения, как я могу это сделать?

EDIT: Обратите внимание, что мне нужно будет запустить это ядро ​​с переменным числом блоков и числом переменных.

Answer 1

Это может не решить вашу проблему параллелизма так, как вы ее намереваетесь, но вы можете использовать двухэтапный подход: сначала вычислите min, а затем найдите парня, у которого есть этот мин.

Кроме того, если несколько tid имеют одно и то же неизменяемое значение, выход может отличаться от одного исполнения к другому, ведь, как указал Таро, порядок выполнения варпсов не подчиняется правилам предикатива. Этот подход, основанный на двух фазах, может помочь вам построить прогнозируемый шаблон для списка минимумов.

В более Hacky подход, если ndistent значение и TID может и помещается в 64 бита, вы можете попытаться иметь биты высокого порядка в 64bits значения, питаемой ndistent и низкого порядка бит проведения TID, и делать атомный Min в одной инструкции.