第一部分 理解CPU的并行性

第1章 CPU并行编程概述

1.1 并行编程的演化

1.2 核心越多，并行性越高

1.3 核心与线程

1.3.1 并行化更多的是线程还是核心

1.3.2 核心资源共享的影响

1.3.3 内存资源共享的影响

1.4 第一个串行程序

1.4.1 理解数据传输速度

1.4.2 imflip.c中的main（）函数

1.4.3 垂直翻转行：FlipImageV（）

1.4.4 水平翻转列：FlipImageH（）

1.5 程序的编辑、编译、运行

1.5.1 选择编辑器和编译器

1.5.2 在Windows 7、8、10平台上开发

1.5.3 在Mac平台上开发

1.5.4 在Unix平台上开发

1.6 Unix速成

1.6.1 与目录相关的Unix命令

1.6.2 与文件相关的Unix命令

1.7 调试程序

1.7.1 gdb

1.7.2 古典调试方法

1.7.3 valgrind

1.8 第一个串行程序的性能

1.8.1 可以估计执行时间吗

1.8.2 代码执行时OS在做什么

1.8.3 如何并行化

1.8.4 关于资源的思考

第2章 开发第一个CPU并行程序

2.1 第一个并行程序

2.1.1 imflipP.c中的main（）函数

2.1.2 运行时间

2.1.3 imflipP.c中main（）函数代码的划分

2.1.4 线程初始化

2.1.5 创建线程

2.1.6 线程启动／执行

2.1.7 线程终止（合并）

2.1.8 线程任务和数据划分

2.2 位图文件

2.2.1 BMP是一种无损／不压缩的文件格式

2.2.2 BMP图像文件格式

2.2.3 头文件ImageStuff.h

2.2.4 ImageStuff.c中的图像操作函数

2.3 执行线程任务

2.3.1 启动线程

2.3.2 多线程垂直翻转函数MTFlipV（）

2.3.3 FlipImageV（）和MTFlipV（）的比较

2.3.4 多线程水平翻转函数MTFlipH（）

2.4 多线程代码的测试／计时

第3章 改进第一个CPU并行程序

3.1 程序员对性能的影响

3.2 CPU对性能的影响

3.2.1 按序核心与乱序核心

3.2.2 瘦线程与胖线程

3.3 imflipP的性能

3.4 操作系统对性能的影响

3.4.1 创建线程

3.4.2 线程启动和执行

3.4.3 线程状态

3.4.4 将软件线程映射到硬件线程

3.4.5 程序性能与启动的线程

3.5 改进imflipP

3.5.1 分析MTFlipH（）中的内存访问模式

3.5.2 MTFlipH（）的多线程内存访问

3.5.3 DRAM访问的规则

3.6 imflipPM：遵循DRAM的规则

3.6.1 imflipP的混乱内存访问模式

3.6.2 改进imflipP的内存访问模式

3.6.3 MTFlipHM（）：内存友好的MTFlipH（）

3.6.4 MTFlipVM（）：内存友好的MTFlipV（）

3.7 imflipPM.C的性能

3.7.1 imflipP.c和imflipPM.c的性能比较

3.7.2 速度提升：MTFlipV（）与MTFlipVM（）

3.7.3 速度提升：MTFlipH（）与MTFlipHM（）

3.7.4 理解加速：MTFlipH（）与MTFlipHM（）

3.8 进程内存映像

3.9 英特尔MIC架构：Xeon Phi

3.10 GPU是怎样的

3.11 本章小结

第4章 理解核心和内存

4.1 曾经的英特尔

4.2 CPU和内存制造商

4.3 动态存储器与静态存储器

4.3.1 静态随机存取存储器（SRAM）

4.3.2 动态随机存取存储器（DRAM）

4.3.3 DRAM接口标准

4.3.4 DRAM对程序性能的影响

4.3.5 SRAM对程序性能的影响

4.4 图像旋转程序：imrotate.c

4.4.1 imrotate.c的说明

4.4.2 imrotate.c：参数限制和简化

4.4.3 imrotate.c：实现原理

4.5 imrotate的性能

4.5.1 线程效率的定性分析

