基于CUDA的GPU并行程序开发指南2025|PDF|Epub|mobi|kindle电子书版本百度云盘下载

基于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 gdb19

1.7.2 古典调试方法20

1.7.3 valgrind22

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.h37

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 改进imflipP61

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 Phi74

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.c81

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/O93

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.c104

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 Sobel113

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 PrGaussianFilter120

5.4.6 PrSobel121

5.4.7 PrThreshold121

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 曾经的Nvidia134

6.1.1 GPU的诞生134

6.1.2 早期的GPU架构136

6.1.3 GPGPU的诞生137

6.1.4 Nvidia、ATI Technologies和Intel138

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.cu144

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

6.4.2 初始化和查询GPU147

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结果传回CPU158

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.x165

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.0167

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

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

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

6.5.5 确保程序的正确性174

6.6 Mac平台上的CUDA程序开发175

6.6.1 在Mac上安装XCode175

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.cu205

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架构和NVLink229

8.4 CUDA边缘检测：imedgeG.cu229

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、CUBIN250

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和峰值DGFLOPS266

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 什么是FLOP278

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

9.3.13 四倍和八倍精度浮点279

9.3.14 Pascal GP104引擎的SM结构279

9.4 imflipGC.cu：核心友好的imflipG280

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：核心友好的imedgeG293

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：核心和内存友好的imflipG301

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：核心和内存友好的imedgeG312

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

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.cu365

11.8.1 同步与异步结果366

11.8.2 结果的随机性366

11.8.3 队列优化366

11.8.4 最佳流式结果367

11.8.5 最差流式结果369

11.9 Nvidia可视化分析器：nvvp370

11.9.1 安装nvvp和nvprof370

11.9.2 使用nvvp370

11.9.3 使用nvprof371

11.9.4 imGStr的同步和单流结果372

11.9.5 imGStr的2流和4流结果373

第三部分 拓展知识376

第12章 CUDA库376

12.1 cuBLAS377

12.1.1 BLAS级别377

12.1.2 cuBLAS数据类型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.11 cuBLAS程序示例：矩阵的标量操作381

12.2 cuFFT382

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 什么是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.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 OpenGL408

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

14.4 Vulkan409

14.5 微软的高级着色语言409

14.5.1 着色410

14.5.2 HLSL410

14.6 Apple的MetalAPI411

14.7 Apple的Swift编程语言411

14.8 OpenCV411

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.5 cuDNN深度学习库416

15.5.1 创建一个层417

15.5.2 创建一个网络418

15.5.3 前向传播418

15.5.4 反向传播419

15.5.5 在网络中使用cuBLAS419

15.6 Keras419

参考文献421

术语表424