在cuda示例中,例如在这里,使用了
__match_all_sync
__match_any_sync
。
下面是一个示例,其中 warp 被分成多个(一个或多个(组,每个组跟踪自己的原子计数器。
// increment the value at ptr by 1 and return the old value
__device__ int atomicAggInc(int* ptr) {
int pred;
//const auto mask = __match_all_sync(__activemask(), ptr, &pred); //error, should be any_sync, not all_sync
const auto mask = __match_any_sync(__activemask(), ptr, &pred);
const auto leader = __ffs(mask) - 1; // select a leader
int res;
const auto lane_id = ThreadId() % warpSize;
if (lane_id == leader) { // leader does the update
res = atomicAdd(ptr, __popc(mask));
}
res = __shfl_sync(mask, res, leader); // get leader’s old value
return res + __popc(mask & ((1 << lane_id) - 1)); //compute old value
}
这里的__match_any_sync
将 warp 中的线程拆分为具有相同ptr
值的组,以便每个组都可以原子地更新自己的 ptr,而不会妨碍其他线程。
我知道 nvcc 编译器(自 cuda 9 以来(会自动在引擎盖下进行这种优化,但这只是关于__match_any_sync
的机制
有没有办法做这个预计算能力7?
编辑:博客文章现在已经修改以反映__match_any_sync()
而不是__match_all_sync()
,所以下面的任何评论都应该被忽略。 下面的答案经过编辑以反映这一点。
根据您的陈述:
这只是关于
__match_any_sync
的机制
我们将专注于替换__match_any_sync
本身,而不是任何其他形式的重写atomicAggInc
函数。 因此,我们必须提供一个掩码,该掩码的值与 cc7.0 或更高版本架构上的__match_any_sync()
返回的值相同。
我相信这将需要一个循环,它广播ptr
值,在最坏的情况下,对 warp 中的每个线程进行一次迭代(因为每个线程可能具有唯一的ptr
值(并测试哪些线程具有相同的值。 有多种方法可以针对此函数"优化"此循环,以便根据每个线程中的实际ptr
值将行程计数从 32 减少到较小的值,但在我看来,这种优化引入了相当大的复杂性,这使得最坏情况的处理时间更长(这是典型的提前退出优化(。 因此,我将演示一个相当简单的方法,而无需进行此优化。
另一个考虑因素是在翘曲不收敛的情况下该怎么办? 为此,我们可以采用__activemask()
来识别这种情况。
这是一个工作示例:
$ cat t1646.cu
#include <iostream>
#include <stdio.h>
// increment the value at ptr by 1 and return the old value
__device__ int atomicAggInc(int* ptr) {
int mask;
#if __CUDA_ARCH__ >= 700
mask = __match_any_sync(__activemask(), (unsigned long long)ptr);
#else
unsigned tmask = __activemask();
for (int i = 0; i < warpSize; i++){
#ifdef USE_OPT
if ((1U<<i) & tmask){
#endif
unsigned long long tptr = __shfl_sync(tmask, (unsigned long long)ptr, i);
unsigned my_mask = __ballot_sync(tmask, (tptr == (unsigned long long)ptr));
if (i == (threadIdx.x & (warpSize-1))) mask = my_mask;}
#ifdef USE_OPT
}
#endif
#endif
int leader = __ffs(mask) - 1; // select a leader
int res;
unsigned lane_id = threadIdx.x % warpSize;
if (lane_id == leader) { // leader does the update
res = atomicAdd(ptr, __popc(mask));
}
res = __shfl_sync(mask, res, leader); // get leader’s old value
return res + __popc(mask & ((1 << lane_id) - 1)); //compute old value
}
__global__ void k(int *d){
int *ptr = d + threadIdx.x/4;
if ((threadIdx.x >= 16) && (threadIdx.x < 32))
atomicAggInc(ptr);
}
const int ds = 32;
int main(){
int *d_d, *h_d;
h_d = new int[ds];
cudaMalloc(&d_d, ds*sizeof(d_d[0]));
cudaMemset(d_d, 0, ds*sizeof(d_d[0]));
k<<<1,ds>>>(d_d);
cudaMemcpy(h_d, d_d, ds*sizeof(d_d[0]), cudaMemcpyDeviceToHost);
for (int i = 0; i < ds; i++)
std::cout << h_d[i] << " ";
std::cout << std::endl;
}
$ nvcc -o t1646 t1646.cu -DUSE_OPT
$ cuda-memcheck ./t1646
========= CUDA-MEMCHECK
0 0 0 0 4 4 4 4 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
========= ERROR SUMMARY: 0 errors
$
(CentOS 7,CUDA 10.1.243,设备 0 是 Tesla V100,设备 1 是 cc3.5 设备(。
我为翘曲发散的情况添加了一个可选的优化(即tmask
不是0xFFFFFFFF
(。 这可以通过定义USE_OPT
来选择。