English version

February 15, 2025 · View on GitHub

[toc]

English version

I am writing a simplified English version (about half size of the Chinese version) of the book:

友情链接

关于本书第一版：

封面

注意：本书仅有 184 页，没有如下封面展示的那么厚。该封面只是编辑发给我的一个示意图，而不是实物图。要点：不要对本书的厚度（及深度）有太高的期望，否则会有点失望。

cover

购买渠道

已于2020年10月由清华大学出版社出版，语言为中文。在京东或者淘宝搜索“CUDA 编程 樊哲勇”可找到本书。

第一版勘误

本仓库的 master 分支将对应开发版本，与第一版对应的源代码见此发布版本。欢迎读者针对本书找错。找到一个其他人没有报告的错误并说服我改正者，我承诺送您此书第二版一本。目前收到的错误报告如下：

报告者错误类型页码信息更正信息
GPUSLady笔误前言“苏州吉浦讯科技有限公司”应改为“苏州吉浦迅科技有限公司”。
EverNorif笔误第 34 页“调用该核函时” 应改为 “调用该核函数时”。
Ebrece笔误第 52-53 页$ nvcc -O3 -arch=sm_75 -arithmetic1cpu.cu 应改为 $ nvcc -O3 -arch=sm_75 arithmetic1cpu.cu。类似地，$ nvcc -O3 -arch=sm_75 -arithmetic2gpu.cu 应改为 $ nvcc -O3 -arch=sm_75 arithmetic2gpu.cu。
Ebrece笔误第 135 页倒数第二行的“函数将退化同步的”应改为“函数将退化为同步的”。
我自己不安全的代码第 143 页程序中第 23 行核函数的第二个参数 size 应改为 size / sizeof(uint64_t)。我已修改本仓库中 对应的程序。
某网友笔误第 106 页“在10.1节”应改为“在第10.2节”。
静听风吟认知错误第 3 页表 1.2 中，Jetson 系列的显卡没有图灵架构的，应把 “AGX Xavier” 改成 “无”。
静听风吟认知错误第 12 章本章中关于统一内存的使用，错误地认为在使用第二代统一内存（在 Linux 系统中使用帕斯卡及以上架构的 GPU）的情况下，在调用核函数之后不需要进行主机与设备的同步即可 1）从主机访问任何统一内存数据；2）并总是得到正确的结果。以上论断中，1）是正确的，而 2）不正确，因为无论是使用第一代统一内存，还是使用第二代统一内存，在从主机访问被核函数修改的统一内存数据之前，都需要某种（显式的或隐式的）同步操作，才能避免读写竞争，从而保证结果的正确性。
Zhenkun Li错误的初始化第 10 章第 108 页尾部，语句 real v = 0; 应改为 real v = s_y[tid];
https://github.com/yuchangminghit错误的初始化第 9 章neighbor2gpu.cu 中，在调用核函数find_neighbor_atomic之前应该将数组 d_NN 的每个元素初始化为零，并去掉核函数中的初始化语句 d_NN[n1] = 0;
https://github.com/yuchangminghit错误的论述第 10 章第10.2节最后一段，从“如果想要在循环内去掉对线程号的约束”开始的论述是错误的，因为此处给的示范代码会导致非法显存访问。

目录和源代码条目

第 1 章：GPU 硬件和 CUDA 工具

本章无源代码。

第 2 章：CUDA 中的线程组织

文件知识点
hello.cpp用 C++ 写一个 Hello World 程序
hello1.cu一个正确的 C++ 程序也是一个正确的 CUDA 程序
hello2.cu写一个打印字符串的 CUDA 核函数并调用
hello3.cu使用含有多个线程的线程块
hello4.cu使用多个线程块
hello5.cu使用两维线程块

第 3 章：CUDA 程序的基本框架

文件知识点
add.cpp数组相加的 C++ 版本
add1.cu数组相加的 CUDA 版本
add2wrong.cu如果数据传输方向搞错了会怎样？
add3if.cu什么时候必须在核函数使用 if 语句？
add4device.cu定义与使用 __device__ 函数

第 4 章：CUDA 程序的错误检测

文件知识点
check1api.cu检测 CUDA 运行时 API 函数的调用
check2kernel.cu检测 CUDA 核函数的调用
memcheck.cu用 cuda-memcheck 检测内存方面的错误
error.cuh本书常用的用于检测错误的宏函数

