CUDA：如何原子地执行这些指令（4个地址）

Question

我正在 Cuda 中实现一个算法，需要执行以下步骤：

给定一个数组 x（在共享内存中）和一些设备函数 f，

Select 一对索引 (i,j) 到 x（随机）。
计算y = f(x[i], x[i - 1], x[j], x[j + 1]).
根据y决定是否交换x[i]和x[j]的位置。

问题是函数f依赖于共享内存中的4个值，所有这些值都必须保证在交换之前保持不变。

有一分钟我认为这可能是 critical section 的主体，但我不明白如何使用单个锁定地址来锁定 4 个变量。我认为，主要问题是，当某些线程在 (i,j) 上工作时，其他线程不允许在任何对 (k,l) 上工作，其中 k 或 l 是任何{i, i-1, j, j+1}。

编辑

刚发帖，我就想到了一个主意……能不能做个级联锁？首先锁定 x[i]，如果成功则锁定 x[i-1]，等等所有 4 个值。只有最终加锁成功，才进行上述步骤。我会去试验并保持这个问题对其他建议开放。

Answer 1

在我看来，你想得太多了。如果所有内存事务都必须序列化以使操作成为线程安全的，那么最简单的解决方案是让每个块一个线程执行该操作。所以像

if (threadIdx.x == 0) // assume 1D block for simplicity
{
    y = f(x[i], x[i - 1], x[j], x[j + 1]);
    compare_and_swap(y, x[i], x[j];
}
__syncthreads();

将正常工作，因为正在操作的数组在共享内存中，因此保证每个块的单个线程正在执行操作，并且不存在先写后读的危险。在实践中，这种方法不应该比让一整块线程争用锁或大量序列化原子内存事务慢。

Answer 2

CUDA 对锁和临界区非常不友好 :) 众多原因之一是它以 32 宽 SIMD 模式运行。这可能会导致意外的死锁。考虑例如：

__shared__ int crit;
crit = 0;
__syncthreads();
int old;
do {
    old = atomicCas(&crit, 0, 1);
} while (old==1);
//critical section
crit = 0;

意图是线程在do-while循环中主动等待。一次只有一个线程存在循环，在临界区执行操作，然后将 crit 重置为 0。但是，在 CUDA 中，warp 调度程序将始终优先于循环中的 1 个线程的 31 个线程出口。因为 warp 在 SIMD 中运行，所以临界区中的线程永远不会执行，您会遇到意外的死锁。

出于这个原因，我强烈建议尽量避免使用临界区。

现在，我不知道你算法的细节。我假设您有一些 "master" for/while 循环，并且在每次迭代中您选择一个随机对进行可能的交换。

你说碰撞不常发生。如果是这样，您是否可以选择完全删除其中一个冲突对，而不是等待它成功？

如果这是您可以接受的，那么问题就出在检测碰撞上，而不是您之后采取的行动上。例如，要检测碰撞，您可以：

在每个线程提出一个候选对后，对索引对进行排序，然后检查邻居持有的值。
有一个与 x 和 atomicCas 大小相同的标志数组 f 4 次，类似于您的建议。如果f在共享内存中，应该不会很昂贵。

现在，当一个线程发现它有冲突时，它什么都不做。只是等待所有其他线程完成他们的工作，__syncthreads，然后进入主循环的下一个迭代 for/while 循环。

与您提出的解决方案的不同之处在于，如果锁定失败，您的线程只会放弃他的工作，而不是尝试等待。

CUDA：如何原子地执行这些指令（4个地址）

CUDA: How to atomically execute these instructions (4 addresses)

c++

cuda

atomic