GPU 编程的起点不是记住某个接口，而是看清一项计算能否拆成大量相似的工作，以及这些工作如何交给设备执行。若只比较核函数的运行时间，而忽略数据准备、传输和结果验证，即使程序确实用了 GPU，也未必改善了整个任务。
本文是 GPU 编程系列的导读。能够读懂 C++ 的变量、循环和数组，就可以从这里开始；指针、编译流程和设备内存会在后续用到时补充。阅读本文不需要安装 CUDA 工具链。我们用一张小图像完成纸面上的线程映射练习，建立 CPU 与 GPU 分工、核函数和线程层级的基本模型。下一篇再编写、编译并运行第一个 CUDA C++ 程序。
1. GPU 计算与任务选择
输入、规则与输出
初学并行计算时，可以先把任务拆成三个问题：输入是什么，每个输出要怎样计算，一个输出是否要等待其他输出。例如，单通道灰度图像可以看作按行排列的一组像素值；本文使用 0 到 255 的整数表示像素亮度。阈值处理只需判断每个像素是否达到指定数值，输出仍是一组像素值。
用 CPU 编写最直接的版本时，可以通过循环依次处理各个像素。这份简单实现不仅能提供正确结果，也为 GPU 版本提供核对依据。需要交给 GPU 的，是循环中反复执行且可以拆开的计算；文件读取、界面显示等外围工作仍要由应用程序组织。
数据并行
考虑一批像素的阈值处理：像素值大于或等于 128 时输出 255，否则输出 0。如果每个输出只依赖对应的输入像素，那么各像素的处理可以独立进行。这种“同一规则作用于许多数据元素”的工作，容易表达为数据并行。
这里的“可以分别进行”描述的是任务之间没有依赖，并不表示四个像素就能体现 GPU 的性能优势。实际程序还要考虑工作数量、线程安排和数据移动。相反，若第 2 项必须用第 1 项刚算出的结果，两项就不能简单地按此方式同时执行。
CPU 和 GPU 都能完成这个任务。CPU 适合组织程序流程、处理复杂控制逻辑，也可以通过多线程或向量指令并行计算。GPU 的优势通常来自让大量工作同时处于可执行状态，以较高吞吐量处理规模足够大的并行任务。它并不保证每个单独操作都比 CPU 更快。
与阈值处理不同，若后一步必须等待前一步的结果，且难以分解出足够多的独立工作，仅增加 GPU 线程数量不会自动消除这种依赖。选择实现方式前，应先找出实际耗时的部分，并确认它能否并行。CUDA C++ 最佳实践指南也将评估现有程序与定位瓶颈放在并行化之前。
整体耗时
在常见的独立显卡配置中，一次任务可以简化为以下过程：
CPU 准备输入
↓
向 GPU 提供数据
↓
CPU 启动核函数 → GPU 计算
↓
确认完成并使用结果
这是一幅用于理解职责的流程图，不表示每一步都必须串行等待。后续会讨论异步执行、数据驻留设备和流水线。集成式设备与统一内存等环境的实际内存关系也可能不同，不能把“每次都必须复制数据”当作所有 GPU 程序的规则。
若任务确实需要输入和输出传输，可先用下面的估算拆分耗时：
总耗时 ≈ 准备 + 输入传输 + 计算
+ 输出传输 + 必要的等待
这个式子用于找出可能的开销，并非精确的性能公式。只有十个像素的练习显然不值得为了加速而搬到 GPU；大批量图像处理则可能有足够多的工作来分摊启动和传输成本。实际结论仍要通过具有代表性的输入与端到端测量获得。
| 任务特征 | 需要考虑的问题 |
|---|---|
| 大量元素执行相似操作 | 元素之间是否独立，能否同时处理？ |
| 少量输入或很短的计算 | 启动、传输和准备开销是否超过计算本身？ |
| 复杂的跨元素依赖 | 是否能重新划分任务或分阶段计算？ |
| 数据已在 GPU 上 | 能否连续完成后续操作，减少往返传输？ |
2. CUDA 编程模型
主机、设备与核函数
本系列选择 CUDA C++ 作为具体实践路线。CUDA 程序由主机端代码发起：主机通常负责准备数据、管理设备资源、启动计算和处理结果；运行在 GPU 上的函数称为核函数（kernel）。一次核函数启动会让许多线程执行同一段核函数代码，各线程通过自己的索引决定处理哪一项数据。CUDA 编程指南的编程模型以主机、设备和核函数解释了这种分工。
可以把核函数理解为“交给设备执行的计算规则”，而不是一份已经包含文件读取、结果显示和所有内存管理的完整程序。主机需要告诉设备执行哪个核函数、安排多少工作，并在使用结果前确认所需计算已经完成。主机内存与设备内存描述的是程序可访问数据的位置和方式；它们在具体硬件上如何关联，不能只凭“CPU 和 GPU 是否装在同一块板上”判断。
CUDA 是 NVIDIA 的平台与编程模型，不等同于全部 GPU 编程技术。理解线程映射、数据移动、同步和性能测量后，再比较 HIP、SYCL、Triton 等路线，会比同时学习多套工具链更容易分清共性与差异。本系列聚焦通用计算；图形渲染和着色器编程属于另一条学习路径。
线程层级
启动核函数时，线程被组织为线程块（block），线程块再组成网格（grid）。网格和线程块可以设为一维、二维或三维，以便对应数组、图像等数据形状。每个线程能取得自身在线程块中的位置、所属线程块在网格中的位置，以及线程块大小；这些信息共同决定它要处理的数据。CUDA 核函数指南说明了这些内置索引及其计算方式。
对刚接触这一模型的读者，可以先记住一条关系：一个网格包含多个线程块，一个线程块包含多个逻辑线程。本系列初期通常让一个线程负责一个数据元素，便于检查是否遗漏或重复；更复杂的任务可能让一个线程处理多个元素。线程编号是程序分工的依据，并不是 CPU 线程编号，也不表示 GPU 必须同时拥有同样多的实体运算单元。
线程块是理解协作边界的重要单位：同一块中的线程可以使用块内共享内存与同步机制；普通核函数不能假定不同块按特定顺序运行，也不能直接等待另一个块先完成。因此，初学时应优先选用各块可以独立处理的任务。硬件还会把块内线程组织成 warp 执行；占用率在第七章初步介绍，warp 行为与资源用量的深入分析留给规划中的“编译结果与硬件资源分析”专题。这里先区分清楚“数据元素”“逻辑线程”和“实际硬件执行资源”，不要将三者视为同一物件。
3. 图像阈值处理案例
初始数据与预期结果
假设有一张宽 5、高 2 的单通道图像，按行存放为十个整数像素。处理规则是：输入值 >= 128 时输出 255，否则输出 0。
表中的行号和后文的下标都从 0 开始。按行存放意味着先放完第 0 行的五个像素，再放第 1 行。因此，行号为 r、列号为 c 的像素，其线性下标是 r × 图像宽度 + c；例如第 0 行第 4 列是下标 4，第 1 行第 0 列是下标 5。
| 行 | 输入像素 | CPU 参考结果 |
|---|---|---|
| 0 | 10, 130, 128, 200, 0 |
0, 255, 255, 255, 0 |
| 1 | 255, 126, 127, 128, 100 |
255, 0, 0, 255, 0 |
CPU 参考实现只需依次访问下标 0 到 9，对每个像素应用同一规则。这个结果是后续检查 GPU 实现的依据。这里选择整数像素，使预期结果可以逐项核对；浮点计算的比较规则会在正确性专题中讨论。
线程与像素的对应关系
为了便于手算，假设每个线程块有 4 个线程。前两个块分别负责 4 个像素，剩余 2 个像素仍需要第三个块，所以处理 10 个像素需要 3 个线程块，共启动 12 个逻辑线程。对于一维网格，线程处理的线性下标可按下式计算：
像素下标 = 线程块编号 × 每块线程数 + 块内线程编号
| 线程块 | 块内线程编号 | 对应像素下标 | 处理结果 |
|---|---|---|---|
| 0 | 0、1、2、3 | 0、1、2、3 | 处理输入行 0 的前四个像素 |
| 1 | 0、1、2、3 | 4、5、6、7 | 处理输入行 0 的末位及行 1 的前三位 |
| 2 | 0、1、2、3 | 8、9、10、11 | 只处理 8、9；10、11 越界，必须跳过 |
这个例子按行将二维图像展开为一维数组，因而下标 4 和 5 分属不同行，却可以由相邻线程处理。线程应先判断下标是否小于像素总数，再读取输入和写入输出；仅在写入时检查，读取仍可能越界。
在 CUDA C++ 中，上述下标对应 blockIdx.x * blockDim.x + threadIdx.x。此处的 4 仅用于手算，不是推荐的性能配置；设备上线程块的实际取值还要考虑硬件限制与资源使用。本节展示的是映射逻辑，不是可直接编译的完整程序。下一篇会补齐环境、内存管理、核函数启动和错误检查。
结果验证
先按块列出有效下标：0–3、4–7、8–9。它们恰好覆盖十个像素，每个下标出现一次。再逐项应用阈值规则，应得到上表中的 CPU 参考结果。如果遗漏边界判断，下标 10 和 11 可能造成非法读取或写入；如果索引公式写错，也可能出现像素被重复处理或无人处理。
这个练习体现了 GPU 编程的第一个闭环：确定工作分配 → 检查边界 → 对照参考结果。运行成功与运行更快是两件事。后续的所有优化都必须保留正确性检查。
4. 系列学习顺序
系列沿同一个图像批处理项目逐步增加能力，矩阵乘法作为计算密集型的对照案例。各阶段先回答一个主要问题，再进入依赖它的专题。
| 阶段 | 主要专题 | 学习成果 |
|---|---|---|
| 入门 | 环境与首个核函数、线程映射、设备内存与数据传输 | 能编写和运行基础 CUDA 程序，验证不同尺寸输入。 |
| 核心 | 内存访问与共享内存、同步与归约、数值正确性 | 能处理数据复用、线程协作和结果检查。 |
| 优化 | 性能测量与瓶颈分析、分块矩阵乘法 | 能依据证据解释优化收益与代价。 |
| 系统 | 异步流水线、计算库与框架集成 | 能把核函数放进完整应用，测量端到端效果。 |
| 进阶 | 多 GPU、编译结果与硬件资源分析 | 能分析扩展限制及源码与资源使用的关系。 |
| 拓展 | GPU 编程生态与可移植性 | 能根据硬件与项目约束选择技术路线。 |
初学者建议顺序阅读前三个阶段。已有并行编程经验的读者可以在完成本节的线程映射练习后，按问题进入相应专题。系列大纲允许随案例验证结果调整，但每篇都会给出独立的初始条件和验收标准。
5. 常见问题
GPU 一定比 CPU 快吗
不一定。工作量、可并行程度、数据位置、传输和启动开销都会影响结果。比较时应固定输入与正确性要求，测量实际任务的端到端耗时，而非只挑选某个核函数的最短运行时间。
一个 GPU 线程等于一个 CPU 线程吗
两者都能表达并发工作，但执行模型和资源成本不同。CUDA 线程按块与网格组织，硬件进一步以 warp 等单位执行。设计程序时应依据 GPU 的线程层级和同步范围推理，不能直接套用 CPU 线程的直觉。
没有 NVIDIA GPU 能否学习
可以先完成本文的任务划分与结果验证练习，并学习并行算法、内存访问和测量方法。后续 CUDA C++ 实操需要受支持的 NVIDIA GPU、驱动与工具链；跨平台路线会在系列末尾比较。是否能在某台设备上运行具体示例，应以该设备和软件环境的实际支持情况为准。
概念速查
| 概念 | 本文中的含义 |
|---|---|
| 像素 | 本文图像中的一个整数亮度值，是练习的一个数据元素。 |
| 数据并行 | 将同一计算规则分别应用于大量可以独立处理的数据元素。 |
| 主机 | 发起计算、组织数据和处理结果的 CPU 端程序。 |
| 设备 | 执行核函数的 GPU 及其可用资源。 |
| 核函数 | 在 GPU 上由许多逻辑线程执行的计算代码。 |
| 线程块 | 一组能够在规定范围内协作的线程。 |
| 网格 | 一次核函数启动所组织的线程块集合。 |
| 参考结果 | 用于检查 GPU 计算是否正确的已知结果。 |
实践练习
- 将案例中的图像改为宽 5、高 3，并自行给第三行添加五个像素。仍按每块 4 个线程计算所需线程块数，列出每个线程块负责的下标。验收标准：有效下标恰好覆盖
0–14，无重复、无遗漏，所有越界下标均被标出。 - 比较“十个像素只处理一次”和“一批大图像连续进行多步处理”两种任务。分别指出可能占主要时间的部分，以及需要测量哪些量才能判断 GPU 是否带来收益。验收标准：至少考虑数据准备、传输、核函数和等待，不仅比较核函数时间。
参考资料
- NVIDIA CUDA Programming Guide：Programming Model
- NVIDIA CUDA Programming Guide：Writing SIMT Kernels
- NVIDIA CUDA C++ Best Practices Guide
系列导航：第一篇 · GPU 编程系列入口 · 下一篇 · 开发环境与首个核函数。