第 5 章：获得 GPU 加速的前提

文件知识点
add1cpu.cu为 C++ 版的数组相加函数计时
add2gpu.cu为数组相加核函数计时
add3memcpy.cu如果把数据传输的时间也包含进来，还有加速吗？
arithmetic1cpu.cu提高算术强度的 C++ 函数
arithmetic2gpu.cu提高算术强度的核函数；GPU/CPU 加速比是不是很高？

第 6 章： CUDA 中的内存组织

文件知识点
static.cu如何使用静态全局内存
query.cu如何在 CUDA 程序中查询所用 GPU 的相关技术指标

第 7 章：全局内存的合理使用

文件知识点
matrix.cu合并与非合并读、写对程序性能的影响

第 8 章：共享内存的合理使用

文件知识点
reduce1cpu.cuC++ 版本的归约函数
reduce2gpu.cu仅使用全局内存和同时使用全局内存和共享内存的归约核函数
bank.cu使用共享内存实现矩阵转置并避免共享内存的 bank 冲突

第 9 章：原子函数的合理使用

文件知识点
reduce.cu在归约核函数中使用原子函数 atomicAdd
neighbor1cpu.cuCPU 版本的邻居列表构建函数
neighbor2gpu.cuGPU 版本的邻居列表构建函数，分使用和不使用原子函数的情况

第 10 章: 线程束内部函数

文件知识点
reduce.cu线程束同步函数、线程束洗牌函数以及协作组的使用
reduce1parallelism.cu提高线程利用率
reduce2static.cu利用静态全局内存加速

第 11 章： CUDA 流

文件知识点
host-kernel.cu重叠主机与设备计算
kernel-kernel.cu重叠核函数之间的计算
kernel-transfer.cu重叠核函数执行与数据传输

第 12 章：统一内存

文件知识点
add.cu使用统一内存可以简化代码
oversubscription1.cu统一内存在初始化时才被分配
oversubscription2.cu用 GPU 先访问统一内存时可以超过显存的容量
oversubscription3.cu用 CPU 先访问统一内存时不可超过主机内存容量
prefetch.cu使用 cudaMemPrefetchAsync 函数

第 13 章：分子动力学模拟（MD）

文件夹知识点
cppC++ 版本的 MD 程序
force-only仅将求力的函数移植到 CUDA
whole-code全部移植到 CUDA

第 14 章：CUDA 库

文件知识点
thrust_scan_vector.cu使用 thrust 中的设备矢量
thrust_scan_pointer.cu使用 thrust 中的设备指针
cublas_gemm.cu用 cuBLAS 实现矩阵相乘
cusolver.cu用 cuSolver 求矩阵本征值
curand_host1.cu用 cuRAND 产生均匀分布的随机数
curand_host2.cu用 cuRAND 产生高斯分布的随机数

我的部分测试结果

我的测试系统

  • Linux: 主机编译器用的 g++。
  • Windows: 仅使用命令行解释器 CMD，主机编译器用 Visual Studio 中的 cl。在用 nvcc 编译 CUDA 程序时，可能需要添加 -Xcompiler "/wd 4819" 选项消除和 unicode 有关的警告。
  • 全书代码可在 CUDA 9.0-10.2 （包含）之间的版本运行。

矢量相加 (第 5 章)

  • 数组元素个数 = 1.0e8。
  • CPU (我的笔记本) 函数的执行时间是 60 ms （单精度）和 120 ms （双精度）。
  • GPU 执行时间见下表：
V100 (S)V100 (D)2080ti (S)2080ti (D)P100 (S)P100 (D)laptop-2070 (S)laptop-2070 (D)K40 (S)K40 (D)
1.5 ms3.0 ms2.1 ms4.3 ms2.2 ms4.3 ms3.3 ms6.8 ms6.5 ms13 ms
  • 如果包含 cudaMemcpy 所花时间，GeForce RTX 2070-laptop 用时 180 ms （单精度）和 360 ms （双精度），是 CPU 版本的三倍慢！

一个高算术强度的函数（第 5 章）

  • CPU 函数（数组长度为 10410^{4}）用时 320 ms （单精度）和 450 ms （双精度）。
  • GPU 函数（数组长度为 10610^{6}）用时情况如下表：
