先把三个概念拆开

“用了 volatile”“调用了 atomic”“在同一 stream”分别只解决部分问题，不能相互替代。正确并发程序需要同时说明参与者、作用域、顺序和对象生命周期。

Block 与 Warp 同步

__shared__ float tile[256];
tile[threadIdx.x] = input[i];
__syncthreads();
float neighbor = tile[(threadIdx.x + 1) % blockDim.x];

__syncthreads() 是 block barrier，并为此前的 shared/global 访问提供相应的内存顺序保证。它必须由 block 内所有未退出线程以一致条件到达：

// 错误：条件可能只对部分线程成立。
if (threadIdx.x < n) __syncthreads();

// 正确模式：谓词保护数据，不保护 barrier。
bool valid = threadIdx.x < n;
tile[threadIdx.x] = valid ? input[threadIdx.x] : 0.0f;
__syncthreads();
if (valid) output[threadIdx.x] = tile[threadIdx.x];

Warp 内使用 __syncwarp(mask) 或带 _sync 的 shuffle/vote API。传入 mask 的线程必须满足该原语的参与规则；分支内随手用 0xffffffff 可能把不会到达的 lane 包含进来。

Thread Scope 与内存模型

CUDA C++ 为原子、barrier 等并发原语提供分层 scope：

thread_scope_thread  ⊂  thread_scope_block  ⊂  thread_scope_device  ⊂  thread_scope_system

同步距离越远，通常成本越高。选择能覆盖通信双方的最小 scope：

• block 内生产者/消费者：block scope；
• 同一 GPU 的不同 block：device scope；
• GPU 与 CPU 或其他设备：system scope，并确认内存类型和设备属性支持所需原子性。

内存顺序可类比标准 C++ 的 relaxed/acquire/release/seq_cst，但作用域是额外维度。常见消息传递模式：生产者先写数据，再 release 发布 flag；消费者 acquire 观察 flag 后再读数据。只把 flag 设为 atomic 而不建立正确的顺序，不保证 payload 可见。

Fence 不是 Barrier

Fence 只约束调用线程的内存操作顺序，不会让其他线程自动走到同一点。多数新代码优先使用作用域和 memory order 明确的 cuda::atomic/cuda::barrier，比“裸 fence + flag”更容易审查。

原子操作

原子保证对某个地址的读改写不可被同 scope 的竞争者拆开，不保证整个算法无竞态，也不保证多个地址组成事务。

__global__ void histogram(const unsigned char* in, int n,
                          unsigned int* global_bins) {
  __shared__ unsigned int local[256];
  for (int b = threadIdx.x; b < 256; b += blockDim.x) local[b] = 0;
  __syncthreads();

  for (int i = blockIdx.x * blockDim.x + threadIdx.x;
       i < n; i += blockDim.x * gridDim.x) {
    atomicAdd(&local[in[i]], 1u);
  }
  __syncthreads();

  for (int b = threadIdx.x; b < 256; b += blockDim.x)
    if (local[b]) atomicAdd(&global_bins[b], local[b]);
}

这里把大量 global atomics 变成 shared atomics，最后每 block 每 bin 至多一次 global atomic。若输入高度偏斜，shared 中同一 bin 仍会争用，还可进一步做 warp-private bins、按 key 聚合后由 leader 原子更新等。

原子浮点加法的执行顺序通常不固定，因此结果可能 run-to-run 有细微差异。需要确定性时，应使用固定归约树、可重现库模式或更高精度累加，并接受性能成本。

Stream：设备工作队列

同一 stream 中操作按提交顺序执行；不同 stream 之间除非建立依赖，否则允许并发且顺序未定义。

stream 0: H2D(A) ─ kernel(A) ─ D2H(A)
stream 1:         H2D(B) ─ kernel(B) ─ D2H(B)

可能的时间线：
copy engine H2D: [ A ][ B ]
compute engine :     [kernel A][kernel B]
copy engine D2H:               [ A ][ B ]

是否真正重叠取决于硬件 copy engine、数据依赖、资源占用、内存是否 page-locked、传输方向和问题规模。Stream 只表达机会，不承诺并发。

Default stream 陷阱

Legacy default stream 会与其他 blocking streams 存在隐式同步；per-thread default stream 表现为每个 host thread 的普通 stream。可通过编译选项或宏选择语义，但第三方库可能有自己的 stream 约定。工程中最好显式传递 stream，减少对默认行为的依赖。

Event：轻量依赖与设备计时

cudaEvent_t ready;
cudaEventCreateWithFlags(&ready, cudaEventDisableTiming);

producer<<<grid, block, 0, stream_a>>>(...);
cudaEventRecord(ready, stream_a);
cudaStreamWaitEvent(stream_b, ready);
consumer<<<grid, block, 0, stream_b>>>(...);

这不会让 host 阻塞。只有 stream_b 中 event 之后的工作等待 producer 之前的工作。

