CUDA学习（八十七）

独立的线程调度：
Volta体系结构在变形中的线程之间引入了独立线程调度功能，可实现先前不可用的内部变形同步模式，并在移植CPU代码时简化代码更改。 但是，如果开发人员对以前的硬件体系结构的同步性做出假设，这可能会导致参与执行的代码的一组相当不同的线程。
以下是关注的代码模式，并提出了针对Volta安全代码的纠正措施：
1.对于使用warp内在函数（__shfl *，__any，__all，__ballot）的应用程序，开发人员必须使用* _sync后缀将代码移植到新的，安全的同步对象。 新的扭曲内在函数需要一个线程掩码来明确定义哪些通道（warp的线程）必须参与warp内在。 有关详细信息，请参阅变形表决函数和变形混合函数。
由于CUDA 9.0+提供了内在函数，因此可以使用以下预处理器宏有条件地执行代码（如有必要）：

#if defined(CUDART_VERSION) && CUDART_VERSION >= 9000
// *_sync intrinsic
#endif

这些内在函数在所有体系结构上都可用，不仅仅是Volta，在大多数情况下，单个代码库就足以适用于所有体系结构。 但是，请注意，对于Pascal和更早的体系结构，掩码中的所有线程必须在收敛中执行相同的warp内部指令，并且掩码中所有值的并集必须与warp的活动掩码相等。 以下代码模式在Volta上有效，但不在Pascal或更早版本的体系结构上。

if (tid % warpSize < 16) {
    ...
        float swapped = __shfl_xor_sync(0xffffffff, val, 16);
    ...
}
else {
    ...
        float swapped = __shfl_xor_sync(0xffffffff, val, 16);
    ...
}

__ballot（1）的替换是__activemask（）。 请注意，即使在单个代码路径内，warp内的线程也可能会发散。 因此，__activemask（）和__ballot（1）可能只返回当前代码路径上的线程子集。 以下无效代码示例在data [i]大于阈值时将输出的位i设置为1。 __activemask（）用于尝试启用dataLen不是32的倍数的情况。

// Sets bit in output[] to 1 if the correspond element in data[i]
// is greater than ‘threshold’, using 32 threads in a warp.
for (int i = warpLane; i<dataLen; i += warpSize) {
    unsigned active = __activemask();
    unsigned bitPack = __ballot_sync(active, data[i] > threshold);
    if (warpLane == 0)
        output[i / 32] = bitPack;
}

此代码无效，因为CUDA不保证仅在循环条件下变形才会发生变化。 当由于其他原因发生分歧时，将通过变形中的不同线程子集针对相同的32位输出元素计算冲突结果。 正确的代码可能会使用非发散循环条件和__ballot_sync（）一起安全地枚举参与阈值计算的warp中的线程集合，如下所示:

for (int i = warpLane; i - warpLane<dataLen; i += warpSize) {
    unsigned active = __ballot_sync(0xFFFFFFFF, i < dataLen);
    if (i < dataLen) {
        unsigned bitPack = __ballot_sync(active, data[i] > threshold);
        if (warpLane == 0)
            output[i / 32] = bitPack;
    }
}

2.如果应用程序具有warp同步代码，则它们将需要在通过全局或共享内存在线程之间交换数据的任何步骤之间插入新的__syncwarp（）全范围屏障同步指令。 代码以锁步执行的假设或从单独线程读取/写入的假设在整个变形中都是可见的，而不同步是无效的。

float s_buff[tid] = val;
__syncthreads();
// Inter-warp reduction
for (int i = BSIZE / 2; i>32; i /= 2) {
    s_buff[tid] += s_buff[tid + i];
    __syncthreads();
}
// Intra-warp reduction
// Butterfly reduction simplifies syncwarp mask
if (tid < 32) {
    float temp;
    temp = s_buff[tid ^ 16]; __syncwarp();
    s_buff[tid] += temp; __syncwarp();
    temp = s_buff[tid ^ 8]; __syncwarp();
    s_buff[tid] += temp; __syncwarp();
    temp = s_buff[tid ^ 4]; __syncwarp();
    s_buff[tid] += temp; __syncwarp();
    temp = s_buff[tid ^ 2]; __syncwarp();
    s_buff[tid] += temp; __syncwarp();
}
if (tid == 0) {
    *output = s_buff[0] + s_buff[1];
}
__syncthreads()