企业官网建设流程全解析

简介：NVIDIA官方发布的《CUDA C++编程指南》PDF文档（Release 12.8）是一份面向C++开发者的专业参考，旨在帮助利用GPU并行计算能力解决大规模计算性能瓶颈。文档系统阐述CUDA编程模型：涵盖内核定义、线程层次及线程块簇的组织方式，详细解释内存层次结构中各类存储器的用途，并说明异构编程与异步SIMT编程模型如何支持并发执行与资源管理；编程接口部分则深入介绍NVCC编译器，解析离线编译与即时编译的区别、二进制/PTX/应用/C++兼容性以及64位系统适配，同时涉及L2缓存管理等性能优化主题。资源包共包含1个PDF文件，大小约4.3MB，为官方原版英文指南，目录结构从GPU优势到高级编程接口层次分明，便于按需查阅。该资源目前已有166人学习浏览，适合想要系统学习CUDA的初学者及需要对照最新版本（12.8）进行实践的开发者。通过完整阅读，读者可以掌握从内核编写、内存访问优化到编译部署的整套方法，为开发高性能GPU加速应用打下坚实基础。

1. 把 CUDA C++ 当“扩展的 C++”来学，是很多人的第一个误区

CUDA C++ 编程指南这份官方文档，我读了三遍才敢说自己摸清了门道。它讲的是如何用 C++ 语法编写在 GPU 上运行的并行程序：Kernel 函数、线程层次、内存层次、NVCC 编译流程、运行时 API。很多人第一次上手时以为 CUDA C++ 就是“C++ 加几个关键字”，实际上它是“同一套源码，两段编译，两种执行域”。这份指南把从线程块簇到 L2 持久化缓存、从流并发到 CUDA Graphs 的整套机制都讲透了。适合正在做并行计算、图像处理、数值模拟的 C++ 开发者，能帮你搞清楚 GPU 程序为什么快、快在哪、哪些写法反而拖慢速度，省掉自己查文档踩坑的时间。

2. 编程模型与编译流程：先搞懂 Kernel 和 NVCC 的两段式编译

2.1 从矢量加法看 Kernel 的线程层次：grid、block、thread 到底怎么映射

文档里编程模型一章的开篇就是 Kernel。Kernel 是用__global__修饰的函数，调用时用<<<grid, block>>>语法指定执行配置。这份指南讲得最清楚的一点是：Kernel 不是普通的 C++ 函数，它由 GPU 上的大量线程同时执行，而线程的编排方式是三层结构——grid 包含多个 block，block 包含多个 thread。

我在模拟项目X里写过一个最简单的矢量加法，这是理解线程索引映射最快的方式：

