基于CUDA的GPU并行程序开发指南PDF电子书下载

第一部分 理解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