CUDA 执行模型与调度 Q&A

Q1：什么是 Warp Divergence（线程束分歧）？如何减少其影响？

什么是 Warp Divergence

GPU 以 Warp（32 个线程）为单位同时发射指令。当 Warp 内的线程由于条件分支（if/else、switch、while）而走向不同的代码路径时，就发生了 Warp Divergence（线程束分歧）。

发生分歧时，GPU 必须串行执行每条分支路径——先执行走某条路径的线程（其余线程被屏蔽），再执行走另一条路径的线程。所有路径都执行完毕后，Warp 才重新汇合（reconverge）继续往下执行。

极端情况： 若 Warp 内 32 个线程走了 32 条完全不同的路径，性能退化至理论峰值的 1/32。

分歧仅在 Warp 内部发生，不同 Warp 之间完全独立。 两个 Warp 各自走不同分支是完全并行的，没有任何性能损失。

CC 7.0 前后的差异

在 CC 7.0（Volta）之前，Warp 内所有 32 个线程共享一个程序计数器，必须严格锁步（lockstep）执行。

CC 7.0 起引入了独立线程调度（Independent Thread Scheduling），每个线程拥有独立的程序计数器和调用栈，可以在子 Warp 粒度上分散和重聚，灵活性更高。但这也意味着：依赖"同一 Warp 内线程一定锁步"隐式假设的老代码，在 CC 7.0+ 上可能行为改变，需要显式加 __syncwarp() 保证正确性。

如何减少 Warp Divergence

1. 让同一 Warp 内的线程走相同分支

最根本的方法：让分支条件依赖于 blockIdx 或 Warp ID，而不是 threadIdx——这样同一 Warp 内的线程会整体走同一条路径。

1
2
3
4
5
// 差：按线程 ID 分支，Warp 内一半线程走 if，一半走 else
if (threadIdx.x % 2 == 0) { /* ... */ } else { /* ... */ }

// 好：按 Warp ID 分支，整个 Warp 走同一路径
if (threadIdx.x / 32 % 2 == 0) { /* ... */ } else { /* ... */ }

2. 将分支粒度从线程级提升到线程块级

1
2
3
4
5
6
// 按块分工，每个块只做一种任务，块内无分歧
if (blockIdx.x < numBlocksForTaskA) {
    doTaskA();
} else {
    doTaskB();
}

3. 利用谓词指令（Predication）

对于短小的 if/else（如单行赋值），编译器有时会把分支转换为谓词指令——两条分支都执行，但通过掩码控制哪条结果生效。这消除了分歧，但代价是执行了多余的指令。当分支体很小时，这比串行执行两条路径更快。

4. 重新组织数据或算法

有时根本解决方案是调整数据布局或算法结构，让处理同类数据的线程自然落在同一个 Warp 里。例如在排序后再并行处理，或对稀疏数据做压缩。

5. 用 Cooperative Groups 的 coalesced_threads

1
2
3
4
5
6
7
8
9
10
11
#include <cooperative_groups.h>
namespace cg = cooperative_groups;

__global__ void kernel() {
    auto warp = cg::tiled_partition<32>(cg::this_thread_block());
    if (condition) {
        // 获取当前 Warp 内满足条件的所有线程
        auto active = cg::coalesced_threads();
        // active 内的线程是汇聚的，可以高效协作
    }
}

Q2：CUDA Kernel 的启动配置（Grid/Block 大小）如何设置？有何原则？

Kernel 的启动配置决定了线程的组织方式，直接影响 SM 利用率、内存访问效率和整体性能。没有万能公式，但有明确的设计原则。

线程块大小（blockDim）的原则

1. 必须是 32 的整数倍

线程块大小必须是 Warp 大小（32）的整数倍，否则最后一个 Warp 有若干 lane 空闲，浪费计算资源。常用值：128、256、512。

2. 每块至少 128 个线程

线程数过少，每个 SM 承载的 Warp 数不足以隐藏内存延迟。

3. 不超过 1024 个线程

这是每个线程块的硬件上限。

4. 考虑共享内存和寄存器的限制

每块使用的共享内存越多、每线程使用的寄存器越多，每个 SM 能并发的线程块数越少，Occupancy 越低。

