7 月 27 日，2026 夏季 InfiniTensor 训练营 CUDA 编程方向第四课《CUDA 基础入门：向量规约》开讲。
前三课依次讲解 CPU / GPU 架构差异、SIMT 基础概念、 Roofline 性能模型、GPU 多层内存架构、Shared Memory Bank Conflict，以及原子操作、Warp / 树状规约、Cooperative Groups 等多级并行规约优化方案。

本节课从 CPU / GPU 硬件差异切入，讲解 CUDA 完整编程模型、向量基础算子与并行向量规约实现，打通 CUDA 底层硬件到并行代码落地全链路。

CPU 与 GPU 硬件核心差异

• CPU：大量缓存、控制单元，计算单元少，适配复杂分支逻辑；
• GPU 海量计算单元，缓存占比极低，面向大批量连续并行数据运算。
• 关键区别：CPU 依赖 L3 缓存复用重复数据，GPU 依靠 Warp 切换掩盖访存延迟，不采用类似 CPU 的大容量共享 L3 Cache 设计。
• Warp 由 32 个线程组成，GPU 为每个 Warp 独立分配寄存器；GPU 一个 SM 可以同时驻留多个 Warp，Warp Scheduler 会在每个周期选择满足执行条件的 Warp 发射指令。当某个 Warp 因访存等原因等待时，调度器可以切换执行其他 ready Warp，从而隐藏访存延迟。

GPU 完整存储层级与性能差距

GPU 存储层级（速度由快至慢）

1. 寄存器（RMEM）：线程私有，常见单线程上限 255 个，变量直接映射寄存器；
2. Shared Memory（SMEM）：单 Block 线程共享，可编程，存在 Bug Conflict 性能隐患；
3. 显存（GMEM）：全局大容量存储。在大模型推理尤其是 Decode 阶段，显存带宽通常是影响性能的重要指标之一。

补充：L2 Cache 硬件自动缓存，不可手动控制，极限优化场景少量利用。

异步访存硬件革新

GPU 支持显存 ↔ Shared Memory 异步传输，访存与计算可并行执行，大幅掩盖内存延迟；新一代 TMA 硬件进一步节省寄存器开销。

GPU 的计算单元

1. SM 基础组成

单 SM 包含多组 Sub Core、Warp 调度器、CUDA Core、Tensor Core，每 Block 固定调度至单一 SM 执行，不可跨 SM。

2. CUDA Core 与 Tensor Core 区分

• CUDA Core：GPU 中负责通用计算的执行单元，支持多种运算类型；
• Tensor Core 专用矩阵乘硬件，一次指令完成批量矩阵计算。

3. 线程数工程规范

单 Block 推荐 128 / 256 个线程（4 个 Warp），匹配 SM 内调度单元，充分利用硬件算力；线程过多会出现寄存器溢出，严重降低性能。

CUDA 编程模型基础

三类函数修饰符

1. __global__：CPU 调用、GPU 执行，标准 Kernel 函数，无返回值；
2. __device__：仅 GPU 内部调用，建议强制内联消除函数调用开销；
3. __host__：普通 CPU 串行函数。

Grid / Block / Thread 任务划分

• Grid：任务总容器，由若干 Block 组成；
• Block：CUDA 编程中的线程组织单位，也是资源分配和调度的基本单位，同一 Block 内线程可以共享 Shared Memory；
• Thread：最小执行单元，通过 BlockIdx、ThreadIdx 计算全局数据索引。

Kernel 完整执行流程

1. cudaMalloc 在显存分配空间；
2. cudaMemcpy 主机 ↔ 显存拷贝数据；
3. 启动 Kernel 执行并行计算；
4. 拷贝计算结果回主机，释放显存。

SIMT 架构特点

所有线程执行同一套代码，通过不同线程索引访问不同数据；同一 Warp 内分支判断会产生线程空转（Warp 发散），尽量避免复杂控制流。

Vector Add 向量加法

任务特性

向量加法属于逐元素无关计算，各线程处理独立数据，无线程间数据依赖，无同步需求。

并行划分思路

• 按线程均分向量元素，每个线程仅计算一组 X + Y = Z；
• 线程处理数据保持内存连续，优化显存读取效率。

Vector Reduction 向量规约

串行 vs 并行规约

• 串行规约单线程循环累加，速度极慢；
• 并行规约拆分多线程分段求和，再逐层聚合得到最终结果。

两层分层规约流程

1. 线程内规约：单线程循环累加分段内全部元素，结果存入寄存器；
2. Block 内树状规约：将寄存器结果写入 Shared Memory，通过 __syncthreads() 同步，对半聚合逐层合并。

规约加速

1. Warp Shuffle：__shfl_down_sync 寄存器内直接交换数据，无需 Shared Memory；
2. CUB 是 NVIDIA 提供的 CUDA 高性能并行算法库，工程开发优先复用。

加入训练营

📺完整课程内容，请观看直播或查看课程回放：

InfiniTensor 官网： https://www.infinitensor.com

B站直播： InfiniTensor 官方直播间

视频号直播： InfiniTensor 视频号预约

答疑交流： 训练营官方社群⬇️

关注与交流