V100 (S)V100 (D)2080ti (S)2080ti (D)laptop-2070 (S)laptop-2070 (D)
11 ms28 ms15 ms450 ms28 ms1000 ms
  • 用 GeForce RTX 2070-laptop 时核函数执行时间与数组元素个数 N 的关系见下表（单精度）：
N时间
10000.91 ms
100000.99 ms
1000003.8 ms
100000028 ms
10000000250 ms
1000000002500 ms

矩阵复制和转置 (第 7-8 章)

  • 矩阵维度为 10000 乘 10000。
  • 核函数执行时间见下表：
计算V100 (S)V100 (D)2080ti (S)2080ti (D)K40 (S)
矩阵复制1.1 ms2.0 ms1.6 ms2.9 ms
读取为合并、写入为非合并的矩阵转置4.5 ms6.2 ms5.3 ms5.4 ms12 ms
写入为合并、读取为非合并的矩阵转置1.6 ms2.2 ms2.8 ms3.7 ms23 ms
在上一个版本的基础上使用 __ldg 函数1.6 ms2.2 ms2.8 ms3.7 ms8 ms
利用共享内存转置，但有 bank 冲突1.8 ms2.6 ms3.5 ms4.3 ms
利用共享内存转置，且无 bank 冲突1.4 ms2.5 ms2.3 ms4.2 ms

数组归约（第 8-10 章）

  • 数组长度为 1.0e8，每个元素为 1.23。
  • 归约的精确结果为 123000000。
  • GPU 为笔记本版本的 GeForce RTX 2070。
  • 下面是用单精度浮点数测试的结果：
计算方法与机器计算时间结果
CPU 中循环累加100 ms33554432 （完全错误）
全局内存+线程块同步函数5.8 ms123633392 （三位正确的有效数字）
静态共享内存+线程块同步函数5.8 ms123633392 （三位正确的有效数字）
动态共享内存+线程块同步函数5.8 ms123633392 （三位正确的有效数字）
共享内存+原子函数+线程块同步函数3.8 ms123633392 （三位正确的有效数字）
共享内存+原子函数+线程束同步函数3.4 ms123633392 （三位正确的有效数字）
共享内存+原子函数+线程束洗牌函数2.8 ms123633392 （三位正确的有效数字）
共享内存+原子函数+协作组2.8 ms123633392 （三位正确的有效数字）
共享内存+协作组+两个核函数2.0 ms123000064 （七位正确的有效数字）
共享内存+协作组+两个核函数+静态全局内存1.5 ms123000064 （七位正确的有效数字）

邻居列表（第 9 章）

  • 原子数为 22464。
  • 使用单精度或双精度时，CPU 都用时约 250 毫秒。
  • GPU 测试结果见下表：
是否使用原子函数V100 (S)V100 (D)RTX 2070 (S)RTX 2070 (D)
否1.9 ms2.6 ms2.8 ms23 ms
是1.8 ms2.6 ms2.5 ms16 ms

分子动力学模拟（第 13 章）

  • 模拟体系为固态氩
  • GPU 为笔记本中的 RTX 2070，使用单精度浮点数
  • CPU 为 Intel i7-8750H 处理器

CPU 版本计算速度测试

  • 原子数 N = 10310^{3} * 4 = 4000
  • 产出步数 = 20000
  • 各个部分所花时间见下表
求力部分运动方程积分部分全部
62 s0.7 s62.7 s

force-only 版本的计算速度测试

原子数产出步数求力和数据传输的时间运动方程积分的时间全部时间整体速度
4000200005.8 s0.7 s6.5 s1.2e7 原子步每秒
32000100005.0 s2.5 s7.5 s4.3e7 原子步每秒
10800040005.4 s3.3 s8.7 s5.0e7 原子步每秒
25600020005.4 s4.6 s10 s5.1e7 原子步每秒

whole-code 版本的计算速度测试

原子数产出步数求力的时间运动方程积分的时间全部时间整体速度
4000200001.5 s0.6 s2.1 s3.8e7 原子步每秒
32000100001.6 s0.3 s1.9 s1.7e8 原子步每秒
10800040002.0 s0.4 s2.4 s1.8e8 原子步每秒
25600020002.2 s0.4 s2.6 s2.0e8 原子步每秒