CUDA C编程权威指南-第二章：CUDA编程模型

系列导航：导读 | 上一篇：第1章 基于CUDA的异构并行计算 | 下一篇：第3章 CUDA执行模型

本篇是系列第 2 篇，对应原书第 2 章。第 1 章搭好了并行与异构的宏观认知，这一章从「概念」落到「代码」：内存管理、线程组织、核函数怎么写、怎么计时和选设备，用向量加法和矩阵加法两个例子串起来。

一、本章在全书中的位置与学习目标

1.1 为什么要学「编程模型」

第 1 章回答了「为什么需要 GPU」和「CUDA 是什么」；本章回答的是：具体怎么写 CUDA 程序？

这一章里强调：编程模型（Programming Model） 提供了计算机架构的抽象，作为应用程序与底层硬件之间的桥梁。它从逻辑上给出一种特定的计算架构视图，让程序员专注于算法逻辑，而不必纠缠硬件细节。CUDA 编程模型在 GPU 架构上提供了两个核心抽象：

  • 一种通过层次结构在 GPU 中组织线程的方法（本章 + 第 3 章）
  • 一种通过层次结构在 GPU 中访问内存的方法（第 4 章 + 第 5 章）

因此，学完本章你就掌握了「如何组织线程」和「如何管理主机/设备内存」的编程层用法；至于「为什么这样组织更快」「如何访问内存更高效」，将在第 3～5 章从执行模型与内存层次展开。

1.2 学完本章，你应该能回答

学习目标 检验方式
区分主机（Host）与设备（Device），以及典型 CUDA 程序流程 能写出「六步法」并说明每步作用
使用 cudaMalloc / cudaMemcpy / cudaFree 管理设备内存与数据传输 能说明 cudaMemcpyKind 的四种方向及同步含义
理解网格（Grid）— 线程块（Block）— 线程（Thread） 两层层次 能根据数据量写出 grid 与 block 的配置
掌握 threadIdx / blockIdx / blockDim / gridDim 的含义与用法 能写出一维/二维下的全局索引公式
理解核函数 异步启动 与 cudaMemcpy 隐式同步 能解释「为何计时时要在 kernel 后加 cudaDeviceSynchronize」
会编写 global 核函数，并从 C 循环转化为「每线程做一事」 能写出向量加法的核函数并说明与串行循环的对应关系
会用错误处理宏封装 CUDA API，并会做核函数结果验证 能写出 CHECK 宏并说明核函数错误检查方式
会用 CPU 计时 与 nvprof 度量核函数时间，理解带宽/计算受限 能解释「通信计算比」对重叠计算与传输的指导意义
对矩阵会用二维/一维/混合方式组织线程，并做边界检查 能写出 ix, iy, idx 的映射及 if (ix < nx && iy < ny) 的原因
会查询 cudaDeviceProp 与 nvidia-smi，理解块/网格限制 能说出每个块最大线程数（如 1024）与网格维度上限

1.3 博客阅读导图（本章架构）

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
第 2 章  CUDA 编程模型
│
├── 二、CUDA 编程模型概述（书 2.1 节）
│   ├── 2.1.1 编程结构：Host/Device、典型流程、nvcc（书 2.1.1）
│   ├── 2.1.2 内存管理：cudaMalloc / cudaMemcpy / 全局与共享内存（书 2.1.2）
│   ├── 2.1.3 线程管理：Grid–Block–Thread、内置变量、网格/块尺寸（书 2.1.3）
│   ├── 2.1.4 启动核函数：<<<grid, block>>>、异步与同步（书 2.1.4）
│   ├── 2.1.5 编写核函数：__global__、从 C 循环到核函数、全局索引（书 2.1.5）
│   ├── 2.1.6 验证核函数（书 2.1.6）
│   ├── 2.1.7 错误处理（书 2.1.7）
│   └── 2.1.8 完整示例：GPU 向量加法（书 2.1.8）
│
├── 三、给核函数计时（书 2.2 节）
│   ├── CPU 计时器与 cudaDeviceSynchronize
│   ├── nvprof 工具
│   └── 理论峰值与通信计算比
│
├── 四、组织并行线程（书 2.3 节）
│   ├── 二维网格 + 二维块（矩阵索引与边界检查）
│   ├── 一维网格 + 一维块（每线程处理一列）
│   ├── 二维网格 + 一维块（混合）
│   └── 执行配置对性能的影响（为第 3 章铺垫）
│
├── 五、设备管理（书 2.4 节）
│   ├── cudaGetDeviceProperties 与常用属性
│   ├── 多 GPU 下选择设备与 nvidia-smi / CUDA_VISIBLE_DEVICES
│   └── 块/网格硬件限制
│
└── 六、本章小结与重难点

