CUDA专家手册

中文版序

推荐序

译者序

前　言

第一部分　基础知识

第1章　简介

1.1　方法

1.2　代码

1.2.1　验证型代码

1.2.2　演示型代码

1.2.3　探究型代码

1.3　资源

1.3.1　开源代码

1.3.2　CUDA专家手册库（chLib）

1.3.3　编码风格

1.3.4　CUDA SDK

1.4　结构

第2章　硬件架构

2.1　CPU配置

2.1.1　前端总线

2.1.2　对称处理器簇

2.1.3　非一致内存访问（NUMA）

2.1.4　集成的PCIe

2.2　集成GPU

2.3　多GPU

2.4　CUDA中的地址空间

2.4.1　虚拟寻址简史

2.4.2　不相交的地址空间

2.4.3　映射锁页内存

2.4.4　可分享锁页内存

2.4.5　统一寻址

2.4.6　点对点映射

2.5　CPU/GPU交互

2.5.1　锁页主机内存和命令缓冲区

2.5.2　CPU/GPU并发

2.5.3　主机接口和内部GPU同步

2.5.4　GPU间同步

2.6　GPU架构

2.6.1　综述

2.6.2　流处理器簇

2.7　延伸阅读

第3章　软件架构

3.1　软件层

3.1.1　CUDA运行时和驱动程序

3.1.2　驱动程序模型

3.1.3　nvcc、PTX和微码

3.2　设备与初始化

3.2.1　设备数量

3.2.2　设备属性

3.2.3　无CUDA支持情况

3.3　上下文

3.3.1　生命周期与作用域

3.3.2　资源预分配

3.3.3　地址空间

3.3.4　当前上下文栈

3.3.5　上下文状态

3.4　模块与函数

3.5　内核（函数）

3.6　设备内存

3.7　流与事件

3.7.1　软件流水线

3.7.2　流回调

3.7.3　NULL流

3.7.4　事件

3.8　主机内存

3.8.1　锁页主机内存

3.8.2　可分享的锁页内存

3.8.3　映射锁页内存

3.8.4　主机内存注册

3.9　CUDA数组与纹理操作

3.9.1　纹理引用

3.9.2　表面引用

3.10　图形互操作性

3.11　CUDA运行时与CUDA驱动程序API

第4章　软件环境

4.1　nvcc--CUDA编译器驱动程序

4.2　ptxas--PTX汇编工具

4.3　cuobjdump

4.4　nvidia-smi

4.5　亚马逊Web服务

4.5.1　命令行工具

4.5.2　EC2和虚拟化

4.5.3　密钥对

4.5.4　可用区域（AZ）和地理区域

4.5.5　S3

4.5.6　EBS

4.5.7　AMI

4.5.8　EC2上的Linux

4.5.9　EC2上的Windows

第二部分　CUDA编程

第5章　内存

5.1　主机内存

5.1.1　分配锁页内存

5.1.2　可共享锁页内存

5.1.3　映射锁页内存

5.1.4　写结合锁页内存

5.1.5　注册锁页内存

5.1.6　锁页内存与统一虚拟寻址

5.1.7　映射锁页内存用法

5.1.8　NUMA、线程亲和性与锁页内存

5.2　全局内存

5.2.1　指针

5.2.2　动态内存分配

5.2.3　查询全局内存数量

5.2.4　静态内存分配

5.2.5　内存初始化API

5.2.6　指针查询

5.2.7　点对点内存访问

5.2.8　读写全局内存

5.2.9　合并限制

5.2.10　验证实验：内存峰值带宽

5.2.11　原子操作

5.2.12　全局内存的纹理操作

5.2.13　ECC（纠错码）

5.3　常量内存

5.3.1　主机与设备常量内存

5.3.2　访问常量内存

5.4　本地内存

5.5　纹理内存

5.6　共享内存

5.6.1　不定大小共享内存声明

5.6.2　束同步编码

5.6.3　共享内存的指针

5.7　内存复制

5.7.1　同步内存复制与异步内存复制

5.7.2　统一虚拟寻址

5.7.3　CUDA运行时

5.7.4　驱动程序API

第6章　流与事件

6.1　CPU/GPU的并发：隐藏驱动程序开销

6.2　异步的内存复制

6.2.1　异步的内存复制：主机端到设备端

6.2.2　异步内存复制：设备端到主机端

6.2.3　NULL流和并发中断

6.3　CUDA事件：CPU/GPU同步

6.3.1　阻塞事件

6.3.2　查询

6.4　CUDA事件：计时

6.5　并发复制和内核处理

6.5.1　concurrencyMemcpyKernel.cu

6.5.2　性能结果

6.5.3　中断引擎间的并发性

6.6　映射锁页内存

6.7　并发内核处理

6.8　GPU/GPU同步：cudaStreamWaitEvent()

6.9　源代码参考

第7章　内核执行

7.1　概况

7.2　语法

7.2.1　局限性

7.2.2　高速缓存和一致性

7.2.3　异步与错误处理

7.2.4　超时

7.2.5　本地内存

7.2.6　共享内存

7.3　线程块、线程、线程束、束内线程

7.3.1　线程块网格

7.3.2　执行保证

7.3.3　线程块与线程ID

7.4　占用率

7.5　动态并行

7.5.1　作用域和同步

7.5.2　内存模型

7.5.3　流与事件

7.5.4　错误处理

7.5.5　编译和链接

7.5.6　资源管理

7.5.7　小结

第8章　流处理器簇

8.1　内存

8.1.1　寄存器

8.1.2　本地内存