1
2
3
4
5
6
// 查询 SM 资源限制
cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, 0);
// prop.maxThreadsPerBlock     // 最大线程/块（1024）
// prop.sharedMemPerBlock      // 最大共享内存/块
// prop.regsPerBlock           // 最大寄存器/块

Grid 大小（gridDim）的原则

1. 覆盖所有数据元素

最基本的要求：Grid 中的总线程数 ≥ 问题规模，且每个线程处理一个或固定数量的元素。

1
2
3
int threads = 256;
int blocks = (N + threads - 1) / threads;  // 向上取整
kernel<<<blocks, threads>>>(data, N);

2. 让所有 SM 都有活可干

Grid 中的线程块数应至少是 GPU 中 SM 数量的若干倍，让每个 SM 都能获得线程块。若线程块数少于 SM 数，部分 SM 空闲。

1
2
3
4
cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, 0);
int smCount = prop.multiProcessorCount;
// 建议 grid 大小至少为 2~4 × smCount，保证足够的并行度

3. 使用 Occupancy API 自动确定最优配置

1
2
3
4
5
6
7
8
9
int numBlocksPerSm = 0;
int numThreads = 128;
cudaOccupancyMaxActiveBlocksPerMultiprocessor(
    &numBlocksPerSm, my_kernel, numThreads, 0);

// 启动时填满所有 SM
dim3 dimBlock(numThreads, 1, 1);
dim3 dimGrid(prop.multiProcessorCount * numBlocksPerSm, 1, 1);
cudaLaunchCooperativeKernel(...);

4. 固定线程块数 + grid-stride 循环（高级模式）

对于大规模数据，可以固定启动数量为 SM 数 × 每 SM 最大块数 的线程块，每个线程块通过循环处理多批数据。这减少了线程块启动开销，并支持更好的负载均衡：

1
2
3
4
5
6
__global__ void gridStrideKernel(float* data, int N) {
    int stride = gridDim.x * blockDim.x;
    for (int i = threadIdx.x + blockDim.x * blockIdx.x; i < N; i += stride) {
        data[i] = process(data[i]);
    }
}

多维 Grid/Block

对于 2D 图像或矩阵处理，使用 2D 配置更直观，但对性能无影响：

1
2
3
dim3 block(16, 16);                        // 16×16=256 线程/块
dim3 grid((W+15)/16, (H+15)/16);           // 覆盖整张图
imageKernel<<<grid, block>>>(img, W, H);

Q3：什么是 CUDA Stream？同步 Stream 与异步 Stream 的区别？

CUDA Stream 的定义

CUDA Stream 是一个操作队列的抽象——程序向 Stream 中提交操作（kernel 启动、内存拷贝等），Stream 按入队顺序依次执行这些操作。

  • 同一 Stream 内：操作严格按顺序执行，前一个操作完成后才执行下一个。
  • 不同 Stream 间：在硬件资源允许的情况下，可以并发执行。

所有提交到 Stream 的操作都是对 CPU 线程异步的——API 调用立即返回，CPU 不等待 GPU 完成。

默认 Stream（Default Stream）

没有指定 Stream 时，操作进入默认 Stream（也称 NULL Stream 或 Stream 0）。

1
2
3
// 这两行等价，都进入默认 Stream
kernel<<<grid, block>>>(args);
kernel<<<grid, block, 0, 0>>>(args);

阻塞 Stream vs 非阻塞 Stream

这里的"阻塞/非阻塞"是指与默认 Stream 的同步关系，而不是对 CPU 是否阻塞：

阻塞 Stream（Blocking Stream）——cudaStreamCreate 默认创建

  • 与默认 Stream 互相等待：向默认 Stream 提交操作时，必须等待所有阻塞 Stream 中的操作完成；阻塞 Stream 中的操作也必须等待默认 Stream 的操作完成。
  • 实际效果：有默认 Stream 操作时，阻塞 Stream 不能真正并发。
1
2
3
4
5
6
7
8
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);  // 阻塞 Stream
cudaStreamCreate(&stream2);  // 阻塞 Stream

kernel1<<<grid, block, 0, stream1>>>(args);
kernel2<<<grid, block>>>(args);        // 默认 Stream，等待 kernel1 完成
kernel3<<<grid, block, 0, stream2>>>(args); // 等待 kernel2 完成
// 三个 kernel 实际上串行执行