二、CUDA 编程模型概述（书 2.1 节）

2.1.1 编程结构：一个 CUDA 程序的生命周期（书 2.1.1）

在异构环境中，CPU 与 GPU 各自拥有独立内存，通过 PCIe 总线相连。需要区分清楚（表 2-1 对应）：

术语 含义
主机（Host） CPU 及其内存（主机内存）
设备（Device） GPU 及其内存（设备内存）

命名约定：书中示例代码用 h_ 前缀表示主机内存中的变量，用 d_ 前缀表示设备内存中的变量。坚持这一习惯有助于阅读和调试。

典型 CUDA 程序流程（建议配合原书图 2-2 理解）：

  1. 把数据从 CPU 内存拷贝到 GPU 内存（Host → Device）
  2. 调用核函数对 GPU 内存中的数据进行操作
  3. 把结果从 GPU 内存传送回 CPU 内存（Device → Host）

下面用流程图概括上述数据流与执行顺序（建议配合原书图 2-2 理解）：

flowchart LR
    subgraph host [主机 Host]
        H2D["cudaMemcpy\nHost→Device"]
        D2H["cudaMemcpy\nDevice→Host"]
    end
    subgraph device [设备 Device]
        Kernel["核函数执行"]
    end
    HostData[(主机内存)] --> H2D
    H2D --> DevData[(设备内存)]
    DevData --> Kernel
    Kernel --> DevData
    DevData --> D2H
    D2H --> HostData

两个重要特性：

  • 异步性：核函数启动后，控制权立刻返回给主机；CPU 可在 GPU 计算的同时执行其他任务，这是 CUDA 编程模型的天然优势。
  • 编译分离：NVIDIA 的 nvcc 将主机代码与设备代码分离——主机代码交标准 C 编译器，设备代码由 CUDA 编译器处理。你可将主机与设备代码放在同一源文件或不同文件中构建应用。

从 CUDA 6.0 起引入了统一寻址（Unified Memory），可用单一指针访问 CPU 与 GPU 内存，无需手动拷贝；细节在第 4 章。本章重点仍是程序员显式管理的主机/设备内存分配与拷贝，以便理解数据流并为后续优化打基础。

理解与体会：把「主机串行 + 设备并行」想成两条流水线：主机负责准备数据、发起 kernel、取回结果；设备负责大规模并行计算。理解「异步」后，后面学流（Stream）重叠计算与传输时会更自然。

2.1.2 内存管理（书 2.1.2）

GPU 拥有独立内存空间，CUDA 运行时提供与标准 C 对应的内存管理函数（原书表 2-1）：

标准 C 函数 CUDA C 函数 功能
malloc() cudaMalloc() 在设备上分配线性内存
memcpy() cudaMemcpy() 主机与设备间拷贝数据
memset() cudaMemset() 设备内存初始化
free() cudaFree() 释放设备内存

cudaMalloc — 在设备上分配一定字节的线性内存，通过指针返回：

1
cudaError_t cudaMalloc(void **devPtr, size_t size);

cudaMemcpy — 从 src 复制 count 字节到 dst，方向由 kind 指定：

