原子操作、L2 cache、sysmem

我在测试
lts__t_sectors_srcnode_gpc_aperture_device_op_atom_dot_alu_lookup_hit lts__t_sectors_srcnode_gpc_aperture_device_op_atom_dot_alu_lookup_miss
lts__t_sectors_srcnode_gpc_aperture_sysmem_op_atom_dot_alu_lookup_hit lts__t_sectors_srcnode_gpc_aperture_sysmem_op_atom_dot_alu_lookup_miss
这四个metrics，都和原子操作有关，前两个和global memory有关，后两个和sysmem原子操作有关，
使用的kernel相同，并且只启动一个warp：
__global__ void LtsTSectorsSrcnodeGpcApertureDeviceOpAtomDotAluLookupHitMissKernel(unsigned int *input) {
    atomicAdd_system(&input[0], threadIdx.x);
}
其中当input是cudaMalloc分配的空间时，
lts__t_sectors_srcnode_gpc_aperture_device_op_atom_dot_alu_lookup_hit 为 31 lts__t_sectors_srcnode_gpc_aperture_device_op_atom_dot_alu_lookup_miss 为 1
这个比较好理解 ，因为一个warp执行相同位置处的原子操作会串行执行并且顺序未知，但是第一次执行完了以后数据会被load到L2 cache中，然后后面的31次操作都会hit cache；
但是当使用cudaHostAlloc分配空间测试后面的两个metrics时，发现情况有所不同：
lts__t_sectors_srcnode_gpc_aperture_sysmem_op_atom_dot_alu_lookup_hit 等于0
lts__t_sectors_srcnode_gpc_aperture_sysmem_op_atom_dot_alu_lookup_miss 等于32
 
然后我google中发现：
Do not try to use atomics on mapped pinned host memory, either for the host
(locked compare-exchange) or the device (atomicAdd()). On the CPU side,
the facilities to enforce mutual exclusion for locked operations are not visible
to peripherals on the PCI Express bus. Conversely, on the GPU side, atomic
operations only work on local device memory locations because they are
implemented using the GPU’s local memory controller.
但是Pascal结构以后，新增了了atomicAdd_system原子操作，所以我换成了这个函数，但是测试的结果依然没变还是全部都是Miss。

此外我还google到了这个描述：
Yes, this works atomically from a single GPU. So if no other CPU or GPU is accessing the memory it will be atomic. Atomics are implemented in the L2 cache and the CROP (on various GPUs), and both can handle system memory accesses.
链接：https://stackoverflow.com/questions/23193151/atomic-operations-in-cuda-kernels-on-mapped-pinned-host-memory-to-do-or-not-to

所以我有一下两个问题：
1. system memory的原子操作是在L2 cache中处理的嘛？为什么会全是miss？为什么使用atomicAdd_system也是全部miss？
2.shared memory的原子操作是在L1中处理的嘛？

我在测试
lts__t_sectors_srcnode_gpc_aperture_device_op_atom_dot_alu_lookup_hit lts__t_sectors_srcnode_gpc_aperture_device_op_atom_dot_alu_lookup_miss
lts__t_sectors_srcnode_gpc_aperture_sysmem_op_atom_dot_alu_lookup_hit lts__t_sectors_srcnode_gpc_aperture_sysmem_op_atom_dot_alu_lookup_miss
这四个metrics，都和原子操作有关，前两个和global memory有关，后两个和sysmem原子操作有关，
使用的kernel相同，并且只启动一个warp：
__global__ void LtsTSectorsSrcnodeGpcApertureDeviceOpAtomDotAluLookupHitMissKernel(unsigned int *input) {
    atomicAdd_system(&input[0], threadIdx.x);
}
其中当input是cudaMalloc分配的空间时，
lts__t_sectors_srcnode_gpc_aperture_device_op_atom_dot_alu_lookup_hit 为 31 lts__t_sectors_srcnode_gpc_aperture_device_op_atom_dot_alu_lookup_miss 为 1
这个比较好理解 ，因为一个warp执行相同位置处的原子操作会串行执行并且顺序未知，但是第一次执行完了以后数据会被load到L2 cache中，然后后面的31次操作都会hit cache；
但是当使用cudaHostAlloc分配空间测试后面的两个metrics时，发现情况有所不同：
lts__t_sectors_srcnode_gpc_aperture_sysmem_op_atom_dot_alu_lookup_hit 等于0
lts__t_sectors_srcnode_gpc_aperture_sysmem_op_atom_dot_alu_lookup_miss 等于32
 
