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 可能会在访问数组和执行原子之间改变对地址 next 的操作导致不正确的结果。

为了更笼统地提出问题，原子操作的参数中的所有操作都是原子执行的吗？

Answer 1

is the memory access of array done atomically along with the access of next?

不，不是。如果你研究对应的SASS代码，你会发现与entry->next关联的读操作在这里：

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

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

这是一个例子：

$ 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（always-zero 寄存器，充当“丢弃”寄存器）因为我们不要求函数的 return 值.寄存器对 R2、R3 包含要自动“更新”（在本例中，替换其值）的位置的 64 位地址，替换值包含在 R4 中。向后工作，我们看到 R4 已加载到此处：

     LDG.E R4, [R4] ;

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

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

CUDA 原子参数竞争条件

CUDA atomic argument race conditions

cuda

atomic

race-condition