GPU & CPU程式設計

     GPU除了用處圖形渲染領域外，還可以用來做大規模的并行運算，這裡我們稱其為GPGPU(General Purpose GPU)；GPGPU計算通常采用CPU+GPU異構模式，由CPU負責執行複雜邏輯處理和事務管理等不适合資料并行的計算，由GPU負責計算密集型的大規模并行計算。比如醫學上對圖像進行重建、解大規模方程組等，接下來讓我們進入GPU高性能運算之CUDA的世界吧！

CUDA程式設計：

     CUDA程式設計中，習慣稱CPU為Host，GPU為Device。Grid、Block和Thread的關系

Kernel ：在GPU上執行的程式，一個Kernel對應一個Grid。

Grid     ：一組Block，有共享全局記憶體

Block   ：由互相合作的一組線程組成。一個block中的thread可以彼此同步，快速交換資料，最多可以同時512個線程。

Thread  ：并行運算的基本機關（輕量級的線程）

其結構如下圖所示：

?

1 2 3 4 5 6 7 8 9 10

?

存儲層次

1 2 3 4 5 6 7

per- thread register 1 cycle per- thread local memory                     slow per-block shared memory                   1 cycle per-grid global memory                       500 cycle,not cached!! constant and texture memories            500 cycle, but cached and read-only 配置設定記憶體：cudaMalloc，cudaFree，它們配置設定的是global memory Hose-Device資料交換：cudaMemcpy

?

變量類型

1 2 3 4 5

__device__   // GPU的global memory空間，grid中所有線程可通路 __constant__ // GPU的constant memory空間，grid中所有線程可通路 __shared__   // GPU上的thread block空間，block中所有線程可通路 local        // 位于SM内，僅本thread可通路 // 在程式設計中，可以在變量名前面加上這些字首以區分。

?

資料類型

1 2 3 4 5 6 7 8 9

// 内建矢量類型： int1，int2，int3，int4，float1，float2， float3，float4 ... // 紋理類型： texture<Type, Dim, ReadMode>texRef; // 内建dim3類型：定義grid和block的組織方法。例如： dim3 dimGrid(2, 2); dim3 dimBlock(4, 2, 2); // CUDA函數CPU端調用方法 kernelFoo<<<dimGrid, dimBlock>>>(argument);

?

函數定義

1 2 3 4 5 6 7 8 9 10

__device__ // 執行于Device，僅能從Device調用。限制，不能用&取位址；不支援遞歸；不支援static variable；不支援可變長度參數 __global__ // void： 執行于Device，僅能從Host調用。此類函數必須傳回void __host__ // 執行于Host，僅能從Host調用，是函數的預設類型 // 在執行kernel函數時，必須提供execution configuration，即<<<....>>>的部分。 //   例如： __global__ void KernelFunc(...); dim3 DimGrid(100, 50); // 5000 thread blocks dim3 DimBlock(4, 8, 8); // 256 threads per block size_t SharedMemBytes = 64; // 64 bytes of shared memory KernelFunc<<< DimGrid, DimBlock, SharedMemBytes >>>(...);

?

數學函數

1 2	CUDA包含一些數學函數，如 sin ， pow 等。每一個函數包含有兩個版本， 例如正弦函數 sin ，一個普通版本 sin ，另一個不精确但速度極快的__sin版本。

?

内置變量

1 2 3 4 5

?

編寫程式

1 2 3 4 5 6 7

?

相關擴充

1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59 60 61 62 63 64 65 66 67 68 69 70 71 72 73 74 75 76 77 78 79 80 81 82 83 84 85 86 87 88 89 90 91 92 93

