使用CUDA Warp-Level级原语
使用CUDA Warp-Level级原语
NVIDIA GPU以SIMT（单指令，多线程）的方式执行称为warps 的线程组。许多CUDA程序通过利用warp执行来实现高性能。本文将展示如何使用cuda9中引入的原语，使warp级编程安全有效。
Figure 1: The Tesla V100 Accelerator with Volta GV100 GPU. SXM2 Form Factor.
图1:Volta GV100 GPU的特斯拉V100加速器。SXM2外形尺寸。
Warp-level Primitives
nvidiagpu和CUDA编程模型采用了一种称为SIMT（单指令多线程）的执行模型。SIMT扩展了Flynn的计算机体系结构分类法，该分类法根据指令流和数据流的数量描述了四类体系结构。Flynn的四个类之一SIMD（单指令多数据）通常用于描述gpu之类的体系结构。但是SIMD和SIMT之间有一个微妙但重要的区别。在SIMD体系结构中，每条指令在多个数据元素上并行应用相同的操作。SIMD通常使用带有向量寄存器和执行单元的处理器来实现；标量线程发出以SIMD方式执行的向量指令。在SIMT体系结构中，多个线程向任意数据发出公共指令，而不是单个线程向数据向量发出向量指令。
SIMT对于可编程性的好处,使NVIDIA的GPU架构师为这种体系结构创造了一个新的名称，而不是将其描述为SIMD。NVIDIA GPU使用SIMT执行32个并行线程的warp，这使每个线程能够访问自己的寄存器，从不同的地址加载和存储，并遵循不同的控制流路径。CUDA编译器和GPU协同工作，以确保warp的线程尽可能频繁地一起执行相同的指令序列，从而最大限度地提高性能。
虽然warp执行所获得的高性能发生在后台，但许多CUDA程序通过使用显式warp级编程可以获得更高的性能。并行程序通常使用选择通信操作，例如并行缩减和扫描。CUDA C++通过提供经编级别的原语和合作组集合来支持这种选择操作。协作组collectives（在前一篇文章中描述）是在本文关注的warp原语之上实现的。
表1显示了一个使用warp-level primitives原语的示例。它用于__shfl_down_sync()执行tree-reduction，以计算val warp中每个线程所持有的变量的总和。在循环结束时，val warp中的第一个线程the warp contains the sum总和。
#define FULL_MASK 0xffffffff
for (int offset = 16; offset > 0; offset /= 2)
val += __shfl_down_sync(FULL_MASK, val, offset);
一个thread包含32个通道，每个线程占用一个通道。对于X thread通道中的线程，请从同一thread通道中的线程__shfl_down_sync(FULL_MASK, val, offset)获取val变量的值X+offset。数据交换是在寄存器之间执行的，并且比通过共享内存更有效，共享内存需要加载，存储和额外的寄存器来保存地址。
CUDA 9引入了三类新的或更新的warp级图元。
- 同步数据交换：在warp中的线程之间交换数据。
- __all_sync，__any_sync，__uni_sync，__ballot_sync
- __shfl_sync，__shfl_up_sync，__shfl_down_sync，__shfl_xor_sync
- __match_any_sync， __match_all_sync
- 活动掩码查询：返回一个32位掩码，该掩码指示warp中的哪些线程在当前执行线程中处于活动状态。
- __activemask
- 线程同步：同步线程束中的线程并提供内存隔离。
- __syncwarp
Synchronized Data Exchange
每个“同步数据交换”原语在线程束中的一组线程之间执行集合操作。例如，表2显示了其中的三个。每个线程在同一warp中从一个线程调用__shfl_sync()或 __shfl_down_sync()接收数据，并且每个调用的线程都__ballot_sync()接收一个位掩码，该掩码表示warp中所有传递谓词参数真值的线程。
int __shfl_sync(unsigned mask, int val, int src_line, int width=warpSize);
int __shfl_down_sync(unsigned mask, int var, unsigned detla,
int width=warpSize);
int __ballot_sync(unsigned mask, int predicate);
参与调用每个原语的线程集使用32位掩码指定，这是这些原语的第一个参数。必须使所有参与线程同步，以使选择操作正常工作。因此，如果这些原语尚未同步，则它们首先将同步线程。
一个常见的问题是“应该用什么作为mask论点？”。可以考虑将掩码表示thread中应参与选择操作的一组thread。这组线程由程序逻辑确定，通常可以通过程序流程中较早的某些分支条件进行计算。以表1中的简化代码为例。假设要计算一个数组的所有元素的总和，该数组input[],的大小NUM_ELEMENTS小于线程块中的线程数。可以使用表3中的方法。
unsigned mask = __ballot_sync(FULL_MASK, threadIdx.x < NUM_ELEMENTS);
if (threadIdx.x < NUM_ELEMENTS) {
val = input[threadIdx.x];
for (int offset = 16; offset > 0; offset /= 2)
val += __shfl_down_sync(mask, val, offset);
…
}
该代码使用条件thread.idx.x < NUM_ELEMENTS来确定线程是否将减少参与。 __ballot_sync()用于计算操作的成员资格掩码__shfl_down_sync()。__ballot_sync()本身使用FULL_MASK（0xffffffff用于32个线程），因为假设所有线程都将执行它。
在Volta和更高版本的GPU架构上，数据交换原语可以在线程分散的分支中使用：分支中，warp中的某些线程所采用的路径与其他线程不同。表4显示了一个示例，其中，warp中的所有线程val从通道0处的线程获取值。偶数和奇数线程采用if语句的不同分支。
if (threadIdx.x % 2) {
val += __shfl_sync(FULL_MASK, val, 0);
…
}
else {
val += __shfl_sync(FULL_MASK, val, 0);
…
}
在最新的Volta（及以后）的GPU上，可以运行使用warp同步原语的库函数，而不必担心该函数是否在线程thread-divergent发散分支中调用。
Active Mask Query
__activemask()返回unsigned int调用warp中所有当前活动线程的32位掩码。换句话说，它显示了调用线程中的哪些线程也在执行相同的线程__activemask()。这对于稍后说明的“opportunistic warp-level programming”技术，调试和理解程序行为很有用。
但是，__activemask()正确使用很重要。表5说明了一种不正确的用法。该代码尝试执行表4所示的相同的总和减少，但是没有__ballot_sync()在分支之前用来计算掩码，而是__activemask()在分支内部使用。这是不正确的，将导致部分和而不是总数。CUDA执行模型不能保证将分支合并在一起的所有线程都将一起执行__activemask()。正如将要解释的，不能保证隐式锁定步骤的执行。
//
// Incorrect use of __activemask()
//
if (threadIdx.x < NUM_ELEMENTS) {
unsigned mask = __activemask();
val = input[threadIdx.x];
for (int offset = 16; offset > 0; offset /= 2)
val += __shfl_down_sync(mask, val, offset);
…
}
Warp Synchronization
当warp中的线程需要执行比数据交换原语所提供的更复杂的通信或选择操作时，可以使用该__syncwarp()原语来同步warp中的线程。它类似于__syncthreads()原语（同步线程块中的所有线程），但粒度更细。
void __syncwarp(unsigned mask=FULL_MASK);
__syncwarp()原语使正在执行的线程等待，直到指定的所有线程mask都已执行了__syncwarp()（具有相同的mask），然后才恢复执行。它还提供了一个 内存屏障， 以允许线程在调用原语之前和之后通过内存进行通信。
表6显示了一个示例，其中将对warp中的线程之间的矩阵元素的所有权进行改组。
float val = get_value(…);
__shared__ float smem[4][8];
// 0 1 2 3 4 5 6 7
// 8 9 10 11 12 13 14 15
// 16 17 18 19 20 21 22 23
// 24 25 26 27 28 29 30 31
int x1 = threadIdx.x % 8;
int y1 = threadIdx.x / 8;
// 0 4 8 12 16 20 24 28
// 1 5 10 13 17 21 25 29
// 2 6 11 14 18 22 26 30
// 3 7 12 15 19 23 27 31
int x2= threadIdx.x / 4;
int y2 = threadIdx.x % 4;
smem[y1][x1] = val;
__syncwarp();
val = smem[y2][x2];
use(val);
假设使用一维线程块（即threadIdx.y始终为0）。在代码的开头，warp中的每个线程都拥有一个4×8矩阵的元素，并具有行优先索引。换句话说，通道0拥有[0][0]，而通道1拥有[0][1]。每个线程将其值存储到共享内存中4×8数组的相应位置。然后__syncwarp()，在每个线程从数组中的转置位置读取之前，使用来确保所有线程都已完成存储。最后，thread中的每个线程都拥有矩阵的一个元素，并且具有列索引索引：lane道0拥有[0][0]，lane道1拥有[1][0]。
确保__syncwarp()将共享内存的读取和写入分开，以避免出现竞争状况。表7说明了在减少共享内存中的树总数方面的错误使用。每两个__syncwarp()调用之间有一个共享内存读取，然后是一个共享内存写入。CUDA编程模型不能保证所有读取都将在所有写入之前执行，因此存在竞争条件。
unsigned tid = threadIdx.x;
// Incorrect use of __syncwarp()
shmem[tid] += shmem[tid+16]; __syncwarp();
shmem[tid] += shmem[tid+8]; __syncwarp();
shmem[tid] += shmem[tid+4]; __syncwarp();
shmem[tid] += shmem[tid+2]; __syncwarp();
shmem[tid] += shmem[tid+1]; __syncwarp();
表8通过插入额外的__syncwarp()调用来解决竞争条件。CUDA编译器可能会根据目标架构（例如，在Vol-pre架构之前），在最终生成的代码中，取消其中一些同步指令。
unsigned tid = threadIdx.x;
int v = 0;
v += shmem[tid+16]; __syncwarp();
shmem[tid] = v; __syncwarp();
v += shmem[tid+8]; __syncwarp();
shmem[tid] = v; __syncwarp();
v += shmem[tid+4]; __syncwarp();
shmem[tid] = v; __syncwarp();
v += shmem[tid+2]; __syncwarp();
shmem[tid] = v; __syncwarp();
v += shmem[tid+1]; __syncwarp();
shmem[tid] = v;
在最新的Volta（及以后）GPU上，还可以__syncwarp()在线程分散的分支中，使用以同步来自两个分支的线程。但是一旦它们从原语返回，线程将再次变得发散。有关示例，请参见表13。
Opportunistic Warp-level Programming
正如我们在“同步数据交换”部分mask所显示的那样，同步数据交换原语中使用的成员资格，通常是在程序流中的分支条件之前计算的。在许多情况下，程序需要沿着程序流传递掩码；例如，当在函数内部使用warp-level基元时，作为函数参数。如果要在库函数中使用warp-level编程，但是不能更改函数接口，可能会很困难。
一些计算可以使用碰巧一起执行的任何线程。可以使用称为机会warp-level编程的技术，如以下示例所示。（有关此算法的更多信息，请参见有关warp聚合原子的文章，以及有关“Cooperative Groups”如何使实现更简单的讨论）。
// increment the value at ptr by 1 and return the old value
__device__ int atomicAggInc(int *ptr) {
int mask = __match_any_sync(__activemask(), (unsigned long long)ptr);
int leader = __ffs(mask) – 1; // select a leader
int res;
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
}
atomicAggInc()以原子方式将指向的值ptr加1，然后返回以前值。它使用该atomicAdd()功能，可能会引起争用。为了减少争用，请atomicAggInc使用atomicAdd()per-warp替换每个线程的操作atomicAdd()。第4行__activemask()在thread中找到将要执行原子操作的线程集。__match_any_sync()返回具有相同值的线程的位掩码ptr，分割输入的线程到其成员具有相同的ptr的值。每个组选择一个lead线程（第5行），该线程atomicAdd()为整个组执行（第8行）。每个线程都从线程返回的leader（第9行）获取以前值。atomicAdd()。第10行计算并返回当前线程（atomicInc()如果调用该函数而不是）将获得的以前值atomicAggInc。
Implicit Warp-Synchronous Programming is Unsafe
9.0版之前的CUDA工具包提供了一个（现在是旧的）warp-level原语版本。与CUDA 9基元相比，旧基元不接受 mask参数。例如，int __any(int predicate)是的旧版本int __any_sync(unsigned mask, int predicate)。
mask如前所述，该参数指定必须参与原语的线程束中的线程集。如果掩码指定的线程在执行过程中尚未同步，则新的原语将执行线程内线程级同步。
传统的warp级原语不允许程序员指定所需的线程，并且不执行同步。因此，CUDA程序未明确表示必须参与warp-level操作的线程。这种程序的正确性取决于隐式的warp同步行为，这种行为可能会从一种硬件体系结构，更改为另一种硬件体系结构，从一种CUDA工具包版本，更改为另一种（例如，由于编译器优化的更改），甚至可能从一个runtime更改。执行到另一个。这样的隐式warp同步编程是不安全的，可能无法正常工作。
例如，在下面的代码中，假设warp中的所有32个线程一起执行第2行。第4行的if语句导致线程发散，奇数线程foo()在第5行调用，偶数线程bar()在第8行调用。
// Assuming all 32 threads in a warp execute line 1 together.
assert(__ballot(1) == FULL_MASK);
int result;
if (thread_id % 2) {
result = foo();
}
else {
result = bar();
}
unsigned ballot_result = __ballot(result);
CUDA编译器和硬件将尝试在第10行重新收敛线程，以提高性能。但是不能保证这种重新收敛。因此，ballot_result可能不包含所有32个线程的投票结果。
如表11所示，__syncwarp()在第10行之前调用新原语__ballot()也无法解决问题。这又是隐式的warp同步编程。假定曾经同步过的同一线程束中的线程将保持同步，直到下一个the next thread-divergent branch线程分歧分支为止。尽管通常是这样，但在CUDA编程模型中并不能保证。
__syncwarp();
unsigned ballot_result = __ballot(result);
The correct fix is to use __ballot_sync() as in Listing 12.
unsigned ballot_result = __ballot_sync(FULL_MASK, result);
一个常见的错误是，假定__syncwarp()在旧的warp-level原语之前，和/或之后进行调用在功能上，等同于调用sync原语的版本。例如，__syncwarp(); v = __shfl(0); __syncwarp();与__shfl_sync(FULL_MASK, 0)？相同吗？答案是否定的，原因有二。首先，如果序列在线程分歧分支中使用，则__shfl(0)不会被所有线程一起执行。表13显示了一个示例。第3行__syncwarp()和第7行将确保foo()在执行第4行或第8行之前，warp中的所有线程都调用了at 。线程退出后__syncwarp()，奇数线程和偶数线程再次变得发散。因此，第4行的__shfl(0)将获得未定义的at值，因为在执行第4行时通道0不活动。__shfl_sync(FULL_MASK, 0) 可以在线程发散分支中使用，而不会出现此问题。
v = foo();
if (threadIdx.x % 2) {
__syncwarp();
v = __shfl(0); // L3 will get undefined result because lane 0
__syncwarp(); // is not active when L3 is executed. L3 and L6
} else { // will execute divergently.
__syncwarp();
v = __shfl(0);
__syncwarp();
}
其次，即使序列被所有线程一起调用，CUDA执行模型也不保证线程离开后将保持收敛__syncwarp()，如表14所示。不保证隐式执行锁步。请记住，只有在显式同步的warp级原语内才能保证线程收敛。
assert(__activemask() == FULL_MASK); // assume this is true
__syncwarp();
assert(__activemask() == FULL_MASK); // this may fail
因为使用它们会导致程序不安全，所以从CUDA 9.0开始不推荐使用旧的扭曲级原语。
Update Legacy Warp-Level Programming
如果程序使用以前的warp级原语或任何形式的隐式warp同步编程（例如，warp的线程之间进行通信而不同步），则应更新代码以使用sync原语的版本。还需要调整代码中使用Cooperative Groups，它提供的抽象以及新功能，如多块同步的更高的水平。
使用warp级基元最棘手的部分是弄清楚要使用的成员资格掩码。希望以上各节提供一个好方法，从何处开始和要注意什么。以下是建议列表：
- 不要只使用 FULL_MASK（即0xffffffff用于32个线程）作为mask值。如果不是所有线程thread都可以根据程序逻辑到达原语，则使用FULL_MASK可能会导致程序挂起。
- 不要只是__activemask()用作掩码值。__activemask()调用该函数时，哪些线程恰好收敛，这可能希望在选择操作中获得的线程不同。
- 请分析程序逻辑并了解成员资格要求。根据您的程序逻辑预先计算掩码。
- 如果程序执行机会warp同步编程，请使用“detective”功能，例如__activemask()和__match_all_sync()查找正确的掩码。
- 使用__syncwarp()与intra-warp依赖性独立的操作。不要假定执行锁定步骤。
最后一招。如果现有的CUDA程序在Volta架构GPU上给出了不同的结果，并且怀疑该差异是由Volta的新独立线程调度 （可以更改warp同步行为）引起的，则可能希望使用nvccoptions重新编译程序-arch=compute_60 -code=sm_70。这样的编译程序选择加入Pascal的线程调度。如果有选择地使用，则可以帮助更快地固定罪魁祸首模块，从而可以更新代码，以避免隐式warp同步编程。
Volta独立线程调度，支持交错执行来自不同分支的语句。这样就可以执行细粒度的并行算法，其中warp内的线程可以同步和通信。