1
cudaError_t cudaMemcpy(void *dst, const void *src, size_t count, cudaMemcpyKind kind);
kind 值 方向
cudaMemcpyHostToDevice 主机 → 设备
cudaMemcpyDeviceToHost 设备 → 主机
cudaMemcpyDeviceToDevice 设备 → 设备
cudaMemcpyHostToHost 主机 → 主机

这里要搞清楚：cudaMemcpy 以同步方式执行——在拷贝完成并返回之前，主机应用程序是阻塞的。而核函数启动是异步的。这一区别在调试和性能分析时非常重要。

GPU 内存结构（简化）（原书图 2-3）：主要包含两层——

  • 全局内存（Global Memory）：类似 CPU 的系统内存，所有线程可访问，容量大、延迟高。
  • 共享内存（Shared Memory）：类似 CPU 缓存，但可由程序员显式控制，容量小、速度高。

内存层次与访问模式的详细内容在第 4、5 章展开。

2.1.3 线程管理：网格、块与线程（书 2.1.3）

当核函数在主机端启动时，执行转移到设备，设备会生成大量线程，每个线程执行核函数中的相同语句。需要留意的是：了解如何组织线程是 CUDA 编程的关键部分。 CUDA 采用两层线程层次结构（原书图 2-5）：

1
2
3
4
5
6
7
网格（Grid）
 ├── 线程块（Block 0）
 │    ├── 线程 0, 1, 2, ...
 │    └── 可通过共享内存与同步协作
 ├── 线程块（Block 1）
 │    └── ...
 └── ...

核心规则（原书表述）：

  • 由一个核函数启动的所有线程统称为一个网格（Grid）；
  • 同一网格中的线程共享全局内存；
  • 一个网格由多个**线程块（Block）**组成；
  • 同一线程块内的线程可通过共享内存和同步协作；
  • 不同线程块内的线程不能协作。

内置坐标变量（书 2.1.3）

线程通过以下预初始化的内置变量区分彼此（原书表 2-2）：

变量 含义 类型
threadIdx 线程在块内的索引 uint3（.x, .y, .z）
blockIdx 线程块在网格内的索引 uint3
blockDim 线程块的维度（每块线程数） dim3
gridDim 网格的维度（网格中块数） dim3

网格和块均可组织为一维、二维或三维；未指定的维度默认为 1。

主机端 vs 设备端：主机端用 dim3 定义网格和块的维度（作为 kernel 调用的一部分）；设备端（核函数内）可访问 uint3 类型的 threadIdx、blockIdx，以及 dim3 类型的 blockDim、gridDim。尤其要注意：主机端手动定义的 dim3 仅在主机可见，设备端预初始化的变量仅在设备端可见。

确定网格和块尺寸的一般步骤（书 2.1.3）

典型调用顺序为：

  1. 确定块的大小——需考虑 GPU 资源限制与内核性能特性（第 3 章详述）；
  2. 在已知数据大小和块大小的基础上计算网格维度。

一维示例：

1
2
dim3 block(256);
dim3 grid((nElem + block.x - 1) / block.x);

向上取整保证在数据量不是块大小整数倍时仍能覆盖全部数据。

重难点：块大小不能任意取，受硬件限制（如 Fermi 上每块最多 1024 线程）；网格维度也受限制（如每维最多 65535）。先掌握「先定块、再算网格」的流程，第 3 章会解释为何某些块大小更优。

2.1.4 启动核函数（书 2.1.4）

核函数通过 <<<>>> 指定执行配置（书中对 C 函数调用语法的扩展）：

1
kernel_name<<<grid, block>>>(argument list);
  • 第一个参数：网格维度（启动的块数）；
  • 第二个参数：块维度（每块线程数）。

例如对 32 个数据元素（书中举例）：

1
2
3
4
5
6
7
8
// 4 个块，每块 8 个线程
kernel<<<4, 8>>>(args);

// 1 个块，32 个线程
kernel<<<1, 32>>>(args);

