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;
    }
}

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

谢谢 ！

Answer 1

1. 这里还有关于互斥锁的其他问题。你可能想看看其中的一些。例如，搜索“cuda 临界区”。

2. 假设一个会起作用而另一个不起作用，因为它似乎适用于您的测试用例是危险的。管理互斥体或临界区，尤其是当协商是在同一经线中的线程之间时，是非常困难和脆弱的。一般建议是避免它。正如其他地方所讨论的，如果您必须使用互斥锁或临界区，请在线程块中为任何需要它的线程协商一个线程，然后使用线程块内同步机制控制线程块内的行为，例如__syncthreads().

3. 如果不查看编译器对各种执行路径进行排序的方式，就无法真正回答这个问题 (IMO)。因此我们需要查看SASS代码（机器码）。您可以使用cuda 二进制实用程序来执行此操作，并且可能需要同时参考PTX 参考和SASS 参考。这也意味着您需要完整的代码，而不仅仅是您提供的代码片段。

这是我的分析代码：

$ 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表示，也是我对您的“工作”代码的表示。当我在 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
$

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

这些实现之间的主要区别在于编译器如何根据条件分支调度atomicExch指令。

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

如果我们然后将其与“工作”代码进行比较，我们会看到一旦ATOM.E.CAS发出指令，在该点和发出指令的点之间就没有条件分支ATOM.E.EXCH，从而释放刚刚获得的锁。由于每个获取锁 (via ATOM.E.CAS) 的线程都会ATOM.E.EXCH在任何条件分支发生之前释放它 (via )，因此之前 (with kernel1)看到的那种死锁不可能发生（鉴于此代码实现）。

（@P0是一种预测形式，您可以在此处的 PTX 参考中阅读有关它的信息，以了解它如何导致条件分支。）

注意： 我认为这两个代码都很危险，并且可能存在缺陷。尽管当前的测试似乎没有发现“工作”代码的问题，但我认为未来的 CUDA 编译器可能会选择以不同的方式安排事情，并破坏该代码。甚至有可能针对不同的机器架构进行编译可能会在此处生成不同的代码。我认为像一个机制，这是更强大的，它完全避免了内部竞争的扭曲。然而，即使这样的机制也可能导致线程块间死锁。必须在特定的编程和使用限制下使用任何互斥锁。