然后我google中发现：
Do not try to use atomics on mapped pinned host memory, either for the host
(locked compare-exchange) or the device (atomicAdd()). On the CPU side,
the facilities to enforce mutual exclusion for locked operations are not visible
to peripherals on the PCI Express bus. Conversely, on the GPU side, atomic
operations only work on local device memory locations because they are
implemented using the GPU’s local memory controller.
但是Pascal结构以后，新增了了atomicAdd_system原子操作，所以我换成了这个函数，但是测试的结果依然没变还是全部都是Miss。

此外我还google到了这个描述：
Yes, this works atomically from a single GPU. So if no other CPU or GPU is accessing the memory it will be atomic. Atomics are implemented in the L2 cache and the CROP (on various GPUs), and both can handle system memory accesses.
链接：https://stackoverflow.com/questions/23193151/atomic-operations-in-cuda-kernels-on-mapped-pinned-host-memory-to-do-or-not-to

所以我有一下两个问题：
1. system memory的原子操作是在L2 cache中处理的嘛？为什么会全是miss？为什么使用atomicAdd_system也是全部miss？
2.shared memory的原子操作是在L1中处理的嘛？

（1）我从来没有成功的使用过atomicAdd_system(), 虽然手册上曾经说过Pascal（6.0+）和更高的设备即可支持，但是我在6.1， 7.5和8.6，配合多种主板和CPU上从来都没有成功过。所以你不要期望你能有一段缓冲区，能同时让CPU和GPU，细粒度的通过CPU上的例如C里面的_Atomic, 和GPU上的*_system结尾的原子操作函数的进行整个系统内部的协作。（如果你能成功，请告诉我）。

（2）考虑到1点从来都不能用，所以只剩下普通的系统内存上的原子操作（即普通的atomic*()）了，这种的确是在L2上完成的。

以上回答你的1-1问题。

关于你的1-2问题：为何使用普通的atomicAdd(), 对于内存映射成的global memory，全部报告miss，这个我不知道答案。

同样关于你的1-2问题：请使用atomicAdd(input, 1); 测试，然后告诉我结果变化（这个是一个已知的现象）。


关于你的2：shared memory上的原子操作分成多种情况，

较低计算能力(<5.0)是模拟的：即用SP来完全实际的计算（例如min，max，add，逻辑等等），shared memory只提供存储单元，然后SP通过锁定读取-计算改变-尝试锁定写入--失败重试的方式来完成的，这个流程和一些ARM的处理器类似，细节在《CUDA Handbook》这本书里（2013年的大约）有详细描述。

较高计算能力(>=5.0)是shared memory单元来完成的，在Maxwell出来的那一年的GTC中有资料，NV叫它“远程原子操作”，即将相应的shared存储单元上的原子操作指令，直接提交给shared memory本身来完成计算，此时shared memory本身在SP之外，具有简单的ALU能力。这一点和GCN+的A卡类似，local data share (LDS), 具有简单的运算能力。

所以综合上面两点，shared上的原子操作不通过L1 cache完成。

（1）我从来没有成功的使用过atomicAdd_system(), 虽然手册上曾经说过Pascal（6.0+）和更高的设备即可支持，但是我在6.1， 7.5和8.6，配合多种主板和CPU上从来都没有成功过。所以你不要期望你能有一段缓冲区，能同时让CPU和GPU，细粒度的通过CPU上的例如C里面的_Atomic, 和GPU上的*_system结尾的原子操作函数的进行整个系统内部的协作。（如果你能成功，请告诉我）。

（2）考虑到1点从来都不能用，所以只剩下普通的系统内存上的原子操作（即普通的atomic*()）了，这种的确是在L2上完成的。

以上回答你的1-1问题。

关于你的1-2问题：为何使用普通的atomicAdd(), 对于内存映射成的global memory，全部报告miss，这个我不知道答案。