1 GPU硬體 // i GPU一個最小單元稱為Streaming Processor(SP)，全流水線單事件無序微處理器， 包含兩個ALU和一個FPU，多組寄存器檔案（ register file，很多寄存器的組合）， 這個SP沒有cache。事實上，現代GPU就是一組SP的array，即SPA。 每一個SP執行一個 thread // ii 多個SP組成Streaming Multiprocessor(SM)。 每一個SM執行一個block。每個SM包含8個SP； 2個special function unit(SFU)： 這裡面有4個FPU可以進行超越函數和插值計算 MultiThreading Issue Unit：分發線程指令 具有指令和常量緩存。 包含shared memory // iii Texture Processor Cluster(TPC) ：包含某些其他單元的一組SM 2 Single-Program Multiple-Data （SPMD）模型  // i CPU以順序結構執行代碼， GPU以threads blocks組織并發執行的代碼，即無數個threads同時執行 // ii 回顧一下CUDA的概念： 一個kernel程式執行在一個grid of threads blocks之中 一個threads block是一批互相合作的threads： 可以用過__syncthreads同步； 通過shared memory共享變量，不同block的不能同步。 // iii Threads block聲明： 可以包含有1到512個并發線程，具有唯一的blockID，可以是1,2,3D 同一個block中的線程執行同一個程式，不同的操作數，可以同步，每個線程具有唯一的ID 3 線程硬體原理 // i GPU通過Global block scheduler來排程block， 根據硬體架構配置設定block到某一個SM。 每個SM最多配置設定8個block，每個SM最多可接受768個 thread （可以是一個block包含512個 thread ， 也可以是3個block每個包含256個 thread （3*256=768！））。 同一個SM上面的block的尺寸必須相同。每個線程的排程與ID由該SM管理。 // ii SM滿負載工作效率最高！考慮某個Block，其尺寸可以為8*8,16*16,32*32 8*8：每個block有64個線程， 由于每個SM最多處理768個線程，是以需要768/64=12個block。 但是由于SM最多8個block，是以一個SM實際執行的線程為8*64=512個線程。 16*16：每個block有256個線程，SM可以同時接受三個block，3*256=768，滿負載 32*32：每個block有1024個線程，SM無法處理！ // iii Block是獨立執行的，每個Block内的threads是可協同的。 // iv 每個線程由SM中的一個SP執行。 當然，由于SM中僅有8個SP，768個線程是以warp為機關執行的， 每個warp包含32個線程，這是基于線程指令的流水線特性完成的。 Warp是SM基本排程機關，實際上，一個Warp是一個32路SIMD指令 。基本機關是half-warp。 如，SM滿負載工作有768個線程，則共有768/32=24個warp ，每一瞬時，隻有一組warp在SM中執行。 Warp全部線程是執行同一個指令， 每個指令需要4個 clock cycle，通過複雜的機制執行。 // v 一個thread的一生： Grid在GPU上啟動； block被配置設定到SM上； SM把線程組織為warp； SM排程執行warp； 執行結束後釋放資源； block繼續被配置設定.... 4 線程存儲模型 // i Register and local memory：線程私有，對程式員透明。 每個SM中有8192個 register ，配置設定給某些block， block内部的 thread 隻能使用配置設定的寄存器。 線程數多，每個線程使用的寄存器就少了。 // ii shared memory：block内共享，動态配置設定。 如__shared__ float region[N]。 shared memory 存儲器是被劃分為16個小單元， 與half-warp長度相同，稱為bank，每個bank可以提供自己的位址服務。 連續的32位word映射到連續的bank。 對同一bank的同時通路稱為bank conflict。 盡量減少這種情形。 // iii Global memory：沒有緩存！容易稱為性能瓶頸，是優化的關鍵！ 一個half-warp裡面的16個線程對global memory的通路可以被coalesce成整塊記憶體的通路，如果： 資料長度為4,8或16bytes；位址連續；起始位址對齊；第N個線程通路第N個資料。 Coalesce可以大大提升性能。 // uncoalesced Coalesced方法：如果所有線程讀取同一位址， 不妨使用constant memory； 如果為不規則讀取可以使用texture記憶體 如果使用了某種結構體，其大小不是4 8 16的倍數， 可以通過__align(X)強制對齊，X=4 8 16