第一部分 理解CPU的并行性 2
第1章 CPU并行编程概述 2
1.1 并行编程的演化 2
1.2 核心越多，并行性越高 3
1.3 核心与线程 4
1.3.1 并行化更多的是线程还是核心 5
1.3.2 核心资源共享的影响 6
1.3.3 内存资源共享的影响 6
1.4 第一个串行程序 7
1.4.1 理解数据传输速度 8
1.4.2 imflip.c中的main（）函数 9
1.4.3 垂直翻转行：FlipImageV（） 10
1.4.4 水平翻转列：FlipImageH（） 11
1.5 程序的编辑、编译、运行 12
1.5.1 选择编辑器和编译器 12
1.5.2 在Windows 7、8、10平台上开发 12
1.5.3 在Mac平台上开发 14
1.5.4 在Unix平台上开发 14
1.6 Unix速成 15
1.6.1 与目录相关的Unix命令 15
1.6.2 与文件相关的Unix命令 16
1.7 调试程序 19
1.7.1 gdb 19
1.7.2 古典调试方法 20
1.7.3 valgrind 22
1.8 第一个串行程序的性能 22
1.8.1 可以估计执行时间吗 23
1.8.2 代码执行时OS在做什么 23
1.8.3 如何并行化 24
1.8.4 关于资源的思考 25
第2章 开发第一个CPU并行程序 26
2.1 第一个并行程序 26
2.1.1 imflipP.c中的main（）函数 27
2.1.2 运行时间 28
2.1.3 imflipP.c中main（）函数代码的划分 28
2.1.4 线程初始化 30
2.1.5 创建线程 31
2.1.6 线程启动／执行 32
2.1.7 线程终止（合并） 33
2.1.8 线程任务和数据划分 34
2.2 位图文件 35
2.2.1 BMP是一种无损／不压缩的文件格式 35
2.2.2 BMP图像文件格式 36
2.2.3 头文件ImageStuff.h 37
2.2.4 ImageStuff.c中的图像操作函数 38
2.3 执行线程任务 40
2.3.1 启动线程 41
2.3.2 多线程垂直翻转函数MTFlipV（） 43
2.3.3 FlipImageV（）和MTFlipV（）的比较 46
2.3.4 多线程水平翻转函数MTFlipH（） 47
2.4 多线程代码的测试／计时 49
第3章 改进第一个CPU并行程序 51
3.1 程序员对性能的影响 51
3.2 CPU对性能的影响 52
3.2.1 按序核心与乱序核心 53
3.2.2 瘦线程与胖线程 55
3.3 imflipP的性能 55
3.4 操作系统对性能的影响 56
3.4.1 创建线程 57
3.4.2 线程启动和执行 57
3.4.3 线程状态 58
3.4.4 将软件线程映射到硬件线程 59
3.4.5 程序性能与启动的线程 60
3.5 改进imflipP 61
3.5.1 分析MTFlipH（）中的内存访问模式 62
3.5.2 MTFlipH（）的多线程内存访问 63
3.5.3 DRAM访问的规则 64
3.6 imflipPM：遵循DRAM的规则 65
3.6.1 imflipP的混乱内存访问模式 65
3.6.2 改进imflipP的内存访问模式 65
3.6.3 MTFlipHM（）：内存友好的MTFlipH（） 66
3.6.4 MTFlipVM（）：内存友好的MTFlipV（） 69
3.7 imflipPM.C的性能 69
3.7.1 imflipP.c和imflipPM.c的性能比较 70
3.7.2 速度提升：MTFlipV（）与MTFlipVM（） 71
3.7.3 速度提升：MTFlipH（）与MTFlipHM（） 71
3.7.4 理解加速：MTFlipH（）与MTFlipHM（） 71
3.8 进程内存映像 72
3.9 英特尔MIC架构：Xeon Phi 74
3.10 GPU是怎样的 75
3.11 本章小结 76
第4章 理解核心和内存 77
4.1 曾经的英特尔 77
4.2 CPU和内存制造商 78
4.3 动态存储器与静态存储器 79
4.3.1 静态随机存取存储器（SRAM） 79
4.3.2 动态随机存取存储器（DRAM） 79
4.3.3 DRAM接口标准 79
4.3.4 DRAM对程序性能的影响 80
4.3.5 SRAM对程序性能的影响 81
4.4 图像旋转程序：imrotate.c 81
4.4.1 imrotate.c的说明 82
4.4.2 imrotate.c：参数限制和简化 82
4.4.3 imrotate.c：实现原理 83
4.5 imrotate的性能 87
4.5.1 线程效率的定性分析 87
4.5.2 定量分析：定义线程效率 87
4.6 计算机的体系结构 89
4.6.1 核心、L1$和L2$ 89
4.6.2 核心内部资源 90
4.6.3 共享L3高速缓存（L3$） 91
4.6.4 内存控制器 92
4.6.5 主存 92
4.6.6 队列、非核心和I/O 93
4.7 imrotateMC：让imrotate更高效 94
4.7.1 Rotate2（）：平方根和浮点除法有多差 96
4.7.2 Rotate3（）和Rotate4（）：sin（）和cos（）有多差 97
4.7.3 Rotate5（）：整数除法／乘法有多差 98
4.7.4 Rotate6（）：合并计算 100
4.7.5 Rotate7（）：合并更多计算 100
4.7.6 imrotateMC的总体性能 101
4.8 本章小结 103
第5章 线程管理和同步 104
5.1 边缘检测程序：imedge.c 104
5.1.1 imedge.c的说明 105
5.1.2 imedge.c：参数限制和简化 106
5.1.3 imedge.c：实现原理 106
5.2 imedge.c：实现 108
5.2.1 初始化和时间戳 109
5.2.2 不同图像表示的初始化函数 110
5.2.3 启动和终止线程 111
5.2.4 高斯滤波 112
5.2.5 Sobel 113
5.2.6 阈值过滤 114
5.3 imedge的性能 115
5.4 imedgeMC：让imedge更高效 116
5.4.1 利用预计算降低带宽 116
5.4.2 存储预计算的像素值 117
5.4.3 预计算像素值 118
5.4.4 读取图像并预计算像素值 119
5.4.5 PrGaussianFilter 120
5.4.6 PrSobel 121
5.4.7 PrThreshold 121
5.5 imedgeMC的性能 122
5.6 imedgeMCT：高效的线程同步 123
5.6.1 屏障同步 124
5.6.2 用于数据共享的MUTEX结构 125
5.7 imedgeMCT：实现 127
5.7.1 使用MUTEX：读取图像、预计算 128
5.7.2 一次预计算一行 130
5.8 imedgeMCT的性能 131
第二部分 基于CUDA的GPU编程 134
第6章 GPU并行性和CUDA概述 134
6.1 曾经的Nvidia 134
6.1.1 GPU的诞生 134
6.1.2 早期的GPU架构 136
6.1.3 GPGPU的诞生 137
6.1.4 Nvidia、ATI Technologies和Intel 138
6.2 统一计算设备架构 140
6.2.1 CUDA、OpenCL和其他GPU语言 140
6.2.2 设备端与主机端代码 140
6.3 理解GPU并行 141
6.3.1 GPU如何实现高性能 142
6.3.2 CPU与GPU架构的差异 143
6.4 图像翻转的CUDA版：imflipG.cu 144
6.4.1 imflipG.cu：将图像读入CPU端数组 146
6.4.2 初始化和查询GPU 147
6.4.3 GPU端的时间戳 148
6.4.4 GPU端内存分配 152
6.4.5 GPU驱动程序和Nvidia运行时引擎 153
6.4.6 CPU到GPU的数据传输 153
6.4.7 用封装函数进行错误报告 154
6.4.8 GPU核函数的执行 155
6.4.9 完成GPU核函数的执行 157
6.4.10 将GPU结果传回CPU 158
6.4.11 完成时间戳 158
6.4.12 输出结果和清理 158
6.4.13 读取和输出BMP文件 159
6.4.14 Vflip（）：垂直翻转的GPU核函数 160
6.4.15 什么是线程ID、块ID和块维度 163
6.4.16 Hflip（）：水平翻转的GPU核函数 165
6.4.17 硬件参数：threadIDx.x、blockIdx.x和blockDim.x 165
6.4.18 PixCopy（）：复制图像的GPU核函数 165
6.4.19 CUDA关键字 166
6.5 Windows中的CUDA程序开发 167
6.5.1 安装MS Visual Studio 2015和CUDA Toolkit 8.0 167
6.5.2 在Visual Studio 2015中创建项目imflipG.cu 168
6.5.3 在Visual Studio 2015中编译项目imflipG.cu 170
6.5.4 运行第一个CUDA应用程序：imflipG.exe 173
6.5.5 确保程序的正确性 174
6.6 Mac平台上的CUDA程序开发 175
6.6.1 在Mac上安装XCode 175
6.6.2 安装CUDA驱动程序和CUDA工具包 176
6.6.3 在Mac上编译和运行CUDA应用程序 177
6.7 Unix平台上的CUDA程序开发 177
6.7.1 安装Eclipse和CUDA工具包 177
6.7.2 使用ssh登录一个集群 178
6.7.3 编译和执行CUDA代码 179
第7章 CUDA主机／设备编程模型 181
7.1 设计程序的并行性 181
7.1.1 任务的并行化 182
7.1.2 什么是Vflip（）的最佳块尺寸 183
7.1.3 imflipG.cu：程序输出的解释 183
7.1.4 imflipG.cu：线程块和图像的大小对性能的影响 184
7.2 核函数的启动 185
7.2.1 网格 185
7.2.2 线程块 187
7.2.3 线程 187
7.2.4 线程束和通道 189
7.3 imflipG.cu：理解核函数的细节 189
7.3.1 在main（）中启动核函数并将参数传递给它们 189
7.3.2 线程执行步骤 190
7.3.3 Vflip（）核函数 191
7.3.4 Vflip（）和MTFlipV（）的比较 192
7.3.5 Hflip（）核函数 194
7.3.6 PixCopy（）核函数 194
7.4 PCIe速度与CPU的关系 196
7.5 PCIe总线对性能的影响 196
7.5.1 数据传输时间、速度、延迟、吞吐量和带宽 196
7.5.2 imflipG.cu的PCIe吞吐量 197
7.6 全局内存总线对性能的影响 200
7.7 计算能力对性能的影响 203
7.7.1 Fermi、Kepler、Maxwell、Pascal和Volta系列 203
7.7.2 不同系列实现的相对带宽 204
7.7.3 imflipG2.cu：计算能力2.0版本的imflipG.cu 205
7.7.4 imflipG2.cu：main（）的修改 206
7.7.5 核函数PxCC20（） 208
7.7.6 核函数VfCC20（） 208
7.8 imflipG2.cu的性能 210
7.9 古典的CUDA调试方法 212
7.9.1 常见的CUDA错误 212
7.9.2 return调试法 214
7.9.3 基于注释的调试 216
7.9.4 printf（）调试 216
7.10 软件错误的生物学原因 217
7.10.1 大脑如何参与编写／调试代码 218
7.10.2 当我们疲倦时是否会写出错误代码 219
第8章 理解GPU的硬件架构 221
8.1 GPU硬件架构 222
8.2 GPU硬件的部件 222
8.2.1 SM：流处理器 222
8.2.2 GPU核心 223
8.2.3 千兆线程调度器 223
8.2.4 内存控制器 225
8.2.5 共享高速缓存（L2$） 225
8.2.6 主机接口 225
8.3 Nvidia GPU架构 226
8.3.1 Fermi架构 227
8.3.2 GT、GTX和计算加速器 227
8.3.3 Kepler架构 228
8.3.4 Maxwell架构 228
8.3.5 Pascal架构和NVLink 229
8.4 CUDA边缘检测：imedgeG.cu 229
8.4.1 CPU和GPU内存中存储图像的变量 229
8.4.2 为GPU变量分配内存 231
8.4.3 调用核函数并对其进行计时 233
8.4.4 计算核函数的性能 234
8.4.5 计算核函数的数据移动量 235
8.4.6 输出核函数性能 236
8.5 imedgeG：核函数 237
8.5.1 BWKernel（） 237
8.5.2 GaussKernel（） 239
8.5.3 SobelKernel（） 240
8.5.4 ThresholdKernel（） 242
8.6 imedgeG.cu的性能 243
8.6.1 imedgeG.cu：PCIe总线利用率 244
8.6.2 imedgeG.cu：运行时间 245
8.6.3 imedgeG.cu：核函数性能比较 247
8.7 GPU代码：编译时间 248
8.7.1 设计CUDA代码 248
8.7.2 编译CUDA代码 250
8.7.3 GPU汇编：PTX、CUBIN 250
8.8 GPU代码：启动 250
8.8.1 操作系统的参与和CUDA DLL文件 250
8.8.2 GPU图形驱动 251
8.8.3 CPU与GPU之间的内存传输 251
8.9 GPU代码：执行（运行时间） 252
8.9.1 获取数据 252
8.9.2 获取代码和参数 252
8.9.3 启动线程块网格 252
8.9.4 千兆线程调度器（GTS） 253
8.9.5 线程块调度 254
8.9.6 线程块的执行 255
8.9.7 透明的可扩展性 256
第9章 理解GPU核心 257
9.1 GPU的架构系列 258
9.1.1 Fermi架构 258
9.1.2 Fermi SM的结构 259
9.1.3 Kepler架构 260
9.1.4 Kepler SMX的结构 260
9.1.5 Maxwell架构 261
9.1.6 MaxwellSMM的结构 262
9.1.7 PascalGP100架构 264
9.1.8 PascalGP100 SM的结构 265
9.1.9 系列比较：峰值GFLOPS和峰值DGFLOPS 266
9.1.10 GPU睿频 267
9.1.11 GPU功耗 268
9.1.12 计算机电源 268
9.2 流处理器的构建模块 269
9.2.1 GPU核心 269
9.2.2 双精度单元（DPU） 270
9.2.3 特殊功能单元（SFU） 270
9.2.4 寄存器文件（RF） 270
9.2.5 读取／存储队列（LDST） 271
9.2.6 L1$和纹理高速缓存 272
9.2.7 共享内存 272
9.2.8 常量高速缓存 272
9.2.9 指令高速缓存 272
9.2.10 指令缓冲区 272
9.2.11 线程束调度器 272
9.2.12 分发单元 273
9.3 并行线程执行（PTX）的数据类型 273
9.3.1 INT8：8位整数 274
9.3.2 INT16：16位整数 274
9.3.3 24位整数 275
9.3.4 INT32：32位整数 275
9.3.5 判定寄存器（32位） 275
9.3.6 INT64：64位整数 276
9.3.7 128位整数 276
9.3.8 FP32：单精度浮点（float） 277
9.3.9 FP64：双精度浮点（double） 277
9.3.10 FP16：半精度浮点（half） 278
9.3.11 什么是FLOP 278
9.3.12 融合乘法累加（FMA）与乘加（MAD） 278
9.3.13 四倍和八倍精度浮点 279
9.3.14 Pascal GP104引擎的SM结构 279
9.4 imflipGC.cu：核心友好的imflipG 280
9.4.1 Hflip2（）：预计算核函数参数 282
9.4.2 Vflip2（）：预计算核函数参数 284
9.4.3 使用线程计算图像坐标 285
9.4.4 线程块ID与图像的行映射 285
9.4.5 Hflip3（）：使用二维启动网格 286
9.4.6 Vflip3（）：使用二维启动网格 287
9.4.7 Hflip4（）：计算2个连续的像素 288
9.4.8 Vflip4（）：计算2个连续的像素 289
9.4.9 Hflip5（）：计算4个连续的像素 290
9.4.10 Vflip5（）：计算4个连续的像素 291
9.4.11 PixCopy2（）、PixCopy3（）：一次分别复制2个和4个连续的像素 292
9.5 imedgeGC.cu：核心友好的imedgeG 293
9.5.1 BWKerne12（）：使用预计算的值和2D块 293
9.5.2 GaussKerne12（）：使用预计算的值和2D块 294
第10章 理解GPU内存 296
10.1 全局内存 297
10.2 L2高速缓存 297
10.3 纹理／L1高速缓存 298
10.4 共享内存 298
10.4.1 分拆与专用共享内存 299
10.4.2 每核心可用的内存资源 299
10.4.3 使用共享内存作为软件高速缓存 300
10.4.4 分配SM中的共享内存 300
10.5 指令高速缓存 300
10.6 常量内存 301
10.7 imflipGCM.cu：核心和内存友好的imflipG 301
10.7.1 Hflip6（）、Vflip6（）：使用共享内存作为缓冲区 301
10.7.2 Hflip7（）：共享内存中连续的交换操作 303
10.7.3 Hflip8（）：使用寄存器交换4个像素 305
10.7.4 Vflip7（）：一次复制4个字节（int） 307
10.7.5 对齐与未对齐的内存数据访问 308
10.7.6 Vflip8（）：一次复制8个字节 308
10.7.7 Vflip9（）：仅使用全局内存，一次复制8个字节 309
10.7.8 PixCopy4（）、PixCopy5（）：使用共享内存复制1个和4个字节 310
10.7.9 PixCopy6（）、PixCopy7（）：使用全局内存复制1个和2个整数 311
10.8 imedgeGCM.cu：核心和内存友好的imedgeG 312
10.8.1 BWKerne13（）：使用字节操作来提取RGB 312
10.8.2 GaussKerne13（）：使用常量内存 314
10.8.3 处理常量的方法 315
10.8.4 GaussKerne14（）：在共享内存中缓冲1个像素的邻居 316
10.8.5 GaussKerne15（）：在共享内存中缓冲4个像素的邻居 318
10.8.6 GaussKerne16（）：将5个垂直像素读入共享内存 320
10.8.7 GaussKerne17（）：去除边界像素的影响 322
10.8.8 GaussKerne18（）：计算8个垂直像素 324
10.9 CUDA占用率计算器 326
10.9.1 选择最佳的线程／块 327
10.9.2 SM级资源限制 328
10.9.3 什么是占用率 329
10.9.4 CUDA占用率计算器：资源计算 330
10.9.5 案例分析：GaussKerne17（） 334
10.9.6 案例分析：GaussKerne18（） 337
第11章 CUDA流 340
11.1 什么是流水线 342
11.1.1 重叠执行 342
11.1.2 暴露与合并的运行时间 343
11.2 内存分配 344
11.2.1 物理与虚拟内存 345
11.2.2 物理地址到虚拟地址的转换 345
11.2.3 固定内存 345
11.2.4 使用cudaMallocHost（）分配固定内存 346
11.3 CPU与GPU之间快速的数据传输 346
11.3.1 同步数据传输 346
11.3.2 异步数据传输 347
11.4 CUDA流的原理 347
11.4.1 CPU到GPU的传输、核函数的执行、GPU到CPU的传输 347
11.4.2 在CUDA中实现流 348
11.4.3 复制引擎 348
11.4.4 核函数执行引擎 349
11.4.5 并发的上行和下行PCIe传输 349
11.4.6 创建CUDA流 350
11.4.7 销毁CUDA流 350
11.4.8 同步CUDA流 350
11.5 imGStr.cu：流式图像处理 351
11.5.1 将图像读入固定内存 351
11.5.2 同步与单个流 353
11.5.3 多个流 354
11.5.4 多流之间的数据依赖 356
11.6 流式水平翻转核函数 360
11.7 imGStr.cu：流式边缘检测 362
11.8 性能对比：imGStr.cu 365
11.8.1 同步与异步结果 366
11.8.2 结果的随机性 366
11.8.3 队列优化 366
11.8.4 最佳流式结果 367
11.8.5 最差流式结果 369
11.9 Nvidia可视化分析器：nvvp 370
11.9.1 安装nvvp和nvprof 370
11.9.2 使用nvvp 370
11.9.3 使用nvprof 371
11.9.4 imGStr的同步和单流结果 372
11.9.5 imGStr的2流和4流结果 373
第三部分 拓展知识 376
第12章 CUDA库 376
12.1 cuBLAS 377
12.1.1 BLAS级别 377
12.1.2 cuBLAS数据类型 377
12.1.3 安装cuBLAS 378
12.1.4 变量声明和初始化 378
12.1.5 设备内存分配 379
12.1.6 创建上下文 379
12.1.7 将数据传输到设备端 379
12.1.8 调用cuBLAS函数 380
12.1.9 将数据传回主机 380
12.1.10 释放内存 381
12.1.11 cuBLAS程序示例：矩阵的标量操作 381
12.2 cuFFT 382
12.2.1 cuFFT库特征 383
12.2.2 复数到复数变换示例 383
12.2.3 实数到复数变换示例 384
12.3 Nvidia性能库（NPP） 384
12.4 Thrust库 386
第13章 OpenCL简介 388
13.1 什么是OpenCL 388
13.1.1 多平台 388
13.1.2 基于队列 389
13.2 图像翻转核函数 389
13.3 运行核函数 390
13.3.1 选择设备 390
13.3.2 运行核函数 392
13.3.3 OpenCL程序的运行时间 396
13.4 OpenCL中的边缘检测 396
第14章 其他GPU编程语言 402
14.1 使用Python进行GPU编程 402
14.1.1 imflip的PyOpenCL版本 403
14.1.2 PyOpenCL的逐元素核函数 406
14.2 OpenGL 408
14.3 OpenGL ES：用于嵌入式系统的OpenGL 409
14.4 Vulkan 409
14.5 微软的高级着色语言 409
14.5.1 着色 410
14.5.2 HLSL 410
14.6 Apple的MetalAPI 411
14.7 Apple的Swift编程语言 411
14.8 OpenCV 411
14.8.1 安装OpenCV和人脸识别 411
14.8.2 移动-微云-云实时人脸识别 412
14.8.3 加速即服务（AXaas） 412
第15章 深度学习中的CUDA 413
15.1 人工神经网络 413
15.1.1 神经元 413
15.1.2 激活函数 414
15.2 全连接神经网络 415
15.3 深度网络／卷积神经网络 415
15.4 训练网络 416
15.5 cuDNN深度学习库 416
15.5.1 创建一个层 417
15.5.2 创建一个网络 418
15.5.3 前向传播 418
15.5.4 反向传播 419
15.5.5 在网络中使用cuBLAS 419
15.6 Keras 419
参考文献 421
术语表 424