同样关于你的1-2问题：请使用atomicAdd(input, 1); 测试，然后告诉我结果变化（这个是一个已知的现象）。


关于你的2：shared memory上的原子操作分成多种情况，

较低计算能力(<5.0)是模拟的：即用SP来完全实际的计算（例如min，max，add，逻辑等等），shared memory只提供存储单元，然后SP通过锁定读取-计算改变-尝试锁定写入--失败重试的方式来完成的，这个流程和一些ARM的处理器类似，细节在《CUDA Handbook》这本书里（2013年的大约）有详细描述。

较高计算能力(>=5.0)是shared memory单元来完成的，在Maxwell出来的那一年的GTC中有资料，NV叫它“远程原子操作”，即将相应的shared存储单元上的原子操作指令，直接提交给shared memory本身来完成计算，此时shared memory本身在SP之外，具有简单的ALU能力。这一点和GCN+的A卡类似，local data share (LDS), 具有简单的运算能力。

所以综合上面两点，shared上的原子操作不通过L1 cache完成。

对于atomicAdd_system()我没使用过，而且现在还没遇到需要时使用它的情况，如果之后遇到我会进行尝试使用和测试；

对于使用cudaHostAlloc（，，cudaHostAllocDefault）分配的空间做atomicAdd(input, 1); ，我测试了一下结果，和预期一致，始终等于开辟的线程数，目前看来应该是可以保证结果的正确性。
然后我又进行了多次测试：
cudaHostAlloc（，，cudaHostAllocDefault）
 thread_num = 32;grid_num = 1;
原子操作的结果是 32，metrics的结果如下：
    Section: Command line profiler metrics
    ---------------------------------------------------------------------- --------------- ------------------------------
    lts__t_sectors_srcnode_gpc_aperture_device_op_atom_dot_alu_lookup_hit.          sector                              0
    sum                                                                                                                 
    lts__t_sectors_srcnode_gpc_aperture_device_op_atom_dot_alu_lookup_miss          sector                              0
    .sum                                                                                                                 
    lts__t_sectors_srcnode_gpc_aperture_sysmem_op_atom_dot_alu_lookup_hit.          sector                              0
    sum                                                                                                                 
    lts__t_sectors_srcnode_gpc_aperture_sysmem_op_atom_dot_alu_lookup_miss          sector                             32
    .sum                                                                                                                 
    ---------------------------------------------------------------------- --------------- ------------------------------
 thread_num = 32;grid_num = 10;
原子操作的结果是 320，metrics的结果如下：
  LtsTSectorsSrcnodeGpcApertureDeviceOpAtomDotAluLookupHitMissKernel(unsigned int *), 2022-Jan-08 10:33:18, Context 1, Stream 7
    Section: Command line profiler metrics
    ---------------------------------------------------------------------- --------------- ------------------------------
    lts__t_sectors_srcnode_gpc_aperture_device_op_atom_dot_alu_lookup_hit.          sector                              0
    sum                                                                                                                 
    lts__t_sectors_srcnode_gpc_aperture_device_op_atom_dot_alu_lookup_miss          sector                              0
    .sum                                                                                                                 
    lts__t_sectors_srcnode_gpc_aperture_sysmem_op_atom_dot_alu_lookup_hit.          sector                              0
    sum                                                                                                                 
    lts__t_sectors_srcnode_gpc_aperture_sysmem_op_atom_dot_alu_lookup_miss          sector                            320
    .sum                                                                                                                 
    ---------------------------------------------------------------------- --------------- ------------------------------

然后我改为cudaHostAllocMapped、cudaHostAllocPortable、cudaHostAllocWriteCombined，与cudaHostAllocDefault结果都一致，

对于atomicAdd_system()我没使用过，而且现在还没遇到需要时使用它的情况，如果之后遇到我会进行尝试使用和测试；

对于使用cudaHostAlloc（，，cudaHostAllocDefault）分配的空间做atomicAdd(input, 1); ，我测试了一下结果，和预期一致，始终等于开辟的线程数，目前看来应该是可以保证结果的正确性。
然后我又进行了多次测试：
cudaHostAlloc（，，cudaHostAllocDefault）
 thread_num = 32;grid_num = 1;
