[CUDA] atomic函数闭坑技巧

cuda atomic函数使用时的一些注意事项

1. 尽量减少atomic的使用频率

• 由于atomic会增加threads之间的同步性，所以在有选择性的atomic操作时，可以考虑用if(condition) atomic； 代替 atomic(addr, condition? a:b); 也就是尽量减少atomic的使用频率；

因为atomic 容易冲突，所以用if包起来要比用三元操作每个kernel都执行要快
// more effective than atomicAdd(sum + i, factor ? 1 : 0).
 if (factor) {
 atomicAdd(sum + i, 1);
 }

2. 小心atomic的不规则读数

• 在kernel进行中有atomic存在时， 有时候float类型的kernel函数执行效率要比half类型的kernel执行效率快
  • 当读取数据之间没有重叠时，half kernel比float kernel快；
  • 当不规则的读取数据进行累加时，则half比float慢；有可能会慢很多。
  • 此外，使用half2优化也不太会达到float的性能。

// 规则读取时， half kernel更快
template <typename T>
__global__ void TestKernel(T* a, T* b, int size, int count) {
  const int idx = blockIdx.x * blockDim.x + threadIdx.x;
  if (idx >= size) {
    return;
  }
  
  AtomicAdd(a[idx], count, &b[idx]);
}
// 不规则读取，多个kernel存在公用一个b数据时，half类型kernel耗时急剧增加；（idx%100会让多个kernel读取同一个数据）
// float kernel更快
template <typename T>
__global__ void TestKernel(T* a, T* b, int size, int count) {
  const int idx = blockIdx.x * blockDim.x + threadIdx.x;
  if (idx >= size) {
    return;
  }
  
  AtomicAdd(a[idx], count, &b[idx%100]);
}

• 规则读取时的nsight 评测
  
• 不规则读取时的nsight 评测， half 比float长很多，耗时增加很大。
  

3. 使用shared memory加速atomic

• 在atomic累加时，可以考虑用shared memory来保证快速的累加，并且需要注意每个block的shared memory是相互隔离的，需要考虑最终进行reduce规约。