CUDA專家手冊(cè)：GPU編程權(quán)威指南

中文版序
推薦序
譯者序
前　言
第一部分　基礎(chǔ)知識(shí)
第1章　簡介 2
1.1　方法 4
1.2　代碼 4
1.2.1　驗(yàn)證型代碼 5
1.2.2　演示型代碼 5
1.2.3　探究型代碼 5
1.3　資源 5
1.3.1　開源代碼 5
1.3.2　CUDA專家手冊(cè)庫（chLib） 6
1.3.3　編碼風(fēng)格 6
1.3.4　CUDA SDK 6
1.4　結(jié)構(gòu) 6
第2章　硬件架構(gòu) 8
2.1　CPU配置 8
2.1.1　前端總線 9
2.1.2　對(duì)稱處理器簇 9
2.1.3　非一致內(nèi)存訪問（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　映射鎖頁內(nèi)存 21
2.4.4　可分享鎖頁內(nèi)存 21
2.4.5　統(tǒng)一尋址 23
2.4.6　點(diǎn)對(duì)點(diǎn)映射 24
2.5　CPU/GPU交互 24
2.5.1　鎖頁主機(jī)內(nèi)存和命令緩沖區(qū) 25
2.5.2　CPU/GPU并發(fā) 26
2.5.3　主機(jī)接口和內(nèi)部GPU同步 29
2.5.4　GPU間同步 31
2.6　GPU架構(gòu) 31
2.6.1　綜述 31
2.6.2　流處理器簇 34
2.7　延伸閱讀 37
第3章　軟件架構(gòu) 39
3.1　軟件層 39
3.1.1　CUDA運(yùn)行時(shí)和驅(qū)動(dòng)程序 40
3.1.2　驅(qū)動(dòng)程序模型 41
3.1.3　nvcc、PTX和微碼 43
3.2　設(shè)備與初始化 45
3.2.1　設(shè)備數(shù)量 46
3.2.2　設(shè)備屬性 46
3.2.3　無CUDA支持情況 48
3.3　上下文 50
3.3.1　生命周期與作用域 51
3.3.2　資源預(yù)分配 51
3.3.3　地址空間 52
3.3.4　當(dāng)前上下文棧 52
3.3.5　上下文狀態(tài) 53
3.4　模塊與函數(shù) 53
3.5　內(nèi)核（函數(shù)） 55
3.6　設(shè)備內(nèi)存 56
3.7　流與事件 57
3.7.1　軟件流水線 57
3.7.2　流回調(diào) 57
3.7.3　NULL流 57
3.7.4　事件 58
3.8　主機(jī)內(nèi)存 59
3.8.1　鎖頁主機(jī)內(nèi)存 60
3.8.2　可分享的鎖頁內(nèi)存 60
3.8.3　映射鎖頁內(nèi)存 60
3.8.4　主機(jī)內(nèi)存注冊(cè) 60
3.9　CUDA數(shù)組與紋理操作 61
3.9.1　紋理引用 61
3.9.2　表面引用 63
3.10　圖形互操作性 63
3.11　CUDA運(yùn)行時(shí)與CUDA驅(qū)動(dòng)程序API 65
第4章　軟件環(huán)境 69
4.1　nvcc——CUDA編譯器驅(qū)動(dòng)程序 69
4.2　ptxas——PTX匯編工具 73
4.3　cuobjdump 76
4.4　nvidia-smi 77
4.5　亞馬遜Web服務(wù) 79
4.5.1　命令行工具 79
4.5.2　EC2和虛擬化 79
4.5.3　密鑰對(duì) 80
4.5.4　可用區(qū)域（AZ）和地理區(qū)域 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章　內(nèi)存 88
5.1　主機(jī)內(nèi)存 89
5.1.1　分配鎖頁內(nèi)存 89
5.1.2　可共享鎖頁內(nèi)存 90
5.1.3　映射鎖頁內(nèi)存 90
5.1.4　寫結(jié)合鎖頁內(nèi)存 91
5.1.5　注冊(cè)鎖頁內(nèi)存 91
5.1.6　鎖頁內(nèi)存與統(tǒng)一虛擬尋址 92
5.1.7　映射鎖頁內(nèi)存用法 92
5.1.8　NUMA、線程親和性與鎖頁內(nèi)存 93
5.2　全局內(nèi)存 95
5.2.1　指針 96
5.2.2　動(dòng)態(tài)內(nèi)存分配 97
5.2.3　查詢?nèi)謨?nèi)存數(shù)量 100
5.2.4　靜態(tài)內(nèi)存分配 101
5.2.5　內(nèi)存初始化API 102
5.2.6　指針查詢 103
5.2.7　點(diǎn)對(duì)點(diǎn)內(nèi)存訪問 104
5.2.8　讀寫全局內(nèi)存 105
5.2.9　合并限制 105
5.2.10　驗(yàn)證實(shí)驗(yàn)：內(nèi)存峰值帶寬 107
5.2.11　原子操作 111
5.2.12　全局內(nèi)存的紋理操作 113
5.2.13　ECC（糾錯(cuò)碼） 113
5.3　常量內(nèi)存 114
5.3.1　主機(jī)與設(shè)備常量內(nèi)存 114
5.3.2　訪問常量內(nèi)存 114
5.4　本地內(nèi)存 115
5.5　紋理內(nèi)存 118
5.6　共享內(nèi)存 118
5.6.1　不定大小共享內(nèi)存聲明 119
5.6.2　束同步編碼 119
5.6.3　共享內(nèi)存的指針 119
5.7　內(nèi)存復(fù)制 119
5.7.1　同步內(nèi)存復(fù)制與異步內(nèi)存復(fù)制 120
5.7.2　統(tǒng)一虛擬尋址 121
5.7.3　CUDA運(yùn)行時(shí) 121
5.7.4　驅(qū)動(dòng)程序API 123
第6章　流與事件 125
6.1　CPU/GPU的并發(fā)：隱藏驅(qū)動(dòng)程序開銷 126
6.2　異步的內(nèi)存復(fù)制 129
6.2.1　異步的內(nèi)存復(fù)制：主機(jī)端到設(shè)備端 130
6.2.2　異步內(nèi)存復(fù)制：設(shè)備端到主機(jī)端 130
6.2.3　NULL流和并發(fā)中斷 131
6.3　CUDA事件：CPU/GPU同步 133
6.3.1　阻塞事件 135
6.3.2　查詢 135
6.4　CUDA事件：計(jì)時(shí) 135
6.5　并發(fā)復(fù)制和內(nèi)核處理 136
6.5.1　concurrencyMemcpyKernel.cu 137
6.5.2　性能結(jié)果 141
6.5.3　中斷引擎間的并發(fā)性 142
6.6　映射鎖頁內(nèi)存 143
6.7　并發(fā)內(nèi)核處理 145
6.8　GPU/GPU同步：cudaStreamWaitEvent() 146
6.9　源代碼參考 147
第7章　內(nèi)核執(zhí)行 148
7.1　概況 148
7.2　語法 149
7.2.1　局限性 150
7.2.2　高速緩存和一致性 151
7.2.3　異步與錯(cuò)誤處理 151
7.2.4　超時(shí) 152
7.2.5　本地內(nèi)存 152
7.2.6　共享內(nèi)存 153
7.3　線程塊、線程、線程束、束內(nèi)線程 153
7.3.1　線程塊網(wǎng)格 153
7.3.2　執(zhí)行保證 156
7.3.3　線程塊與線程ID 156
7.4　占用率 159
7.5　動(dòng)態(tài)并行 160
7.5.1　作用域和同步 161
7.5.2　內(nèi)存模型 162
7.5.3　流與事件 163
7.5.4　錯(cuò)誤處理 163
7.5.5　編譯和鏈接 164
7.5.6　資源管理 164
7.5.7　小結(jié) 165
第8章　流處理器簇 167
8.1　內(nèi)存 168
8.1.1　寄存器 168
8.1.2　本地內(nèi)存 169
8.1.3　全局內(nèi)存 170
8.1.4　常量內(nèi)存 171
8.1.5　共享內(nèi)存 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　浮點(diǎn)支持 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的轉(zhuǎn)換 182
8.3.6　數(shù)學(xué)函數(shù)庫 185
8.3.7　延伸閱讀 190
8.4　條件代碼 191
8.4.1　斷定 191
8.4.2　分支與匯聚 191
8.4.3　特殊情況：最小值、最大值和絕對(duì)值 192
8.5　紋理與表面操作 193
8.6　其他指令 193
8.6.1　線程束級(jí)原語 193
8.6.2　線程塊級(jí)原語 194
8.6.3　性能計(jì)數(shù)器 195
8.6.4　視頻指令 195
8.6.5　特殊寄存器 196
8.7　指令集 196
第9章　多GPU 203
9.1　概述 203
9.2　點(diǎn)對(duì)點(diǎn)機(jī)制 204
9.2.1　點(diǎn)對(duì)點(diǎn)內(nèi)存復(fù)制 204
9.2.2　點(diǎn)對(duì)點(diǎn)尋址 205
9.3　UVA：從地址推斷設(shè)備 206
9.4　多GPU間同步 207
9.5　單線程多GPU方案 208
9.5.1　當(dāng)前上下文棧 208
9.5.2　N-體問題 210
9.6　多線程多GPU方案 212
第10章　紋理操作 216
10.1　簡介 216
10.2　紋理內(nèi)存 217
10.2.1　設(shè)備內(nèi)存 217
10.2.2　CUDA數(shù)組與塊的線性尋址 218
10.2.3　設(shè)備內(nèi)存與CUDA數(shù)組對(duì)比 222
10.3　一維紋理操作 223
10.4　紋理作為數(shù)據(jù)讀取方式 226
10.4.1　增加有效地址范圍 226
10.4.2　主機(jī)內(nèi)存紋理操作 228
10.5　使用非歸一化坐標(biāo)的紋理操作 230
10.6　使用歸一化坐標(biāo)的紋理操作 237
10.7　一維表面內(nèi)存的讀寫 238
10.8　二維紋理操作 240
10.9　二維紋理操作：避免復(fù)制 242
10.9.1　設(shè)備內(nèi)存上的二維紋理操作 242
10.9.2　二維表面內(nèi)存的讀寫 243
10.10　三維紋理操作 244
10.11　分層紋理 245
　10.11.1　一維分層紋理 246
　10.11.2　二維分層紋理 246
10.12　最優(yōu)線程塊大小選擇以及性能 246
10.13　紋理操作快速參考 248
　10.13.1　硬件能力 248
　10.13.2　CUDA運(yùn)行時(shí) 249
　10.13.3　驅(qū)動(dòng)API 250
第三部分　實(shí)例
第11章　流式負(fù)載 254
11.1　設(shè)備內(nèi)存 255
11.2　異步內(nèi)存復(fù)制 258
11.3　流 259
11.4　映射鎖頁內(nèi)存 260
11.5　性能評(píng)價(jià)與本章小結(jié) 261
第12章　歸約算法 263
12.1　概述 263
12.2　兩遍歸約 265
12.3　單遍歸約 269
12.4　使用原子操作的歸約 271
12.5　任意線程塊大小的歸約 272
12.6　適應(yīng)任意數(shù)據(jù)類型的歸約 273
12.7　基于斷定的歸約 276
12.8　基于洗牌指令的線程束歸約 277
第13章　掃描算法 278
13.1　定義與變形 278
13.2　概述 279
13.3　掃描和電路設(shè)計(jì) 281
13.4　CUDA實(shí)現(xiàn) 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　指令數(shù)對(duì)比 298
13.6　流壓縮 300
13.7　參考文獻(xiàn)（并行掃描算法） 302
13.8　延伸閱讀（并行前綴求和電路） 303
第14章　N-體問題 304
14.1　概述 305
14.2　簡單實(shí)現(xiàn) 309
14.3　基于共享內(nèi)存實(shí)現(xiàn) 312
14.4　基于常量內(nèi)存實(shí)現(xiàn) 313
14.5　基于線程束洗牌實(shí)現(xiàn) 315
14.6　多GPU及其擴(kuò)展性 316
14.7　CPU的優(yōu)化 317
14.8　小結(jié) 321
14.9　參考文獻(xiàn)與延伸閱讀 323
第15章　圖像處理的歸一化相關(guān)系數(shù)計(jì)算 324
15.1　概述 324
15.2　簡單的紋理實(shí)現(xiàn) 326
15.3　常量內(nèi)存中的模板 329
15.4　共享內(nèi)存中的圖像 331
15.5　進(jìn)一步優(yōu)化 334
15.5.1　基于流處理器簇的實(shí)現(xiàn)代碼 334
15.5.2　循環(huán)展開 335
15.6　源代碼 336
15.7　性能評(píng)價(jià) 337
15.8　延伸閱讀 339
附錄A　CUDA專家手冊(cè)庫 340
術(shù)語表 347