原子操作的结果是 32，metrics的结果如下：
    Section: Command line profiler metrics
    ---------------------------------------------------------------------- --------------- ------------------------------
    lts__t_sectors_srcnode_gpc_aperture_device_op_atom_dot_alu_lookup_hit.          sector                              0
    sum                                                                                                                 
    lts__t_sectors_srcnode_gpc_aperture_device_op_atom_dot_alu_lookup_miss          sector                              0
    .sum                                                                                                                 
    lts__t_sectors_srcnode_gpc_aperture_sysmem_op_atom_dot_alu_lookup_hit.          sector                              0
    sum                                                                                                                 
    lts__t_sectors_srcnode_gpc_aperture_sysmem_op_atom_dot_alu_lookup_miss          sector                             32
    .sum                                                                                                                 
    ---------------------------------------------------------------------- --------------- ------------------------------
 thread_num = 32;grid_num = 10;
原子操作的结果是 320，metrics的结果如下：
  LtsTSectorsSrcnodeGpcApertureDeviceOpAtomDotAluLookupHitMissKernel(unsigned int *), 2022-Jan-08 10:33:18, Context 1, Stream 7
    Section: Command line profiler metrics
    ---------------------------------------------------------------------- --------------- ------------------------------
    lts__t_sectors_srcnode_gpc_aperture_device_op_atom_dot_alu_lookup_hit.          sector                              0
    sum                                                                                                                 
    lts__t_sectors_srcnode_gpc_aperture_device_op_atom_dot_alu_lookup_miss          sector                              0
    .sum                                                                                                                 
    lts__t_sectors_srcnode_gpc_aperture_sysmem_op_atom_dot_alu_lookup_hit.          sector                              0
    sum                                                                                                                 
    lts__t_sectors_srcnode_gpc_aperture_sysmem_op_atom_dot_alu_lookup_miss          sector                            320
    .sum                                                                                                                 
    ---------------------------------------------------------------------- --------------- ------------------------------

然后我改为cudaHostAllocMapped、cudaHostAllocPortable、cudaHostAllocWriteCombined，与cudaHostAllocDefault结果都一致，

那样的话，就没法解释了，因为atomicAdd(&addr, 1);是一个特定的行为，会被自动识别并且编译为如下步骤（NV所称的，自动的为了减少原子操作的竞争所进行的优化），

（1）在atomicAdd所在的行进行warp投票(vote.all), 选出当前的warp中的所有存活线程,
（2）统计warp中所有激活的线程的数量（值：1-32），记录为V
（3）判断当前线程是否为warp中第一个激活的线程，如果是，执行1次old = atomicAdd(..., V);
------到此你的测试流程结束-----
（4）如果atomicAdd的返回值还需要继续使用（即这是一个ATOM，而不是RED），将激活的线程的返回值传递给每个人，每个人进行warp内SP计算之前的激活线程数量，获取到正确分布的（而且是固定顺序分布的）old + 1... old + 32

也就是这种情况下，32个线程实际上只会执行1次L2上的原子操作。并没有32次。这个特性是从CUDA 8.0+引入的软件实现的特性（和硬件无关）。

所以我很奇怪你的结果，这将是无法解释的。所以，也许nsight报告的相关这个值，看看就好，不反应实际情况了。（你可以用nvcc -arch sm_86 your.cu -cubin -o your.cubin, cubojdump --dump-sass your.cubin, 反编译，来验证这个编译器的“模拟一次原子操作”的行为）。

那样的话，就没法解释了，因为atomicAdd(&addr, 1);是一个特定的行为，会被自动识别并且编译为如下步骤（NV所称的，自动的为了减少原子操作的竞争所进行的优化），

（1）在atomicAdd所在的行进行warp投票(vote.all), 选出当前的warp中的所有存活线程,
（2）统计warp中所有激活的线程的数量（值：1-32），记录为V
（3）判断当前线程是否为warp中第一个激活的线程，如果是，执行1次old = atomicAdd(..., V);
------到此你的测试流程结束-----
（4）如果atomicAdd的返回值还需要继续使用（即这是一个ATOM，而不是RED），将激活的线程的返回值传递给每个人，每个人进行warp内SP计算之前的激活线程数量，获取到正确分布的（而且是固定顺序分布的）old + 1... old + 32

