●譯者序
前言
關於作者
部分理解CPU的並行性
章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.2imflip.c中的main()函數9
1.4.3垂直翻轉行：FlipImageV()10
1.4.4水平翻轉列：FlipImageH()11
1.5程序的編輯、編譯、運行12
1.5.1選擇編輯器和編譯器12
1.5.2在Windows7、8、10平臺上開發12
1.5.3在Mac平臺上開發14
1.5.4在Unix平臺上開發14
1.6Unix速成15
1.6.1與目錄相關的Unix命令15
1.6.2與文件相關的Unix命令16
1.7調試程序19
1.7.1gdb19
1.7.2古典調試方法20
1.7.3valgrind22
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.1imflipP.c中的main()函數27
2.1.2運行時間28
2.1.3imflipP.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.1BMP是一種無損/不壓縮的文件格式35
2.2.2BMP圖像文件格式36
2.2.3頭文件ImageStuff.h37
2.2.4ImageStuff.c中的圖像操作函數38
2.3執行線程任務40
2.3.1啟動線程41
2.3.2多線程垂直翻轉函數MTFlipV()43
2.3.3FlipImageV()和MTFlipV()的比較46
2.3.4多線程水平翻轉函數MTFlipH(?)47
2.4多線程代碼的測試/計時49
第3章改進個CPU並行程序51
3.1程序員對性能的影響51
3.2CPU對性能的影響52
3.2.1按序核心與亂序核心53
3.2.2瘦線程與胖線程55
3.3imf?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?lipP61
3.5.1分析MTFlipH()中的內存訪問模式62
3.5.2MTFlipH()的多線程內存訪問63
3.5.3DRAM訪問的規則64
3.6imflipPM：遵循DRAM的規則65
3.6.1imflipP的混亂內存訪問模式65
3.6.2改進imflipP的內存訪問模式65
3.6.3MTFlipHM()：內存友好的MTFlipH()66
3.6.4MTFlipVM()：內存友好的MTFlipV()69
3.7imflipPM.C的性能69
3.7.1imflipP.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架構：XeonPhi74
3.10GPU是怎樣的75
3.11本章小結76
第4章理解核心和內存77
4.1曾經的英特爾77
4.2CPU和內存制造商78
4.3動態存儲器與靜態存儲器79
4.3.1靜態隨機存取存儲器（SRAM）79
4.3.2動態隨機存取存儲器（DRAM）79
4.3.3DRAM接口標準79
4.3.4DRAM對程序性能的影響80
4.3.5SRAM對程序性能的影響81
4.4圖像旋轉程序：imrotate.c81
4.4.1imrotate.c的說明82
4.4.2imrotate.c：參數限制和簡化82
4.4.3imrotate.c：實現原理83
4.5imrotate的性能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.7imrotateMC：讓imrotate更高效94
4.7.1Rotate2()：平方根和浮點除法有多差96
4.7.2Rotate3()和Rotate4()：sin()和cos()有多差97
4.7.3Rotate5()：整數除法/乘法有多差98
4.7.4Rotate6()：合並計算100
4.7.5Rotate7()：合並更多計算100
4.7.6imrotateMC的總體性能101
4.8本章小結103
第5章線程管理和同步104
5.1邊緣檢測程序：imedge.c104
5.1.1imedge.c的說明105
5.1.2imedge.c：參數限制和簡化106
5.1.3imedge.c：實現原理106
5.2imedge.c：實現108
5.2.1初始化和時間戳109
5.2.2不同圖像表示的初始化函數110
5.2.3啟動和終止線程111
5.2.4高斯濾波112
5.2.5Sobel113
5.2.6閾值過濾114
5.3imedge的性能115
5.4imedgeMC：讓imedge更高效116
5.4.1利用預計算降低帶寬116
5.4.2存儲預計算的像素值117
5.4.3預計算像素值118
5.4.4讀取圖像並預計算像素值119
5.4.5PrGaussianFilter120
5.4.6PrSobel121
5.4.7PrThreshold121
5.5imedgeMC的性能122
5.6imedgeMCT：高效的線程同步123
5.6.1屏障同步124
5.6.2用於數據共享的MUTEX結構125
5.7imedgeMCT：實現127
5.7.1使用MUTEX：讀取圖像、預計算128
5.7.2一次預計算一行130
5.8imedgeMCT的性能131
第二部分基於CUDA的GPU編程
第6章GPU並行性和CUDA概述134
6.1曾經的Nvidia134
6.1.1GPU的誕生134
6.1.2早期的GPU架構136
6.1.3GPGPU的誕生137
6.1.4Nvidia、ATITechnologies和Intel138
6.2統一計算設備架構140
6.2.1CUDA、OpenCL和其他GPU語言140
6.2.2設備端與主機端代碼140
6.3理解GPU並行141
6.3.1GPU如何實現高性能142
6.3.2CPU與GPU架構的差異143
6.4圖像翻轉的CUDA版：imflipG.cu144
6.4.1imflipG.cu：將圖像讀入CPU端數組146
6.4.2初始化和查詢GPU147
6.4.3GPU端的時間戳148
6.4.4GPU端內存分配152
6.4.5GPU驅動程序和Nvidia運行時引擎153
6.4.6CPU到GPU的數據傳輸153
6.4.7用封裝函數進行錯誤報告154
6.4.8GPU核函數的執行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.14Vflip()：垂直翻轉的GPU核函數160
6.4.15什麼是線程ID、塊ID和塊維度163
6.4.16Hflip()：水平翻轉的GPU核函數165
6.4.17硬件參數：threadIDx.x、blockIdx.x和blockDim.x165
6.4.18PixCopy()：復制圖像的GPU核函數165
6.4.19CUDA關鍵字166
6.5Windows中的CUDA程序開發167
6.5.1安裝MSVisualStudio2015和CUDAToolkit8.0167
6.5.2在VisualStudio2015中創建項目imflipG.cu168
6.5.3在VisualStudio2015中編譯項目imflipG.cu170
6.5.4運行個CUDA應用程序：imflipG.exe173
6.5.5確保程序的正確性174
6.6Mac平臺上的CUDA程序開發175
6.6.1在Mac上安裝XCode175
6.6.2安裝CUDA驅動程序和CUDA工具包176
6.6.3在Mac上編譯和運行CUDA應用程序177
6.7Unix平臺上的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.3imflipG.cu：程序輸出的解釋183
7.1.4imflipG.cu：線程塊和圖像的大小對性能的影響184
7.2核函數的啟動185
7.2.1網格185
7.2.2線程塊187
7.2.3線程187
7.2.4線程束和通道189
7.3imf?lipG.cu：理解核函數的細節189
7.3.1在main()中啟動核函數並將參數傳遞給它們189
7.3.2線程執行步驟190
7.3.3Vflip()核函數191
7.3.4Vflip()和MTFlipV()的比較192
7.3.5Hflip()核函數194
7.3.6PixCopy()核函數194
7.4PCIe速度與CPU的關繫196
7.5PCIe總線對性能的影響196
7.5.1數據傳輸時間、速度、延遲、吞吐量和帶寬196
7.5.2imflipG.cu的PCIe吞吐量197
7.6全局內存總線對性能的影響200
7.7計算能力對性能的影響203
7.7.1Fermi、Kepler、Maxwell、Pascal和Volta繫列203
7.7.2不同繫列實現的相對帶寬204
7.7.3imflipG2.cu：計算能力2.0版本的imflipG.cu205
7.7.4imflipG2.cu：main()的修改206
7.7.5核函數PxCC20()208
7.7.6核函數VfCC20()208
7.8imflipG2.cu的性能210
7.9古典的CUDA調試方法212
7.9.1常見的CUDA錯誤212
7.9.2return調試法214
7.9.3基於注釋的調試216
7.9.4printf(?)調試216
7.10軟件錯誤的生物學原因217
7.10.1大腦如何參與編寫/調試代碼218
7.10.2當我們疲倦時是否會寫出錯誤代碼219
第8章理解GPU的硬件架構221
8.1GPU硬件架構222
8.2GPU硬件的部件222
8.2.1SM：流處理器222
8.2.2GPU核心223
8.2.3千兆線程調度器223
8.2.4內存控制器225
8.2.5共享高速緩存（L2$）225
8.2.6主機接口225
8.3NvidiaGPU架構226
8.3.1Fermi架構227
8.3.2GT、GTX和計算加速器227
8.3.3Kepler架構228
8.3.4Maxwell架構228
8.3.5Pascal架構和NVLink229
8.4CUDA邊緣檢測：imedgeG.cu229
8.4.1CPU和GPU內存中存儲圖像的變量229
8.4.2為GPU變量分配內存231
8.4.3調用核函數並對其進行計時233
8.4.4計算核函數的性能234
8.4.5計算核函數的數據移動量235
8.4.6輸出核函數性能236
8.5imedgeG：核函數237
8.5.1BWKernel()237
8.5.2GaussKernel()239
8.5.3SobelKernel()240
8.5.4ThresholdKernel()242
8.6imedgeG.cu的性能243
8.6.1imedgeG.cu：PCIe總線利用率244
8.6.2imedgeG.cu：運行時間245
8.6.3imedgeG.cu：核函數性能比較247
8.7GPU代碼：編譯時間248
8.7.1設計CUDA代碼248
8.7.2編譯CUDA代碼250
8.7.3GPU彙編：PTX、CUBIN250
8.8GPU代碼：啟動250
8.8.1操作繫統的參與和CUDADLL文件250
8.8.2GPU圖形驅動251
8.8.3CPU與GPU之間的內存傳輸251
8.9GPU代碼：執行（運行時間）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.1GPU的架構繫列258
9.1.1Fermi架構258
9.1.2FermiSM的結構259
9.1.3Kepler架構260
9.1.4KeplerSMX的結構260
9.1.5Maxwell架構261
9.1.6MaxwellSMM的結構262
9.1.7PascalGP100架構264
9.1.8PascalGP100SM的結構265
9.1.9繫列比較：峰值GFLOPS和峰值DGFLOPS266
9.1.10GPU睿頻267
9.1.11GPU功耗268
9.1.12計算機電源268
9.2流處理器的構建模塊269
9.2.1GPU核心269
9.2.2雙（DPU）270
9.2.3特殊（SFU）270
9.2.4寄存器文件（RF）270
9.2.5讀取/存儲隊列（LDST）271
9.2.6L1$和紋理高速緩存272
9.2.7共享內存272
9.2.8常量高速緩存272
9.2.9指令高速緩存272
9.2.10指令緩衝區272
9.2.11線程束調度器272
9.2.12273
9.3並行線程執行（PTX）的數據類型273
9.3.1INT8：8位整數274
9.3.2INT16：16位整數274
9.3.324位整數275
9.3.4INT32：32位整數275
9.3.5判定寄存器（32位）275
9.3.6INT64：64位整數276
9.3.7128位整數276
9.3.8FP32：單精度浮點（float）277
9.3.9FP64：雙精度浮點（double）277
9.3.10FP16：半精度浮點（half）278
9.3.11什麼是FLOP278
9.3.12融合乘法累加（FMA）與乘加（MAD）278
9.3.13四倍和八倍精度浮點279
9.3.14PascalGP104引擎的SM結構279
9.4imflipGC.cu：核心友好的imflipG280
9.4.1Hflip2()：預計算核函數參數282
9.4.2Vflip2()：預計算核函數參數284
9.4.3使用線程計算圖像坐標285
9.4.4線程塊ID與圖像的行映射285
9.4.5Hflip3()：使用二維啟動網格286
9.4.6Vflip3()：使用二維啟動網格287
9.4.7Hflip4()：計算2個連續的像素288
9.4.8Vflip4()：計算2個連續的像素289
9.4.9Hflip5()：計算4個連續的像素290
9.4.10Vflip5()：計算4個連續的像素291
9.4.11PixCopy2()、PixCopy3()：一次分別復制2個和4個連續的像素292
9.5imedgeGC.cu：核心友好的imedgeG293
9.5.1BWKernel2()：使用預計算的值和2D塊293
9.5.2GaussKernel2()：使用預計算的值和2D塊294
0章理解GPU內存296
10.1全局內存297
10.2L2高速緩存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.7imflipGCM.cu：核心和內存友好的imflipG301
10.7.1Hflip6()、Vflip6()：使用共享內存作為緩衝區301
10.7.2Hflip7()：共享內存中連續的交換操作303
10.7.3Hflip8()：使用寄存器交換4個像素305
10.7.4Vflip7()：一次復制4個字節（int）307
10.7.5對齊與未對齊的內存數據訪問308
10.7.6Vflip8()：一次復制8個字節308
10.7.7Vflip9()：僅使用全局內存，一次復制8個字節309
10.7.8PixCopy4()、PixCopy5()：使用共享內存復制1個和4個字節310
10.7.9PixCopy6()、PixCopy7()：使用全局內存復制1個和2個整數311
10.8imedgeGCM.cu：核心和內存友好的imedgeG312
10.8.1BWKernel3()：使用字節操作來提取RGB312
10.8.2GaussKernel3()：使用常量內存314
10.8.3處理常量的方法315
10.8.4GaussKernel4()：在共享內存中緩衝1個像素的鄰居316
10.8.5GaussKernel5()：在共享內存中緩衝4個像素的鄰居318
10.8.6GaussKernel6()：將5個垂直像素讀入共享內存320
10.8.7GaussKernel7()：去除邊界像素的影響322
10.8.8GaussKernel8()：計算8個垂直像素324
10.9CUDA占用率計算器326
10.9.1選擇最佳的線程/塊327
10.9.2SM級資源限制328
10.9.3什麼是占用率329
10.9.4CUDA占用率計算器：資源計算330
10.9.5案例分析：GaussKernel7(?)334
10.9.6案例分析：GaussKernel8(?)337
1章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.3CPU與GPU之間快速的數據傳輸346
11.3.1同步數據傳輸346
11.3.2異步數據傳輸347
11.4CUDA流的原理347
11.4.1CPU到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.5imGStr.cu：流式圖像處理351
11.5.1將圖像讀入固定內存351
11.5.2同步與單個流353
11.5.3多個流354
11.5.4多流之間的數據依賴356
11.6流式水平翻轉核函數360
11.7imGStr.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.9Nvidia可視化分析器：nvvp370
11.9.1安裝nvvp和nvprof370
11.9.2使用nvvp370
11.9.3使用nvprof371
11.9.4imGStr的同步和單流結果372
11.9.5imGStr的2流和4流結果373
第三部分拓展知識
2章CUDA庫376
12.1cuBLAS377
12.1.1BLAS級別377
12.1.2cuBLAS數據類型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.11cuBLAS程序示例：矩陣的標量操作381
12.2cuFFT382
12.2.1cuFFT庫特征383
12.2.2復數到復數變換示例383
12.2.3實數到復數變換示例384
12.3Nvidia性能庫（NPP）384
12.4Thrust庫386
3章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.3OpenCL程序的運行時間396
13.4OpenCL中的邊緣檢測396
4章其他GPU編程語言402
14.1使用Python進行GPU編程402
14.1.1imflip的PyOpenCL版本403
14.1.2PyOpenC素核函數406
14.2OpenGL408
14.3OpenGLES：用於嵌入式繫統的OpenGL409
14.4Vulkan409
14.5微軟的不錯著色語言409
14.5.1著色410
14.5.2HLSL410
14.6Apple的MetalAPI411
14.7Apple的Swift編程語言411
14.8OpenCV411
14.8.1安裝OpenCV和人臉識別411
14.8.2移動–微雲–雲實時人臉識別412
14.8.3加速即服務（AXaas）412
5章深度學習中的CUDA413
15.1人工神經網絡413
15.1.413
15.1.2激活函數414
15.2全連接神經網絡415
15.3深度網絡/卷積神經網絡415
15.4訓練網絡416
15.5cuDNN深度學習庫416
15.5.1創建一個層417
15.5.2創建一個網絡418
15.5.3前向傳播418
15.5.4反向傳播419
15.5.5在網絡中使用cuBLAS419
15.6Keras419
參考文獻421
術語表424