4.5.2 定量分析：定义线程效率

4.6 计算机的体系结构

4.6.1 核心、L1$和L2$

4.6.2 核心内部资源

4.6.3 共享L3高速缓存（L3$）

4.6.4 内存控制器

4.6.5 主存

4.6.6 队列、非核心和I/O

4.7 imrotateMC：让imrotate更高效

4.7.1 Rotate2（）：平方根和浮点除法有多差

4.7.2 Rotate3（）和Rotate4（）：sin（）和cos（）有多差

4.7.3 Rotate5（）：整数除法／乘法有多差

4.7.4 Rotate6（）：合并计算

4.7.5 Rotate7（）：合并更多计算

4.7.6 imrotateMC的总体性能

4.8 本章小结

第5章 线程管理和同步

5.1 边缘检测程序：imedge.c

5.1.1 imedge.c的说明

5.1.2 imedge.c：参数限制和简化

5.1.3 imedge.c：实现原理

5.2 imedge.c：实现

5.2.1 初始化和时间戳

5.2.2 不同图像表示的初始化函数

5.2.3 启动和终止线程

5.2.4 高斯滤波

5.2.5 Sobel

5.2.6 阈值过滤

5.3 imedge的性能

5.4 imedgeMC：让imedge更高效

5.4.1 利用预计算降低带宽

5.4.2 存储预计算的像素值

5.4.3 预计算像素值

5.4.4 读取图像并预计算像素值

5.4.5 PrGaussianFilter

5.4.6 PrSobel

5.4.7 PrThreshold

5.5 imedgeMC的性能

5.6 imedgeMCT：高效的线程同步

5.6.1 屏障同步

5.6.2 用于数据共享的MUTEX结构

5.7 imedgeMCT：实现

5.7.1 使用MUTEX：读取图像、预计算

5.7.2 一次预计算一行

5.8 imedgeMCT的性能

第二部分 基于CUDA的GPU编程

第6章 GPU并行性和CUDA概述

6.1 曾经的Nvidia

6.1.1 GPU的诞生

6.1.2 早期的GPU架构

6.1.3 GPGPU的诞生

6.1.4 Nvidia、ATI Technologies和Intel

6.2 统一计算设备架构

6.2.1 CUDA、OpenCL和其他GPU语言

6.2.2 设备端与主机端代码

6.3 理解GPU并行

6.3.1 GPU如何实现高性能

6.3.2 CPU与GPU架构的差异

6.4 图像翻转的CUDA版：imflipG.cu

6.4.1 imflipG.cu：将图像读入CPU端数组

6.4.2 初始化和查询GPU

6.4.3 GPU端的时间戳

6.4.4 GPU端内存分配

6.4.5 GPU驱动程序和Nvidia运行时引擎

6.4.6 CPU到GPU的数据传输

6.4.7 用封装函数进行错误报告

6.4.8 GPU核函数的执行

6.4.9 完成GPU核函数的执行

6.4.10 将GPU结果传回CPU

6.4.11 完成时间戳

6.4.12 输出结果和清理

6.4.13 读取和输出BMP文件

6.4.14 Vflip（）：垂直翻转的GPU核函数

6.4.15 什么是线程ID、块ID和块维度

6.4.16 Hflip（）：水平翻转的GPU核函数

6.4.17 硬件参数：threadIDx.x、blockIdx.x和blockDim.x

6.4.18 PixCopy（）：复制图像的GPU核函数

6.4.19 CUDA关键字

6.5 Windows中的CUDA程序开发

6.5.1 安装MS Visual Studio 2015和CUDA Toolkit 8.0

6.5.2 在Visual Studio 2015中创建项目imflipG.cu

6.5.3 在Visual Studio 2015中编译项目imflipG.cu

6.5.4 运行第一个CUDA应用程序：imflipG.exe

6.5.5 确保程序的正确性

6.6 Mac平台上的CUDA程序开发

6.6.1 在Mac上安装XCode

6.6.2 安装CUDA驱动程序和CUDA工具包

6.6.3 在Mac上编译和运行CUDA应用程序

6.7 Unix平台上的CUDA程序开发

