CUDA並行計算架構（二） 執行個體相關

從這部分開始 結合蟲子的demo程式給大家分析下cuda的效能與可行性。

一。先概述下實現流程。

CUDA在執行的時候是讓host裡面的一個一個的kernel按照線程網格（Grid）的概念在顯卡硬體（GPU）上執行。每一個線程網格又可以包含多個線程塊（block），每一個線程塊中又可以包含多個線程（thread）。

每一個kernel交給每一個Grid來完成。當要執行這些任務的時候，每一個Grid又把任務分成一部分一部分的block，block再分線程來完成。每個Grid中的任務是一定的。二維線程塊的索引關係為如下：
unsigned int xIndex = blockDim.x * blockIdx.x + threadIdx.x;

unsigned int yIndex = blockDim.y * blockIdx.y + threadIdx.y;

block中的每個線程都有自己的寄存器和local memory，block中的所有線程共用一個shared memory，一個grid共用一個global memory。

每一個刻度內，warp（一個block裡面一起啟動並執行thread）包含的thread數量是有限的，現在的規定是32個。一個block中含有16個warp。所以一個block中最多含有512個線程

每次Device（就是顯卡）只處理一個grid。

下面說明一下硬體的執行模型。    

   假如出於某種原因，公司的辦公室被徵用了搞活動。只留下一個小房間來給Team Dev。每一個刻度內按照wrap（就理解為啟動並執行時候，一個block裡面一起啟動並執行thread，例如block裡面有512個thread，但是每次只有32個thread在運行，那麼這32個thread就是一個啟動並執行warp組 線程束）。每一個warp裡麵包含的thread數量是有限的，現在的規定是32個。將來不知道會不會有變化，這個只有CUDA開發人員知道了。每次Device（就是顯卡）只處理一個grid(在未來支援directX11 的硬體中這一限制可能被解除)。假如我們一個部門有x個人，辦公室內有N個桌子，每張桌子可以坐32個人。然後輪流來開發….。這裡的桌子可以理解成multiprocessor（多處理器）。每個sm中包含8個標量流處理器（sp）。GPU所謂的多核中核的概念就是sp的數量。Cuda中的kernel函數實質上是以block為單位執行的。同一block中需要共用資料，因此他們必須在同一個SM中發射，而block中的每一個線程則被發射到sp上去執行。      疑點：既然有這樣的線程簇限制、為何還要設定高於warp線程數的線程。

二。demo

安裝部署方面driver、toolkit、sdk順序安裝好。Cuda的項目支援4種調試方式release、debug、emurelease、emudubug。前2個是需要gpu真正的支援cuda後者是cpu類比gpu。至於你的電腦能否支援cuda 可以運行下deviceQuery.exe程式

圖中我們關注一下幾點就可以了，首先 有一個支援cuda的裝置。計算能力1，局儲存空間的大小，核的數量，多處理器的數量，常量儲存空間的大小、每個block的共用儲存空間的大小、wrap的線程數等等。
想看cuda在圖形領域的應用可以運行這個smokeParticles.exe程式哦。

在我的demo中，cpp檔案主要是處理一些cpu端的處理、cu檔案通常是與gpu核函數和cuda api的一些內容。其中My_kernel封裝了具體的核函數實現方法。Cudatool項目就是cuda的應用程式，CudaProviders是我串連C#與cuda之間的驅動、CudaWeb就是我們平常的web項目。CUDAWinApp這個就是一些小的功能示範。
下面介紹下cuda的函數類型限定符。
__device__ 在裝置上執行、只能在裝置上調用。
__global__ 用於聲明核心功能、在裝置上執行只能從主機端調用。
__host__ 在主機端執行，只能從主機端調用，預設。
__device__與__global__不支援遞迴，函數體內不能聲明靜態變數、參數數目不可變化，不能對device取指標。__global__與__host__不能連用。__global__只能返回空，調用__global__函數必須聲明其執行配置、__global__函數的調用是非同步、__global__參數的值目前是通過共用儲存空間傳遞，總的大小不能超過256byte。
變數類型限定符分為__device__（變數存在裝置端上）、__constant__（存在常數儲存空間空間）、__share__（block的共用儲存空間）、volatile關鍵字 當線程間資料可能互相影響變換時使用。