也就是这种情况下，32个线程实际上只会执行1次L2上的原子操作。并没有32次。这个特性是从CUDA 8.0+引入的软件实现的特性（和硬件无关）。

所以我很奇怪你的结果，这将是无法解释的。所以，也许nsight报告的相关这个值，看看就好，不反应实际情况了。（你可以用nvcc -arch sm_86 your.cu -cubin -o your.cubin, cubojdump --dump-sass your.cubin, 反编译，来验证这个编译器的“模拟一次原子操作”的行为）。

我按照上述步骤进行了操作，得到的结果如下：
   .headerflags    @"EF_CUDA_SM86 EF_CUDA_PTX_SM(EF_CUDA_SM86)"
        /*0000*/                   IMAD.MOV.U32 R1, RZ, RZ, c[0x0][0x28] ;   /* 0x00000a00ff017624 */
                                                                             /* 0x000fc400078e00ff */
        /*0010*/                   S2R R0, SR_LANEID ;                       /* 0x0000000000007919 */
                                                                             /* 0x000e220000000000 */
        /*0020*/                   VOTEU.ALL UR4, UPT, PT ;                  /* 0x0000000000047886 */
                                                                             /* 0x000fe200038e0000 */
        /*0030*/                   MOV R2, c[0x0][0x160] ;                   /* 0x0000580000027a02 */
                                                                             /* 0x000fe20000000f00 */
        /*0040*/                   UFLO.U32 UR5, UR4 ;                       /* 0x00000004000572bd */
                                                                             /* 0x000fe200080e0000 */
        /*0050*/                   POPC R5, UR4 ;                            /* 0x0000000400057d09 */
                                                                             /* 0x000e620008000000 */
        /*0060*/                   IMAD.MOV.U32 R3, RZ, RZ, c[0x0][0x164] ;  /* 0x00005900ff037624 */
                                                                             /* 0x000fc800078e00ff */
        /*0070*/                   ISETP.EQ.U32.AND P0, PT, R0, UR5, PT ;    /* 0x0000000500007c0c */
                                                                             /* 0x001fe2000bf02070 */
        /*0080*/                   ULDC.64 UR4, c[0x0][0x118] ;              /* 0x0000460000047ab9 */
                                                                             /* 0x000fd80000000a00 */
        /*0090*/               @P0 RED.E.ADD.STRONG.GPU [R2.64], R5 ;        /* 0x000000050200098e */
                                                                             /* 0x002fe2000c10e184 */
        /*00a0*/                   EXIT ;                                    /* 0x000000000000794d */
                                                                             /* 0x000fea0003800000 */
        /*00b0*/                   BRA 0xb0;                                 /* 0xfffffff000007947 */
                                                                             /* 0x000fc0000383ffff */
        /*00c0*/                   NOP;                                      /* 0x0000000000007918 */
                                                                             /* 0x000fc00000000000 */
        /*00d0*/                   NOP;                                      /* 0x0000000000007918 */
                                                                             /* 0x000fc00000000000 */
        /*00e0*/                   NOP;                                      /* 0x0000000000007918 */
                                                                             /* 0x000fc00000000000 */
        /*00f0*/                   NOP;                                      /* 0x0000000000007918 */
                                                                             /* 0x000fc00000000000 */
        /*0100*/                   NOP;                                      /* 0x0000000000007918 */
                                                                             /* 0x000fc00000000000 */
        /*0110*/                   NOP;                                      /* 0x0000000000007918 */
                                                                             /* 0x000fc00000000000 */
        /*0120*/                   NOP;                                      /* 0x0000000000007918 */
                                                                             /* 0x000fc00000000000 */
        /*0130*/                   NOP;                                      /* 0x0000000000007918 */
                                                                             /* 0x000fc00000000000 */
        /*0140*/                   NOP;                                      /* 0x0000000000007918 */
                                                                             /* 0x000fc00000000000 */
        /*0150*/                   NOP;                                      /* 0x0000000000007918 */
                                                                             /* 0x000fc00000000000 */
        /*0160*/                   NOP;                                      /* 0x0000000000007918 */
                                                                             /* 0x000fc00000000000 */
        /*0170*/                   NOP;                                      /* 0x0000000000007918 */
                                                                             /* 0x000fc00000000000 */
