1

2

3

4

5

6

7

8

9

10

/*

另外：Block和Thread都有各自的ID，記作blockIdx（1D，2D），threadIdx（1D，2D，3D）

Block和Thread還有Dim，即blockDim與threadDim. 他們都有三個分量x，y，z

執行緒同步：void __syncthreads(); 可以同步一個Block內的所有執行緒

總結來說，每個 thread 都有自己的一份 register 和 local memory 的空間。

一組thread構成一個 block，這些 thread 則共享有一份shared memory。

此外，所有的 thread(包括不同 block 的 thread)都共享一份

global memory、constant memory、和 texture memory。

不同的 grid 則有各自的 global memory、constant memory 和 texture memory。

*/

1

2

3

4

5

6

7

per-

per-threadlocal 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__ voidKernelFunc(...);

dim3 DimGrid(100, 50);// 5000 thread blocks

dim3 DimBlock(4, 8, 8);// 256 threads per block

size_tSharedMemBytes = 64; // 64 bytes of shared memory

KernelFunc<<< DimGrid, DimBlock, SharedMemBytes >>>(...);

1

2

CUDA包含一些數學函式，如sin，pow等。每一個函式包含有兩個版本，

例如正弦函式sin，一個普通版本sin，另一個不精確但速度極快的__sin版本。

1

2

3

4

5

/*

gridDim, blockIdx, blockDim,

threadIdx, wrapsize.

這些內建變數不允許賦值的

*/

1

2

3

4

5

6

7

/*

目前CUDA僅能良好的支援C，在編寫含有CUDA程式碼的程式時，

首先要匯入標頭檔案cuda_runtime_api.h。檔名字尾為.cu，使用nvcc編譯器編譯。

目前最新的CUDA版本為5.0，可以在官方網站下載最新的工具包，網址為：

該工具包內包含了ToolKit、樣例等，安裝起來比原先的版本也方便了很多。

*/

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，多組暫存器檔案（registerfile，很多暫存器的組合），

這個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個clockcycle，通過複雜的機制執行。

// 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



轉自http://luofl1992.is-programmer.com/posts/38830.html