Cuda Mutex，为什么会死锁？

Question

我正在尝试实现基于原子的互斥量。

我成功了，但我有一个关于扭曲/死锁的问题。

此代码运行良好。

bool blocked = true;

while(blocked) {
    if(0 == atomicCAS(&mLock, 0, 1)) {
        index = mSize++;

        doCriticJob();

        atomicExch(&mLock, 0);
        blocked = false;
    }
}

但是这个没有...

while(true) {
    if(0 == atomicCAS(&mLock, 0, 1)) {
        index = mSize++;

        doCriticJob();

        atomicExch(&mLock, 0);
        break;
    }
}

我认为这是退出循环的位置。在第一个中，exit 发生在条件所在的地方，在第二个中它发生在 if 的末尾，所以线程等待其他 warp 完成循环，但其他线程也等待第一个线程......但我想我我错了，所以如果你能解释我:)。

谢谢！

Answer 1

这里还有其他关于互斥体的问题。你可能想看看其中的一些。例如，在 "cuda critical section" 上搜索。
假设一个会工作而另一个不会因为它似乎对您的测试用例工作是危险的。管理互斥体或关键部分，尤其是在同一个 warp 中的线程之间进行协商时 是出了名的困难和脆弱。一般的建议是避免它。正如其他地方所讨论的，如果您必须使用互斥体或关键部分，让线程块中的单个线程协商需要它的任何线程，然后使用线程块内同步机制控制线程块内的行为，例如 __syncthreads().
如果不查看编译器对各种执行路径进行排序的方式，就无法真正回答这个问题 (IMO)。因此，我们需要查看 SASS 代码（机器代码）。你可以使用cuda binary utilities to do this, and will probably want to refer to both the PTX reference as well as the SASS reference。这也意味着您需要 完整代码，而不仅仅是您提供的代码片段。

下面是我的分析代码：

$ cat t830.cu
#include <stdio.h>


__device__ int mLock = 0;

__device__ void doCriticJob(){

}

__global__ void kernel1(){

  int index = 0;
  int mSize = 1;
  while(true) {
    if(0 == atomicCAS(&mLock, 0, 1)) {
        index = mSize++;

        doCriticJob();

        atomicExch(&mLock, 0);
        break;
    }
  }
}

__global__ void kernel2(){

  int index = 0;
  int mSize = 1;
  bool blocked = true;

  while(blocked) {
    if(0 == atomicCAS(&mLock, 0, 1)) {
        index = mSize++;

        doCriticJob();

        atomicExch(&mLock, 0);
        blocked = false;
    }
  }
}
int main(){

 kernel2<<<4,128>>>();
 cudaDeviceSynchronize();
}

kernel1 是我对你的死锁代码的表示，kernel2 是我对你的 "working" 代码的表示。当我在 CUDA 7 下的 linux 和 cc2.0 设备（Quadro5000）上的运行上编译它时，如果我调用 kernel1 代码将死锁，如果我调用 kernel2（如图所示）它没有。

我用cuobjdump -sass转储机器码：

$ cuobjdump -sass ./t830

Fatbin elf code:
================
arch = sm_20
code version = [1,7]
producer = <unknown>
host = linux
compile_size = 64bit

        code for sm_20

