0

我遇到了一个性能问题atomicAdd，float在int使用nv-nsight-cu-cli. 检查生成的 SASS 后，我发现共享内存的生成 SASS 是打开的atomicAdd，float并且int完全不相似。

在这里，我以最小的 cuda 代码展示了一个示例：

$ cat test.cu                                                                                                                                                                                                                                                   
__global__ void testAtomicInt() {
    __shared__ int SM_INT;
    SM_INT = 0;
    __syncthreads();
    atomicAdd(&(SM_INT), ((int)1));
}

__global__ void testAtomicFloat() {
    __shared__ float SM_FLOAT;
    SM_FLOAT = 0.0;
    __syncthreads();
    atomicAdd(&(SM_FLOAT), ((float)1.1));
}

$ nvcc -arch=sm_86 -c test.cu 
$ cuobjdump -sass test.o                                                                                                                                                                                                                                        
Fatbin elf code:
================
arch = sm_86
code version = [1,7]
producer = <unknown>
host = linux
compile_size = 64bit

    code for sm_86
        Function : _Z15testAtomicFloatv
    .headerflags    @"EF_CUDA_SM86 EF_CUDA_PTX_SM(EF_CUDA_SM86)"
        /*0000*/                   MOV R1, c[0x0][0x28] ;                  /* 0x00000a0000017a02 */
                                                                           /* 0x000fc40000000f00 */
        /*0010*/                   STS [RZ], RZ ;                          /* 0x000000ffff007388 */
                                                                           /* 0x000fe80000000800 */
        /*0020*/                   BAR.SYNC 0x0 ;                          /* 0x0000000000007b1d */
                                                                           /* 0x000fec0000000000 */
        /*0030*/                   LDS R2, [RZ] ;                          /* 0x00000000ff027984 */
                                                                           /* 0x000e240000000800 */
        /*0040*/                   FADD R3, R2, 1.1000000238418579102 ;    /* 0x3f8ccccd02037421 */
                                                                           /* 0x001fcc0000000000 */
        /*0050*/                   ATOMS.CAST.SPIN R3, [RZ], R2, R3 ;      /* 0x00000002ff03738d */
                                                                           /* 0x000e240001800003 */
        /*0060*/                   ISETP.EQ.U32.AND P0, PT, R3, 0x1, PT ;  /* 0x000000010300780c */
                                                                           /* 0x001fda0003f02070 */
        /*0070*/              @!P0 BRA 0x30 ;                              /* 0xffffffb000008947 */
                                                                           /* 0x000fea000383ffff */
        /*0080*/                   EXIT ;                                  /* 0x000000000000794d */
                                                                           /* 0x000fea0003800000 */
        /*0090*/                   BRA 0x90;                               /* 0xfffffff000007947 */
                                                                           /* 0x000fc0000383ffff */
        /*00a0*/                   NOP;                                    /* 0x0000000000007918 */
        ..........


        Function : _Z13testAtomicIntv
    .headerflags    @"EF_CUDA_SM86 EF_CUDA_PTX_SM(EF_CUDA_SM86)"
        /*0000*/                   MOV R1, c[0x0][0x28] ;         /* 0x00000a0000017a02 */
                                                                  /* 0x000fc40000000f00 */
        /*0010*/                   STS [RZ], RZ ;                 /* 0x000000ffff007388 */
                                                                  /* 0x000fe80000000800 */
        /*0020*/                   BAR.SYNC 0x0 ;                 /* 0x0000000000007b1d */
                                                                  /* 0x000fec0000000000 */
        /*0030*/                   ATOMS.POPC.INC.32 RZ, [URZ] ;  /* 0x00000000ffff7f8c */
                                                                  /* 0x000fe2000d00003f */
        /*0040*/                   EXIT ;                         /* 0x000000000000794d */
                                                                  /* 0x000fea0003800000 */
        /*0050*/                   BRA 0x50;                      /* 0xfffffff000007947 */
                                                                  /* 0x000fc0000383ffff */
        /*0060*/                   NOP;                           /* 0x0000000000007918 */
        ..........



Fatbin ptx code:
================
arch = sm_86
code version = [7,5]
producer = <unknown>
host = linux
compile_size = 64bit
compressed

从上面生成的 SASS 代码中，我们可以清楚地得知，共享内存atomicAdd在int生成单个轻量级ATOMS.POPC.INC.32 RZ, [URZ]，而在float生成一堆 SASS 时具有重量级ATOMS.CAST.SPIN R3, [RZ], R2, R3 。

CUDA Binary Utilities没有告诉我CASTor的含义SPIN。但是，我猜想这意味着共享内存地址上的独占自旋锁。（纠正我，如果我猜错了。）在我的真实代码中，没有一个 SASSatomicAdd有int热点。但是，这比ofATOMS.CAST.SPIN生成的其他 SASS 代码要热得多。atomicAddfloat

