本书旨在帮助读者了解与基于CUDA的并行编程技术有关的基本概念，并掌握实用c语言进行GPU高性能编程的相关技巧。本书第一部分通过CPU多线程编程解释了并行计算，使得没有太多并行计算基础的读者也能毫无阻碍地进入CUDA天地；第二部分重点介绍了基于CUDA的GPU大规模并行程序的开发与实现，并通过大量的性能分析帮助读者理解如何开发一个好的GPU并行程序以及GPU架构对程序性能的影响；本书的第三部分介绍了一些常用的CUDA库。
本书内容详尽、实例丰富，可作为高等院校相关专业高年级本科生和研究生课程的教材，也可作为有关专业研究生或计算机技术人员的参考书。

作者结合其多年CUDA教学经验以及与Nvidia公司多年合作积累的工程经验精心撰写
紧扣深度学习热点，包含cuDNN技术
通俗易懂、循序渐进，是学习CUDA编程的最佳选择
GPU Parallel Program Development Using CUDA
基于CUDA的
GPU并行程序开发指南
[ 美 ]　托尔加·索亚塔(Tolga Soyata)　　 著
唐杰　 译

我经历过在IBM大型机上编写汇编语言来开发高性能程序的日子。用穿孔卡片编写程序，编译需要一天时间；你要留下在穿孔卡片上编写的程序，第二天再来拿结果。如果出现错误，你需要重复这些操作。在那些日子里，一位优秀的程序员必须理解底层的机器硬件才能编写出好的代码。当我看到现在的计算机科学专业的学生只学习抽象层次较高的内容以及像Ruby这样的语言时，我总会感到有些焦虑。尽管抽象是一件好事，因为它可以避免由于不必要的细节而使程序开发陷入困境，但当你尝试开发高性能代码时，抽象就变成了一件坏事。
自第一个CPU出现以来，计算机架构师在CPU硬件中添加了令人难以置信的功能来“容忍”糟糕的编程技巧。20年前，你必须手动设置机器指令的执行顺序，而如今在硬件中CPU会为你做这些（例如，乱序执行）。在GPU世界中也能清晰地看到类似的趋势。由于GPU架构师正在改进硬件功能，5年前我们在GPU编程中学习的大多数性能提升技术（例如，线程发散、共享存储体冲突以及减少原子操作的使用）正变得与改进的GPU架构越来越不相关，甚至5～10年后，即使是一名非常马虎的程序员，这些因素也会变得无关紧要。当然，这只是一个猜测。GPU架构师可以做的事取决于晶体管总数及客户需求。当说晶体管总数时，是指GPU制造商可以将多少个晶体管封装到集成电路（IC）即“芯片”中。当说客户需求时，是指即使GPU架构师能够实现某个功能，但如果客户使用的应用程序不能从中受益，就意味着浪费了部分的晶体管数量。
从编写教科书的角度出发，我考虑了所有的因素，逐渐明确讲授GPU编程的最佳方式是说明不同系列GPU（如Fermi、Kepler、Maxwell和Pascal）之间的不同并指明发展趋势，这可以让读者准备好迎接即将到来的下一代GPU，再下一代，……我会重点强调那些相对来说会长期存在的概念，同时也关注那些与平台相关的概念。也就是说，GPU编程完全关乎性能，如果你了解程序运行的平台架构，编写出了与平台相关的代码，就可以获得更高的性能。所以，提供平台相关的解释与通用的GPU概念一样有价值。本书内容的设计方式是，越靠后的章节，内容越具有平台特定性。
我认为本书最独特的地方就是通过第一部分中的CPU多线程来解释并行。第二部分介绍了GPU的大规模并行（与CPU的并行不同）。由于第一部分解释了CPU并行的方式，因此读者在第二部分中可以较为容易地理解GPU的并行。在过去的6年中，我设计了这种方法来讲授GPU编程，认识到从未学过并行编程课程的学生并不是很清楚大规模并行的概念。与GPU相比，“并行化任务”的概念在CPU架构中更容易理解。
本书的组织如下。第一部分（第1章至第5章）使用一些简单的程序来演示如何将大任务分成多个并行的子任务并将它们映射到CPU线程，分析了同一任务的多种并行实现方式，并根据计算核心和存储单元操作来研究这些方法的优缺点。本书的第二部分（第6章至第11章）将同一个程序在多个Nvidia GPU平台（Fermi、Kepler、Maxwell和Pascal）上并行化，并进行性能分析。由于CPU和GPU的核心和内存结构不同，分析结果的差异有时很有趣，有时与直觉相反。本书指出了这些结果的不同之处，并讨论了如何让GPU代码运行得更快。本书的最终目标是让程序员了解所有的做法，这样他们就可以应用好的做法，并避免将不好的做法应用到项目中。
尽管第一部分和第二部分已经完全涵盖了编写一个好的CUDA程序需要的所有内容，但总会有更多需要了解的东西。本书的第三部分为希望拓宽视野的读者指明了方向。第三部分并不是相关主题的详细参考文档，只是给出了一些入门介绍，读者可以从中获得学习这些内容的动力。这部分主要介绍了一些流行的CUDA库，比如cuBLAS、cuFFT、Nvidia Performance Primitives和Thrust（第12章）；OpenCL编程语言（第13章）；使用其他编程语言和API库进行GPU编程，包括Python、Metal、Swift、OpenGL、OpenGL ES、OpenCV和微软HLSL（第14章）；深度学习库cuDNN（第15章）。
书中代码的下载地址为：https://www.crcpress.com/GPU-Parallel-ProgramDevelopment-Using- CUDA /Soyata/p/book/9781498750752。