Fatbin elf code:
================
arch = sm_20
code version = [1,7]
producer = cuda
host = linux
compile_size = 64bit

        code for sm_20
                Function : _Z7kernel1v
        .headerflags    @"EF_CUDA_SM20 EF_CUDA_PTX_SM(EF_CUDA_SM20)"
        /*0000*/         MOV R1, c[0x1][0x100];            /* 0x2800440400005de4 */
        /*0008*/         MOV32I R4, 0x1;                   /* 0x1800000004011de2 */
        /*0010*/         SSY 0x48;                         /* 0x60000000c0000007 */
        /*0018*/         MOV R2, c[0xe][0x0];              /* 0x2800780000009de4 */
        /*0020*/         MOV R3, c[0xe][0x4];              /* 0x280078001000dde4 */
        /*0028*/         ATOM.E.CAS R0, [R2], RZ, R4;      /* 0x54080000002fdd25 */
        /*0030*/         ISETP.NE.AND P0, PT, R0, RZ, PT;  /* 0x1a8e0000fc01dc23 */
        /*0038*/     @P0 BRA 0x18;                         /* 0x4003ffff600001e7 */
        /*0040*/         NOP.S;                            /* 0x4000000000001df4 */
        /*0048*/         ATOM.E.EXCH RZ, [R2], RZ;         /* 0x547ff800002fdd05 */
        /*0050*/         EXIT;                             /* 0x8000000000001de7 */
                ............................


                Function : _Z7kernel2v
        .headerflags    @"EF_CUDA_SM20 EF_CUDA_PTX_SM(EF_CUDA_SM20)"
        /*0000*/         MOV R1, c[0x1][0x100];            /* 0x2800440400005de4 */
        /*0008*/         MOV32I R0, 0x1;                   /* 0x1800000004001de2 */
        /*0010*/         MOV32I R3, 0x1;                   /* 0x180000000400dde2 */
        /*0018*/         MOV R4, c[0xe][0x0];              /* 0x2800780000011de4 */
        /*0020*/         MOV R5, c[0xe][0x4];              /* 0x2800780010015de4 */
        /*0028*/         ATOM.E.CAS R2, [R4], RZ, R3;      /* 0x54061000004fdd25 */
        /*0030*/         ISETP.NE.AND P1, PT, R2, RZ, PT;  /* 0x1a8e0000fc23dc23 */
        /*0038*/    @!P1 MOV R0, RZ;                       /* 0x28000000fc0025e4 */
        /*0040*/    @!P1 ATOM.E.EXCH RZ, [R4], RZ;         /* 0x547ff800004fe505 */
        /*0048*/         LOP.AND R2, R0, 0xff;             /* 0x6800c003fc009c03 */
        /*0050*/         I2I.S32.S16 R2, R2;               /* 0x1c00000008a09e84 */
        /*0058*/         ISETP.NE.AND P0, PT, R2, RZ, PT;  /* 0x1a8e0000fc21dc23 */
        /*0060*/     @P0 BRA 0x18;                         /* 0x4003fffec00001e7 */
        /*0068*/         EXIT;                             /* 0x8000000000001de7 */
                ............................



Fatbin ptx code:
================
arch = sm_20
code version = [4,2]
producer = cuda
host = linux
compile_size = 64bit
compressed
$

考虑到单个 warp，无论使用哪种代码，所有线程都必须获取锁（通过 atomicCAS）一次，以便代码成功完成。使用任何一种代码，在任何给定时间，warp 中只有一个线程可以获取锁，并且为了让 warp 中的其他线程（稍后）获取锁，该线程必须有机会释放它（通过 atomicExch）。

这些实现之间的主要区别在于编译器如何根据条件分支.

调度atomicExch指令
让我们考虑一下 "deadlock" 代码 (kernel1)。在这种情况下，ATOM.E.EXCH 指令不会出现，直到在一个（也是唯一的）条件分支 (@P0 BRA 0x18;) 指令之后。 CUDA 代码中的条件分支表示一个可能的 warp 发散点，并且在 warp 发散之后的执行在某种程度上是未指定的，取决于机器的具体情况。但是鉴于这种不确定性，获得锁的线程可能会等待其他线程完成它们的分支，在执行atomicExch指令之前，这意味着另一个线程将没有机会获得锁，我们有死锁。

如果我们将其与 "working" 代码进行比较，我们会发现一旦发出 ATOM.E.CAS 指令，它们之间就有 no 条件分支那个点和发出 ATOM.E.EXCH 指令的点，从而释放刚刚获得的锁。由于每个获得锁（通过ATOM.E.CAS）的线程都会在任何条件分支发生之前释放它（通过ATOM.E.EXCH），所以没有任何可能（鉴于此代码实现）见证这种死锁之前（kernel1）发生。

(@P0 是 predication 的一种形式，您可以在 PTX 参考文献 here 中阅读它以了解它如何导致条件判断分支。)

注意： 我认为这两种代码都很危险，而且可能存在缺陷。尽管当前的测试似乎没有发现 "working" 代码的问题，但我认为未来的 CUDA 编译器可能会选择以不同的方式安排事情，并破坏该代码。甚至有可能为不同的机器架构编译可能会在这里产生不同的代码。我认为像 this 这样的机制更健壮，它可以完全避免内部 warp 争用。然而，即使是这样的机制也可能导致线程块间死锁。必须在特定的编程和使用限制下使用任何互斥量。

Cuda Mutex，为什么会死锁？

Cuda Mutex, why deadlock?

cuda