用于计时：

cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, stream);
for (int r = 0; r < repeats; ++r) kernel<<<g, b, 0, stream>>>(...);
cudaEventRecord(stop, stream);
cudaEventSynchronize(stop);
float ms = 0.0f;
cudaEventElapsedTime(&ms, start, stop);

先预热；报告 median/percentile 或总时长除以重复数；不要在每次迭代中 cudaDeviceSynchronize()，否则测到的是强制串行和同步成本。

双缓冲数据流水线

把数据拆为 chunk，为每条 stream 准备互不重叠的 device buffer：

for (int i = 0; i < chunks; ++i) {
  int s = i % stream_count;
  size_t offset = static_cast<size_t>(i) * chunk_elems;
  size_t count = min(chunk_elems, n - offset);
  size_t bytes = count * sizeof(float);

  cudaMemcpyAsync(d_in[s], h_in + offset, bytes,
                  cudaMemcpyHostToDevice, streams[s]);
  transform<<<grid(count), block, 0, streams[s]>>>(d_in[s], d_out[s], count);
  cudaMemcpyAsync(h_out + offset, d_out[s], bytes,
                  cudaMemcpyDeviceToHost, streams[s]);
}

Host buffer 通常需 page-locked 才能实现预期的异步 DMA。Buffer 在对应 stream 完成之前不能被 host 或其他无依赖操作复用。

Stream-ordered memory allocator

同步 cudaMalloc/cudaFree 可能引入昂贵的全局协调。内存池 API 把分配/释放放进 stream 顺序：

void* ptr = nullptr;
cudaMallocAsync(&ptr, bytes, stream);
kernel<<<grid, block, 0, stream>>>(ptr);
cudaFreeAsync(ptr, stream);

ptr 只能在分配之后且释放之前、通过正确 stream/event 依赖访问。池的释放阈值、跨 stream 重用和跨设备访问可配置；性能关键代码应长期复用内存，而不是在热路径反复分配。

CUDA Graph：复用整张任务图

普通 stream 提交时，host 为每个 kernel/memcpy 分别付出 launch 设置成本。CUDA Graph 将节点和依赖先定义、实例化，再反复提交 executable graph：

定义/捕获一次： H2D → kernel A → kernel B → D2H
                     └──── kernel C ────┘
                         ↓ instantiate
重复执行：             graph launch × N

Stream capture 最容易接入已有代码：

cudaGraph_t graph;
cudaGraphExec_t exec;

cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);
kernel_a<<<ga, ba, 0, stream>>>(...);
kernel_b<<<gb, bb, 0, stream>>>(...);
cudaStreamEndCapture(stream, &graph);

cudaGraphInstantiate(&exec, graph, nullptr, nullptr, 0);
for (int i = 0; i < iterations; ++i)
  cudaGraphLaunch(exec, stream);

最适合：重复执行、拓扑较稳定、大量短 kernel、CPU launch overhead 明显的工作流。Graph 不是让单个 kernel 计算更快；它减少提交开销并给 Runtime 整体依赖信息。

注意：

• capture 期间并非所有 API 都允许；错误操作会使 capture invalidated；
• 跨 stream capture 要用 event fork/join 并回到 origin stream；
• graph 与 executable graph 的线程安全/并发约束需查 API 文档；
• 参数或拓扑变化可尝试 node update/graph update，但不是任意变化都能原地更新；
• 捕获的资源必须在 replay 期间保持有效。

Device 内异步 pipeline

Host 端 stream pipeline 与 device 端 cuda::pipeline 是两个层次：

Host 层：H2D chunk i  || kernel chunk i-1 || D2H chunk i-2
Kernel 层：load tile k+1 到 shared || compute tile k

现代架构可用 cuda::memcpy_async 配合 cuda::pipeline/cuda::barrier 构造多 stage buffer。核心状态机：

1. producer 获得空 stage；
2. 发起异步 copy；
3. commit stage；
4. consumer 等待 stage ready；
5. 计算；
6. release stage 供下一轮覆盖。

过多 stages 会消耗 shared memory 并降低 occupancy；过少则覆盖不了内存延迟。需要用 pipeline stall 和吞吐实测。

并发设计检查表

• 每个依赖由同 stream 顺序、event、graph edge、barrier 或 atomic 中的哪一个建立？
• 同步原语的 scope 是否覆盖通信双方，又没有过宽？
• 所有 barrier 参与线程是否必定到达？
• 异步操作完成前，输入、输出、参数和 graph 资源是否仍存活？
• pinned memory、stream 数、chunk 大小是否经过时间线验证？
• 是否无意使用 legacy default stream 引入全局串行？
• 原子热点是否能分层聚合或改变算法？

官方参考

• Asynchronous Execution
• CUDA Graphs
• Pipelines
• CUDA C++ Memory Model
• Stream-Ordered Memory Allocator