非阻塞 Stream（Non-blocking Stream）——需要显式指定标志

1
2
3
4
5
6
7
8
cudaStreamCreateWithFlags(&stream1, cudaStreamNonBlocking);
cudaStreamCreateWithFlags(&stream2, cudaStreamNonBlocking);

kernel1<<<grid, block, 0, stream1>>>(args);
kernel2<<<grid, block>>>(args);        // 默认 Stream，不等待 stream1/2
kernel3<<<grid, block, 0, stream2>>>(args);
// 三个 kernel 可能并发执行
cudaDeviceSynchronize();  // 需要显式同步

Per-thread Default Stream（每线程独立默认 Stream）

从 CUDA 7 起，通过编译选项 --default-stream per-thread，每个 CPU 线程有独立的默认 Stream，不再共享同一个默认 Stream，解决了多线程编程中默认 Stream 的同步问题。

Stream 同步 API

1
2
3
4
5
6
7
8
// 阻塞等待某个 Stream 中的所有操作完成
cudaStreamSynchronize(stream);

// 非阻塞检查 Stream 是否完成（返回 cudaSuccess 或 cudaErrorNotReady）
cudaError_t status = cudaStreamQuery(stream);

// 等待设备上所有 Stream 的所有操作完成（最重量级）
cudaDeviceSynchronize();

Q4：如何利用多 Stream 实现计算与数据传输的重叠？

现代 GPU 有独立的 DMA 引擎负责数据传输，可以与 SM 上的 kernel 执行并发进行。利用多 Stream 可以实现"计算 + 传输"的流水线，消除数据传输开销。

前提条件

  1. CPU 端内存必须是 Pinned Memory（通过 cudaMallocHost 分配），否则 cudaMemcpyAsync 会退化为同步行为。
  2. 数据传输使用 cudaMemcpyAsync（而非 cudaMemcpy）。
  3. 传输和计算必须在不同的 Stream 中。

基本流水线模式

将大数据集分成若干 chunk，在处理第 i 个 chunk 的同时传输第 i+1 个 chunk：

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
32
33
34
35
36
37
38
39
const int NUM_STREAMS = 2;
const int CHUNK = N / NUM_STREAMS;

cudaStream_t streams[NUM_STREAMS];
for (int i = 0; i < NUM_STREAMS; i++)
    cudaStreamCreate(&streams[i]);

float *h_A, *h_B, *h_C;
cudaMallocHost(&h_A, N * sizeof(float));  // Pinned Memory，必须！
cudaMallocHost(&h_B, N * sizeof(float));
cudaMallocHost(&h_C, N * sizeof(float));

float *d_A, *d_B, *d_C;
cudaMalloc(&d_A, N * sizeof(float));
cudaMalloc(&d_B, N * sizeof(float));
cudaMalloc(&d_C, N * sizeof(float));

for (int i = 0; i < NUM_STREAMS; i++) {
    int offset = i * CHUNK;

    // H→D 传输，异步
    cudaMemcpyAsync(d_A + offset, h_A + offset,
                    CHUNK * sizeof(float), cudaMemcpyHostToDevice, streams[i]);
    cudaMemcpyAsync(d_B + offset, h_B + offset,
                    CHUNK * sizeof(float), cudaMemcpyHostToDevice, streams[i]);

    // Kernel 计算，在同一 Stream 中（自动等待上面两个传输完成）
    int blocks = (CHUNK + 255) / 256;
    vecAdd<<<blocks, 256, 0, streams[i]>>>(d_A + offset, d_B + offset,
                                            d_C + offset, CHUNK);

    // D→H 传输（等 kernel 完成后执行）
    cudaMemcpyAsync(h_C + offset, d_C + offset,
                    CHUNK * sizeof(float), cudaMemcpyDeviceToHost, streams[i]);
}

// 等待所有 Stream 完成
for (int i = 0; i < NUM_STREAMS; i++)
    cudaStreamSynchronize(streams[i]);

用 Event 精确控制跨 Stream 依赖

当不同 Stream 之间存在数据依赖时，使用 Event 建立精确的同步点：