// 32 个块，每块 1 个线程
kernel<<<32, 1>>>(args);

核函数调用与主机线程异步：调用后控制权立刻返回主机。若需等待所有设备任务完成，可调用：

1
cudaError_t cudaDeviceSynchronize(void);

需要留意的是：部分 CUDA 运行时 API 是隐式同步的。例如 cudaMemcpy 在主机与设备间拷贝时，主机会阻塞直到拷贝完成；且在此拷贝开始前，之前所有已启动的核函数都需执行完毕。

2.1.5 编写核函数（书 2.1.5）

核函数用 global 修饰符声明，必须返回 void（语法规定）：

1
2
3
__global__ void kernel_name(argument list) {
    // 每个线程执行的代码
}

CUDA C 三种函数类型限定符（原书表 2-2）：

限定符 执行位置 调用方 说明
__global__ 设备 主机（或设备，compute capability ≥ 3.2） 核函数，必须返回 void
__device__ 设备 设备 仅被核函数或其他 device 函数调用
__host__ 主机 主机 默认行为，可省略

__device__ 与 __host__ 可同时使用，使同一函数在主机与设备上各有一份实现；__global__ 不能与二者组合。

从 C 循环到 CUDA 核函数（书 2.1.5）

书中用向量加法对比：主机端 C 代码是迭代 N 次的串行循环；核函数中循环消失，用线程坐标替代循环变量。

C 串行版本：

1
2
3
4
5
void sumArraysOnHost(float *A, float *B, float *C, const int N) {
    for (int i = 0; i < N; i++) {
        C[i] = A[i] + B[i];
    }
}

CUDA 核函数版本（仅当 N = 块内线程数时可直接用 threadIdx.x）：

1
2
3
4
__global__ void sumArraysOnGPU(float *A, float *B, float *C) {
    int i = threadIdx.x;
    C[i] = A[i] + B[i];
}

当使用一维网格 + 一维块的通用配置时，全局唯一索引由下式给出：

1
int i = blockIdx.x * blockDim.x + threadIdx.x;

理解与体会：可以把「核函数」理解为：把原来的循环体抽出来，让每个线程执行一次；循环变量 i 换成「当前线程的全局索引」。这个公式是后续所有一维数据映射的基础，务必熟练。

2.1.6 验证核函数（书 2.1.6）

书中建议的验证方式包括：

  1. 主机端验证函数：用 CPU 实现相同算法，逐元素比较结果（如 checkResult）。
  2. 核函数内 printf：Fermi 及以上架构支持在核函数内使用 printf 输出调试信息（需 -arch=sm_20 等）。
  3. 单线程执行：将配置设为 <<<1, 1>>>，相当于单线程串行执行，便于对比数值正确性。

2.1.7 错误处理（书 2.1.7）

由于许多 CUDA 调用是异步的，错误可能不会在发生处立即暴露。书中建议：定义一个错误处理宏，封装所有 CUDA API 调用，便于定位出错步骤与可读错误信息。示例（与书中一致）：

1
2
3
4
5
6
7
8
9
#define CHECK(call)                                                        \
{                                                                          \
    const cudaError_t error = call;                                        \
    if (error != cudaSuccess) {                                            \
        printf("Error: %s:%d, ", __FILE__, __LINE__);                       \
        printf("code: %d, reason: %s\n", error, cudaGetErrorString(error)); \
        exit(1);                                                           \
    }                                                                      \
}

使用示例：

1
2
CHECK(cudaMalloc((void **)&d_A, nBytes));
CHECK(cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice));

核函数错误的检查：核函数本身不返回错误码，需在 kernel 调用后通过同步再检查：

1
2
kernel<<<grid, block>>>(args);
CHECK(cudaDeviceSynchronize());

这一章里强调：cudaDeviceSynchronize() 会阻塞主机直到设备上所有请求完成，因此仅建议在调试时使用，以避免掩盖异步带来的性能优势。

2.1.8 完整示例：GPU 向量加法（书 2.1.8 / 代码清单 2-4）

