CUDA · parallel_processing

而在device 上執行的function（device 和global）有一些基本的限制：. 不支援遞迴; 不能有static 變數; 不能使用variable number of arguments. 呼叫kernel 函式的語法比 ... parallel_processing Introduction 程序/行程(Process) 平行處理系統類別 MPI(Messagepassinginterface) MPICH安裝與設定 點對點通訊 集合通訊 Communicator One-sidedcommunication TORQUE OpenMP OpenMP設定 pthread ApacheSpark Spark安裝與設定 RDD介紹 基本程式設計 HDFS MapReduce介紹 CUDA 安裝(installation) 第一個程式 PyCUDA Theano Tensorflow XORproblem Cython 編譯 平行處理 PEP numpy Pandas C語言語法 Clibrary Pointer Doublepointer Functionpointer Struct Memoryleak C++參考 C++物件 C++樣版 x86函式調用協定 編譯器參數 cmake Pythonclass設計 nosetest py.test BLASlibrary Linpackbenchmark PoweredbyGitBook CUDA NvidiaCUDA CUDA是Nvidia的平行運算架構，可運用繪圖處理單元(GPU)的強大處理能力，大幅增加運算效能。

OpenCL與CUDA是功能相似的技術，但由於CUDA提供的工具與解決方案較完整，因此將以CUDA為主做GPU計算的功能筆記。

不同的Nvidia顯示卡支援的CUDA版本功能不同，詳見Nvidia提供的列表。

CUDAZONE CUDA流程圖。

現代的顯示晶片已經具有高度的可程式化能力，由於顯示晶片通常具有相當高的記憶體頻寬，以及大量的執行單元，因此開始有利用顯示晶片來幫助進行一些計算工作的想法，即GPGPU(generalpurposegraphicalprocessingunit)。

CUDA即是NVIDIA的GPGPU模型。

GPGPU的優點 顯示晶片通常具有更大的記憶體頻寬。

顯示晶片具有更大量的執行單元。

和高階CPU相比，顯示卡的價格較為低廉。

GPGPU的缺點 顯示晶片的運算單元數量很多，因此對於不能高度平行化的工作，所能帶來的幫助就不大。

顯示晶片目前通常在浮點數單精度(32-bit)的性能較佳，且多半不能完全支援IEEE754規格，有些運算的精確度可能較低。

目前許多顯示晶片並沒有分開的整數運算單元，因此整數運算的效率較差。

顯示晶片通常不具有分支預測等複雜的流程控制單元，因此對於具有高度分支的程式，效率會比較差。

整體來說，顯示晶片的性質類似streamprocessor，適合一次進行大量相同的工作(如SIMD的流程)。

CPU則比較有彈性，能同時進行變化較多的工作。

　CUDA架構 ＊　在CUDA的架構下，一個程式分為兩個部份：host端和device端。

Host端是指在CPU上執行的部份，而device端則是在顯示晶片上執行的部份。

Device端的程式又稱為"kernel"。

通常host端程式會將資料準備好後，複製到顯示卡的記憶體中，再由顯示晶片執行device端程式，完成後再由host端程式將結果從顯示卡的記憶體中取回。

CUDAHostdevice架構。

由於CPU存取顯示記憶體時只能透過PCIExpress介面(目前顯示卡均使用PCIexpress架構，以後也許會使用更快的介面)，因此速度較慢（PCIExpressx16的理論頻寬是雙向各4GB/s），因此不能太常進行這類動作，以免降低效率。

在CUDA架構下，顯示晶片執行時的最小單位是thread。

數個thread可以組成一個block。

一個block中的thread能存取同一塊共用的記憶體，而且可以快速進行同步的動作。

(類似同一個process下的thread可以共享記憶體)。

每一個block所能包含的thread數目是有限的。

不過，執行相同程式的block，可以組成grid。

不同block中的thread無法存取同一個共用的記憶體，因此無法直接互通或進行同步。

因此，不同block中的thread能合作的程度是比較低的。

不過，利用這個模式，可以讓程式不用擔心顯示晶片實際上能同時執行的thread數目限制。

例如，一個具有很少量執行單元的顯示晶片，可能會把各個block中的thread循序執行，而非同時執行。

不同的grid則可以執行不同的程式（即kernel）。

