CUDA 原子参数竞争条件

Question

如果存在如下结构的设备代码

Item* prev_entry = array[entry->prev];
prev_entry->next = entry->next;

并且将其重写为原子操作

atomicExch(&(array[entry->prev]->next),entry->next);

是array 的内存访问与 next 的访问一起以原子方式完成？可能有其他线程修改 entry->prev （因为它们可能是另一个 Item 的下一个值），并且如果数组访问是非原子地完成的，则 entry->prev code> 可能会在访问数组和对地址 next 执行原子操作之间发生变化，从而导致不正确的结果。

更笼统地概括这个问题，原子操作参数中的所有操作都是原子执行的吗？

Answer 1

数组的内存访问是否与下一个的访问一起以原子方式完成？

不，不是。如果你研究一下相应的SASS代码，你会发现这里的entry->next相关的读操作

atomicExch(&(array[entry->prev]->next), entry->next);

是一个普通的读操作，没有受到任何保护。该读取操作将原子“更新值”放入寄存器中。另一个寄存器保存要更新的地址。原子操作对这些寄存器起作用（如果相关，则将其结果返回到另一个寄存器中）。

这是一个例子：

$ cat t1983.cu
__global__ void k(int *al, int *d){

  atomicExch(al, d[threadIdx.x]);
}


$ nvcc -c t1983.cu
$ cuobjdump -sass ./t1983.o

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

        code for sm_52
                Function : _Z1kPiS_
        .headerflags    @"EF_CUDA_SM52 EF_CUDA_PTX_SM(EF_CUDA_SM52)"
                                                                           /* 0x001c7c00fe0007f6 */
        /*0008*/                   MOV R1, c[0x0][0x20] ;                  /* 0x4c98078000870001 */
        /*0010*/         {         MOV R2, c[0x0][0x140] ;                 /* 0x4c98078005070002 */
        /*0018*/                   S2R R4, SR_TID.X         }
                                                                           /* 0xf0c8000002170004 */
                                                                           /* 0x001fc800fec20ff1 */
        /*0028*/                   SHR.U32 R0, R4.reuse, 0x1e ;            /* 0x3828000001e70400 */
        /*0030*/                   ISCADD R4.CC, R4, c[0x0][0x148], 0x2 ;  /* 0x4c18810005270404 */
        /*0038*/                   IADD.X R5, R0, c[0x0][0x14c] ;          /* 0x4c10080005370005 */
                                                                           /* 0x041fc400fe8007b1 */
        /*0048*/                   LDG.E R4, [R4] ;                        /* 0xeed4200000070404 */
        /*0050*/                   MOV R3, c[0x0][0x144] ;                 /* 0x4c98078005170003 */
        /*0058*/                   ATOM.E.EXCH RZ, [R2], R4 ;              /* 0xed810000004702ff */
                                                                           /* 0x001ffc00ffe007ed */
        /*0068*/                   NOP ;                                   /* 0x50b0000000070f00 */
        /*0070*/                   EXIT ;                                  /* 0xe30000000007000f */
        /*0078*/                   BRA 0x78 ;                              /* 0xe2400fffff87000f */
                ..........

首先，我们注意到原子操作纯粹基于寄存器来工作：

   ATOM.E.EXCH RZ, [R2], R4 ; 

寄存器RZ是“目的地”，它是RZ（始终为零的寄存器，充当“丢弃”寄存器），因为我们不要求函数的返回值。寄存器对R2、R3包括要自动“更新”的位置的64位地址（在这种情况下，替换其值），并且替换值包含在R4中。向后看，我们看到 R4 被加载在这里：

     LDG.E R4, [R4] ; 

它与原子完全分开。这是“普通”负载。您可以向后推，发现包含加载地址的寄存器对 R4、R5 填充了常量内存中的内核参数，使用 SR_TID.X 寄存器进行偏移（对应于 threadIdx .x）这是有道理的。

同样，包含原子更新位置地址的 R2、R3 寄存器对直接从内核参数加载，没有偏移量，这也是有意义的。