6.7.1 安装Eclipse和CUDA工具包

6.7.2 使用ssh登录一个集群

6.7.3 编译和执行CUDA代码

第7章 CUDA主机／设备编程模型

7.1 设计程序的并行性

7.1.1 任务的并行化

7.1.2 什么是Vflip（）的最佳块尺寸

7.1.3 imflipG.cu：程序输出的解释

7.1.4 imflipG.cu：线程块和图像的大小对性能的影响

7.2 核函数的启动

7.2.1 网格

7.2.2 线程块

7.2.3 线程

7.2.4 线程束和通道

7.3 imflipG.cu：理解核函数的细节

7.3.1 在main（）中启动核函数并将参数传递给它们

7.3.2 线程执行步骤

7.3.3 Vflip（）核函数

7.3.4 Vflip（）和MTFlipV（）的比较

7.3.5 Hflip（）核函数

7.3.6 PixCopy（）核函数

7.4 PCIe速度与CPU的关系

7.5 PCIe总线对性能的影响

7.5.1 数据传输时间、速度、延迟、吞吐量和带宽

7.5.2 imflipG.cu的PCIe吞吐量

7.6 全局内存总线对性能的影响

7.7 计算能力对性能的影响

7.7.1 Fermi、Kepler、Maxwell、Pascal和Volta系列

7.7.2 不同系列实现的相对带宽

7.7.3 imflipG2.cu：计算能力2.0版本的imflipG.cu

7.7.4 imflipG2.cu：main（）的修改

7.7.5 核函数PxCC20（）

7.7.6 核函数VfCC20（）

7.8 imflipG2.cu的性能

7.9 古典的CUDA调试方法

7.9.1 常见的CUDA错误

7.9.2 return调试法

7.9.3 基于注释的调试

7.9.4 printf（）调试

7.10 软件错误的生物学原因

7.10.1 大脑如何参与编写／调试代码

7.10.2 当我们疲倦时是否会写出错误代码

第8章 理解GPU的硬件架构

8.1 GPU硬件架构

8.2 GPU硬件的部件

8.2.1 SM：流处理器

8.2.2 GPU核心

8.2.3 千兆线程调度器

8.2.4 内存控制器

8.2.5 共享高速缓存（L2$）

8.2.6 主机接口

8.3 Nvidia GPU架构

8.3.1 Fermi架构

8.3.2 GT、GTX和计算加速器

8.3.3 Kepler架构

8.3.4 Maxwell架构

8.3.5 Pascal架构和NVLink

8.4 CUDA边缘检测：imedgeG.cu

8.4.1 CPU和GPU内存中存储图像的变量

8.4.2 为GPU变量分配内存

8.4.3 调用核函数并对其进行计时

8.4.4 计算核函数的性能

8.4.5 计算核函数的数据移动量

8.4.6 输出核函数性能

8.5 imedgeG：核函数

8.5.1 BWKernel（）

8.5.2 GaussKernel（）

8.5.3 SobelKernel（）

8.5.4 ThresholdKernel（）

8.6 imedgeG.cu的性能

8.6.1 imedgeG.cu：PCIe总线利用率

8.6.2 imedgeG.cu：运行时间

8.6.3 imedgeG.cu：核函数性能比较

8.7 GPU代码：编译时间

8.7.1 设计CUDA代码

8.7.2 编译CUDA代码

8.7.3 GPU汇编：PTX、CUBIN

8.8 GPU代码：启动

8.8.1 操作系统的参与和CUDA DLL文件

8.8.2 GPU图形驱动

8.8.3 CPU与GPU之间的内存传输

8.9 GPU代码：执行（运行时间）

8.9.1 获取数据

8.9.2 获取代码和参数

8.9.3 启动线程块网格

8.9.4 千兆线程调度器（GTS）

8.9.5 线程块调度

8.9.6 线程块的执行

8.9.7 透明的可扩展性

第9章 理解GPU核心

9.1 GPU的架构系列

9.1.1 Fermi架构

9.1.2 Fermi SM的结构

9.1.3 Kepler架构

9.1.4 Kepler SMX的结构