每個thread都有自己的一份register和localmemory的空間。

同一個block中的每個thread則有共用的一份sharememory。

此外，所有的thread（包括不同block的thread）都共用一份globalmemory、constantmemory、和texturememory。

不同的grid則有各自的globalmemory、constantmemory和texturememory。

CUDAgrid,block,thread結構圖。

執行模式 由於顯示晶片大量平行計算的特性，它處理一些問題的方式，和一般CPU是不同的。

主要的特點包括： 記憶體存取latency的問題：CPU通常使用cache來減少存取主記憶體的次數，以避免記憶體latency影響到執行效率。

顯示晶片則多半沒有cache（或很小），而利用平行化執行的方式來隱藏記憶體的latency（即，當第一個thread需要等待記憶體讀取結果時，則開始執行第二個thread，依此類推）。

分支指令的問題：CPU通常利用分支預測等方式來減少分支指令造成的pipelinebubble。

顯示晶片則多半使用類似處理記憶體latency的方式。

不過，通常顯示晶片處理分支的效率會比較差。

因此，最適合利用CUDA處理的問題，是可以大量平行化的問題，才能有效隱藏記憶體的latency，並有效利用顯示晶片上的大量執行單元。

使用CUDA時，同時有上千個thread在執行是很正常的。

因此，如果不能大量平行化的問題，使用CUDA就沒辦法達到最好的效率了。

主機與裝置(hostanddevice) 區分主機和裝置的不同： 主機就是PC。

裝置就是顯示卡。

兩者皆有中央處理器，主機上為CPU，裝置上為GPU，指令集不同： 主機上的程式碼使用傳統C/C++語法撰寫成，實作與呼叫和一般函式無異， 裝置上的程式碼稱為【核心】(kernel)，需使用CUDA的延伸語法(函式前加global等標籤)來撰寫，並於呼叫時指定執行緒群組大小。

兩者皆有各自的記憶體(DRAM)，擁有獨立的定址空間： 主機上的透過malloc()、free()、new、delete等函式配置與釋放， 裝置上的透過cudaMalloc()、cudaFree()等API配置與釋放， 因為主機和裝置的不同，C/C++的標準函式庫不能在kernel中直接使用，例如要秀出計算結果，必需使用cudaMemcpy()先將資料搬移至主機，再呼叫printf或cout等標準輸出函式。

主機和裝置之間的資料搬移，使用cudaMemcpy()這個API。

使用時先在主機記憶體設好資料的初始值，然後傳入裝置記憶體，接著執行核心，如果可以的話就儘量讓資料保留在裝置中，進行一連串的kernel操作，避免透過PCI-E搬移造成效能下降，最後再將結果傳回主機中顯示。

使用API配置裝置記憶體與主機和裝置間資料搬移 最基本的API有5個(incuda.h): 配置裝置記憶體:cudaMalloc() 釋放裝置記憶體:cudaFree() 記憶體複制:cudaMemcpy() 錯誤字串解譯:cudaGetErrorString() 同步化:cudaThreadSynchronize() 配置顯示記憶體cudaMalloc() cudaError_tcudaMalloc(void**ptr,size_tcount); ptr指向目的指位器之位址m，count欲配置的大小(單位bytes) 傳回值cudaError_t是個enum,執行成功時傳回0,其他的錯誤代號可用cudaGetErrorString()來解譯. 釋放顯示記憶體cudaFree() cudaError_tcudaFree(void*ptr); ptr指向欲釋放的位址(devicememory) 記憶體複制cudaMemcpy() cudaError_tcudaMemcpy(void*dst,constvoid*src,size_tcount, enumcudaMemcpyKindkind); dst:指向目的位址，src:指向來源位址，count拷貝區塊大小(單位bytes) kind有四種拷貝流向 cudaMemcpyHostToHost主機->主機 cudaMemcpyHostToDevice主機->裝置 cudaMemcpyDeviceToHost裝置->主機 cudaMemcpyDeviceToDevice裝置->裝置 錯誤字串解譯cudaGetErrorString() constchar*cudaGetErrorString(cudaError_terror); 傳回錯誤代號(error)所代表的字串 同步化cudaThreadSynchronize() cudaError_tcudaThreadSynchronize(void); 用來進行核心和主機程序的同步。