虽然我看不懂这个汇编代码，不过似乎的确是有 VOTEU.ALL这种指令，我修改为atomicAdd(&addr, 10)似乎还是同样的流程，什么情况下会编译成这样的汇编？一个warp内的线程对同样位置的显存数据做常量原子加、减？这块内容有官方文档嘛？想去学习下

我按照上述步骤进行了操作，得到的结果如下：
   .headerflags    @"EF_CUDA_SM86 EF_CUDA_PTX_SM(EF_CUDA_SM86)"
        /*0000*/                   IMAD.MOV.U32 R1, RZ, RZ, c[0x0][0x28] ;   /* 0x00000a00ff017624 */
                                                                             /* 0x000fc400078e00ff */
        /*0010*/                   S2R R0, SR_LANEID ;                       /* 0x0000000000007919 */
                                                                             /* 0x000e220000000000 */
        /*0020*/                   VOTEU.ALL UR4, UPT, PT ;                  /* 0x0000000000047886 */
                                                                             /* 0x000fe200038e0000 */
        /*0030*/                   MOV R2, c[0x0][0x160] ;                   /* 0x0000580000027a02 */
                                                                             /* 0x000fe20000000f00 */
        /*0040*/                   UFLO.U32 UR5, UR4 ;                       /* 0x00000004000572bd */
                                                                             /* 0x000fe200080e0000 */
        /*0050*/                   POPC R5, UR4 ;                            /* 0x0000000400057d09 */
                                                                             /* 0x000e620008000000 */
        /*0060*/                   IMAD.MOV.U32 R3, RZ, RZ, c[0x0][0x164] ;  /* 0x00005900ff037624 */
                                                                             /* 0x000fc800078e00ff */
        /*0070*/                   ISETP.EQ.U32.AND P0, PT, R0, UR5, PT ;    /* 0x0000000500007c0c */
                                                                             /* 0x001fe2000bf02070 */
        /*0080*/                   ULDC.64 UR4, c[0x0][0x118] ;              /* 0x0000460000047ab9 */
                                                                             /* 0x000fd80000000a00 */
        /*0090*/               @P0 RED.E.ADD.STRONG.GPU [R2.64], R5 ;        /* 0x000000050200098e */
                                                                             /* 0x002fe2000c10e184 */
        /*00a0*/                   EXIT ;                                    /* 0x000000000000794d */
                                                                             /* 0x000fea0003800000 */
        /*00b0*/                   BRA 0xb0;                                 /* 0xfffffff000007947 */
                                                                             /* 0x000fc0000383ffff */
        /*00c0*/                   NOP;                                      /* 0x0000000000007918 */
                                                                             /* 0x000fc00000000000 */
        /*00d0*/                   NOP;                                      /* 0x0000000000007918 */
                                                                             /* 0x000fc00000000000 */
        /*00e0*/                   NOP;                                      /* 0x0000000000007918 */
                                                                             /* 0x000fc00000000000 */
        /*00f0*/                   NOP;                                      /* 0x0000000000007918 */
                                                                             /* 0x000fc00000000000 */
        /*0100*/                   NOP;                                      /* 0x0000000000007918 */
                                                                             /* 0x000fc00000000000 */
        /*0110*/                   NOP;                                      /* 0x0000000000007918 */
                                                                             /* 0x000fc00000000000 */
        /*0120*/                   NOP;                                      /* 0x0000000000007918 */
                                                                             /* 0x000fc00000000000 */
        /*0130*/                   NOP;                                      /* 0x0000000000007918 */
                                                                             /* 0x000fc00000000000 */
        /*0140*/                   NOP;                                      /* 0x0000000000007918 */
                                                                             /* 0x000fc00000000000 */
        /*0150*/                   NOP;                                      /* 0x0000000000007918 */
                                                                             /* 0x000fc00000000000 */
        /*0160*/                   NOP;                                      /* 0x0000000000007918 */
                                                                             /* 0x000fc00000000000 */
        /*0170*/                   NOP;                                      /* 0x0000000000007918 */
                                                                             /* 0x000fc00000000000 */