9.1.5 Maxwell架构

9.1.6 MaxwellSMM的结构

9.1.7 PascalGP100架构

9.1.8 PascalGP100 SM的结构

9.1.9 系列比较：峰值GFLOPS和峰值DGFLOPS

9.1.10 GPU睿频

9.1.11 GPU功耗

9.1.12 计算机电源

9.2 流处理器的构建模块

9.2.1 GPU核心

9.2.2 双精度单元（DPU）

9.2.3 特殊功能单元（SFU）

9.2.4 寄存器文件（RF）

9.2.5 读取／存储队列（LDST）

9.2.6 L1$和纹理高速缓存

9.2.7 共享内存

9.2.8 常量高速缓存

9.2.9 指令高速缓存

9.2.10 指令缓冲区

9.2.11 线程束调度器

9.2.12 分发单元

9.3 并行线程执行（PTX）的数据类型

9.3.1 INT8：8位整数

9.3.2 INT16：16位整数

9.3.3 24位整数

9.3.4 INT32：32位整数

9.3.5 判定寄存器（32位）

9.3.6 INT64：64位整数

9.3.7 128位整数

9.3.8 FP32：单精度浮点（float）

9.3.9 FP64：双精度浮点（double）

9.3.10 FP16：半精度浮点（half）

9.3.11 什么是FLOP

9.3.12 融合乘法累加（FMA）与乘加（MAD）

9.3.13 四倍和八倍精度浮点

9.3.14 Pascal GP104引擎的SM结构

9.4 imflipGC.cu：核心友好的imflipG

9.4.1 Hflip2（）：预计算核函数参数

9.4.2 Vflip2（）：预计算核函数参数

9.4.3 使用线程计算图像坐标

9.4.4 线程块ID与图像的行映射

9.4.5 Hflip3（）：使用二维启动网格

9.4.6 Vflip3（）：使用二维启动网格

9.4.7 Hflip4（）：计算2个连续的像素

9.4.8 Vflip4（）：计算2个连续的像素

9.4.9 Hflip5（）：计算4个连续的像素

9.4.10 Vflip5（）：计算4个连续的像素

9.4.11 PixCopy2（）、PixCopy3（）：一次分别复制2个和4个连续的像素

9.5 imedgeGC.cu：核心友好的imedgeG

9.5.1 BWKerne12（）：使用预计算的值和2D块

9.5.2 GaussKerne12（）：使用预计算的值和2D块

第10章 理解GPU内存

10.1 全局内存

10.2 L2高速缓存

10.3 纹理／L1高速缓存

10.4 共享内存

10.4.1 分拆与专用共享内存

10.4.2 每核心可用的内存资源

10.4.3 使用共享内存作为软件高速缓存

10.4.4 分配SM中的共享内存

10.5 指令高速缓存

10.6 常量内存

10.7 imflipGCM.cu：核心和内存友好的imflipG

10.7.1 Hflip6（）、Vflip6（）：使用共享内存作为缓冲区

10.7.2 Hflip7（）：共享内存中连续的交换操作

10.7.3 Hflip8（）：使用寄存器交换4个像素

10.7.4 Vflip7（）：一次复制4个字节（int）

10.7.5 对齐与未对齐的内存数据访问

10.7.6 Vflip8（）：一次复制8个字节

10.7.7 Vflip9（）：仅使用全局内存，一次复制8个字节

10.7.8 PixCopy4（）、PixCopy5（）：使用共享内存复制1个和4个字节

10.7.9 PixCopy6（）、PixCopy7（）：使用全局内存复制1个和2个整数

10.8 imedgeGCM.cu：核心和内存友好的imedgeG

10.8.1 BWKerne13（）：使用字节操作来提取RGB

10.8.2 GaussKerne13（）：使用常量内存

10.8.3 处理常量的方法

10.8.4 GaussKerne14（）：在共享内存中缓冲1个像素的邻居

10.8.5 GaussKerne15（）：在共享内存中缓冲4个像素的邻居

10.8.6 GaussKerne16（）：将5个垂直像素读入共享内存

10.8.7 GaussKerne17（）：去除边界像素的影响