将上述内容串联，完整流程如下（书中 sumArraysOnGPU-small-case.cu 对应）：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
#include <stdio.h>
#include <stdlib.h>

__global__ void sumArraysOnGPU(float *A, float *B, float *C) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    C[i] = A[i] + B[i];
}

int main(int argc, char **argv) {
    int nElem = 32;
    size_t nBytes = nElem * sizeof(float);

    // 1. 主机端分配并初始化
    float *h_A = (float *)malloc(nBytes);
    float *h_B = (float *)malloc(nBytes);
    float *h_C = (float *)malloc(nBytes);
    // ... 初始化 h_A, h_B ...

    // 2. 设备端分配
    float *d_A, *d_B, *d_C;
    CHECK(cudaMalloc((void **)&d_A, nBytes));
    CHECK(cudaMalloc((void **)&d_B, nBytes));
    CHECK(cudaMalloc((void **)&d_C, nBytes));

    // 3. 主机 → 设备
    CHECK(cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice));

    // 4. 启动核函数
    dim3 block(32);
    dim3 grid(nElem / block.x);
    sumArraysOnGPU<<<grid, block>>>(d_A, d_B, d_C);

    // 5. 设备 → 主机
    CHECK(cudaMemcpy(h_C, d_C, nBytes, cudaMemcpyDeviceToHost));

    // 6. 验证、释放
    cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
    free(h_A); free(h_B); free(h_C);
    cudaDeviceReset();
    return 0;
}

小结：六步法——主机分配 → 设备分配 → H2D 传输 → 核函数执行 → D2H 传输 → 释放。这是绝大多数 CUDA 程序的主干，建议熟记。

三、给核函数计时（书 2.2 节）

性能度量是优化的前提。常用有两种方式。

3.1 CPU 计时器

在主机端用系统时钟测量包含 kernel 在内的代码段。注意：必须在核函数调用后调用 cudaDeviceSynchronize()，否则测到的只是「启动 kernel」的开销，而非实际执行时间（核函数启动是异步的）。

1
2
3
4
double iStart = cpuSecond();
kernel<<<grid, block>>>(args);
cudaDeviceSynchronize();
double iElaps = cpuSecond() - iStart;

3.2 nvprof 工具

从 CUDA 5.0 起，NVIDIA 提供命令行性能分析工具 nvprof，可得到 CPU/GPU 活动时间线、各 kernel 与内存传输耗时等（书中 2.2 节）：

1
nvprof ./sumArraysOnGPU

3.3 理论峰值与通信计算比（书 2.2 节）

这一章里强调：将应用测量值与理论峰值对比很重要。例如以 Tesla K10 为例：单精度峰值性能、内存带宽峰值、以及指令与字节比（如 13.6 : 1）。若应用每访问一字节产生的指令数大于该比，则受限于计算；否则受限于内存带宽。多数 HPC 工作负载是带宽受限的。

理解与体会：若计算时间 > 数据传输时间，可通过流（Stream）重叠计算与传输，隐藏传输延迟（第 6 章）；若计算时间 < 传输时间，则应尽量减少主机–设备间的数据搬运。这就是「通信计算比」对设计的指导意义。

四、组织并行线程（书 2.3 节）

本节通过矩阵加法展示如何用不同方式组织线程，并说明执行配置对性能有显著影响（书中用不同块大小做实验）。

4.1 用块和线程建立矩阵索引（书 2.3.1）

矩阵在全局内存中以行优先线性存储。对 nx × ny 矩阵，需要建立：

  • 线程/块索引 → 矩阵坐标 (ix, iy) → 线性全局偏移 idx

二维网格、二维块时（书中公式）：

1
2
3
unsigned int ix = threadIdx.x + blockIdx.x * blockDim.x;
unsigned int iy = threadIdx.y + blockIdx.y * blockDim.y;
unsigned int idx = iy * nx + ix;

4.2 二维网格 + 二维块（书 2.3.2）