1
2
3
4
5
6
7
8
9
10
cudaEvent_t event;
cudaEventCreate(&event);

// Stream1 完成某段计算后打标记
kernel1<<<grid, block, 0, stream1>>>(output);
cudaEventRecord(event, stream1);

// Stream2 必须等 Event 触发后才能开始（Stream2 中的后续操作等待 event）
cudaStreamWaitEvent(stream2, event, 0);
kernel2<<<grid, block, 0, stream2>>>(output);  // 依赖 kernel1 的输出

cudaStreamWaitEvent 是非阻塞的 CPU 调用——它只是告诉 GPU 调度器在 event 完成前不要启动 stream2 里的后续工作，CPU 线程本身不等待。

Q5：CUDA Event 的用途是什么？如何用于计时和同步？

CUDA Event 是插入到 Stream 中的时间戳标记。当 Event 到达 Stream 队列的队首并被执行时，GPU 记录一个时间戳，并将 Event 标记为已完成（triggered）。

Event 有两大用途：精确计时和跨 Stream 同步。

用途一：精确计时

CPU 端的 clock() 或 chrono 计时会把 CPU 等待时间也算进去，无法准确反映 GPU kernel 的真实执行时间。CUDA Event 是在 GPU 端打时间戳，精度可达微秒级：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
cudaStream_t stream;
cudaStreamCreate(&stream);

cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);

// 在 kernel 前后各插入一个 Event
cudaEventRecord(start, stream);
kernel<<<grid, block, 0, stream>>>(args);
cudaEventRecord(stop, stream);

// 等待 stop Event 完成（阻塞 CPU）
cudaStreamSynchronize(stream);

// 计算耗时（单位：毫秒）
float ms;
cudaEventElapsedTime(&ms, start, stop);
printf("Kernel time: %.3f ms\n", ms);

cudaEventDestroy(start);
cudaEventDestroy(stop);
cudaStreamDestroy(stream);

用途二：跨 Stream 细粒度同步

Event 可以让一个 Stream 等待另一个 Stream 中某个特定操作完成，而不必等待整个 Stream 清空：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
cudaEvent_t event;
cudaStream_t stream1, stream2;
cudaEventCreate(&event);
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);

// stream1 依次执行 kernel1 → event 标记 → kernel2
kernel1<<<grid, block, 0, stream1>>>(output);
cudaEventRecord(event, stream1);         // kernel1 完成后打标记
kernel2<<<grid, block, 0, stream1>>>(output); // 继续执行

// stream2 等待 event（即等 kernel1 完成），不必等 kernel2
cudaStreamWaitEvent(stream2, event, 0);
dependentKernel<<<grid, block, 0, stream2>>>(output); // 依赖 kernel1 输出

// CPU 端若需要等待 kernel1 完成才能继续
cudaEventSynchronize(event);  // 阻塞 CPU 直到 event 触发
doCPUWork(output);            // 此时 kernel1 保证已完成

// 最后等待所有工作完成
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);

非阻塞查询 Event 状态

1
2
3
4
5
6
7
// 非阻塞查询，不等待 GPU
cudaError_t status = cudaEventQuery(event);
if (status == cudaSuccess) {
    // event 已触发，kernel 已完成
} else if (status == cudaErrorNotReady) {
    // 还没完成，可以继续做其他工作
}

这种轮询模式适合在等待 GPU 的同时，持续推进 CPU 端的工作（见 Q4 中的示例）。

Event 的创建与销毁

1
2
3
cudaEventCreate(&event);                                   // 默认标志
cudaEventCreateWithFlags(&event, cudaEventDisableTiming);  // 禁用计时（性能更好，用于纯同步场景）
cudaEventDestroy(event);                                   // 用完必须销毁

若只用于同步而不需要计时，使用 cudaEventDisableTiming 可以减少 GPU 端的时间戳记录开销，略微提升性能。

Q6：__syncthreads() 的作用是什么？在哪些场景下必须使用？

作用

__syncthreads() 是块内的**屏障同步（Barrier Synchronization）**原语。调用时：

  1. 块内所有线程必须先到达这个调用点，才能继续往后执行。
  2. 保证在调用之前所有线程对共享内存和全局内存的写操作，在调用之后对块内所有线程可见。

