Cuda原子操作

Cuda atomic operations

本文关键字：原子操作 Cuda      更新时间：2023-10-16

我有以下内核：

__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]更新为所有ndistent的最小值。

我使用实现了这一点

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为true，但在它执行prob_statement 1之前，让我们假设将给出最小值的线程，比如线程Y执行prob_sttatement 1并存储它的tid。现在线程X将恢复并存储它的tid，因此min-tid丢失。

因此，我希望prob_condition 1和prob_statement 1以原子方式执行。

或者，我需要原子地执行以下3个操作：

  1. 检查ndistent<ndist[vtex.head]

  2. 更新ndist[vtex.head]

  3. 将tid存储在父中

有人对我该怎么做有什么建议吗？

编辑：请注意，我将不得不用可变数量的块和可变数量的线程来运行这个内核。

它可能无法按照您的意愿解决并发问题，但您可以有两个阶段的方法：首先，计算最小值，然后找到具有最小值的人。

此外，如果几个tid具有相同的ndistent值，则输出可能会因一次执行而不同，事实上，正如Taro所指出的，warps的执行顺序不符合可预测规则。这种分两个阶段的方法可以帮助您为最小值列表建立一个可预测的模式。

在一种更为简单的方法中，如果ndistent值和tid都可以容纳64位，则可以尝试将64位值的高阶位与ndistent一起馈送，并将低阶位保存tid，然后在一条指令中执行atomicMin。

相关文章：