CUDA中如何高效实现int3类型的AtomicAdd原子加法操作

和double类型自定义原子加的CAS自旋逻辑同源，但int3的实现需要结合硬件原子操作的位宽支持、业务对原子性的粒度要求选方案，不要盲目套while+CAS的模板。

首先明确两个核心前提：

• CUDA原生32位atomicAdd对int类型的支持从sm_1x架构就存在，硬件实现开销远低于软件CAS自旋循环
• sm_70（伏特架构，2017年及之后发布的消费级/计算卡）及以上硬件原生支持128位宽度的atomicCAS操作，要求操作地址16字节对齐

按业务原子性要求选实现方案 #

• 场景1：仅要求每个分量独立累加正确，不要求三个分量的更新是整体原子
  这是性能最高的实现，不需要任何CAS自旋，直接对三个分量分别调用原生32位atomicAdd即可：

  __device__ int3 atomicAdd_int3(int3* addr, int3 val) {
      int3 old;
      old.x = atomicAdd(&addr->x, val.x);
      old.y = atomicAdd(&addr->y, val.y);
      old.z = atomicAdd(&addr->z, val.z);
      return old;
  }
  

  注意：这个实现返回的旧值不是某一时刻int3的完整原子快照，执行过程中其他线程可能读到「x已更新、y/z未更新」的中间态，仅适合不需要整体一致性的场景，比如三个分量独立统计互不相关的计数。

• 场景2：要求三个分量的加法是整体原子（不存在部分更新的中间态，返回值是更新前的完整快照）
  优先选128位CAS自旋方案，性能远高于拆多次32位CAS、自旋锁等方案：

  1. 首先保证存储int3的内存地址是16字节对齐：堆内存分配用对齐分配接口，全局/共享内存变量加__align__(16)修饰即可，int3本身占12字节，剩余4字节留空不影响使用。
  2. 自旋逻辑和double原子加完全一致：循环中先读取当前128位值，计算加完增量后的新值，用CAS尝试写回，若写回过程中值被其他线程修改则重试。
     参考实现：

  // 注意：传入的addr必须为16字节对齐地址
  __device__ int3 atomicAdd_int3(int3* __restrict__ addr, int3 val) {
      uint4* addr_u4 = reinterpret_cast<uint4*>(addr);
      uint4 old = *addr_u4, assumed;
      int3 ret;
      do {
          assumed = old;
          uint4 new_u4 = make_uint4(
              assumed.x + val.x,
              assumed.y + val.y,
              assumed.z + val.z,
              assumed.w // 预留的4字节保持原值不变
          );
          old = atomicCAS(addr_u4, assumed, new_u4);
      } while (__vcmpeq4(old, assumed) == 0); // 用向量比较指令加速判断，比逐分量比较快
      ret = make_int3(old.x, old.y, old.z);
      return ret;
  }
  

  如果需要兼容sm_70以下的老架构（不支持128位原子CAS），没有特别高效的通用方案，建议通过块内/全局自旋锁保护int3的读写更新，但是高冲突场景下性能会下降明显，非必要不推荐。

避坑提示 #

不要为了实现整体原子性拆成三次32位CAS做自旋，单次循环就需要三次原子操作，访存冲突概率是128位CAS方案的3倍以上，性能极差。

内容的提问来源于stack exchange，提问作者Nyaruko