CUDA 编程模型与 Kernel 生命周期

Hostprepare + launch

async enqueue

StreamH2D → kernel → D2H

executed by

Devicegrid / blocks / threads

completion

Host boundarysync + check

1. 最小 vector-add kernel

__global__ void add(const float* a, const float* b, float* c, int n) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) c[i] = a[i] + b[i];
}

int threads = 256;
int blocks = (n + threads - 1) / threads;
add<<<blocks, threads>>>(a, b, c, n);

__global__ 表示 host launch/device execute；grid 覆盖 n elements，最后 block 可能部分越界，所以 guard 必不可少。

2. 一次完整生命周期

1. 选 device/context；
2. 分配 host/device memory；
3. 初始化并 H2D copy；
4. 配置 grid/block/shared memory/stream；
5. launch kernel；
6. 检查 launch + asynchronous execution errors；
7. 必要时同步并 D2H copy；
8. 释放资源。

framework 隐藏了多数步骤，但 custom op 仍必须服从当前 device、stream、allocator 和 tensor lifetime。

3. Runtime API 与 Driver API

Runtime API（cuda*）更高层，自动管理 primary context/编译链接；Driver API（cu*）显式 module/context/function，适合 JIT、runtime loading 与更细控制。两者能互操作，但 context ownership/version 边界要清楚。

4. Memory choices

方式	用途与边界
cudaMalloc	device allocation；同步 allocator 可有 overhead
cudaMallocAsync	stream-ordered pool allocation，便于复用
Pinned host memory	支持真正 async DMA；过量会伤系统
Unified Memory	统一地址与按需迁移；方便但 page fault/locality 要管理
Mapped/zero-copy	GPU 直接访问 host memory；离散 GPU latency/bandwidth 有限

5. Synchronization scopes

• __syncthreads()：block 内 barrier + shared/global visibility 规则；
• warp primitives：shuffle/ballot 与显式 active mask；
• stream event：跨 streams 建 dependency；
• cudaStreamSynchronize：等待单 stream；
• cudaDeviceSynchronize：粗粒度等待整个 device，调试方便、性能代价大。

6. Error handling

kernel<<<...>>>(...);
check(cudaGetLastError());       // launch/config error
check(cudaStreamSynchronize(s)); // surfaces async execution error at boundary

illegal memory access 可能在后续无关 API 才显现，且 CUDA errors 常具有 sticky/context-corrupting 特性；production 不能假设 catch exception 后继续安全复用同一进程。

7. Grid-stride loop

for (int i = blockIdx.x*blockDim.x + threadIdx.x;
     i < n;
     i += blockDim.x*gridDim.x) { ... }

它让有限 threads 覆盖任意 n，便于限制 grid、复用 threads 和调试；仍需保证 index type 足够大、访问 pattern 合并。

8. 编译链

nvcc 分离 host/device code，device side 可生成 architecture-specific cubin/SASS 和/或 virtual PTX。PTX 不是最终硬件指令；driver JIT 兼容性、目标 sm_xx、fat binary 与 toolkit/driver version 决定部署。

9. 最小 correctness checklist

• 所有 shapes、strides、dtype、device、alignment 与 bounds；
• 零长度、非整除、极大 index、non-contiguous tensors；
• 同 stream ordering 与跨 stream event；
• CPU/reference 多随机 shape 对比；
• compute-sanitizer、racecheck 与 deterministic seed；
• release build 仍保留 error recorder。

10. 四个 CUDA lifecycle 手算

自测

1. 为什么 kernel launch 后立即计 CPU 时间不对？

launch 对 host 通常异步，计到的是 enqueue 时间；需 CUDA events 或同步边界测 device execution。

2. __syncthreads 能放在部分 threads 才进入的分支吗？

只有当整个 block 对该条件一致时安全；否则部分 threads 不到 barrier 会死锁/未定义。

3. PTX 等于 GPU machine code 吗？

不等于；PTX 是 virtual ISA，最终由工具链/driver 生成目标 architecture 的 SASS。