範列程式 #include #include intmain(){ constintnum=100; int*g; cudaError_tr; //主機陣列&初始化 inta[num],b[num]; for(intk=0;k裝置記憶體g[] r=cudaMemcpy(g,a,sizeof(int)*num,cudaMemcpyHostToDevice); printf("cudaMemcpya=>g:%s\n",cudaGetErrorString(r)); //複制記憶體:裝置記憶體g[]------>主機記憶體b[] r=cudaMemcpy(b,g,sizeof(int)*num,cudaMemcpyDeviceToHost); printf("cudaMemcpyg=>b:%s\n",cudaGetErrorString(r)); //結果比對 boolooo=true; for(intk=0;kg:noerror cudaMemcpyg=>b:noerror checka==b?:pass cudaFree:noerror 函式與呼叫(主機、裝置) CUDA中，主機函式的寫法與呼叫和傳統C/C++無異，而裝置核心(kernel)要使用延伸語法：__global__void函式名稱(函式引數...){ ...函式內容... }; 多了global這標籤來標明這道函式是核心程式碼，要編譯器特別照顧一下，注意事項如下： 傳回值隻能是void(要傳東西出來請透過引數) 裏面不能呼叫主機函式或global函式(這兩者皆是主機用的) 輸入的資料若是位址或參考時，必需指向裝置記憶體。

用來指定function是要在host或device上執行，以及是用來被host或device呼叫。

他的類別有三種： +device -在device上執行，且只能被device呼叫。

-同時，他永遠是inlinefunction。

+global -將function宣告成一個kernel，在device上執行，只能被host呼叫。

-他的returntype必須要是void；傳入的參數會是透過sharedmemory給device host 在host上執行，且只能被host呼叫。

（相當於一般的function） 而在device上執行的function（device和global）有一些基本的限制： 不支援遞迴 不能有static變數 不能使用variablenumberofarguments 呼叫kernel函式的語法比一般C函式多了指定網格和區塊大小的序： 函式名稱<<>>(函式引數...); #include #include //裝置函式(核心)在顯示卡記憶體中填入helloCUDA字串 __global__voidhello(char*s){ charw[50]="helloCUDA~~~=^.^="; intk; for(k=0;w[k]!=0;k++) s[k]=w[k]; s[k]=0; }; //主機函式 intmain(){ char*d; charh[100]; //配置裝置記憶體 cudaMalloc((void**)&d,100); //呼叫裝置核心(隻使用單一執行緒) hello<<<1,1>>>(d); //下載裝置記憶體內容到主機上 cudaMemcpy(h,d,100,cudaMemcpyDeviceToHost); //顯示內容 printf("%s\n",h); //釋放裝置記憶體 cudaFree(d); return0; } 變數修飾字 在變數類型方面，是用來指定記憶體的類型。

分成三種： device 宣告變數存在device上；可以和下面兩者同時使用，來做更進一步的設定。

如果沒有額外指定的話，那這個變數 會存在globalmemory空間 生命週期和程式相同 可以被grid中的所有thread透過runtimelibrary存取。

constant 可和device同時使用，會將變數宣告成 存在constantmemory空間 生命週期和程式相同 可以被grid中的所有thread透過runtimelibrary存取。

shared 可和device同時使用，會將變數宣告成： 存在threadblock的sharedmemory空間 生命週期和threadblock相同 只能被block中的thread存取 網格、區塊、執行緒(grid,blockandthread) GPU是具備超多核心，能行大量平行化運算的晶片，執行緒眾多，要分群組管理： 最基本的執行單位是執行緒(thread)， 數個執行緒組成區塊(block)， 數個區塊組成網格(grid)， 整個網格就是所謂的核心(kernel)。

執行緒是最基本的執行單位，程式設計師站在執行緒的角度，透過內建變數，定出執行緒的位置，對工作進行主動切割。

區塊為執行緒的群組，一個區塊可包含1~512個執行緒， 每個執行緒在區塊中擁有唯一的索引編號，記錄於內建變數threadIdx。

每個區塊中包含的執行緒數目，記錄於內建變數blockDim。

相同區塊內的執行緒可同步化，而且可透過共享記憶體交換資料 網格為區塊的群組，一個網格可包含1~65535個區塊， 每個區塊在網格中擁有唯一的索引編號，記錄於內建變數blockIdx。

