先分清四个版本

nvidia-smi 显示的 “CUDA Version” 通常表示驱动最高兼容能力，不等于本机安装了同版本 nvcc。能运行 PyTorch CUDA 也不代表能编译 .cu 文件。

基础自检：

nvidia-smi
nvcc --version
cmake --version

Linux 还可检查动态库；Windows 需要受当前 Toolkit 支持的 MSVC。具体组合不要凭经验猜，直接查对应 Toolkit 的 Installation Guide 与 Release Notes。

编译链到底做了什么

                       ┌─ host C++ ── host compiler ── CPU object ─┐
hello.cu ── nvcc 拆分 ─┤                                         ├─ linker ─ executable
                       └─ device C++ ─ PTX ─ ptxas ─ cubin/fatbin ┘

运行时：fatbin 中有匹配的 SASS/cubin → 直接加载
        只有兼容 PTX                → Driver JIT 为当前 GPU 编译

• PTX 是面向虚拟 GPU 的中间指令集，便于向未来架构 JIT。
• SASS/cubin 是面向具体 sm_xx 的机器码，启动快且可针对目标架构优化。
• fatbin 可同时携带多个 cubin 和一份或多份 PTX。
• compute_xx 表示虚拟架构能力，sm_xx 表示真实架构目标。

官方 NVCC 文档建议把“最低需要的虚拟能力”与“实际部署的真实架构”分开考虑。生产发行常带目标 GPU 的 cubin，并保留一份合适的 PTX 作为前向兼容路径；但 PTX JIT 不是无条件保证，过老 Toolkit 的 PTX 可能无法利用新特性，驱动也必须能理解该 PTX 版本。

示例：

# 当前机器快速实验；生成本机可见 GPU 的机器码，不额外保留 PTX
nvcc -std=c++17 -O3 -arch=native vector_add.cu -o vector_add

# 例：为已知目标生成 cubin，并保留同一虚拟架构 PTX
nvcc -std=c++17 -O3 \
  -gencode arch=compute_90,code=sm_90 \
  -gencode arch=compute_90,code=compute_90 \
  vector_add.cu -o vector_add

不要照抄 sm_90：用 nvidia-smi --query-gpu=name,compute_cap 或 Runtime API 查询本机 Compute Capability，并按部署矩阵选择目标。

第一个完整程序

#include <cuda_runtime.h>

#include <cmath>
#include <cstdlib>
#include <iostream>
#include <vector>

#define CUDA_CHECK(call)                                                     \
  do {                                                                       \
    cudaError_t err__ = (call);                                              \
    if (err__ != cudaSuccess) {                                              \
      std::cerr << "CUDA error: " << cudaGetErrorString(err__)              \
                << " at " << __FILE__ << ':' << __LINE__ << '\n';           \
      std::exit(EXIT_FAILURE);                                               \
    }                                                                        \
  } while (0)

__global__ void vector_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 main() {
  constexpr int n = 1 << 20;
  const size_t bytes = n * sizeof(float);
  std::vector<float> a(n, 1.25f), b(n, 2.5f), c(n);

  float *d_a = nullptr, *d_b = nullptr, *d_c = nullptr;
  CUDA_CHECK(cudaMalloc(&d_a, bytes));
  CUDA_CHECK(cudaMalloc(&d_b, bytes));
  CUDA_CHECK(cudaMalloc(&d_c, bytes));
  CUDA_CHECK(cudaMemcpy(d_a, a.data(), bytes, cudaMemcpyHostToDevice));
  CUDA_CHECK(cudaMemcpy(d_b, b.data(), bytes, cudaMemcpyHostToDevice));

  constexpr int threads = 256;
  const int blocks = (n + threads - 1) / threads;
  vector_add<<<blocks, threads>>>(d_a, d_b, d_c, n);

  // 检查 launch 配置类错误；真正的执行错误常在同步点才暴露。
  CUDA_CHECK(cudaGetLastError());
  CUDA_CHECK(cudaMemcpy(c.data(), d_c, bytes, cudaMemcpyDeviceToHost));

  for (int i = 0; i < n; ++i) {
    if (std::abs(c[i] - 3.75f) > 1e-6f) return EXIT_FAILURE;
  }

  CUDA_CHECK(cudaFree(d_c));
  CUDA_CHECK(cudaFree(d_b));
  CUDA_CHECK(cudaFree(d_a));
  std::cout << "ok\n";
}

四元启动语法：

kernel<<<grid_dim, block_dim, dynamic_shared_bytes, stream>>>(args...);

前两个参数可以是 int 或 dim3；第三项只计算 extern __shared__ 的动态共享内存；第四项决定操作进入哪条 stream。

为什么错误检查必须分两层

Kernel launch 对 host 通常是异步的：

CPU:  launch ── cudaGetLastError ───────── memcpy/synchronize ──►
GPU:             [排队] [真正执行 kernel] ─────────────────────►

• cudaGetLastError() 检查非法 grid/block、参数等 launch 错误。
• cudaDeviceSynchronize()、同步 memcpy 或 event/stream 同步会暴露异步执行错误。
• 调试阶段可暂时设置 CUDA_LAUNCH_BLOCKING=1 让定位更直接，但不要拿这种模式测性能。

更稳妥的调试模板：

kernel<<<grid, block>>>(...);
CUDA_CHECK(cudaGetLastError());
#ifndef NDEBUG
CUDA_CHECK(cudaDeviceSynchronize());
#endif

CMake 最小工程

cmake_minimum_required(VERSION 3.24)
project(cuda_tutorial LANGUAGES CXX CUDA)

add_executable(vector_add vector_add.cu)
target_compile_features(vector_add PRIVATE cxx_std_17 cuda_std_17)
set_target_properties(vector_add PROPERTIES
  CUDA_SEPARABLE_COMPILATION ON
  CUDA_ARCHITECTURES native)

发布项目不应长期写 native，应显式维护部署 GPU 架构列表。跨 .cu 文件调用 device function 时需要可重定位设备代码/设备链接；CUDA_SEPARABLE_COMPILATION 会让 CMake 管理这一过程，但它可能妨碍部分跨文件优化，应按工程结构与性能测量决定。

常用编译选项

Runtime API 与 Driver API

二者可以互操作，但必须清楚当前 device、context、stream 和资源所有权。初学阶段先使用 Runtime API。

本章实验

• 编译并运行向量加法；把 n 改成非 block size 整数倍验证边界检查。
• 删除边界检查，在 Compute Sanitizer 下观察越界报告。
• 使用 -Xptxas=-v 记录寄存器数量。
• 分别生成 PTX 和 cubin，用 cuobjdump --list、cuobjdump --dump-sass 检查产物。
• 写 deviceQuery 小程序打印名称、Compute Capability、SM 数、warp size、最大 block threads、shared memory 和是否支持并发 copy/execute。

官方参考

• NVCC: The NVIDIA CUDA Compiler
• NVIDIA CUDA Compiler Driver
• CUDA Runtime API
• CUDA Samples