c++ - 如何实现像 a>b?c :d? 这样的原子条件运算符函数

Question

我是 CUDA 程序的新手，我需要实现一个像这样的原子函数，atomicMyFunc(int *address, int compare, int val1, int val2)，“old”是值位于地址。

我已经尝试了好几天了，但都失败了。应该是这样的:

__device__ int atomicMyFunc(int * address, int compare, int val1, int val2)
{
    if(old>compare?val1:val2) // compare and then store val1 or val2 in the address
    {
        // do something
    }
    return old;
}

Answer 1

您需要使用 CUDA 中可用的原子函数来构建原子函数。我建议检查 CUDA programming guide有关如何构建原子函数的示例。以那里提供的示例为起点，我认为以下功能可以完成您想要做的事情

__device__ int atomicMyFunc(int* address, int compare, int val1, int val2) {
    int old = *address;
    int assumed;

    do {
        assumed = old;
        old = atomicCAS(address, assumed, assumed > compare ? val1: val2);
    }
    while (assumed != old);

    return old;
}

它基本上在循环中使用 atomicCAS() 以便仅在使用当前存储在所提供地址的相同值计算运算结果时存储运算结果。