用一句话概括："所有线程都做完了前面的事，才能一起做后面的事。"

__syncthreads() 只同步同一线程块内的线程，不能同步不同线程块，也不能同步不同 SM 上的线程。CUDA 编程模型不支持跨块同步（Cooperative Groups 在一定程度上扩展了这一限制）。

必须使用的场景

场景 1：写入共享内存后供其他线程读取

这是最典型的使用场景。线程 A 写入共享内存，线程 B 读取同一块数据——必须保证 A 写完后 B 才能读：

1
2
3
4
5
6
7
8
9
10
11
__global__ void kernel(int* input, int* output) {
    __shared__ int smem[256];

    // 阶段 1：每个线程把数据写入共享内存
    smem[threadIdx.x] = input[threadIdx.x];

    __syncthreads();  // 必须：确保所有写操作完成，再进行读取

    // 阶段 2：每个线程读取其他位置的数据（依赖其他线程写入的结果）
    output[threadIdx.x] = smem[255 - threadIdx.x];
}

不加 __syncthreads()，线程读到的可能是其他线程还未写入的旧值，结果未定义。

场景 2：规约（Reduction）操作

经典的并行规约需要多轮迭代，每轮都必须等所有线程完成当前轮次的计算再进入下一轮：

1
2
3
4
5
6
7
8
9
10
11
12
13
__global__ void reduce(float* data, float* result) {
    __shared__ float smem[256];
    smem[threadIdx.x] = data[threadIdx.x + blockDim.x * blockIdx.x];
    __syncthreads();

    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
        if (threadIdx.x < stride) {
            smem[threadIdx.x] += smem[threadIdx.x + stride];
        }
        __syncthreads();  // 每轮规约后必须同步
    }
    if (threadIdx.x == 0) result[blockIdx.x] = smem[0];
}

场景 3：共享内存中的矩阵转置

先把数据从全局内存加载进共享内存（写），同步后再以不同的访问模式写出：

1
2
3
4
__shared__ float tile[32][33];
tile[threadIdx.x][threadIdx.y] = input[...];
__syncthreads();  // 等所有线程完成写入
output[...] = tile[threadIdx.y][threadIdx.x];

使用注意事项

__syncthreads() 必须被块内所有线程执行，且必须执行相同次数。 如果放在条件分支内，且不同线程可能走不同路径，就会导致死锁或未定义行为：

1
2
3
4
5
// 错误：只有部分线程调用 __syncthreads()，死锁！
if (threadIdx.x < 16) {
    smem[threadIdx.x] = compute();
    __syncthreads();  // ← 只有 0~15 号线程执行，其他线程永远不到达，死锁
}

Q7：__syncwarp() 与 __syncthreads() 的区别？

对比总结

维度 __syncthreads() __syncwarp()
同步范围 整个线程块（所有 Warp） 单个 Warp（32 个线程）
适用场景 共享内存读写屏障、块内归约 Warp 内细粒度同步、Warp 级规约
开销 较高（等待块内所有线程） 极低（只等 Warp 内线程）
引入版本 所有版本 CC 9.0 之前就有，CC 7.0 起因独立线程调度变得必要

__syncwarp() 的使用场景

__syncwarp() 的出现主要是为了应对 CC 7.0 引入独立线程调度后的代码兼容性问题。

背景： CC 7.0 之前，Warp 内线程锁步执行，写完共享/寄存器数据后，同 Warp 内其他线程立即可见。很多代码依赖这一"隐式同步"行为。CC 7.0 后，线程可以在 Warp 内独立调度，这种隐式保证不再成立，需要显式调用 __syncwarp() 恢复。

1
2
3
4
5
6
7
// Warp 级规约：只需同步 Warp 内的线程，不必同步整个块
__device__ float warpReduce(float val) {
    for (int offset = 16; offset > 0; offset >>= 1) {
        val += __shfl_down_sync(0xffffffff, val, offset);
    }
    return val;  // 线程 0 持有规约结果
}

注意：__shfl_down_sync、__shfl_xor_sync 等 Warp-level 原语自带同步语义（通过 mask 参数），在大多数现代 Warp 级规约中不再需要单独调用 __syncwarp()。

__syncwarp() 的实际使用场景：

