cuda编程笔记（8）--线程束warp

CUDA 中的 Warp（线程束） 是 GPU 并行计算的最小执行调度单位

Warp 是由 32 个并行线程组成的执行单元，这些线程将以 SIMT（Single Instruction, Multiple Thread） 的方式同步执行同一条指令。

Grid
└── Block（多个）└── Warp（每 32 个线程组成一个）└── Thread（0 ~ 31）

SIMT 执行模型（Single Instruction Multiple Threads）

• Warp 内 32 个线程 同时执行一条相同的指令

• 但每个线程可以对 不同数据 操作（类似 SIMD）

• 每次由 warp scheduler 发出一个 warp 执行一个指令周期

分支发散（Divergence）问题

如果 warp 内线程执行了不同的分支语句，会导致性能下降：

if (threadIdx.x % 2 == 0)a[threadIdx.x] = 1;
elsea[threadIdx.x] = 2;

这时 warp 会分两轮执行，实际上还是串行，称为 分支发散（Branch Divergence）。

优化建议：

• warp 内线程尽量走相同路径（控制流一致）

• 使用 warpSize 辅助判断

• 结构上分组 warp 级逻辑（例如每个 warp 处理一个任务）

为什么有Warp这个概念？

在之前，我们介绍GPU的线程模型的时候，只有Grid-Block-Thread。这个Warp是从何而来的呢？

Warp 是为了描述 GPU 的底层“调度单位和执行行为”。 

Warp 让你理解性能本质：GPU 是按 32 个线程批处理的 

一个 block 可能有 128 或 1024 个线程，但 GPU 是以 32 个线程为单位调度和执行 的：

• CUDA 内核 launch 时，会将一个 block 拆成若干个 warp

• GPU 的 warp scheduler 一次调度一整个 warp

• warp 内线程是 SIMT（Single Instruction Multiple Thread） 执行的

所以如果你不了解 warp：

你无法解释：为什么某些 if 分支会让 kernel 执行变慢（→ branch divergence）
你也不会理解：为什么 __shfl_sync() 比 shared memory 更快（→ warp 内通信）

 warp-level 的通信操作比 shared memory 更快、更轻

例子：warp 级归约只用 5~6 次寄存器交换；shared memory 要写回再同步再读取。

 warp-level 原语

__shfl_sync

int __shfl_sync(unsigned mask, int var, int srcLane, int width = warpSize);

允许 warp 内线程之间直接交换变量值，用于做 warp 内归约、broadcast 等。 所有线程调用这个函数后，会拿到 srcLane 上线程的 var 值。

int val = threadIdx.x;
int broadcasted = __shfl_sync(0xffffffff, val, 0); // 所有线程获取 lane 0 的值

其他变种：

 __shfl_sync系列的val是向外传递的数据，同时该函数返回值是从别的线程获取的数据

__ballot_sync

unsigned int __ballot_sync(unsigned mask, int predicate);

返回值是一个 32-bit 的整数，每一位表示该 lane 上线程是否为 true。 

int x = threadIdx.x % 2;
unsigned int bitmask = __ballot_sync(0xFFFFFFFF, x == 0);
// 如果 0、2、4 线程满足，则 bitmask = 0b...101010

 __any_sync() / __all_sync()

在 warp 内判断：是否有 / 是否所有线程满足某条件

int __any_sync(unsigned mask, int predicate);
int __all_sync(unsigned mask, int predicate);

 若 warp 中 至少有一个(any)/所有(all) 线程 predicate 为 true返回非零，否则返回 0。

int is_positive = (threadIdx.x > 0);if (__all_sync(0xFFFFFFFF, is_positive)) {// 所有线程都满足条件
}if (__any_sync(0xFFFFFFFF, is_positive)) {// 至少有一个线程满足条件
}

 __activemask

unsigned int __activemask(void);

返回当前 warp 中的 有效线程掩码，用于配合 __shfl_sync 等 warp 操作。

 示例程序

warp实现规约 vs 共享内存

#ifndef __CUDACC__
#define __CUDACC__
#endif
#include <cuda_runtime.h>
#include <device_launch_parameters.h>#include <iostream>
#include<cstdio>#define N 256void error_handling(cudaError_t res) {if (res !=cudaSuccess) {std::cout << "error!" << std::endl;}
}
__global__ void reduce_with_shared(float* in, float* out) {__shared__ float smen[N];int tid = threadIdx.x;smen[tid] = in[tid];__syncthreads();for (int stride = blockDim.x / 2; stride >= 1; stride >>=  1) {if (tid < stride)smen[tid] += smen[tid + stride];__syncthreads();}if (tid == 0) *out = smen[0];
}
__inline__ __device__ float warpReduce(float val) {// 每次从右边获取值，做加法//先广播自己当前的 val 值//然后从其他线程接收一个 val 值，再累加到自身的valfor (int offset = warpSize / 2; offset >= 1; offset >>= 1)val += __shfl_down_sync(0xFFFFFFFF, val, offset);return val;
}
__global__ void reduce_with_warp(float* in, float* out) {int tid = threadIdx.x;float val = in[tid];float sum = warpReduce(val);//WarpSize是cuda提供的线程束大小的值if (tid % warpSize == 0)         // 只有每个 warp 的 lane 0 写结果out[tid / warpSize] = sum;
}
int main() {// 生成 N 个值初始化为 1.0f，期望和为 Nfloat* h_in, * d_in, * d_out;cudaMallocHost(&h_in, N * sizeof(float));for (int i = 0; i < N; ++i) h_in[i] = 1.0f;cudaMalloc(&d_in, N * sizeof(float));cudaMalloc(&d_out, N * sizeof(float));cudaMemcpy(d_in, h_in, N * sizeof(float), cudaMemcpyHostToDevice);reduce_with_warp << <1, N >> > (d_in, d_out);cudaFreeHost(h_in);cudaFree(d_in);cudaFree(d_out);
}

