[正版]基于CUDA的GPU并行程序开发指南 (美)托尔加·索亚塔(Tolga Soyata) 著 唐杰 译 程序设计(

●译者序
前言
关于作者
第一部分理解CPU的并行性
第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.2imflip.c中的main()函数9
1.4.3垂直翻转行：FlipImageV()10
1.4.4水平翻转列：FlipImageH()11
1.5程序的编辑、编译、运行12
1.5.1选择编辑器和编译器12
1.5.2在Windows7、8、10平台上开发12
1.5.3在Mac平台上开发14
1.5.4在Unix平台上开发14
1.6Unix速成15
1.6.1与目录相关的Unix命令15
1.6.2与文件相关的Unix命令16
1.7调试程序19
1.7.1gdb19
1.7.2古典调试方法20
1.7.3valgrind22
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.1imflipP.c中的main()函数27
2.1.2运行时间28
2.1.3imflipP.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.1BMP是一种无损/不压缩的文件格式35
2.2.2BMP图像文件格式36
2.2.3头文件ImageStuff.h37
2.2.4ImageStuff.c中的图像操作函数38
2.3执行线程任务40
2.3.1启动线程41
2.3.2多线程垂直翻转函数MTFlipV()43
2.3.3FlipImageV()和MTFlipV()的比较46
2.3.4多线程水平翻转函数MTFlipH(?)47
2.4多线程代码的测试/计时49
第3章改进第一个CPU并行程序51
3.1程序员对性能的影响51
3.2CPU对性能的影响52
3.2.1按序核心与乱序核心53
3.2.2瘦线程与胖线程55
3.3imf?lipP的性能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改进imf?lipP61
3.5.1分析MTFlipH()中的内存访问模式62
3.5.2MTFlipH()的多线程内存访问63
3.5.3DRAM访问的规则64
3.6imflipPM：遵循DRAM的规则65
3.6.1imflipP的混乱内存访问模式65
3.6.2改进imflipP的内存访问模式65
3.6.3MTFlipHM()：内存友好的MTFlipH()66
3.6.4MTFlipVM()：内存友好的MTFlipV()69
3.7imflipPM.C的性能69
3.7.1imflipP.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架构：XeonPhi74
3.10GPU是怎样的75
3.11本章小结76
第4章理解核心和内存77
4.1曾经的英特尔77
4.2CPU和内存制造商78
4.3动态存储器与静态存储器79
4.3.1静态随机存取存储器（SRAM）79
4.3.2动态随机存取存储器（DRAM）79
4.3.3DRAM接口标准79
4.3.4DRAM对程序性能的影响80
4.3.5SRAM对程序性能的影响81
4.4图像旋转程序：imrotate.c81
4.4.1imrotate.c的说明82
4.4.2imrotate.c：参数和简化82
4.4.3imrotate.c：实现原理83
4.5imrotate的性能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/O93
4.7imrotateMC：让imrotate更高效94
4.7.1Rotate2()：平方根和浮点除法有多差96
4.7.2Rotate3()和Rotate4()：sin()和cos()有多差97
4.7.3Rotate5()：整数除法/乘法有多差98
4.7.4Rotate6()：合并计算100
4.7.5Rotate7()：合并更多计算100
4.7.6imrotateMC的总体性能101
4.8本章小结103
第5章线程管理和同步104
5.1边缘检测程序：imedge.c104
5.1.1imedge.c的说明105
5.1.2imedge.c：参数和简化106
5.1.3imedge.c：实现原理106
5.2imedge.c：实现108
5.2.1初始化和时间戳109
5.2.2不同图像表示的初始化函数110
5.2.3启动和终止线程111
5.2.4高斯滤波112
5.2.5Sobel113
5.2.6阈值过滤114
5.3imedge的性能115
5.4imedgeMC：让imedge更高效116
5.4.1利用预计算降低带宽116
5.4.2存储预计算的像素值117
5.4.3预计算像素值118
5.4.4读取图像并预计算像素值119
5.4.5PrGaussianFilter120
5.4.6PrSobel121
5.4.7PrThreshold121
5.5imedgeMC的性能122
5.6imedgeMCT：高效的线程同步123
5.6.1屏障同步124
5.6.2用于数据共享的MUTEX结构125
5.7imedgeMCT：实现127
5.7.1使用MUTEX：读取图像、预计算128
5.7.2一次预计算一行130
5.8imedgeMCT的性能131
第二部分基于CUDA的GPU编程
第6章GPU并行性和CUDA概述134
6.1曾经的Nvidia134
6.1.1GPU的诞生134
6.1.2早期的GPU架构136
6.1.3GPGPU的诞生137
6.1.4Nvidia、ATITechnologies和Intel138
6.2统一计算设备架构140
6.2.1CUDA、OpenCL和其他GPU语言140
6.2.2设备端与主机端代码140
6.3理解GPU并行141
6.3.1GPU如何实现高性能142
6.3.2CPU与GPU架构的差异143
6.4图像翻转的CUDA版：imflipG.cu144
6.4.1imflipG.cu：将图像读入CPU端数组146
6.4.2初始化和查询GPU147
6.4.3GPU端的时间戳148
6.4.4GPU端内存分配152
6.4.5GPU驱动程序和Nvidia运行时引擎153
6.4.6CPU到GPU的数据传输153
6.4.7用封装函数进行错误报告154
6.4.8GPU核函数的执行155
6.4.9完成GPU核函数的执行157
6.4.10将GPU结果传回CPU158
6.4.11完成时间戳158
6.4.12输出结果和清理158
6.4.13读取和输出BMP文件159
6.4.14Vflip()：垂直翻转的GPU核函数160
6.4.15什么是线程ID、块ID和块维度163
6.4.16Hflip()：水平翻转的GPU核函数165
6.4.17硬件参数：threadIDx.x、blockIdx.x和blockDim.x165
6.4.18PixCopy()：复制图像的GPU核函数165
6.4.19CUDA关键字166
6.5Windows中的CUDA程序开发167
6.5.1安装MSVisualStudio2015和CUDAToolkit8.0167
6.5.2在VisualStudio2015中创建项目imflipG.cu168
6.5.3在VisualStudio2015中编译项目imflipG.cu170
6.5.4运行第一个CUDA应用程序：imflipG.exe173
6.5.5确保程序的正确性174
6.6Mac平台上的CUDA程序开发175
6.6.1在Mac上安装XCode175
6.6.2安装CUDA驱动程序和CUDA工具包176
6.6.3在Mac上编译和运行CUDA应用程序177
6.7Unix平台上的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.3imflipG.cu：程序输出的解释183
7.1.4imflipG.cu：线程块和图像的大小对性能的影响184
7.2核函数的启动185
7.2.1网格185
7.2.2线程块187
7.2.3线程187
7.2.4线程束和通道189
7.3imf?lipG.cu：理解核函数的细节189
7.3.1在main()中启动核函数并将参数传递给它们189
7.3.2线程执行步骤190
7.3.3Vflip()核函数191
7.3.4Vflip()和MTFlipV()的比较192
7.3.5Hflip()核函数194
7.3.6PixCopy()核函数194
7.4PCIe速度与CPU的关系196
7.5PCIe总线对性能的影响196
7.5.1数据传输时间、速度、延迟、吞吐量和带宽196
7.5.2imflipG.cu的PCIe吞吐量197
7.6全局内存总线对性能的影响200
7.7计算能力对性能的影响203
7.7.1Fermi、Kepler、Maxwell、Pascal和Volta系列203
7.7.2不同系列实现的相对带宽204
7.7.3imflipG2.cu：计算能力2.0版本的imflipG.cu205
7.7.4imflipG2.cu：main()的修改206
7.7.5核函数PxCC20()208
7.7.6核函数VfCC20()208
7.8imflipG2.cu的性能210
7.9古典的CUDA调试方法212
7.9.1常见的CUDA错误212
7.9.2return调试法214
7.9.3基于注释的调试216
7.9.4printf(?)调试216
7.10软件错误的生物学原因217
7.10.1大脑如何参与编写/调试代码218
7.10.2当我们疲倦时是否会写出错误代码219
第8章理解GPU的硬件架构221
8.1GPU硬件架构222
8.2GPU硬件的部件222
8.2.1SM：流处理器222
8.2.2GPU核心223
8.2.3千兆线程调度器223
8.2.4内存控制器225
8.2.5共享高速缓存（L2$）225
8.2.6主机接口225
8.3NvidiaGPU架构226
8.3.1Fermi架构227
8.3.2GT、GTX和计算加速器227
8.3.3Kepler架构228
8.3.4Maxwell架构228
8.3.5Pascal架构和NVLink229
8.4CUDA边缘检测：imedgeG.cu229
8.4.1CPU和GPU内存中存储图像的变量229
8.4.2为GPU变量分配内存231
8.4.3调用核函数并对其进行计时233
8.4.4计算核函数的性能234
8.4.5计算核函数的数据移动量235
8.4.6输出核函数性能236
8.5imedgeG：核函数237
8.5.1BWKernel()237
8.5.2GaussKernel()239
8.5.3SobelKernel()240
8.5.4ThresholdKernel()242
8.6imedgeG.cu的性能243
8.6.1imedgeG.cu：PCIe总线利用率244
8.6.2imedgeG.cu：运行时间245
8.6.3imedgeG.cu：核函数性能比较247
8.7GPU代码：编译时间248
8.7.1设计CUDA代码248
8.7.2编译CUDA代码250
8.7.3GPU汇编：PTX、CUBIN250
8.8GPU代码：启动250
8.8.1操作系统的参与和CUDADLL文件250
8.8.2GPU图形驱动251
8.8.3CPU与GPU之间的内存传输251
8.9GPU代码：执行（运行时间）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.1GPU的架构系列258
9.1.1Fermi架构258
9.1.2FermiSM的结构259
9.1.3Kepler架构260
9.1.4KeplerSMX的结构260
9.1.5Maxwell架构261
9.1.6MaxwellSMM的结构262
9.1.7PascalGP100架构264
9.1.8PascalGP100SM的结构265
9.1.9系列比较：峰值GFLOPS和峰值DGFLOPS266
9.1.10GPU睿频267
9.1.11GPU功耗268
9.1.12计算机电源268
9.2流处理器的构建模块269
9.2.1GPU核心269
9.2.2双精度单元（DPU）270
9.2.3特殊功能单元（SFU）270
9.2.4寄存器文件（RF）270
9.2.5读取/存储队列（LDST）271
9.2.6L1$和纹理高速缓存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.1INT8：8位整数274
9.3.2INT16：16位整数274
9.3.324位整数275
9.3.4INT32：32位整数275
9.3.5判定寄存器（32位）275
9.3.6INT64：64位整数276
9.3.7128位整数276
9.3.8FP32：单精度浮点（float）277
9.3.9FP64：双精度浮点（double）277
9.3.10FP16：半精度浮点（half）278
9.3.11什么是FLOP278
9.3.12融合乘法累加（FMA）与乘加（MAD）278
9.3.13四倍和八倍精度浮点279
9.3.14PascalGP104引擎的SM结构279
9.4imflipGC.cu：核心友好的imflipG280
9.4.1Hflip2()：预计算核函数参数282
9.4.2Vflip2()：预计算核函数参数284
9.4.3使用线程计算图像坐标285
9.4.4线程块ID与图像的行映射285
9.4.5Hflip3()：使用二维启动网格286
9.4.6Vflip3()：使用二维启动网格287
9.4.7Hflip4()：计算2个连续的像素288
9.4.8Vflip4()：计算2个连续的像素289
9.4.9Hflip5()：计算4个连续的像素290
9.4.10Vflip5()：计算4个连续的像素291
9.4.11PixCopy2()、PixCopy3()：一次分别复制2个和4个连续的像素292
9.5imedgeGC.cu：核心友好的imedgeG293
9.5.1BWKernel2()：使用预计算的值和2D块293
9.5.2GaussKernel2()：使用预计算的值和2D块294
第10章理解GPU内存296
10.1全局内存297
10.2L2高速缓存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.7imflipGCM.cu：核心和内存友好的imflipG301
10.7.1Hflip6()、Vflip6()：使用共享内存作为缓冲区301
10.7.2Hflip7()：共享内存中连续的交换操作303
10.7.3Hflip8()：使用寄存器交换4个像素305
10.7.4Vflip7()：一次复制4个字节（int）307
10.7.5对齐与未对齐的内存数据访问308
10.7.6Vflip8()：一次复制8个字节308
10.7.7Vflip9()：仅使用全局内存，一次复制8个字节309
10.7.8PixCopy4()、PixCopy5()：使用共享内存复制1个和4个字节310
10.7.9PixCopy6()、PixCopy7()：使用全局内存复制1个和2个整数311
10.8imedgeGCM.cu：核心和内存友好的imedgeG312
10.8.1BWKernel3()：使用字节操作来提取RGB312
10.8.2GaussKernel3()：使用常量内存314
10.8.3处理常量的方法315
10.8.4GaussKernel4()：在共享内存中缓冲1个像素的邻居316
10.8.5GaussKernel5()：在共享内存中缓冲4个像素的邻居318
10.8.6GaussKernel6()：将5个垂直像素读入共享内存320
10.8.7GaussKernel7()：去除边界像素的影响322
10.8.8GaussKernel8()：计算8个垂直像素324
10.9CUDA占用率计算器326
10.9.1选择很好的线程/块327
10.9.2SM级资源328
10.9.3什么是占用率329
10.9.4CUDA占用率计算器：资源计算330
10.9.5案例分析：GaussKernel7(?)334
10.9.6案例分析：GaussKernel8(?)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.3CPU与GPU之间快速的数据传输346
11.3.1同步数据传输346
11.3.2异步数据传输347
11.4CUDA流的原理347
11.4.1CPU到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.5imGStr.cu：流式图像处理351
11.5.1将图像读入固定内存351
11.5.2同步与单个流353
11.5.3多个流354
11.5.4多流之间的数据依赖356
11.6流式水平翻转核函数360
11.7imGStr.cu：流式边缘检测362
11.8性能对比：imGStr.cu365
11.8.1同步与异步结果366
11.8.2结果的随机性366
11.8.3队列优化366
11.8.4很好流式结果367
11.8.5最差流式结果369
11.9Nvidia可视化分析器：nvvp370
11.9.1安装nvvp和nvprof370
11.9.2使用nvvp370
11.9.3使用nvprof371
11.9.4imGStr的同步和单流结果372
11.9.5imGStr的2流和4流结果373
第三部分拓展知识
第12章CUDA库376
12.1cuBLAS377
12.1.1BLAS级别377
12.1.2cuBLAS数据类型377
12.1.3安装cuBLAS378
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.11cuBLAS程序示例：矩阵的标量操作381
12.2cuFFT382
12.2.1cuFFT库特征383
12.2.2复数到复数变换示例383
12.2.3实数到复数变换示例384
12.3Nvidia性能库（NPP）384
12.4Thrust库386
第13章OpenCL简介388
13.1什么是OpenCL388
13.1.1多平台388
13.1.2基于队列389
13.2图像翻转核函数389
13.3运行核函数390
13.3.1选择设备390
13.3.2运行核函数392
13.3.3OpenCL程序的运行时间396
13.4OpenCL中的边缘检测396
第14章其他GPU编程语言402
14.1使用Python进行GPU编程402
14.1.1imflip的PyOpenCL版本403
14.1.2PyOpenCL的逐元素核函数406
14.2OpenGL408
14.3OpenGLES：用于嵌入式系统的OpenGL409
14.4Vulkan409
14.5微软的高级着色语言409
14.5.1着色410
14.5.2HLSL410
14.6Apple的MetalAPI411
14.7Apple的Swift编程语言411
14.8OpenCV411
14.8.1安装OpenCV和人脸识别411
14.8.2移动–微云–云实时人脸识别412
14.8.3加速即服务（AXaas）412
第15章深度学习中的CUDA413
15.1人工神经网络413
15.1.1神经元413
15.1.2激活函数414
15.2全连接神经网络415
15.3深度网络/卷积神经网络415
15.4训练网络416
15.5cuDNN深度学习库416
15.5.1创建一个层417
15.5.2创建一个网络418
15.5.3前向传播418
15.5.4反向传播419
15.5.5在网络中使用cuBLAS419
15.6Keras419
参考文献421
术语表424