执行模型

MT GPU 采用单指令多线程 SIMT（Single Instruction, Multiple Threads）执行模型：

SIMT 核心特征：
├── 一组线程（线程束）同时执行同一条指令
├── 每个线程具有独立的寄存器状态
└── 当线程执行路径出现分支时，线程束序列化执行各分支

与 SIMD 的区别

特性	SIMD	SIMT
控制流	所有数据执行相同路径	每个线程可独立分支
编程模型	需显式使用向量指令	可编写直观标量代码
分支处理	必须执行所有分支	硬件自动处理分歧
灵活性	低，适合规则数据	高，适合不规则任务

---

线程束基础

线程束是 GPU 硬件的基本执行单元，包含固定数量的线程：

GPU 架构	线程束大小
MP31	32 线程
MP21/MP22	128 线程

线程束执行过程

---

分支分化（Branch Divergence）

问题描述

当线程束内线程执行路径出现分支时，硬件必须序列化执行各分支：

__global__ void branchKernel(float* data) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (data[idx] > 0) {
        // 路径 A：正数处理
        data[idx] = data[idx] * 2.0f;
    } else {
        // 路径 B：非正数处理
        data[idx] = 0.0f;
    }
}

硬件处理方式

性能影响：分支分化导致线程束内部分线程空闲等待，总执行时间 = 各分支执行时间之和，硬件利用率下降。

优化建议

// ❌ 可能引起分支分化
if (threadIdx.x % 2 == 0) {
    result = value * 2.0f;
} else {
    result = 0.0f;
}

// ✅ 减少分支分化：使用无分支操作
float sign = (threadIdx.x % 2 == 0) ? 1.0f : 0.0f;
result = value * sign;

---

延迟隐藏（Latency Hiding）

问题

GPU 内存访问延迟高达 100+ 周期，但线程束执行只需几周期：

内存访问延迟：~400ns（100+ 周期）
ALU 操作延迟：~1-4 周期

解决方案：多 线程束切换

GPU 通过在线程间切换隐藏内存延迟：

时间轴：
│ 线程束0: [ALU] [ALU] [等待内存] [ALU] [ALU]
│ 线程束1: [ALU] [ALU] [ALU]        [ALU] [ALU]
│ 线程束2: [ALU] [ALU] [ALU]        [ALU] [等待内存]
│ 线程束3: [ALU] [等待内存] [ALU]   [ALU] [ALU]
         ↑
    当 线程束0 等待内存时，切换到 线程束1/2/3 执行

占用率（Occupancy）

占用率 = 活跃 线程束数 / 每个 MP 最大 线程束数

高占用率的好处：
├── 更好的延迟隐藏
├── 更高的吞吐量
└── 更充分的硬件利用

---

同步原语

线程块内同步

__syncthreads();  // 等待线程块内所有线程到达同步点

使用场景：

• 共享内存读写同步
• 确保数据可见性

__shared__ float shared_data[256];

// 步骤 1：写入共享内存
shared_data[threadIdx.x] = input[threadIdx.x];

// 步骤 2：同步，确保所有线程完成写入
__syncthreads();

// 步骤 3：读取其他线程的数据
float neighbor = shared_data[threadIdx.x + 1];

线程束级同步

__syncwarp();  // 等待 线程束内所有线程到达同步点
__syncwarp(mask);  // 等待指定线程掩码的线程

---

线程束级原语

投票操作

// __all_sync：所有线程满足条件返回 1
if (__all_sync(0xffffffff, predicate)) {
    // 所有线程都满足条件
}

// __any_sync：任意线程满足条件返回 1
if (__any_sync(0xffffffff, predicate)) {
    // 至少一个线程满足条件
}

// __ballot_sync：返回每个线程的谓词结果
unsigned mask = __ballot_sync(0xffffffff, predicate);
// mask 的第 i 位表示第 i 个线程的 predicate 值

Shuffle 操作

// __shfl_sync：从指定 lane 获取值
float val = __shfl_sync(0xffffffff, src, srcLane);

// __shfl_down_sync：从下方 lane 获取值
float val = __shfl_down_sync(0xffffffff, src, delta);

// __shfl_up_sync：从上方 lane 获取值
float val = __shfl_up_sync(0xffffffff, src, delta);

// __shfl_xor_sync：XOR 索引获取值
float val = __shfl_xor_sync(0xffffffff, src, laneMask);

线程束归约示例

__global__ void warpReduce(float* data) {
    float val = data[threadIdx.x];
    
    // 线程束内归约
    val += __shfl_down_sync(0xffffffff, val, 16);
    val += __shfl_down_sync(0xffffffff, val, 8);
    val += __shfl_down_sync(0xffffffff, val, 4);
    val += __shfl_down_sync(0xffffffff, val, 2);
    val += __shfl_down_sync(0xffffffff, val, 1);
    
    if (threadIdx.x % 32 == 0) {
        data[threadIdx.x] = val;
    }
}

---

原子操作（Atomics）

原子操作是不可中断的内存操作，在多线程并发访问时保证数据一致性。

原子函数

// 原子加
__global__ void atomicAddKernel(float* data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    // 全局原子操作
    atomicAdd(&data[0], 1.0f);
}

// 原子 CAS (Compare-And-Swap)
__global__ void atomicCAS_Kernel(int* data, int compare, int val) {
    int old = atomicCAS(&data[0], compare, val);
    // 如果 data[0] == compare，则设置为 val，返回原值
}

原子操作类型

操作	函数	支持类型
原子加	atomicAdd()	int, float, double
原子减	atomicSub()	int
原子交换	atomicExch()	int, float
原子递增	atomicInc()	unsigned int
原子递减	atomicDec()	unsigned int
原子与	atomicAnd()	int
原子或	atomicOr()	int
原子异或	atomicXor()	int
原子 CAS	atomicCAS()	int

原子操作性能优化

// 方案：使用共享内存减少全局原子操作
__global__ void optimizedAtomicKernel(float* output, float* data, int n) {
    __shared__ float sharedSum;
    int tid = threadIdx.x;
    
    // 初始化共享内存
    if (tid == 0) sharedSum = 0.0f;
    __syncthreads();
    
    // 每个线程累加到共享内存
    float localSum = 0.0f;
    for (int i = tid; i < n; i += blockDim.x) {
        localSum += data[i];
    }
    
    // 原子加到共享内存
    atomicAdd(&sharedSum, localSum);
    __syncthreads();
    
    // 单个线程写入全局内存
    if (tid == 0) {
        atomicAdd(output, sharedSum);
    }
}

原子操作使用场景

// 场景 1: 计数器
__global__ void counterKernel(int* counter) {
    atomicInc(counter, UINT_MAX);
}

// 场景 2: 直方图
__global__ void histogramKernel(float* data, int* hist, int n, int bins) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        int bin = min((int)(data[idx] * bins), bins - 1);
        atomicAdd(&hist[bin], 1);
    }
}

// 场景 3: 锁
__global__ void lockKernel(int* lock, int* data) {
    // 获取锁
    while (atomicCAS(lock, 0, 1) != 0);
    
    // 临界区
    *data = *data + 1;
    
    // 释放锁
    atomicExch(lock, 0);
}

---

相关文档

• GPU 并行计算基础：GPU 架构基础，SIMT 基本概念
• 线程模型：线程层次，Grid/Block/Thread 组织
• 内存模型：内存层次，寄存器/共享内存/全局内存
• Host/Device 架构：Host/Device，流与并发