文章目录
- 1. 尽量减少atomic的使用频率
- 2. 小心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长很多,耗时增加很大。