CUDA學習筆記（一）【轉】

CUDA編程中，習慣稱CPU爲Host，GPU爲Device。編程中最開始接觸的東西恐怕是並行架構，諸如Grid、Block的區別會讓人一頭霧水，我所看的書上所講述的內容比較抽象，對這些概念的內容沒有細講，因而在這裏做一個整理。

Grid、Block和Thread的關係

Thread  ：並行運算的基本單位（輕量級的線程）
Block   ：由相互合做的一組線程組成。一個block中的thread能夠彼此同步，快速交換數據，最多能夠同時512個線程。
Grid     ：一組Block，有共享全局內存
Kernel ：在GPU上執行的程序，一個Kernel對應一個Grid。

其結構以下圖所示：

?

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

/* gridDim, blockIdx, blockDim, threadIdx, wrapsize. 這些內置變量不容許賦值的 */

?

編寫程序

1 2 3 4 5 6 7

/* 目前CUDA僅能良好的支持C，在編寫含有CUDA代碼的程序時， 首先要導入頭文件cuda_runtime_api.h。文件名後綴爲.cu，使用nvcc編譯器編譯。 目前最新的CUDA版本爲5.0，能夠在官方網站下載最新的工具包，網址爲： https://developer.nvidia.com/cuda-downloads 該工具包內包含了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，多組寄存器文件（ 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