每个线程处理矩阵的一个元素，核函数内必须做边界检查：因为网格维度按「向上取整」计算，可能产生超出 nx×ny 的线程，若不判断会导致非法访问。

1
2
3
4
5
6
7
8
__global__ void sumMatrixOnGPU2D(float *A, float *B, float *C, int nx, int ny) {
    unsigned int ix = threadIdx.x + blockIdx.x * blockDim.x;
    unsigned int iy = threadIdx.y + blockIdx.y * blockDim.y;
    unsigned int idx = iy * nx + ix;
    if (ix < nx && iy < ny) {
        C[idx] = A[idx] + B[idx];
    }
}

执行配置示例：

1
2
dim3 block(32, 32);   // 每块 32×32 = 1024 线程
dim3 grid((nx + block.x - 1) / block.x, (ny + block.y - 1) / block.y);

书中实验结果：仅改变块大小，性能就有明显差异（如 32×32、32×16、16×16 等）。块数更多并不一定更快——例如 32×16 可能优于 16×16。原因与 GPU 的**线程束（warp）**调度和资源占用有关，第 3 章会详细解释。

4.3 一维网格 + 一维块（书 2.3.3）

每个线程处理一列（ny 个元素），核函数内用循环遍历行：

1
2
3
4
5
6
7
8
9
__global__ void sumMatrixOnGPU1D(float *A, float *B, float *C, int nx, int ny) {
    unsigned int ix = threadIdx.x + blockIdx.x * blockDim.x;
    if (ix < nx) {
        for (int iy = 0; iy < ny; iy++) {
            int idx = iy * nx + ix;
            C[idx] = A[idx] + B[idx];
        }
    }
}

配置为一维：dim3 block(256); dim3 grid((nx + block.x - 1) / block.x);。书中实验表明，这种「一维块」配置有时反而更快。

4.4 二维网格 + 一维块（书 2.3.4）

网格二维、块一维，每线程仍处理一个元素，其中 iy = blockIdx.y：

1
2
3
unsigned int ix = threadIdx.x + blockIdx.x * blockDim.x;
unsigned int iy = blockIdx.y;
unsigned int idx = iy * nx + ix;

4.5 线程组织的三点启示（书 2.3 节总结）

书中总结：

  1. 改变执行配置对内核性能有显著影响；
  2. 直觉上「最自然」的实现不一定最优；
  3. 对同一核函数，尝试不同网格与块大小可获更好性能。

这些现象的本质将在第 3 章（执行模型、warp、占用率）中揭示。

五、设备管理（书 2.4 节）

合理设置执行配置需要了解设备能力与限制，可以用运行时 API 或命令行工具查询。

5.1 运行时 API：cudaGetDeviceProperties（书 2.4.1）

1
2
cudaDeviceProp deviceProp;
cudaGetDeviceProperties(&deviceProp, 0);

cudaDeviceProp 常用字段包括：name、totalGlobalMem、sharedMemPerBlock、maxThreadsPerBlock（如 1024）、maxThreadsDim[3]、maxGridSize[3]、warpSize（通常 32）、multiProcessorCount、major/minor（计算能力）等。Fermi 上每块最大线程数为 1024，网格每维最大 65535，这些直接限制执行配置的设计。

5.2 多 GPU 下选择设备（书 2.4.2）

可通过遍历设备、比较 multiProcessorCount 等选择「最优」GPU，并用 cudaSetDevice() 设定当前设备。

5.3 nvidia-smi 与 CUDA_VISIBLE_DEVICES（书 2.4.3～2.4.4）

  • nvidia-smi：查看 GPU 状态、列表、详细信息（如 nvidia-smi -q -i 0）。
  • CUDA_VISIBLE_DEVICES：在不改代码的前提下指定可见 GPU（如 export CUDA_VISIBLE_DEVICES=2 仅使用 GPU 2）。

六、本章小结与重难点

6.1 知识结构回顾