__global__ void vecAdd(float *A, float *B, float *C, int N) { // 全局线程索引：blockIdx.x 是 block 编号，blockDim.x 是 block 内线程数 int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < N) { // 边界检查，防止越界访问 C[i] = A[i] + B[i]; } } // 调用侧：启动 N 个线程，按 256 个线程一组划分 int threadPerBlock = 256; int blockPerGrid = (N + threadPerBlock - 1) / threadPerBlock; vecAdd<<<blockPerGrid, threadPerBlock>>>(d_A, d_B, d_C, N);

逻辑说明：每个线程处理数组中的一个元素，全局索引i由blockIdx.x * blockDim.x + threadIdx.x计算，这是 CUDA 里最基础的线程到数据的映射方式。边界检查if (i < N)是必须的，因为blockPerGrid向上取整后实际启动的线程数会大于等于 N，多出来的线程不做事但要保证不越界。

参数说明：threadPerBlock取 256 是经验值，文档的性能指南章节提到 block 大小建议取 128 到 512 之间，是占用率和调度开销的折中。blockPerGrid的计算方式是(N + threadPerBlock - 1) / threadPerBlock，这是整数向上取整的标准写法。文档里特别强调线程块簇（Thread Block Clusters）是较新架构引入的概念，允许 block 之间通过分布式共享内存在簇内协作，老架构上要用这块得先查计算能力。我当时在模拟项目X上跑的时候忽略了一个细节：<<<>>>语法里其实还有个可选参数是流（stream），不写默认用的是默认流，后面讲并发时这个参数就非常重要了。

2.2 NVCC 编译流程：离线编译与 JIT 编译的边界

NVCC 的编译流程是这份指南里值得反复读的部分。它不像 g++ 一次生成可执行文件，而是先把设备端代码编译成 PTX（一种中间汇编），再通过 PTX 生成机器码。文档把这个过程分成两条路径：离线编译和即时编译。

离线编译时，NVCC 在编译阶段就把 CUDA 源码分离成主机代码和设备代码，设备代码生成 cubin 文件嵌入主机可执行文件。JIT 编译则是运行时从 PTX 生成设备机器码，好处是 PTX 具有向前兼容性——新驱动可以用旧 PTX 生成适合新架构的机器码，代价是首次加载会有额外的编译延迟。

# 两步走的编译命令：先生成 PTX，再生成 cubin nvcc -arch=compute_80 -ptx vecAdd.cu -o vecAdd.ptx nvcc -arch=compute_80 -cubin vecAdd.ptx -o vecAdd.cubin

逻辑说明：-arch=compute_80指定虚拟架构，-ptx生成中间表示，-cubin把 PTX 汇编成具体的机器码。实际项目里大多数时候不需要手动执行这两步，直接nvcc vecAdd.cu -o vecAdd会默认走完整流程。

参数说明：虚拟架构compute_XY和真实架构sm_XY的对应关系是新手最容易绕晕的。虚拟架构决定 PTX 兼容性，真实架构决定机器码优化。我一般用-arch=sm_80直接指定，这样既生成 PTX 又生成 cubin，兼顾兼容性和性能。文档里提到二进制兼容性只对同一架构族内有效，比如sm_80的 cubin 不能跑在sm_70上，但 PTX 可以。所以发布程序时保留 PTX 作为后备是很重要的习惯，否则换到新架构的机器上就直接起不来了。

2.3 计算能力（Compute Capability）与硬件特性对照

这份指南后面用了一整章讲 CUDA-Enabled GPUs，并反复提及计算能力这个数字。它本质上是一张硬件功能对照表：某个功能从哪个架构开始支持，某个限制值是多少。比如原子操作对不同数据类型的支持、线程束（warp）大小是不是固定 32、最大 block 尺寸是多少，这些都和计算能力绑定。

我在实际开发中遇到的一个典型场景：开发机器是较新的架构，部署机器是老架构，我在代码里用了__shfl_sync这个 warp 级洗牌指令，结果在老架构上编译报错。查了指南才发现这个 API 在更早的计算能力上根本不存在。从那以后我的做法是：编码之前先查目标机器的cudaDeviceProp，在代码里做版本分支。

cudaDeviceProp prop; cudaGetDeviceProperties(&prop, 0); // 查询当前设备的属性 printf("Compute Capability: %d.%d\n", prop.major, prop.minor); printf("Max threads per block: %d\n", prop.maxThreadsPerBlock); printf("Shared memory per block: %zu\n", prop.sharedMemPerBlock);

逻辑说明：cudaGetDeviceProperties可以在运行时查询设备能力。major和minor组成了计算能力编号，maxThreadsPerBlock是单个 block 的线程上限，sharedMemPerBlock是共享内存上限。这些参数决定代码怎么写，不能拍脑袋。

参数说明：打印出来对照文档里的功能支持表，就能确定当前设备能用哪些特性。比如共享内存上限，有的架构是 48KB，有的能到 163KB，如果代码里声明了过大的共享内存数组，启动时会直接报错，这个我在后面避坑章节会专门讲。

3. 运行时 API 实战：从内存管理到流并发的完整链路

3.1 设备内存与 L2 持久化缓存：把高频访问数据留在缓存里

指南在编程接口部分花了大量篇幅讲设备内存，其中 L2 缓存 Set-Aside 机制是较新版本才提到的重要内容。L2 缓存是 GPU 上的最后一级缓存，默认情况下所有内存访问都经过它，但命中率不受控制。L2 Set-Aside 允许把一部分 L2 容量专门留给持久化访问（persisting access），适合那些会被反复读取的数据。

我处理过一个图像处理 Demo，卷积核的权重数组每次都从全局内存读，虽然数据量不大，但访问频率极高。用了 L2 持久化之后，内核执行时间明显下降。

// 设置 32KB L2 为持久化缓存 cudaDeviceSetLimit(cudaLimitPersistingL2CacheSize, 32 * 1024); // 在访问权重数组前标记访问属性 cudaStreamAttrValue attr; attr.accessPolicyWindow.base_ptr = (void*)d_weights; attr.accessPolicyWindow.num_bytes = weightSize; attr.accessPolicyWindow.hitRatio = 1.0f; attr.accessPolicyWindow.hitProp = cudaAccessPropertyPersisting; attr.accessPolicyWindow.missProp = cudaAccessPropertyStreaming; cudaStreamSetAttribute(stream, cudaStreamAttributeAccessPolicyWindow, &attr);

逻辑说明：cudaDeviceSetLimit设置 L2 中持久化部分的大小，cudaStreamAttrValue定义一个访问策略窗口，把权重数组的地址范围和访问策略绑定到指定流上。hitRatio指定窗口内有多少比例的访问走持久化路径，设为 1.0 表示全部持久化。

参数说明：cudaLimitPersistingL2CacheSize的上限取决于设备，超过上限会被截断。cudaAccessPropertyPersisting标记该窗口内的数据尽量留在 L2 中，cudaAccessPropertyStreaming则相反，表示数据用过即走，避免污染缓存。我一般把高频复用的小数组设为 persistent，流式数据设为 streaming，混用效果比全 persistent 好。完全不懂 L2 机制的人听到这个会觉得很玄学，其实它就是让你手动控制缓存策略，效果立竿见影。

3.2 共享内存与同步：Block 内协作的正确姿势

共享内存在 CUDA 里是真正在硬件上存在的、block 内所有线程共享的高速存储，访问速度比全局内存快一个数量级左右。文档对共享内存的讲解包括静态分配和动态分配两种方式，以及__syncthreads()的同步语义。

静态共享内存的数组大小在编译期就要确定，这是最常见的限制。动态共享内存则是在启动内核时用第三个执行配置参数指定大小，适合数组大小依赖运行期数据的场景。

__global__ void blockReduce(const float *input, float *output, int N) { // 静态共享内存，编译期必须知道大小 __shared__ float sdata[256]; int tid = threadIdx.x; int idx = blockIdx.x * blockDim.x + tid; sdata[tid] = (idx < N) ? input[idx] : 0.0f; __syncthreads(); // 确保所有线程的写入都已完成 // 树形归约：每次步长减半 for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) { if (tid < stride) { sdata[tid] += sdata[tid + stride]; } __syncthreads(); // 每轮迭代后都要同步，否则会读到半更新状态 } if (tid == 0) output[blockIdx.x] = sdata[0]; }

逻辑说明：这是典型的 block 内归约范式。每个线程把自己的值写入共享内存，然后通过__syncthreads()保证可见性。归约循环中每次都只让一半线程工作，stride从blockDim.x/2递减到 1。每轮迭代后的__syncthreads()都必不可少，否则某个线程可能读到其他线程尚未更新的旧值。

参数说明：__shared__数组大小取 256 和 block 大小一致，方便对齐且不浪费。__syncthreads()是所有线程都必须到达的同步点，如果放在分支里会导致死锁——线程 A 在分支里等待，线程 B 没进分支直接往下走，程序直接卡死。我在这上面吃过亏，后面排查章节仔细说。

动态共享内存的写法不同：声明时用extern __shared__，启动内核时多传一个大小参数。我一般在 block 内数据量随输入变化时用动态分配，比如不同大小的矩阵分块乘法。

3.3 页锁定主机内存与映射内存：传输带宽的隐形瓶颈

指南在设备内存之后专门讲了 Page-Locked Host Memory，这是一个容易忽略但影响很大的点。普通的主机内存（malloc 分配的）是可分页的，GPU 直接访问时可能被操作系统换出到磁盘，所以cudaMemcpy需要先锁定页面再拷贝，整个过程有额外开销。

页锁定内存用cudaHostAlloc分配，持有该内存的页面不会被操作系统换页，因此主机与设备之间的拷贝可以走更高带宽的 DMA 路径。文档特别提到 Write-Combining Memory 和 Mapped Memory 两种变体，前者适合只写不回读的场景，后者允许 GPU 直接读写主机内存，省掉显式拷贝。

float *h_A, *d_A; // 分配页锁定主机内存 cudaHostAlloc((void**)&h_A, size, cudaHostAllocDefault); // 分配设备内存 cudaMalloc((void**)&d_A, size); // 页锁定内存配合异步拷贝，可以做到传输和计算重叠 cudaMemcpyAsync(d_A, h_A, size, cudaMemcpyHostToDevice, stream);

逻辑说明：cudaHostAlloc替代日常的malloc分配主机内存，配合cudaMemcpyAsync使用才能真正发挥异步传输的优势。普通内存的拷贝是同步的，拷贝期间主机线程会被阻塞，而页锁定内存配合流时不需要阻塞。

参数说明：cudaHostAllocDefault是默认模式，cudaHostAllocWriteCombining是写合并内存，适合主机只写、设备只读的场景，比如把参数一次性传给 GPU。cudaHostAllocMapped则把主机内存映射到设备地址空间，配合 UVA（统一虚拟地址）直接访问。我一般只在频繁传输的小块数据上用页锁定，比如每帧更新的参数缓冲区，大块数据还是走普通拷贝，因为页锁定内存太多会挤占系统可用内存，反而拖慢整个系统。文档里建议页锁定内存的分配量要克制，这是个很容易被忽视的系统级问题。

3.4 流与事件：用重叠隐藏传输延迟

文档里有一节专门讲多流并发执行，这是提升整体吞吐的关键手段。CUDA 的操作默认在默认流上执行，默认流和其他流之间的行为比较特殊——如果用了非默认流，需要特别注意默认流是否被禁用。指南提到的核心思路是：不同流上的操作可以重叠执行，包括数据拷贝和设备内核计算。

cudaStream_t stream1, stream2; cudaStreamCreate(&stream1); cudaStreamCreate(&stream2); // 把数据拷入设备后立刻做计算，同时主机继续准备下一块数据 cudaMemcpyAsync(d_A, h_A, size, cudaMemcpyHostToDevice, stream1); cudaMemcpyAsync(d_B, h_B, size, cudaMemcpyHostToDevice, stream2); kernel_A<<<grid, block, 0, stream1>>>(d_A); kernel_B<<<grid, block, 0, stream2>>>(d_B); cudaStreamDestroy(stream1); cudaStreamDestroy(stream2);

逻辑说明：cudaMemcpyAsync的最后一个参数指定流，同一个流内的操作串行执行，不同流之间可以并行。示例中两个流分别做拷贝和计算，两个内核可以同时跑，前提是它们访问的数据互不依赖，GPU 有足够的资源容纳两个内核。

参数说明：cudaMemcpyAsync要求主机内存是页锁定的，否则调用会退化为同步行为，这是文档明确指出的坑。kernel<<<grid, block, 0, stream>>>的第三个参数是动态共享内存大小，这里为 0，第四个参数指定流。

流的数量不是越多越好。我在某跨平台系统开发时曾经开了 8 个流想提速，结果发现内核切得太碎，调度开销比收益还大，后来改为 2 个流配合 CUDA Events 做时序控制，效果反而更好。指南里也提到流太多会导致无法充分利用硬件资源，合理数量是让 GPU 的各个计算单元都能吃满。

4. 避坑与排查：五个翻车点，每条都是实际踩过的

4.1 内核启动报错却在主机代码里看不见：异步调用的错误码陷阱

现象：kernel<<<grid, block>>>(...)启动后紧接着的cudaMemcpy返回cudaErrorInvalidValue，但内核代码我检查过逻辑没问题。

原因：CUDA 的内核启动是异步的，错误不会在启动时立刻抛出，而是排队到下一个同步操作才返回。<<<>>>本身不返回错误码，如果 grid 或 block 维度超过设备上限，错误会延迟到后续 API 调用才暴露。

解决：每次内核启动后紧跟一行cudaGetLastError()检查，这是我在所有 CUDA 代码里的固定动作。配合cudaPeekAtLastError()可以在不改变错误状态的情况下查看。

4.2 共享内存声明过大导致内核启动失败

现象：把__shared__数组改大后，内核启动返回cudaErrorInvalidValue，把数组缩小就正常。

原因：静态共享内存的大小在编译期已经决定，超过了设备的sharedMemPerBlock上限。文档里各架构的共享内存上限不同，还有需要动态配置才能使用的额外共享内存。

解决：查询cudaDeviceProp.sharedMemPerBlock和cudaDeviceProp.sharedMemPerBlockOptin。需要超过默认值时，用cudaFuncSetAttribute指定内核函数使用额外的共享内存，且声明动态共享内存并用extern __shared__。

4.3 旧架构上运行时报“不支持的指令”：PTX 兼容不等于机器码兼容

现象：在一台较新的架构上编译生成的程序，拷到老架构机器上运行，直接报错说某个指令不支持。

原因：编译时如果用-arch=compute_XY -code=sm_XY只生成特定架构的 cubin，而没有保留 PTX，驱动在目标架构不匹配时就没有后备方案。二进制兼容性只在相同计算能力族内有效。

解决：编译时用-arch=sm_XX让 NVCC 同时生成 PTX 和 cubin，或者在发布时跟随 cubin 一起附上 PTX。我在某图像处理 Demo 里遇到过同样的翻车，加一行编译参数就好了。

4.4 页锁定内存没传够，异步拷贝偷偷变成同步

现象：用了cudaMemcpyAsync以为传输能和计算重叠，结果 profiler 显示拷贝时间还是阻塞的，整体耗时没变化。

原因：cudaMemcpyAsync只有在主机内存是页锁定内存时才是真异步。如果传入的内存来自普通的malloc，驱动要先做一次隐式的页锁定，这个过程会阻塞主机线程。

解决：主机侧缓冲区用cudaHostAlloc分配，或者用cudaHostRegister把已有的malloc内存注册为页锁定。注意cudaHostRegister有额外开销，适合复用缓冲区的场景，不适合一次性的小数据。

4.5 多流并发不生效：默认流在中间捣乱

现象：写了两个流分别跑内核，用 profiler 看时间线发现它们还是串行的，没有重叠。

原因：代码里调用了某些默认流的操作，它和所有非默认流之间存在隐式同步。文档里提到默认流是同步流，除非编译时开启--default-stream per-thread，否则在默认流上的操作会阻塞其他流。

解决：要么所有流操作都显式指定非默认流，避免碰默认流；要么用nvcc --default-stream per-thread编译，让默认流变成每线程的独立流。我在模拟项目X的并发重构里就是被默认流卡了半天，用 per-thread 编译选项后并发立刻生效。

5. 性能优化指南：从占用率到内存吞吐的量化调优

5.1 Occupancy 计算：为什么 block 调到 256 性能反而最好

指南的性能优化章节一上来就讲总体优化策略，核心是最大化利用率（Maximize Utilization）。占用率（Occupancy）是衡量 GPU 计算单元被填满程度的指标：每个流多处理器能同时驻留的活动 block 数和线程数的比例。文档强调占用率和性能不是简单正比关系，约高不一定越快，但占用率太低的程序几乎不可能跑满。

// 用 occupancy calculator 的运行时版本检查内核占用率 cudaOccupancyMaxActiveBlocksPerMultiprocessor(&numBlocks, kernel, 256, 0); cudaOccupancyAvailableDynamicSMemPerBlock(&dynSmemSize, kernel, numBlocks);

逻辑说明：cudaOccupancyMaxActiveBlocksPerMultiprocessor查询给定 block 大小下每个流多处理器最多能驻留多少个 block，cudaOccupancyAvailableDynamicSMemPerBlock查询在这个占用率下还能用多少动态共享内存。这是不用打开计算器页面就能在代码里做的占用率检查。

参数说明：第一个函数的参数中，256 是 block 线程数，0 是动态共享内存大小。占用率受寄存器数量和共享内存双重限制，寄存器用多了也会降低占用率。我在调一个矩阵乘法的内核时，发现一个线程用 32 个寄存器和 40 个寄存器的占用率完全不同，后者直接少了一个 block 的驻留量。文档里给的指导是：限制内核的寄存器使用量，或者用__launch_bounds__提示编译器最大线程数和最小 block 数，让它自动分配寄存器。

5.2 全局内存访问：合并访问是带宽的生命线

指南在 Maximize Memory Throughput 一节反复强调合并访问（coalesced access）：一个 warp 的 32 个线程访问全局内存时，连续地址的访问能合并成少数几个内存事务，不连续的访问会拆成多个事务，带宽利用率直线下降。

我写的矢量加法是天然合并的，因为线程i访问C[i]，相邻线程访问相邻地址。但转置矩阵这类操作就不合并了，这时候要用共享内存做分块转置：先把不连续的数据读入共享内存，再从共享内存中按连续顺序写出。指南对这类模式的讲解非常细，包括行主序和列主序的影响。

__global__ void transpose(float *in, float *out, int width) { __shared__ float tile[TILE][TILE]; int x = blockIdx.x * TILE + threadIdx.x; int y = blockIdx.y * TILE + threadIdx.y; // 合并读入共享内存 tile[threadIdx.y][threadIdx.x] = in[y * width + x]; __syncthreads(); // 从共享内存转置写出，注意下标交换 out[x * width + y] = tile[threadIdx.x][threadIdx.y]; }

逻辑说明：分块转置的核心是让共享内存作为中间缓冲，读入和写出都做到合并访问。读入时线程按行连续访问，写出时目标地址实际是列方向，但由于共享内存的存在，写出的out[x * width + y]中线程访问的地址仍然是线性的。

参数说明：TILE一般取 16 或 32，太大的话共享内存不够用，太小则无法有效合并访问。转置是著名的非合并访问问题，不少框架里都有类似的分块优化。

文档还提到 L2 缓存的访问策略对性能的影响，以及读写不一致时用__ldg走只读数据路径，这些都属于内存吞吐优化的细节。我在实际项目中用一个 profiling 工具对比过优化前后的内核时间，合并访问优化带来的提升常常是数量级的，比调任何计算指令都快。

5.3 指令吞吐：小心隐式整数除法和双精度

指南在 Maximize Instruction Throughput 一节讲算术指令的代价比较。整数除法是最昂贵的操作之一，即使除数是常量，也会被展开成一系列乘法和移位指令。双精度浮点的吞吐通常比单精度低很多，除非硬件专门优化过。

// 避免整数除法：用位运算替代 2 的幂除法 int a = i / 4; // 除法，较慢 int b = i >> 2; // 右移，快得多 // 双精度开方 vs 单精度 double d = sqrt((double)x); // 双精度路径 float f = sqrtf(x); // 单精度路径，吞吐更高

逻辑说明：除数和被除数如果有一个不是编译期常量，编译器就没法优化成乘法近似，指令数会直线上升。代码里的循环都有个i / 4这样的表达式，改成移位后整个内核快了一截。

参数说明：sqrtf对应单精度版本，sqrt是双精度版本。指南里有一张表列出各算术运算的相对吞吐，单精度加法和乘法是最高吞吐的，整数除法和双精度运算都在低吞吐一档。在不需要高精度的场景，用float替代double、用查表替代三角函数，是性价比很高的优化手段。

5.4 控制流与线程发散：warp 里的分支是双刃剑

指南讲控制流指令时提到了线程发散（divergence）：一个 warp 里 32 个线程执行同一指令，如果出现了分支，硬件会把分支两侧都串行执行一遍，不执行的线程被掩蔽掉。所以if分支两侧的执行时间会叠加，而不是并行。

// 发散写法：相邻线程走不同分支，整个 warp 两边都执行 if (threadIdx.x % 2 == 0) { result = processEven(x); } else { result = processOdd(x); } // 更好的写法：按 warp 对齐的场景分支，让整个 warp 走同一边 if (threadIdx.x < 16) { result = processFirstHalf(x); } else { result = processSecondHalf(x); }

逻辑说明：第一种写法里threadIdx.x % 2会让相邻线程进不同分支，同一个 warp 内 32 个线程一半走 A 一半走 B，两个分支串行执行。第二种写法按连续区间划分，warp 更可能整体走同一分支，发散程度大幅降低。

参数说明：写代码时把分支条件设计成基于 warp 边界对齐的方式，比如按threadIdx.x / 32分段，就能显著减少发散代价。这个优化不是让代码更快，而是避免让 GPU 做无用功。

6. CUDA Graphs：把内核启动改成图执行，减少 launch 开销

CUDA Graphs 是文档里比较靠后的高级特性，但它对短内核、频繁启动的程序的提升非常明显。内核启动本身有 CPU 侧的开销，每次<<<>>>都要做参数解析和驱动调用，当内核执行时间只有几十微秒时，启动开销可能占掉一半性能。CUDA Graphs 把一系列内核启动和依赖关系捕获成一张图，然后一次性提交执行，把启动开销摊薄。

// 用流捕获构建一个 CUDA Graph cudaGraph_t graph; cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal); kernelA<<<grid, block, 0, stream>>>(d_A); kernelB<<<grid, block, 0, stream>>>(d_B); // B 依赖 A 的结果 cudaStreamEndCapture(stream, &graph); // 实例化并执行 cudaGraphExec_t graphExec; cudaGraphInstantiate(&graphExec, graph, NULL, NULL, 0); cudaGraphLaunch(graphExec, stream); cudaGraphExecDestroy(graphExec); cudaGraphDestroy(graph);

逻辑说明：cudaStreamBeginCapture开始捕获指定流上的操作，之后在该流上启动的 kernel 会被记录成图节点而不是真实执行。kernelB因为和kernelA在同一个流上，被自动记录为依赖关系。cudaGraphInstantiate把图编译成可执行的句柄，cudaGraphLaunch执行整个图。

参数说明：cudaStreamCaptureModeGlobal是捕获模式，表示允许跨流依赖。cudaGraphInstantiate的第四个参数可以传错误节点用于调试。图执行的一大好处是 CPU 侧的启动开销只发生一次，之后每次执行都可以复用graphExec，这个句柄在 GPU 侧已经做完了所有调度准备工作。

我在某跨平台系统上试过用 CUDA Graphs 优化一个每帧需要启动 30 多个短内核的渲染管线，常见做法是先把整条管线捕获成图，帧循环里只调cudaGraphLaunch，帧耗时从 8 毫秒降到了 5 毫秒左右。不过图也不是银弹：如果内核启动参数每帧都变，就得重新捕获或部分更新，反而更慢；如果图里的操作之间有数据依赖，必须保证依赖通过实际的 API 调用表达，不能靠时机巧合。从那次以后，我每次写高频小内核都会先问一句：能不能把这段录成图？希望这个思路帮到你，少走我走过的弯路。

本文还有配套的精品资源，点击获取