虽然我看不懂这个汇编代码，不过似乎的确是有 VOTEU.ALL这种指令，我修改为atomicAdd(&addr, 10)似乎还是同样的流程，什么情况下会编译成这样的汇编？一个warp内的线程对同样位置的显存数据做常量原子加、减？这块内容有官方文档嘛？想去学习下

这个是在CUDA 8.0出来的那年(2016?), NV作为一个特性提升（针对密集的在global memory上执行原子操作的，且冲突项较多的时候）。

已知每个线程使用不同的地址，或者使用不同的值去进行，则不会触发这种优化。

这个是在CUDA 8.0出来的那年(2016?), NV作为一个特性提升（针对密集的在global memory上执行原子操作的，且冲突项较多的时候）。

已知每个线程使用不同的地址，或者使用不同的值去进行，则不会触发这种优化。

以及，汇编无官方资料。

以及，这里涉及的vote, 和popc操作都是常见的操作（例如常见于prefix sum之类的运算）。CUDA C里有部分导出，例如__ballot_sync(), 例如__popc(), 有的时候很好用。例如使用相同地址和相同值，会自动触发一次走这种变通方式的减少竞争冲突的原子操作（你上文的帖子）。

以及，汇编无官方资料。

以及，这里涉及的vote, 和popc操作都是常见的操作（例如常见于prefix sum之类的运算）。CUDA C里有部分导出，例如__ballot_sync(), 例如__popc(), 有的时候很好用。例如使用相同地址和相同值，会自动触发一次走这种变通方式的减少竞争冲突的原子操作（你上文的帖子）。

嗯嗯好的，十分感谢

那样的话，就没法解释了，因为atomicAdd(&addr, 1);是一个特定的行为，会被自动识别并且编译为如下步骤（NV所称的，自动的为了减少原子操作的竞争所进行的优化），

（1）在atomicAdd所在的行进行warp投票(vote.all), 选出当前的warp中的所有存活线程,
（2）统计warp中所有激活的线程的数量（值：1-32），记录为V
（3）判断当前线程是否为warp中第一个激活的线程，如果是，执行1次old = atomicAdd(..., V);
------到此你的测试流程结束-----
（4）如果atomicAdd的返回值还需要继续使用（即这是一个ATOM，而不是RED），将激活的线程的返回值传递给每个人，每个人进行warp内SP计算之前的激活线程数量，获取到正确分布的（而且是固定顺序分布的）old + 1... old + 32

也就是这种情况下，32个线程实际上只会执行1次L2上的原子操作。并没有32次。这个特性是从CUDA 8.0+引入的软件实现的特性（和硬件无关）。

所以我很奇怪你的结果，这将是无法解释的。所以，也许nsight报告的相关这个值，看看就好，不反应实际情况了。（你可以用nvcc -arch sm_86 your.cu -cubin -o your.cubin, cubojdump --dump-sass your.cubin, 反编译，来验证这个编译器的“模拟一次原子操作”的行为）。

我发现问题在哪里了，我编译的时候加了-G选项，导致结果和你的解释不一样，当我去掉了-G选项时，结果就一致了， 另外我想问下ATOM和RED，为什么不使用返回值就会使用RED，我查了下官方文档：

ATOMG   Atomic Operation on Global Memory
RED   Reduction Operation on Generic Memory

我发现问题在哪里了，我编译的时候加了-G选项，导致结果和你的解释不一样，当我去掉了-G选项时，结果就一致了， 另外我想问下ATOM和RED，为什么不使用返回值就会使用RED，我查了下官方文档：

RED 和ATOM指令的区别经过查阅文档，我已经大概清楚，RED是Reduction+atomic，而ATOM仅仅是Atomic，RED运行耗时更短

RED 和ATOM指令的区别经过查阅文档，我已经大概清楚，RED是Reduction+atomic，而ATOM仅仅是Atomic，RED运行耗时更短

都是原子操作。前者不需要返回值，从而可能减少了一些执行上的成本，例如不需要设定对指令的写入结果值（例如旧值）的寄存器的结果写入完成的跟踪之类的（scoreboarding).