c++ - CUDA原子操作

我有以下内核：

__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

但是现在我想存储给出最小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为true，但是在执行prob_statement 1之前，让我们说要赋予最小值的线程说线程Y执行prob_statement 1并存储其tid。现在线程X将恢复并存储它的tid，因此min tid丢失了。

所以我希望prob_condition 1和prob_statement 1被原子执行。

或者，我需要自动执行以下3种操作：


检查ndistent 更新ndist [vtex.head]
将tid存储在parent [vtex.head]中


任何人有任何建议我该怎么做？

编辑：
请注意，我将必须使用可变数量的块和可变数量的线程来运行此内核。

最佳答案

它可能无法按您预期的方式解决并发问题，但是您可以采用两个阶段的方法：首先计算最小值，然后找到具有该最小值的人员。

同样，如果多个tid具有相同的ndistent值，则输出可能会因一次执行而异，实际上，正如Taro所指出的那样，扭曲的执行顺序不遵循可预测的规则。这两个阶段的方法可以帮助您为最小值列表建立可预测的模式。

在一种更狡猾的方法中，如果ndistent值和tid都可以容纳64位，则可以尝试让64bits值的高阶位输入ndistent和低阶位来保存tid，然后在一条指令中执行atomicMin。