另外，我用编译器标志-arch=sm_86和. 在那些 CC 下，生成的 SPSS 代码非常相似。另一个事实是，毫不奇怪， 的生成的SPSS 与.-arch=sm_80-arch=sm_75atomicAddfloatatomicAdddoublefloat

这个观察比问题更让我困惑。我会从我的分析经验中提出一些简单的问题，并希望我们能进行愉快的讨论。

  • 具体是做什么ATOMS.CAST.SPIN的？我知道的唯一 SASS 文档是CUDA Binary Utilities。
  • 为什么atomicAddof应该float生成更多的 SASS 代码并做更多的工作int？我知道这是一个普遍的问题，很难回答。也许ATOMS.POPC.INC根本不适用于数据类型float或double？
  • 如果它更容易受到更多共享内存加载和存储冲突的影响，因此of比of 的停顿时间更长？前者显然有更多要执行的指令和不同的分支。我的项目中有以下代码片段，其中两个函数的函数调用次数相同。然而，of在开启时不会造成运行时瓶颈。atomicAddfloatatomicAddintatomicAddfloatint
atomicAdd(&(SM_INT), ((int)1));  // no hotspot
atomicAdd(&(SM_FLOAT), ((float)1.1)); // a hotspot
4

1 回答 1

2

我可能无法提供解决所有可能问题的答案。CUDA SASS 确实没有记录到解释这些事情的级别。

ATOMS.CAST.SPIN 到底是做什么的？我知道的唯一 SASS 文档是 CUDA Binary Utilities。

ATOMS.CAST.SPIN
^^^^^ ^^^     
   ||   |  
   ||   compare and swap
   |shared
   atomic

编程指南给出了如何使用原子 CAS（比较和交换）实现“任意”原子操作的指示。您应该首先熟悉原子 CAS 的工作原理。

关于“任意原子”示例，需要注意的是，它显然可用于为“本机”原子指令不支持的数据类型（例如原子添加）提供原子操作。另一件需要注意的是，它本质上是一个围绕原子 CAS 指令的循环，循环检查操作是否“成功”。如果它“不成功”，则循环继续。如果它“成功”，则循环退出。

这实际上是我们在您的float示例中的 SASS 代码中看到的内容：

/*0030*/  LDS R2, [RZ] ;  // get the current value in the location
FADD R3, R2, 1.1000000238418579102 ; // perform ordinary floating-point add
ATOMS.CAST.SPIN R3, [RZ], R2, R3 ;  //  attempt to atomically replace the result in the location
ISETP.EQ.U32.AND P0, PT, R3, 0x1, PT ; // check if replacement was successful
@!P0 BRA 0x30   // if not, loop and try again

这些基本上是编程指南中“任意原子”示例中概述的步骤。基于此，我得出以下结论：

  • 您编译的架构实际上没有您请求的类型的“本机”原子操作
  • 您请求的原子操作可以使用循环方法完成
  • 编译器工具链（通常ptxas，但也可能是 JIT 系统）作为一个便利功能，会自动为您实现此循环方法，而不是抛出编译错误

为什么 float 的 atomicAdd 会生成更多的 SASS 代码，并且比 int 做更多的工作？

显然，您正在为其编译的体系结构没有原子 add for 的“本机”实现float，因此编译器工具链选择为您实现这种循环方法。由于循环有效地涉及成功/失败的可能性，这将确定此循环是否继续，并且成功/失败取决于其他线程的行为（争用执行原子），循环方法可能比原生单线程做更多的“工作”指令会。

如果它更容易出现更多共享内存加载和存储冲突，因此 float 的 atomicAdd 比 int 的 atomicAdd 的停顿时间更长？

是的，我个人会得出结论，原生 atomic 方法效率更高，而 looping 方法可能效率较低，这可以在 profiler 中以多种方式表示，例如 warp stalls。

编辑：

  1. 在一种 GPU 架构中实现/可用的东西是可能的，但在另一种架构中则不然。这当然适用于原子，如果您阅读编程指南中先前链接的有关原子的部分，您可以看到这方面的示例。我不知道今天有任何架构比 cc8.0 或 cc8.6（安培）“更新”，但未来（或任何其他）GPU 的行为在这里可能会有所不同。

  2. 这种循环原子CAS方法不同于以前的方法（lock/update/unlock，它还涉及锁定协商的循环），在Kepler和以前的架构上使用的编译器工具链在不存在正式的SASS指令时在共享内存上提供原子这样做。

于 2022-01-30T14:55:10.193 回答