warp规约有点难懂，需要解释一下

首先要理解warp的规约是在一个线程束内进行的规约（一般是32个线程），以及一个线程束内的线程是同时执行同一条指令的

也就是说32个线程同时走到了WarpReduce调用并进入，同时走到了__shfl_down_sync

val += __shfl_down_sync(0xFFFFFFFF, val, offset);

在这一行代码中，每个线程：

1. 先广播自己当前的 val 值（传出去）；

2. 同时从 lane = 自己的 lane + offset 那里接收一个 val 值；

3. 然后将这个值加到自己的 val 上。

有个细节：如果 __shfl_down_sync 的目标 lane 超过了 31，CUDA 会自动忽略该调用

牢记所有线程都是同时执行该语句的，所以不会有 有的线程先执行__shfl_down_sync修改了自己的val，再把val传出去的情况。

最后线程0会拿到32个线程值的总和

warp实现softmax

对于一组数，softmax 定义为：

                ​​​​​​​        ​​​​​​​        ​​​​​​​        

不过如果x很大，那么对应的e^x也会很大，一般会让所有x减去xi中的最大值再取e的指数幂（这样不会影响比例）

下面是一个warp内实现softmax的代码

__device__ float warp_softmax(float val) {// Step 1: 先减去最大值以提升数值稳定性float max_val = val;// 在warp内找最大值（规约）for (int offset = 16; offset > 0; offset /= 2) {float temp = __shfl_down_sync(0xffffffff, max_val, offset);max_val = fmaxf(max_val, temp);}// 广播最大值给所有线程（lane 0保存了最终结果）max_val = __shfl_sync(0xffffffff, max_val, 0);// Step 2: 计算指数float exp_val = expf(val - max_val);// Step 3: 求和所有expfloat sum_exp = exp_val;for (int offset = 16; offset > 0; offset /= 2) {sum_exp += __shfl_down_sync(0xffffffff, sum_exp, offset);}// 广播sum给所有线程（lane 0保存了最终结果）sum_exp = __shfl_sync(0xffffffff, sum_exp, 0);// Step 4: softmax输出return exp_val / sum_exp;
}

下面是block级别的warp实现，不同warp之间还是需要共享内存来同步；求的是一个block内的softmax

#include <cstdio>
#include <cmath>#define WARP_SIZE 32
//warp内规约求和（lane 0获取总和）
__inline__ __device__
float warpReduceSum(float val) {for (int offset = WARP_SIZE / 2; offset > 0; offset /= 2)val += __shfl_down_sync(0xffffffff, val, offset);return val;
}
//warp内规约求最大值（lane 0获取最大值）
__inline__ __device__
float warpReduceMax(float val) {for (int offset = WARP_SIZE / 2; offset > 0; offset /= 2)val = fmaxf(val, __shfl_down_sync(0xffffffff, val, offset));return val;
}__global__ void block_softmax(float* input, float* output, int N) {int tid = threadIdx.x + blockIdx.x * blockDim.x;int lane = threadIdx.x % WARP_SIZE;//该线程在本warp内的编号int warp_id = threadIdx.x / WARP_SIZE;//该warp在整个block里的序号//每个warp对应的总和和最大值__shared__ float warp_max[32];__shared__ float warp_sum[32];__shared__ float smem_block_max;__shared__ float smem_block_sum;float x = (tid < N) ? input[tid] : -INFINITY;// Step 1: Warp内最大值float max_val = warpReduceMax(x);// Step 2: 线程0收集每个warp的最大值if (lane == 0)warp_max[warp_id] = max_val;__syncthreads();// Step 3: block范围内最大值（使用前WARP_SIZE线程处理）//整个block只有前32个线程干活了，根据warp_max进行规约求出block内的最大值float block_max = -INFINITY;if (threadIdx.x < WARP_SIZE) {block_max = warp_id < (blockDim.x + WARP_SIZE - 1) / WARP_SIZE ? warp_max[threadIdx.x] : -INFINITY;block_max = warpReduceMax(block_max);}// 广播block_max，整个block内的线程都会收到// thread 0 写入共享内存广播if (threadIdx.x == 0) smem_block_max = block_max;__syncthreads();block_max = smem_block_max;// Step 4: 减去最大值后求expfloat exp_x = (tid < N) ? expf(x - block_max) : 0.0f;// Step 5: warp内规约求和float local_sum = warpReduceSum(exp_x);// Step 6: 每个warp写入warp_sumif (lane == 0)warp_sum[warp_id] = local_sum;__syncthreads();// Step 7: block范围内总和（again用前WARP_SIZE线程规约）float block_sum = 0.0f;if (threadIdx.x < WARP_SIZE) {block_sum = warp_id < (blockDim.x + WARP_SIZE - 1) / WARP_SIZE ? warp_sum[threadIdx.x] : 0.0f;block_sum = warpReduceSum(block_sum);}// 广播block_sumif (threadIdx.x == 0) smem_block_sum = block_sum;__syncthreads();block_sum = smem_block_sum;// Step 8: 输出softmaxif (tid < N)output[tid] = exp_x / block_sum;
}

如果要整个 grid 中的所有线程输入数据 进行 softmax，由于block之间无法简便通信，只能将block级别的最大值、和先传回主机，主机先对每个block之间的最大值、和求出最大值、和，再启动新的核函数，也即

全局 Softmax 涉及三次 kernel 启动：

• 第一次归约 max。

• 第二次归约 exp 和。

• 第三次归一化。