1
2
3
4
5
6
// 场景：Warp 内一个线程写共享内存，其他线程读
if (threadIdx.x % 32 == 0) {
    warpShared[warpId] = computeResult();
}
__syncwarp();  // 确保写操作对 Warp 内所有线程可见
float result = warpShared[warpId];  // 所有线程安全读取

__syncwarp(mask) 接受一个掩码参数，指定哪些线程参与同步，支持 Warp 内部分线程的细粒度同步。若要同步整个 Warp，传入 0xffffffff。

选择建议

  • 涉及共享内存在线程块范围内的读写：用 __syncthreads()
  • 只涉及同一 Warp 内的数据交换或规约：用 __syncwarp() 或带 mask 的 Warp 级 intrinsic
  • 在 CC 7.0+ 上维护旧代码中隐式的 Warp 同步逻辑：必须补上 __syncwarp()

Q8：什么是 Dynamic Parallelism（动态并行）？

定义

CUDA 动态并行（CDP，CUDA Dynamic Parallelism） 是 CUDA 编程模型的一项特性，允许 GPU 上正在运行的 kernel（设备代码）直接启动新的子 kernel，而无需返回 CPU。

1
2
3
4
5
6
7
8
9
10
11
12
13
__global__ void childKernel(float* data, int N) {
    int i = threadIdx.x + blockDim.x * blockIdx.x;
    if (i < N) data[i] = process(data[i]);
}

__global__ void parentKernel(float* data, int N) {
    // GPU 线程直接启动子 kernel！
    if (threadIdx.x == 0) {
        int childBlocks = (N + 255) / 256;
        childKernel<<<childBlocks, 256>>>(data, N);
    }
    // 父 kernel 等待所有子 kernel 完成后才算执行完
}

解决了什么问题

在 CDP 出现之前，所有 kernel 都必须从 CPU 启动。对于数据依赖的、工作量在运行时才能确定的算法（如递归分治、自适应网格细化、不规则树遍历），必须在 CPU 端反复做"GPU 计算 → CPU 分析结果 → 再启动 GPU"的循环，CPU-GPU 之间的往返开销严重影响性能。

CDP 让工作量的动态分配决策可以在 GPU 端运行时作出，避免了 CPU-GPU 控制流切换，适合：

  • 递归算法（如快排、树遍历）
  • 不规则并行（如稀疏矩阵、自适应细化网格）
  • 工作量只能在运行时确定的场景

父子 Grid 的关系

  • 子 kernel 由父 Grid 中的线程启动，父 Grid 只有等所有子 Grid 都完成后才算结束（运行时保证隐式同步）。
  • 父 Grid 和子 Grid 共享全局内存和常量内存，但各自有独立的局部内存和共享内存。
  • 不能把父 kernel 的局部变量（栈上变量）指针传给子 kernel，必须使用全局内存分配（cudaMalloc / __device__ 全局变量）。

内存一致性

内存类型 父 Grid 可见 子 Grid 可见
全局内存 ✓ ✓（同一指针可用）
常量内存 ✓ ✓
共享内存 ✓（自己的） ✗（无法访问父的）
局部内存 ✓（自己的） ✗（禁止传递指针）

使用限制

  • 需要 CC 3.5 及以上（CDP1）或 CC 9.0 及以上（CDP2）
  • CDP2 是 CUDA 12.0 起的默认版本，仅 CC 9.0+ 支持；对 CC < 9.0 需用 CDP1（将来会移除）
  • 设备端的 CUDA Runtime API 是主机端的子集，部分 API 不可用
  • 动态并行的启动开销大于从 CPU 端启动，不适合频繁启动小 kernel 的场景

Q9：CUDA 中的原子操作（Atomic Operations）有哪些？使用时需注意什么？

什么是原子操作

原子操作（Atomic Operation）是不可被其他线程打断的读-改-写操作。GPU 上多个线程可能同时读写同一个内存地址，普通操作会产生数据竞争（Race Condition），结果不确定。原子操作保证了整个"读取旧值 → 计算 → 写回新值"的过程是原子的，不会被其他线程的操作插入。

原子函数只能在设备代码（kernel）中使用，作用于全局内存或共享内存中的 32/64/128 位数据。

完整的原子操作列表

算术类：