每個網格中包含的區塊數目，記錄於內建變數gridDim。

網格中的區塊可能會同時或分散在不同時間執行，視硬體情況而定。

內建唯讀變數gridDim,blockDim,blockIdx,threadIdx皆是3D正整數的結構體 uint3gridDim：網格大小(網格中包含的區塊數目) uint3blockIdx：區塊索引(網格中區塊的索引) uint3blockDim：區塊大小(區塊中包含的執行緒數目) uint3threadIdx：執行緒索引(區塊中執行緒的索引) 其中uint3為3D的正整數型態，定義如下structuint3{ unsignedintx,y,z; }; 這些唯讀變數只能在核心中使用。

GRID的大小雖然是uint3的結構，但只能使用2D而已(其z成員隻能是1)，BLOCK才能完整支援3D結構。

核心呼叫時指定的網格和區塊大小對應的就是其中gridDim和blockDim兩變數，網格和區塊大小在設定時有一定的限制。

網格:max(gridDim)=65535 區塊:max(blockDim)=512 實際在用的時候blockDim還會有資源上的限制,主要是暫存器數目,所以有時達不到512這個數量,在3維的情況還會有其他的限制,建議使用1維的方式呼叫,到核心中再去切,執行緒組態比較簡單,而且bug和限制也會比較少。

#include #include //索引用到的緒構體 structIndex{ intblock,thread; }; //核心:把索引寫入裝置記憶體 __global__voidprob_idx(Indexid[]){ intb=blockIdx.x;//區塊索引 intt=threadIdx.x;//執行緒索引 intn=blockDim.x;//區塊中包含的執行緒數目 intx=b*n+t;//執行緒在陣列中對應的位置 //每個執行緒寫入自己的區塊和執行緒索引. id[x].block=b; id[x].thread=t; }; //主函式 intmain(){ Index*d; Indexh[100]; //配置裝置記憶體 cudaMalloc((void**)&d,100*sizeof(Index)); //呼叫裝置核心 intg=3,b=4,m=g*b; prob_idx<<>>(d); //下載裝置記憶體內容到主機上 cudaMemcpy(h,d,100*sizeof(Index),cudaMemcpyDeviceToHost); //顯示內容 for(inti=0;i32~64KB/BLOCK，看是那個世代的GPU)，系統會自動把一些資料置換到全域記憶體中，導致執行緒變多，但效率反而變慢(類似作業系統虛擬記憶體的swap)；另一個會引發這種swap的情況是在使用動態索引存取陣列，因為此時需要陣列的順序性，而暫存器本身是沒有所謂的順序的，所以系統會自動把陣列置於全域記憶體中，再按索引存取，這種情況建議使用共享記憶體手動避免。

執行緒同步(網格、區塊) 同步化函式 使用地點 功能 __syncthreads() 核心程序中 同步化區塊內的執行緒 cudaThreadSynchronize() 主機程序中 同步化核心和主機程序 在kernel中，使用__syncthreads()來進行區塊內的執行緒的同步，避免資料時序上的問題(來自不同threads)，時常和共享記憶體一起使用。

在主機程序中，使用cudaThreadSynchronize()來進行核心和主機程序的同步，它來避免量到不正確的主機時間(kernel仍未完成就量時間)，因為主機的程序和裝置程序預設是不同步的(直到下載結果資料之前)，這個API可以強迫它們同步。

合併存取 合併存取(coalescedI/O)是CUDA中最基本且最重要的最佳化手段,因為GPU的計算能力太強,使得效能瓶頸卡在顯示記憶體到GPU之間的I/O上,合併存取可讓多個顯示記憶體的交易合併成一次,而加速記憶體的存取. 現階段GPU在合併存取上是自動發生的,以半個warp為單位(16個相鄰的執行緒),如果它們的資料位址是連續的,就會被合併,所以使用上很簡單,只要threadIdx對齊即可,它可以合併4,8,16bytes的資料,成為一次or兩次的交易,合併成的最大封包長度可為32,64,128bytes(其中32bytes的封包只有在版本1.2以後支援,避免位址分散情況下的overhead)，以下為連續的資料位址的合併情況。

16*4bytes->64bytes 16*8bytes->128bytes 16*16bytes->256bytes resultsmatching"" Noresultsmatching""