atomicMax + AtomicCAS(atomicExch)

Question

想请教各位大佬有没有更好的方法把2个原子结合起来

我的目标是在 J 参数值列表（非常类似于 2 路输入）下找到一组 K 方程（超过 32 个）的最高结果，并保存该值和 j 索引。

if (atomicMax(&max_k[id], t_max) < t_max) atomicExch(&indexMax[id],t_pos);

最初我们使用了上述方法，但是，由于我们确实希望每个线程都有更高的值，因此在同一个 warp 中可以有线程 B > C > A（线程 B 具有最高值并且线程 C 的值高于 A)。 我不确定，但 atomicExch 可以按不同于 atomicMax 的线程顺序执行（正确吗？），所以我们尝试了临界区，但它导致了死锁。毕竟下面的所有解决方案似乎都有效。

有没有更好的方法或者下面的代码有什么问题？

__device__ int atomicMaxCAS(int* addressMax, int valMax, int* addressCAS, int valCas) {
        int oldCas = *addressCAS, assumedCas;
        int oldMax = *addressMax, assumedMax;
        do {
            assumedCas = oldCas;
            assumedMax = oldMax;
            oldMax = atomicMax(addressMax, valMax);
            if (oldMax < valMax) oldCas = atomicCAS(addressCAS, assumedCas, valCas);
        } while (assumedCas != oldCas || assumedMax != oldMax);
        return (oldMax);
    }

提前致谢！由于所有这些帖子，我能够开始编写 CUDA！

Answer 1

there is any issue in the following code?

是的，您不能像那样使用两个原子并期望得到连贯的结果。您已经设置了一个可能的竞争条件。

假设线程 A 执行 atomicMax 并将旧值替换为 100。然后线程 B 执行 atomicMax 并将 100 值替换为 110。然后假设线程 B 执行 atomicCAS，并替换其索引。然后线程 A 执行 atomicCAS，并用线程 A 索引替换线程 B 索引。您现在的最大值为 110，索引对应于线程 A。

即使在单个 warp 中，也没有规定的原子操作执行顺序。

Is there a better way?

因为您的值都是 32 位数量，您可能有兴趣使用像 this 这样的自定义 64 位原子操作来同时更新值和索引, 原子地.
对于大规模使用（大量线程），您可能想要探索 classical paraellel reduction. There are questions here on the CUDA tag such as and 讨论如何进行索引+值缩减。

Kepler 上的全局原子非常快，因此根据您的确切代码和缩减 "density" 全局原子缩减在性能方面可能不是大问题。

atomicMax + AtomicCAS(atomicExch)

atomicMax + AtomicCAS(atomicExch)

cuda

atomic