Tolga Soyata

计算机科学及应用

本书详细介绍GPU编程方法，展示了不同的GPU系列之间的区别。这样的讲授方式能使读者更好地应对下一代或未来的GPU系列。本书重点关注基础性概念，同时也关注与平台相关的概念。提供平台相关的内容与通用的GPU概念一样有价值。
本书分为三个部分，第一部分利用CPU的多线程概念介绍了并行性的基础知识，通过一些简单的代码示例展示了如何将一个大任务分成若干并行的子任务，然后将它们映射到CPU线程。对于同一个任务，尝试了多种不同的并行实现方法，并针对核心和内存操作分析了这些方法的优缺点。
第二部分介绍了GPU的大规模并行性。在多个Nvidia GPU平台上对相同的程序进行了并行化实现，并分析了它们的性能。由于CPU和GPU的核心与内存结构不同，运行结果也显示出了很多有趣的差异。最终目标是希望程序员既能掌握有益的思想，也能了解那些不好的做法，这样在自己开发的程序中能采用有用的策略，同时避免犯错。
第三部分为想要拓展知识的读者提供了建议。该部分简要介绍了一些常用的CUDA库（例如cuBLAS、cuFFT、NPP和Thrust），OpenCL编程语言，其他GPU编程语言和API（如Python、OpenCV、OpenGL、Swift与Metal），以及深度学习库cuDNN。

[美]托尔加·索亚塔(Tolga Soyata) 著：托尔加·索亚塔(Tolga Soyata) 是纽约州立大学奥尔巴尼分校电气与计算机工程系副教授。主要教授大规模集成电路、GPU编程课程，曾担任罗切斯特大学CUDA教学中心和CUDA研究中心负责人。