8.1.3　全局内存

8.1.4　常量内存

8.1.5　共享内存

8.1.6　栅栏和一致性

8.2　整型支持

8.2.1　乘法

8.2.2　其他操作（位操作）

8.2.3　漏斗移位（SM 3.5）

8.3　浮点支持

8.3.1　格式

8.3.2　单精度（32位）

8.3.3　双精度（64位）

8.3.4　半精度（16位）

8.3.5　案例分析：float到half的转换

8.3.6　数学函数库

8.3.7　延伸阅读

8.4　条件代码

8.4.1　断定

8.4.2　分支与汇聚

8.4.3　特殊情况：最小值、最大值和绝对值

8.5　纹理与表面操作

8.6　其他指令

8.6.1　线程束级原语

8.6.2　线程块级原语

8.6.3　性能计数器

8.6.4　视频指令

8.6.5　特殊寄存器

8.7　指令集

第9章　多GPU

9.1　概述

9.2　点对点机制

9.2.1　点对点内存复制

9.2.2　点对点寻址

9.3　UVA：从地址推断设备

9.4　多GPU间同步

9.5　单线程多GPU方案

9.5.1　当前上下文栈

9.5.2　N-体问题

9.6　多线程多GPU方案

第10章　纹理操作

10.1　简介

10.2　纹理内存

10.2.1　设备内存

10.2.2　CUDA数组与块的线性寻址

10.2.3　设备内存与CUDA数组对比

10.3　一维纹理操作

10.4　纹理作为数据读取方式

10.4.1　增加有效地址范围

10.4.2　主机内存纹理操作

10.5　使用非归一化坐标的纹理操作

10.6　使用归一化坐标的纹理操作

10.7　一维表面内存的读写

10.8　二维纹理操作

10.9　二维纹理操作：避免复制

10.9.1　设备内存上的二维纹理操作

10.9.2　二维表面内存的读写

10.10　三维纹理操作

10.11　分层纹理

10.11.1　一维分层纹理

10.11.2　二维分层纹理

10.12　最优线程块大小选择以及性能

10.13　纹理操作快速参考

10.13.1　硬件能力

10.13.2　CUDA运行时

10.13.3　驱动API

第三部分　实例

第11章　流式负载

11.1　设备内存

11.2　异步内存复制

11.3　流

11.4　映射锁页内存

11.5　性能评价与本章小结

第12章　归约算法

12.1　概述

12.2　两遍归约

12.3　单遍归约

12.4　使用原子操作的归约

12.5　任意线程块大小的归约

12.6　适应任意数据类型的归约

12.7　基于断定的归约

12.8　基于洗牌指令的线程束归约

第13章　扫描算法

13.1　定义与变形

13.2　概述

13.3　扫描和电路设计

13.4　CUDA实现

13.4.1　先扫描再扇出

13.4.2　先归约再扫描（递归）

13.4.3　先归约再扫描（两阶段）

13.5　线程束扫描

13.5.1　零填充

13.5.2　带模板的版本

13.5.3　线程束洗牌

13.5.4　指令数对比

13.6　流压缩

13.7　参考文献（并行扫描算法）

13.8　延伸阅读（并行前缀求和电路）

第14章　N-体问题

14.1　概述

14.2　简单实现

14.3　基于共享内存实现

14.4　基于常量内存实现

14.5　基于线程束洗牌实现

14.6　多GPU及其扩展性

14.7　CPU的优化

14.8　小结

14.9　参考文献与延伸阅读

第15章　图像处理的归一化相关系数计算

15.1　概述

15.2　简单的纹理实现

15.3　常量内存中的模板

15.4　共享内存中的图像

15.5　进一步优化

15.5.1　基于流处理器簇的实现代码

15.5.2　循环展开

15.6　源代码

15.7　性能评价

15.8　延伸阅读

附录A　CUDA专家手册库

术语表

《CUDA专家手册：GPU编程权威指南》详细讨论CUDA的硬件和软件，包括CUDA 5.0和开普勒架构的最新特性。每个CUDA开发人员，不论新手还是高手，都可以在这里找到感兴趣的内容并即时上手。新晋的CUDA开发者将理解硬件如何处理命令以及驱动程序如何检查状态；更有经验者，将会在驱动程序API、上下文迁移以及如何让CPU/GPU最有效率地进行数据交换和同步等骨灰级的主题上得到指导。
　　本书所附的开源代码有25000多行，欢迎开发者自由重用。
　　本书不仅是权威手册，也是实用代码大全。全书分为以下三个部分：
　　第一部分是基础知识概述，对支持CUDA的硬件和软件进行高屋建瓴的描述。
　　第二部分是CUDA编程细节，对CUDA进行全方位的描述，包括内存，流和事件，执行模型（包括动态并行特性以及CUDA 5.0和SM 3.5的新特性），流处理器簇（包括SM 3.5的所有功能介绍），多GPU编程，纹理操作。这部分附带的源代码作为可重用的验证型代码和演示型代码，旨在展示特殊的硬件特性或强调特定的应用方法。
　　第三部分是案例剖析，深入分析精选的CUDA应用场景以及关键的并行算法，包括流式负载、归约、扫描（并行前缀求和）、N-体问题和图像处理，这些算法全方位涵盖各种CUDA应用场景。

英伟达公司Nicholas Wilt亲笔撰写，英伟达中国首批CUDA官方认证工程师翻译；
全面而系统地讲解CUDA编程的各方面知识，深度解析CUDA各种优化技术，包含大量实用代码示例，是深入掌握主流异构并行计算技术的指南。