模块 要点
内存管理 cudaMalloc / cudaFree；cudaMemcpy 四种方向与同步；全局内存与共享内存的抽象
线程管理 Grid → Block → Thread；threadIdx / blockIdx / blockDim / gridDim；先定块、再算网格
核函数 global / device / host；<<<grid, block>>>；异步启动；全局索引 i = blockIdx.x*blockDim.x + threadIdx.x
验证与错误 主机验证、printf、单线程运行；CHECK 宏与 cudaDeviceSynchronize 查 kernel 错误
性能 CPU 计时 + 同步；nvprof；理论峰值与带宽/计算受限、通信计算比
线程组织 二维/一维/混合；矩阵索引 ix, iy, idx；边界检查；执行配置影响性能
设备 cudaDeviceProp、maxThreadsPerBlock、warpSize；nvidia-smi、CUDA_VISIBLE_DEVICES

6.2 重难点与易错点

  1. 异步 vs 同步：kernel 启动是异步的；cudaMemcpy 是同步的，且会等待之前所有 kernel 完成。计时时若不加 cudaDeviceSynchronize()，测到的不是 kernel 真实时间。
  2. 全局索引公式：一维时 i = blockIdx.x * blockDim.x + threadIdx.x；二维矩阵时 ix, iy 再算 idx = iy * nx + ix。务必熟练，并始终做边界检查（ix < nx && iy < ny）。
  3. 块与网格限制：块最大线程数（如 1024）、网格维度上限（如 65535），设计配置时不能超出。
  4. 从循环到核函数：循环变量 → 线程全局索引；循环体 → 核函数体。核函数里不要写「循环遍历所有数据」（除非刻意让每线程处理多元素，如 2.3.3 的按列处理）。

6.3 与前后章的联系

  • 第 1 章：Host/Device、异构分工、第一个 kernel；本章在此基础上展开「完整程序」与线程组织。
  • 第 3 章：为何不同执行配置性能差异大（warp、SM、占用率、延迟隐藏）；为何 32 是「神奇的数字」。
  • 第 4～5 章：全局内存与共享内存的访问模式、对齐与合并、bank 冲突等，都建立在「网格–块–线程」与「全局索引」之上。

下一章预告

第 3 章 CUDA 执行模型 将从硬件视角回答本章留下的问题：GPU 如何调度线程、什么是线程束（Warp）、为何 32×16 的块可能优于 16×16、分支分化的危害与规避、占用率与延迟隐藏、循环展开与动态并行等。从「怎么写」到「怎么跑」——理解执行模型是写出高性能 CUDA 代码的必经之路。

本章自测

  1. 写出典型 CUDA 程序的「六步」流程（从分配主机数据到取回结果）。
  2. 为何给核函数计时时要在 kernel 调用后加 cudaDeviceSynchronize()？
  3. 一维向量加法中，若数组长度为 1000，块大小为 256，网格维度应如何取？写出全局索引公式。

答案与解析

  1. 分配主机/设备内存 → 初始化主机数据 → 将数据从主机拷到设备（H2D）→ 配置并启动核函数 → 将结果从设备拷回主机（D2H）→ 释放设备与主机内存。
  2. 核函数启动是异步的，调用立即返回，若不同步就计时会把「启动到返回」的时间算进去而非实际执行时间；cudaDeviceSynchronize() 会阻塞直到设备上所有操作完成，之后计时才准确。
  3. 网格块数 = (1000 + 256 - 1) / 256 = 4；即 gridDim.x = 4, blockDim.x = 256。全局索引公式：int i = blockIdx.x * blockDim.x + threadIdx.x，核函数内需加边界判断 if (i < n)。

系列导航：导读 | 上一篇：第1章 基于CUDA的异构并行计算 | 下一篇：第3章 CUDA执行模型

本文为「CUDA C编程权威指南」系列博客第 2 篇，共 10 章。基于《Professional CUDA C Programming》by John Cheng, Max Grossman, Ty McKercher。