10.8.8 GaussKerne18（）：计算8个垂直像素

10.9 CUDA占用率计算器

10.9.1 选择最佳的线程／块

10.9.2 SM级资源限制

10.9.3 什么是占用率

10.9.4 CUDA占用率计算器：资源计算

10.9.5 案例分析：GaussKerne17（）

10.9.6 案例分析：GaussKerne18（）

第11章 CUDA流

11.1 什么是流水线

11.1.1 重叠执行

11.1.2 暴露与合并的运行时间

11.2 内存分配

11.2.1 物理与虚拟内存

11.2.2 物理地址到虚拟地址的转换

11.2.3 固定内存

11.2.4 使用cudaMallocHost（）分配固定内存

11.3 CPU与GPU之间快速的数据传输

11.3.1 同步数据传输

11.3.2 异步数据传输

11.4 CUDA流的原理

11.4.1 CPU到GPU的传输、核函数的执行、GPU到CPU的传输

11.4.2 在CUDA中实现流

11.4.3 复制引擎

11.4.4 核函数执行引擎

11.4.5 并发的上行和下行PCIe传输

11.4.6 创建CUDA流

11.4.7 销毁CUDA流

11.4.8 同步CUDA流

11.5 imGStr.cu：流式图像处理

11.5.1 将图像读入固定内存

11.5.2 同步与单个流

11.5.3 多个流

11.5.4 多流之间的数据依赖

11.6 流式水平翻转核函数

11.7 imGStr.cu：流式边缘检测

11.8 性能对比：imGStr.cu

11.8.1 同步与异步结果

11.8.2 结果的随机性

11.8.3 队列优化

11.8.4 最佳流式结果

11.8.5 最差流式结果

11.9 Nvidia可视化分析器：nvvp

11.9.1 安装nvvp和nvprof

11.9.2 使用nvvp

11.9.3 使用nvprof

11.9.4 imGStr的同步和单流结果

11.9.5 imGStr的2流和4流结果

第三部分 拓展知识

第12章 CUDA库

12.1 cuBLAS

12.1.1 BLAS级别

12.1.2 cuBLAS数据类型

12.1.3 安装cuBLAS

12.1.4 变量声明和初始化

12.1.5 设备内存分配

12.1.6 创建上下文

12.1.7 将数据传输到设备端

12.1.8 调用cuBLAS函数

12.1.9 将数据传回主机

12.1.10 释放内存

12.1.11 cuBLAS程序示例：矩阵的标量操作

12.2 cuFFT

12.2.1 cuFFT库特征

12.2.2 复数到复数变换示例

12.2.3 实数到复数变换示例

12.3 Nvidia性能库（NPP）

12.4 Thrust库

第13章 OpenCL简介

13.1 什么是OpenCL

13.1.1 多平台

13.1.2 基于队列

13.2 图像翻转核函数

13.3 运行核函数

13.3.1 选择设备

13.3.2 运行核函数

13.3.3 OpenCL程序的运行时间

13.4 OpenCL中的边缘检测

第14章 其他GPU编程语言

14.1 使用Python进行GPU编程

14.1.1 imflip的PyOpenCL版本

14.1.2 PyOpenCL的逐元素核函数

14.2 OpenGL

14.3 OpenGL ES：用于嵌入式系统的OpenGL

14.4 Vulkan

14.5 微软的高级着色语言

14.5.1 着色

14.5.2 HLSL

14.6 Apple的MetalAPI

14.7 Apple的Swift编程语言

14.8 OpenCV

14.8.1 安装OpenCV和人脸识别

14.8.2 移动-微云-云实时人脸识别

14.8.3 加速即服务（AXaas）

第15章 深度学习中的CUDA

15.1 人工神经网络

15.1.1 神经元

15.1.2 激活函数

15.2 全连接神经网络

15.3 深度网络／卷积神经网络

15.4 训练网络

15.5 cuDNN深度学习库

15.5.1 创建一个层

15.5.2 创建一个网络

15.5.3 前向传播

15.5.4 反向传播

15.5.5 在网络中使用cuBLAS

15.6 Keras

参考文献

术语表