函数 操作 支持类型
atomicAdd(addr, val) *addr += val，返回旧值 int, uint, ulong long, float, double, half, half2 等
atomicSub(addr, val) *addr -= val，返回旧值 int, uint
atomicInc(addr, val) *addr = (*addr >= val) ? 0 : (*addr + 1)，返回旧值 uint
atomicDec(addr, val) *addr = (old==0 \|\| old>val) ? val : (old-1)，返回旧值 uint
atomicMin(addr, val) *addr = min(*addr, val)，返回旧值 int, uint, ulong long, long long
atomicMax(addr, val) *addr = max(*addr, val)，返回旧值 int, uint, ulong long, long long

位操作类：

函数 操作 支持类型
atomicAnd(addr, val) *addr &= val，返回旧值 int, uint, ulong long
atomicOr(addr, val) *addr \|= val，返回旧值 int, uint, ulong long
atomicXor(addr, val) *addr ^= val，返回旧值 int, uint, ulong long

交换类：

函数 操作 支持类型
atomicExch(addr, val) *addr = val，返回旧值 int, uint, ulong long, float
atomicCAS(addr, cmp, val) *addr = (*addr==cmp) ? val : *addr，返回旧值 int, uint, ulong long, uint short

作用域后缀

原子操作有三个作用域级别，通过后缀区分：

后缀 作用域 示例
无后缀 设备级（所有 SM 的线程可见） atomicAdd(addr, 1)
_block 块级（同一线程块内可见） atomicAdd_block(addr, 1)
_system 系统级（CPU + 所有 GPU 可见） atomicAdd_system(addr, 1)

块级原子在共享内存上性能最好；设备级原子作用于全局内存；系统级原子用于 CPU-GPU 协同场景（需要支持并发访问的统一内存）。

atomicCAS：实现自定义原子操作的万能工具

任何原子操作都可以基于 atomicCAS（Compare and Swap）实现。例如，float 类型的 atomicAdd（老版本 CUDA 不支持 float 原子加）可以这样手动实现：

1
2
3
4
5
6
7
8
9
10
__device__ float atomicAddFloat(float* addr, float val) {
    unsigned* addr_u = (unsigned*)addr;
    unsigned old = *addr_u, assumed;
    do {
        assumed = old;
        old = atomicCAS(addr_u, assumed,
                        __float_as_uint(__uint_as_float(assumed) + val));
    } while (assumed != old);
    return __uint_as_float(old);
}

CAS 循环的逻辑：读取当前值 → 计算期望的新值 → 尝试用 CAS 写入（若当前值没变则写入成功，否则重试）。

使用注意事项

1. 原子操作会序列化，是性能瓶颈

对同一地址的多个原子操作在硬件上仍然是串行的，大量线程竞争同一地址时性能很差。

经典优化：先在共享内存中做局部规约，再对全局内存做一次原子操作：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
__global__ void histogram(int* input, int* globalHist, int N) {
    __shared__ int localHist[256];

    // 初始化本地直方图
    if (threadIdx.x < 256) localHist[threadIdx.x] = 0;
    __syncthreads();

    // 对共享内存做原子加（块内串行，但比全局内存快得多）
    for (int i = threadIdx.x + blockDim.x * blockIdx.x; i < N;
             i += blockDim.x * gridDim.x) {
        atomicAdd(&localHist[input[i]], 1);
    }
    __syncthreads();

    // 只做一次全局内存原子加（每块只有 256 次，大幅减少竞争）
    if (threadIdx.x < 256) {
        atomicAdd(&globalHist[threadIdx.x], localHist[threadIdx.x]);
    }
}

2. 原子操作无返回顺序保证

多个线程对同一地址的原子操作是串行的，结果确定，但执行顺序不确定。

3. 浮点原子加的精度问题

atomicAdd 的浮点版本执行顺序不确定，而浮点加法不满足结合律，因此结果可能因运行顺序不同而有微小差异（非确定性）。对精度敏感的科学计算需注意这一点。

4. 共享内存原子操作要注意 Bank 冲突

对共享内存的原子加如果多个线程命中同一 Bank，会同时产生 Bank 冲突和原子序列化的双重开销。尽量分散访问地址（如使用线程 ID 取模分配 Bank）。