CUDA专家手册：GPU编程权威指南

中文版序
推荐序
译者序
前　言
第一部分　基础知识
第1章　简介 2
1.1　方法 4
1.2　代码 4
1.2.1　验证型代码 5
1.2.2　演示型代码 5
1.2.3　探究型代码 5
1.3　资源 5
1.3.1　开源代码 5
1.3.2　CUDA专家手册库（chLib） 6
1.3.3　编码风格 6
1.3.4　CUDA SDK 6
1.4　结构 6
第2章　硬件架构 8
2.1　CPU配置 8
2.1.1　前端总线 9
2.1.2　对称处理器簇 9
2.1.3　非一致内存访问（NUMA） 10
2.1.4　集成的PCIe 12
2.2　集成GPU 13
2.3　多GPU 14
2.4　CUDA中的地址空间 17
2.4.1　虚拟寻址简史 17
2.4.2　不相交的地址空间 20
2.4.3　映射锁页内存 21
2.4.4　可分享锁页内存 21
2.4.5　统一寻址 23
2.4.6　点对点映射 24
2.5　CPU/GPU交互 24
2.5.1　锁页主机内存和命令缓冲区 25
2.5.2　CPU/GPU并发 26
2.5.3　主机接口和内部GPU同步 29
2.5.4　GPU间同步 31
2.6　GPU架构 31
2.6.1　综述 31
2.6.2　流处理器簇 34
2.7　延伸阅读 37
第3章　软件架构 39
3.1　软件层 39
3.1.1　CUDA运行时和驱动程序 40
3.1.2　驱动程序模型 41
3.1.3　nvcc、PTX和微码 43
3.2　设备与初始化 45
3.2.1　设备数量 46
3.2.2　设备属性 46
3.2.3　无CUDA支持情况 48
3.3　上下文 50
3.3.1　生命周期与作用域 51
3.3.2　资源预分配 51
3.3.3　地址空间 52
3.3.4　当前上下文栈 52
3.3.5　上下文状态 53
3.4　模块与函数 53
3.5　内核（函数） 55
3.6　设备内存 56
3.7　流与事件 57
3.7.1　软件流水线 57
3.7.2　流回调 57
3.7.3　NULL流 57
3.7.4　事件 58
3.8　主机内存 59
3.8.1　锁页主机内存 60
3.8.2　可分享的锁页内存 60
3.8.3　映射锁页内存 60
3.8.4　主机内存注册 60
3.9　CUDA数组与纹理操作 61
3.9.1　纹理引用 61
3.9.2　表面引用 63
3.10　图形互操作性 63
3.11　CUDA运行时与CUDA驱动程序API 65
第4章　软件环境 69
4.1　nvcc——CUDA编译器驱动程序 69
4.2　ptxas——PTX汇编工具 73
4.3　cuobjdump 76
4.4　nvidia-smi 77
4.5　亚马逊Web服务 79
4.5.1　命令行工具 79
4.5.2　EC2和虚拟化 79
4.5.3　密钥对 80
4.5.4　可用区域（AZ）和地理区域 81
4.5.5　S3 81
4.5.6　EBS 81
4.5.7　AMI 82
4.5.8　EC2上的Linux 82
4.5.9　EC2上的Windows 83
第二部分　CUDA编程
第5章　内存 88
5.1　主机内存 89
5.1.1　分配锁页内存 89
5.1.2　可共享锁页内存 90
5.1.3　映射锁页内存 90
5.1.4　写结合锁页内存 91
5.1.5　注册锁页内存 91
5.1.6　锁页内存与统一虚拟寻址 92
5.1.7　映射锁页内存用法 92
5.1.8　NUMA、线程亲和性与锁页内存 93
5.2　全局内存 95
5.2.1　指针 96
5.2.2　动态内存分配 97
5.2.3　查询全局内存数量 100
5.2.4　静态内存分配 101
5.2.5　内存初始化API 102
5.2.6　指针查询 103
5.2.7　点对点内存访问 104
5.2.8　读写全局内存 105
5.2.9　合并限制 105
5.2.10　验证实验：内存峰值带宽 107
5.2.11　原子操作 111
5.2.12　全局内存的纹理操作 113
5.2.13　ECC（纠错码） 113
5.3　常量内存 114
5.3.1　主机与设备常量内存 114
5.3.2　访问常量内存 114
5.4　本地内存 115
5.5　纹理内存 118
5.6　共享内存 118
5.6.1　不定大小共享内存声明 119
5.6.2　束同步编码 119
5.6.3　共享内存的指针 119
5.7　内存复制 119
5.7.1　同步内存复制与异步内存复制 120
5.7.2　统一虚拟寻址 121
5.7.3　CUDA运行时 121
5.7.4　驱动程序API 123
第6章　流与事件 125
6.1　CPU/GPU的并发：隐藏驱动程序开销 126
6.2　异步的内存复制 129
6.2.1　异步的内存复制：主机端到设备端 130
6.2.2　异步内存复制：设备端到主机端 130
6.2.3　NULL流和并发中断 131
6.3　CUDA事件：CPU/GPU同步 133
6.3.1　阻塞事件 135
6.3.2　查询 135
6.4　CUDA事件：计时 135
6.5　并发复制和内核处理 136
6.5.1　concurrencyMemcpyKernel.cu 137
6.5.2　性能结果 141
6.5.3　中断引擎间的并发性 142
6.6　映射锁页内存 143
6.7　并发内核处理 145
6.8　GPU/GPU同步：cudaStreamWaitEvent() 146
6.9　源代码参考 147
第7章　内核执行 148
7.1　概况 148
7.2　语法 149
7.2.1　局限性 150
7.2.2　高速缓存和一致性 151
7.2.3　异步与错误处理 151
7.2.4　超时 152
7.2.5　本地内存 152
7.2.6　共享内存 153
7.3　线程块、线程、线程束、束内线程 153
7.3.1　线程块网格 153
7.3.2　执行保证 156
7.3.3　线程块与线程ID 156
7.4　占用率 159
7.5　动态并行 160
7.5.1　作用域和同步 161
7.5.2　内存模型 162
7.5.3　流与事件 163
7.5.4　错误处理 163
7.5.5　编译和链接 164
7.5.6　资源管理 164
7.5.7　小结 165
第8章　流处理器簇 167
8.1　内存 168
8.1.1　寄存器 168
8.1.2　本地内存 169
8.1.3　全局内存 170
8.1.4　常量内存 171
8.1.5　共享内存 171
8.1.6　栅栏和一致性 173
8.2　整型支持 174
8.2.1　乘法 174
8.2.2　其他操作（位操作） 175
8.2.3　漏斗移位（SM 3.5） 175
8.3　浮点支持 176
8.3.1　格式 176
8.3.2　单精度（32位） 180
8.3.3　双精度（64位） 181
8.3.4　半精度（16位） 181
8.3.5　案例分析：float到half的转换 182
8.3.6　数学函数库 185
8.3.7　延伸阅读 190
8.4　条件代码 191
8.4.1　断定 191
8.4.2　分支与汇聚 191
8.4.3　特殊情况：最小值、最大值和绝对值 192
8.5　纹理与表面操作 193
8.6　其他指令 193
8.6.1　线程束级原语 193
8.6.2　线程块级原语 194
8.6.3　性能计数器 195
8.6.4　视频指令 195
8.6.5　特殊寄存器 196
8.7　指令集 196
第9章　多GPU 203
9.1　概述 203
9.2　点对点机制 204
9.2.1　点对点内存复制 204
9.2.2　点对点寻址 205
9.3　UVA：从地址推断设备 206
9.4　多GPU间同步 207
9.5　单线程多GPU方案 208
9.5.1　当前上下文栈 208
9.5.2　N-体问题 210
9.6　多线程多GPU方案 212
第10章　纹理操作 216
10.1　简介 216
10.2　纹理内存 217
10.2.1　设备内存 217
10.2.2　CUDA数组与块的线性寻址 218
10.2.3　设备内存与CUDA数组对比 222
10.3　一维纹理操作 223
10.4　纹理作为数据读取方式 226
10.4.1　增加有效地址范围 226
10.4.2　主机内存纹理操作 228
10.5　使用非归一化坐标的纹理操作 230
10.6　使用归一化坐标的纹理操作 237
10.7　一维表面内存的读写 238
10.8　二维纹理操作 240
10.9　二维纹理操作：避免复制 242
10.9.1　设备内存上的二维纹理操作 242
10.9.2　二维表面内存的读写 243
10.10　三维纹理操作 244
10.11　分层纹理 245
　10.11.1　一维分层纹理 246
　10.11.2　二维分层纹理 246
10.12　最优线程块大小选择以及性能 246
10.13　纹理操作快速参考 248
　10.13.1　硬件能力 248
　10.13.2　CUDA运行时 249
　10.13.3　驱动API 250
第三部分　实例
第11章　流式负载 254
11.1　设备内存 255
11.2　异步内存复制 258
11.3　流 259
11.4　映射锁页内存 260
11.5　性能评价与本章小结 261
第12章　归约算法 263
12.1　概述 263
12.2　两遍归约 265
12.3　单遍归约 269
12.4　使用原子操作的归约 271
12.5　任意线程块大小的归约 272
12.6　适应任意数据类型的归约 273
12.7　基于断定的归约 276
12.8　基于洗牌指令的线程束归约 277
第13章　扫描算法 278
13.1　定义与变形 278
13.2　概述 279
13.3　扫描和电路设计 281
13.4　CUDA实现 284
13.4.1　先扫描再扇出 284
13.4.2　先归约再扫描（递归） 288
13.4.3　先归约再扫描（两阶段） 291
13.5　线程束扫描 294
13.5.1　零填充 295
13.5.2　带模板的版本 296
13.5.3　线程束洗牌 297
13.5.4　指令数对比 298
13.6　流压缩 300
13.7　参考文献（并行扫描算法） 302
13.8　延伸阅读（并行前缀求和电路） 303
第14章　N-体问题 304
14.1　概述 305
14.2　简单实现 309
14.3　基于共享内存实现 312
14.4　基于常量内存实现 313
14.5　基于线程束洗牌实现 315
14.6　多GPU及其扩展性 316
14.7　CPU的优化 317
14.8　小结 321
14.9　参考文献与延伸阅读 323
第15章　图像处理的归一化相关系数计算 324
15.1　概述 324
15.2　简单的纹理实现 326
15.3　常量内存中的模板 329
15.4　共享内存中的图像 331
15.5　进一步优化 334
15.5.1　基于流处理器簇的实现代码 334
15.5.2　循环展开 335
15.6　源代码 336
15.7　性能评价 337
15.8　延伸阅读 339
附录A　CUDA专家手册库 340
术语表 347