近10年来，随着大数据、深度学习等相关领域的发展，对计算能力的需求呈几何级数增长。与此同时，大规模集成电路的发展却受到功耗、散热、晶体管尺寸等客观因素的限制，难以继续维持摩尔定律。因此，人们逐渐把目光转向了并行系统。GPU自诞生之日起就是为计算机的图形图像渲染等大规模并行处理任务而服务的，因而越来越受到研究界和企业界的关注。随着CUDA等计算架构模型的出现，这一趋势更加明显。
CUDA（Compute Unified Device Architecture，统一计算设备架构）是Nvidia（英伟达）提出的并行计算架构，它可以结合CPU和GPU的优点，处理大规模的计算密集型任务。同时，它采用了基于C语言风格的语法，又将CPU端和GPU端的开发有效地集成到了同一环境中，对于大多数C程序员来说，使用十分方便，因而一经推出就迅速占领了GPU开发环境的市场。
然而，会写CUDA程序与会写好的CUDA程序相差甚远！
阻碍CUDA程序获得高性能的原因有很多。首先，GPU属于单指令多数据类型的并行计算，因而任务切分方式非常关键，既要充分挖掘线程级的并行性，也要充分利用流来实现任务级的并行。其次，GPU的存储类型和访问模式比CPU的要丰富得多，一个成功的CUDA程序要能充分利用不同类型的存储。再次，Nvidia GPU的架构还处于高速发展期，新一代GPU所推出的新功能也能够有效地提升计算效率。最后，万丈高楼平地起并不是CUDA开发的最佳方式，Nvidia和一些第三方机构都开发了很多基于CUDA的支撑库，利用好这些第三方库可以让你的开发过程事半功倍。
Tolga Soyata结合他10多年的CUDA教学经验以及与Nvidia多年合作的经历精心撰写了本书，针对上述问题进行了详细而生动的阐述。本书最独特的地方是它在第一部分中通过CPU多线程解释并行计算，使没有太多并行计算基础的读者也能毫无阻碍地进入CUDA天地。第二部分重点介绍了基于CUDA的GPU大规模并行程序的开发与实现。与现有的同类书籍相比，本书的特点是在多个Nvidia GPU平台（Fermi、Kepler、Maxwell和Pascal）上并行化，并进行性能分析，帮助读者理解GPU架构对程序性能的影响。第三部分介绍了一些重要的CUDA库，比如cuBLAS、cuFFT、NPP和Thrust（第12章）；OpenCL编程语言（第13章）；使用其他编程语言和API库进行GPU编程，包括Python、Metal、Swift、OpenGL、OpenGL ES、OpenCV和微软HLSL（第14章）；当下流行的深度学习库cuDNN（第15章）。
本书通过生动的类比、大量的代码和详细的解释向读者循序渐进地介绍了基于CUDA编程开发的GPU并行计算方法，内容丰富翔实，适合所有具备基本的C语言知识的程序员阅读，也适合作为GPU并行计算相关课程的教材。
在本书的翻译过程中得到了机械工业出版社华章分社朱捷先生的大力支持，在此表示由衷的感谢！
限于水平，翻译中难免有错误或不妥之处，真诚希望各位读者批评指正。

唐　杰
2019年1月

译者序
前言
关于作者
第一部分　理解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.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　imf?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?lipP 61
3.5.1　 分析MTFlipH( )中的内存访问模式 62
3.5.2 　MTFlipH( )的多线程内存访问 63
3.5.3 　DRAM访问的规则 64
3.6　imf?lipPM：遵循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编程
第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　imf?lipG.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　Maxwell SMM的结构 262
9.1.7　Pascal GP100架构 264
9.1.8　Pascal GP100 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　BWKernel2( )：使用预计算的值和2D块 293
9.5.2　GaussKernel2( )：使用预计算的值和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　BWKernel3( )：使用字节操作来提取RGB 312
10.8.2　GaussKernel3( )：使用常量内存 314
10.8.3　处理常量的方法 315
10.8.4　GaussKernel4( )：在共享内存中缓冲1个像素的邻居 316
10.8.5　GaussKernel5( )：在共享内存中缓冲4个像素的邻居 318
10.8.6　GaussKernel6( )：将5个垂直像素读入共享内存 320
10.8.7　GaussKernel7( )：去除边界像素的影响 322
10.8.8　GaussKernel8( )：计算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　案例分析：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.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
第三部分　拓展知识
第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的Metal API 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