CUDA学习1
CUDA編目錄：
1.什么是CUDA
2.為什么要用到CUDA
3.CUDA環境搭建
4.第一個CUDA程序
5. CUDA編程
5.1. 基本概念
5.2. 線程層次結構
5.3. 存儲器層次結構
5.4. 運行時API
5.4.1. 初始化
5.4.2. 設備管理
5.4.3. 存儲器管理
5.4.3.1. 共享存儲器
5.4.3.2. 常量存儲器
5.4.3.3. 線性存儲器
5.4.3.4. CUDA數組
5.4.4. 流管理
5.4.5. 事件管理
5.4.6. 紋理參考管理
5.4.6.1. 紋理聲明
5.4.6.2. 紋理綁定
5.4.6.3. 紋理獲取
5.4.7. OpenGL互操作
5.4.8. Direct3D互操作
5.5. 驅動API
5.5.1. 初始化
5.5.2. 設備管理
5.5.3. 上下文管理
5.5.4. 模塊管理
5.5.5. 執行控制
5.5.6. 存儲器管理
5.5.7. 流管理
5.5.8. 事件管理
5.5.9. 紋理參考管理
5.5.10. OpenGL互操作
5.5.11. Direct3D互操作
5.6. 性能優化
5.7. NVCC編譯器
5.8. 設備模擬
5.9. 其他
參考文獻
?
1.什么是CUDA
?????? CUDA(Compute Unified Device Architecture)，統一計算架構，是NVidia推出的并行計算平臺。NVidia官方對其的解釋是：一個并行計算平臺和簡單（簡潔）地使用圖像處理單元（GPU）進行通用計算的編程模型。利用GPU的能力在計算性能上有驚人的提升。
?????? 簡單地說CUDA是便于程序員利用NVidia GPU進行通用計算的開發環境及工具，目前支持C/C++語言，將來還會支持Fortran語言。
?
2.為什么要用到CUDA
?????? CPU主頻要比GPU高2-3倍左右，但是通常情況下GPU核心的數量要比CPU多2-3個數量級以上。因此GPU的計算能力要遠大于CPU，充分發揮GPU的計算能力，可以有成倍的性能提升。???
?????? 早期利用GPU的計算能力是使用著色器和著色語言（GLSL等）。目前廣泛使用的是CUDA和OpenCL。CUDA是針對NVidia GPU硬件設備設計的，而 OpenCL是針對跨平臺設計的。因此CUDA可充分發揮NVidia GPU的計算性能。
?????? CUDA可以直接使用C/C++語言來開發GPU程序，省去了程序員重新學一種新語言的麻煩。
?
3.CUDA環境搭建
?????? CUDA環境主要分為四點：硬件（GPU設備）、操作系統、C/C++編譯器和CUDA工具包。
?????? 硬件（GPU設備），必須是支持CUDA的GPU。可到NVidia官網查詢支持CUDA的GPU設備，具體地址為：http://www.nvidia.com/object/cuda_home_new.html?。
?????? 操作系統，支持Microsoft Windows、Mac OS X和Linux。
?????? C/C++編譯器，對不同的操作系統有不同的要求。
?????? CUDA工具包，NVidia提供了不同操作系統對應的CUDA Toolkit，可從https://developer.nvidia.com/cuda-downloads?下載對應的版本。
?????? 本文只以Microsoft Windows為例介紹如何搭建CUDA環境。
?????? 準備材料：
?????? ·一臺裝有支持CUDA GPU的電腦。
?????? ·Microsoft Windows操作系統（Microsoft Windows XP,Vista,7,or 8 or Windows Server 2003 or 2008）。
?????? ·CUDA工具包（相應操作系統）。下載地址：https://developer.nvidia.com/cuda-downloads
?????? ·C/C++編譯器：Microsoft Visual Studio 2008 或 2010，或者對應版本的Microsoft Visual C++ Express產品。
?????? 安裝步驟：
?????? ·在裝有支持CUDA GPU的電腦上安裝Microsoft Windows操作系統（一般情況下都已經完成這步驟）。
?????? ·安裝C/C++編譯器，可只安裝其中的C++編譯器部分。
?????? ·安裝CUDA工具包。（CUDA工具包中有NVidia GPU的驅動程序，尚未安裝的請選擇安裝。）
?????? 安裝驗證：
?????? Windows XP系統：進入 C:\Documents and Settings\All Users\Application Data\NVIDIA Corporation\CUDA Samples\v5.0\bin\win32\Release 目錄運行deviceQuery.exe文件。
?????? Windows Vista, Windows 7, Windows 8, Windows Server 2003, and Windows Server 2008系統：進入 C:\ProgramData\NVIDIA Corporation\CUDA Samples\v5.0\bin\win32\Release 目錄運行deviceQuery.exe文件。
?????? 如果安裝正確，執行deviceQuery.exe文件會得到GPU設備的相應信息。如果沒有安裝支持CUDA的GPU也會得出GPU的信息，其中CUDA Capability Major/Minor version number信息為9999.9999。
?????? Microsoft Windows上更詳細的安裝信息請查看：
http://docs.nvidia.com/cuda/cuda-getting-started-guide-for-microsoft-windows/index.html?。
?????? Mac OS X的安裝：
http://docs.nvidia.com/cuda/cuda-getting-started-guide-for-mac-os-x/index.html?。
?????? Linux的安裝：
http://docs.nvidia.com/cuda/cuda-getting-started-guide-for-linux/index.html?。
?
4.第一個CUDA程序
?????? 在Microsoft Windows系統上，如果成功搭建了CUDA環境，則在Microsoft Visual Studio中已經集成了CUDA的開發組件。
?????? 以下以Windows 7 + Microsoft Visual Studio 2008為例，創建第一個CUDA程序。
?????? 打開Microsoft Visual Studio 2008，依次：File->New->Project->NVIDIA->CUDA->CUDA 5.0 Runtime，輸入相應的項目名稱確定即可。
?????? 默認會生成一個kernel.cu文件，內容如下：
#include "cuda_runtime.h"#include "device_launch_parameters.h"#include <stdio.h>void addWithCuda(int *c, const int *a, const int *b, size_t size);__global__ void addKernel(int *c, const int *a, const int *b){int i = threadIdx.x;c[i] = a[i] + b[i];}int main(){const int arraySize = 5;const int a[arraySize] = { 1, 2, 3, 4, 5 };const int b[arraySize] = { 10, 20, 30, 40, 50 };int c[arraySize] = { 0 };// Add vectors in parallel.addWithCuda(c, a, b, arraySize);printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n",c[0], c[1], c[2], c[3], c[4]);// cudaThreadExit must be called before exiting in order for profiling and// tracing tools such as Nsight and Visual Profiler to show complete traces.cudaThreadExit();return 0;}// Helper function for using CUDA to add vectors in parallel.void addWithCuda(int *c, const int *a, const int *b, size_t size){int *dev_a = 0;int *dev_b = 0;int *dev_c = 0;// Choose which GPU to run on, change this on a multi-GPU system.cudaSetDevice(0);// Allocate GPU buffers for three vectors (two input, one output) .cudaMalloc((void**)&dev_c, size * sizeof(int));cudaMalloc((void**)&dev_a, size * sizeof(int));cudaMalloc((void**)&dev_b, size * sizeof(int));// Copy input vectors from host memory to GPU buffers.cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);// Launch a kernel on the GPU with one thread for each element.addKernel<<<1, size>>>(dev_c, dev_a, dev_b);// cudaThreadSynchronize waits for the kernel to finish, and returns// any errors encountered during the launch.cudaThreadSynchronize();// Copy output vector from GPU buffer to host memory.cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);cudaFree(dev_c);cudaFree(dev_a);cudaFree(dev_b);}代碼1
?????? 這是一個將兩個一維數組相加的例子。
?????? 其中addKernel是內核函數，它的計算過程是在GPU上實現的，用函數類型限定符__global__限制，且函數類型為void型。
?????? cuda_runtime.h頭文件包括了運行時API和其參數的定義。（如果使用驅動API則使用cuda.h頭文件）。
?????? device_launch_parameters.h頭文件包含了內核函數的5個變量threadIdx、blockDim、blockIdx、gridDim和wrapSize。
?????? 對其中CUDA運行時API函數的解釋：
?????? ·cudaSetDevice()：選擇設備（GPU）。（可以不使用，不使用的情況下，默認選擇設備0）
?????? ·cudaMalloc()：動態分配顯存。
?????? ·cudaMemcpy()：設備與主機之內的數據拷貝。
?????? ·cudaThreadSynchronize()：同步所有設備上的線程，等待所有線程結束。
?????? ·cudaFree():釋放由cudaMalloc分配的顯存。
?????? ·cudaThreadExit():結束CUDA上下文環境，釋放其中的資源。
?????? 這些函數的具體介紹在?http://docs.nvidia.com/cuda/cuda-runtime-api/index.html?中。
?
5. CUDA編程
5.1. 基本概念
?????? CUDA編程中需要注意一些基本概念，分別為：主機、設備、運行時API、驅動API、warp、bank、函數類型限定符、變量類型限定符、thread、block、grid、計算能力、SIMT、內置變量、紋理、CUDA數組等。
?????? 主機：可理解為CPU與內存的組合。
?????? 設備：可理解為GPU與顯存的組合。
?????? 運行時API：是指CUDA運行時API是在驅動API的基礎上封裝而成的，簡化了CUDA的開發。
?????? 驅動API：是指CUDA驅動API，相比運行時API更接近于設備，可靈活運用設備的特性開發CUDA，可實現運行時API無法實現的功能。
?????? warp：多處理器激活、管理、調度和執行并行任務的單位。計算能力2.x的設備warp為32個線程。未來的設備可能不同，可以通過內置變量warpSize查詢。
?????? bank：為了獲得較高的存儲器帶寬，共享存儲器被劃分為多個大小相等的存儲器模塊，稱為存儲體，這些存儲體就叫bank，可同步訪問。
?????? 函數類型限定符：是CUDA C中特有的，用來修飾是主機函數，設備調用的設備函數，還是主機調用的設備函數。有__device__、__global__、__host__。
?????? 變量類型限定符：是用來修飾設備變量的。有__device__、__constant__、__shared__。
?????? thread：設備中的線程，與主機中的線程是同一個概念。
?????? block：線程塊，由一組線程組成。一個線程塊中的所以線程會在同一個多處理器上執行，一個多處理器上可同時執行多個線程塊。
?????? grid：有所有線程塊組成的網格。
?????? 計算能力：是NVidia GPU不同架構的計算能力。
?????? SIMT：單指令多線程，與單指令多數據（SIMD）類似。一條指令多個線程一同執行，實現程序的并行化。
?????? 內置變量：有threadIdx、blockDim、blockIdx、gridDim、warpSize。其中threadIdx指此線程在線程塊中的位置；blockDim指線程塊維度；blockIdx指該線程塊在網格中的位置；gridDim指線程塊網格維度；warpSize指一個warp多少個線程。
?????? 紋理：本文主要涉及到的是紋理參考、紋理綁定、紋理獲取。
?????? CUDA數組：區別于線性存儲器，對數據進行了對齊等的處理，包括一維、二維和三維。其中的數據為：一元、二元或四元組。
?
5.2. 線程層次結構
?????? CUDA線程的層次結構，由小到大依次為線程、線程塊、線程塊網格。一維、二維或三維的線程組組成一個線程塊，一維、二維或三維的線程塊組組成一個線程塊網格。
?????? 下圖是由二維的線程塊組組成的線程塊網絡，其中線程塊是由二維的線程組組成。
?????? 圖1
?????? NVidia GPU的硬件結構是，一組流處理器組成一個多處理器，一個或多個多處理器組成一個GPU。其中流處理器，可以理解為處理計算的核心單元。多處理器類似于多核CPU。NVidia GPU從DX10（DirectX10）開始出現了Tesla、Fermi、Kepler架構，不同的架構多處理器中流處理器數量都有差別。
?
5.3. 存儲器層次結構
?????? CUDA存儲器有：寄存器、共享存儲器、常量存儲器、本地存儲器、全局存儲器、紋理存儲器等。其中寄存器和本地存儲器是線程私有的，共享存儲器是對線程塊中的所有線程可見，常量存儲器、全局存儲器和紋理存儲器是對網格中所有線程可見。
?????? 下圖解釋了存儲器的層次結構：
?????? 圖2
?
5.4. 運行時API
?????? 運用運行時API開發CUDA程序需要了解：初始化、設備管理、存儲器管理、流管理、事件管理、紋理參考管理、OpenGL互操作和Direct3D互操作。
?????? 運行時API文檔地址為：http://docs.nvidia.com/cuda/cuda-runtime-api/index.html?。
?
5.4.1. 初始化
?????? 運行時API不存在顯示初始化函數，初始化會在首次調用運行時函數時完成。雖然不需要調用初始化函數進行初始化，但是退出時需要調用退出函數cudaThreadExit()釋放資源。
?
5.4.2. 設備管理
?????? 有些電腦上可能有多塊設備，因此對于不同的要求選擇合適的設備。設備管理主要是獲取設備信息和選擇執行設備。
?????? 主要有三個函數：
?????? ·cudaGetDeviceCount()：得到電腦上設備的個數。
?????? ·cudaGetDeviceProperties()：獲得對應設備的信息。
?????? ·cudaSetDevice()：設置CUDA上下文對應的設備。
?????? 運行__global__函數前需要提前選擇設備，如果不調用cudaSetDevice()函數，則默認使用0號設備。
?????? 上面三個函數的具體用法請查看CUDA運行時API文檔。
?
5.4.3. 存儲器管理
?????? 共享存儲器、常量存儲器、線性存儲器和CUDA數組的使用是存儲器管理的主要部分。
?
5.4.3.1. 共享存儲器
?????? 共享存儲器，使用__shared__變量限定符修飾，可靜態或動態分配共享存儲器。
?????? 靜態分配共享存儲器，是在設備代碼中直接分配共享存儲器的大小，如下代碼：
#define SHARED_MEM 16__global__ void kernel(…){__shared__ int shared[SHARED_MEM];}void main(){kernel<<<nBlock, nThread>>>(…);}代碼2
?????? 動態分配共享存儲器，是在主機代碼中使用內核函數的第三個特定參數傳入分配共享存儲器的大小，如下代碼：
#define SHARED_MEM 16__global__ void kernel(…){extern __shared__ int shared[];}void main(){int nSharedMem = (int)SHARED_MEM;kernel<<<nBlock, nThread, nSharedMem*sizeof(int)>>>(…);}代碼3
?
5.4.3.2. 常量存儲器
?????? 常量存儲器，使用__constant__變量限定符修飾。使用常量存儲器，是由于其在設備上有片上緩存，比全局存儲器讀取效率高很多。
?????? 使用常量存儲器時會涉及的運行時API函數主要有：
?????? ·cudaMemcpyToSymbol()
?????? ·cudaMemcpyFromSymbol()
?????? ·cudaGetSymbolAddress()
?????? ·cudaGetSymbolSize()
?????? 主機代碼中使用cudaGetSymbolAddress()獲取__constant__或__device__定義的變量地址。設備代碼中可通過提取__device__、__shared__或__constant__變量的指針獲取變量地址。
?
5.4.3.3. 線性存儲器
?????? 線性存儲器是使用cudaMalloc()、cudaMallocPitch()或cudaMalloc3D()分配的，使用cudaFree()釋放。二維的時候建議使用cudaMallocPitch()分配，cudaMallocPitch()函數對對齊進行了調整。這三個分配函數對應cudaMemset()、cudaMemset2D()、cudaMemset3D()三個memset函數和cudaMemcpy()、cudaMemcpy2D()、cudaMemcpy3D()三個memcpy函數。
?
5.4.3.4. CUDA數組
?????? CUDA數組是使用cudaMallocArray()、cudaMalloc3DArray()分配的，使用cudaFreeArray()釋放。
?????? 相關memcpy函數請查閱CUDA運行時API文檔。
?????? 具體使用可查閱CUDA編程指南：
http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html?。
?
5.4.4. 流管理
?????? 主機設備之間的內存拷貝與內核在設備上執行是異步的。在不使用流的情況下，是這樣執行的：設備先從主機上拷貝內存，拷貝完成之后，再在設備上執行內核代碼計算，最后當內核執行完畢，再把設備上的內存拷貝到主機上。當使用兩個流的情況下，0號流執行內核代碼的同時1號流拷貝主機內存到設備，1號流執行的同時0號流拷貝設備內存到主機（具體的實現并不一定如此，這里是為了說明流的作用簡單做了假設）。兩個流的情況下，部分內存拷貝和內置執行是同時進行的（異步的），比同步的內存拷貝和內核執行節省了時間。
?????? 與流有關的函數有：
?????? ·cudaStreamCreate()：流的創建；
?????? ·cudaStreamDestroy()：流的銷毀；
?????? ·cudaStreamSynchronize()：流同步；
?????? ·*Async：與流相關的其他函數。
?????? 內核<<<…>>>的第四個參數為哪個流。
?????? CUDA編程指南中有對流具體實現的講解。
?
5.4.5. 事件管理
?????? 由于部分CUDA運行時函數的執行與主機代碼是異步的。在一塊代碼中，CUDA運行時函數執行沒有結束就直接執行其后的主機代碼了，主機并不知道已經執行到哪個CUDA運行時函數了。事件的引入就是為了解決這一問題，在CUDA運行時函數已經執行完畢后記入下事件，查詢此事件是否記錄就能知道那個CUDA運行時函數已經執行完畢。在CUDA運行時函數前后記入事件就能獲得此函數執行的時間。
?????? 與事件有關的函數有：
?????? ·cudaEventCreate()：事件的創建；
?????? ·cudaEventDestroy()：事件的銷毀；
?????? ·cudaEventRecord()；記錄事件；
?????? ·cudaEventSynchronize()：事件同步；
?????? ·cudaEventElapsedTime()：計算兩事件的時間差。
?????? 具體的實現請查詢CUDA編程指南。
?
5.4.6. 紋理參考管理
?????? 紋理參考的實現是由紋理聲明、紋理綁定、紋理獲取完成的。
?
5.4.6.1. 紋理聲明
?????? 紋理聲明是在文件域中聲明紋理變量，供主機使用CUDA函數綁定紋理和設備獲取紋理。紋理聲明為：
?????? texture<DataType, Type, ReadMode> texRef;
?????? 其中：
?????? ·DateType：紋理元的格式，有float、unsigned char、signed char、unsigned short、signed short及它們的2元和4元組。
?????? ·Type：紋理參考格式，有cudaTextureType1D、cudaTextureType2D、cudaTextureType3D、cudaTextureType1DLayered、cudaTextureType2DLayered。是可選參數，默認為cudaTextureType1D。
?????? ·ReadMode：讀取模式，有cudaReadModeElementType、cudaReadModeNormalizedFloat。是可選參數，默認為cudaReadModeElementType。為cudaReadModeNormalizedFloat時，紋理元數據使用了單位化的映射，映射到了[0.0, 1.0]或[-1.0, 1.0]。為cudaReadModeElementType時，不進行任何映射變換。
?
5.4.6.2. 紋理綁定
?????? 紋理綁定是將分配的線性存儲器或CUDA數組綁定到紋理存儲器。CUDA運行時紋理綁定API分為高級和低級兩種綁定類型。
?????? 下面是CUDA編程文檔上綁定線性存儲器的例子：
?????? 低級API：
texture<float, cudaTextureType2D, cudaReadModeElementType> texRef;textureReference* texRefPtr;cudaGetTextureReference(&texRefPtr, texRef);cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();size_t offset;cudaBindTexture2D(&offset, texRefPtr, devPtr, &channelDesc, width, height, pitch);代碼4
?????? 高級API：
texture<float, cudaTextureType2D, cudaReadModeElementType> texRef;cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();size_t offset;cudaBindTexture2D(&offset, texRef, devPtr, channelDesc, width, height, pitch);代碼5
?????? 其中devPtr是由cudaMallocPitch()分配的線性存儲器指針；width、height、pitch是cudaMallocPitch()使用或獲得的變量。
?????? 綁定線程存儲器還可以使用cudaBindTexture()。
?????? 綁定CUDA數組使用cudaBindTextureToArray()，CUDA編程文檔上有如何使用的具體介紹。
?
5.4.6.3. 紋理獲取
?????? 紋理獲取是在內核中獲得紋理中某一個坐標對應的紋理元值。
?????? 紋理獲取的函數有：tex1Dfetch()、tex1D()、tex2D()、tex3D()、tex1Dlayered()、tex2Dlayered()、texCubemap()、texCubemapLayered()、tex2Dgather()。具體使用請查詢CUDA編程文檔。
?
5.4.7. OpenGL互操作
?????? OpenGL（Open Graphics Library），一個圖形硬件API。
?????? OpenGL與CUDA互操作，主要是緩沖對象的注冊與取消注冊、映射與取消映射。對應的函數有：
?????? ·cudaGLRegisterBufferObject()：緩沖對象注冊；
?????? ·cudaGLUnregisterBufferObject()：取消緩沖對象注冊；
?????? ·cudaGLMapBufferObject()：映射緩沖對象；
?????? ·cudaGLUnmapBufferObject()：取消映射。
?????? cudaGLMapBufferObject()映射緩沖對象后，CUDA可以使用其返回的設備存儲器地址讀取和寫入緩沖對象。
?????? CUDA關于OpenGL互操作的具體介紹請查詢CUDA編程文檔及運行時API。
?????? OpenGL部分的知識，請查看：
http://www.opengl.org/wiki/Getting_started
http://www.opengl.org/sdk/docs/
?
5.4.8. Direct3D互操作
?????? Direct3D是Microsoft自己的3D圖形API。
?????? Direct3D與CUDA互操作，主要是Direct3D設備的設置、資源的注冊、資源映射、映射后信息獲取、取消映射、取消注冊。對應的函數有（以Direct3D9為例）：
?????? ·cudaD3D9SetDirect3DDevice()：Direct3D設備的設置；
?????? ·cudaD3D9RegisterResource ()：注冊資源；
?????? ·cudaD3D9MapResources()：資源映射；
?????? ·cudaD3D9ResourceGetMappedPointer()：獲取資源映射后的CUDA設備存儲器地址；
?????? ·cudaD3D9ResourceGetMappedSize()：獲取大小；
?????? ·cudaD3D9ResourceGetMappedPitch()：獲取間隔；
?????? ·cudaD3D9UnmapResources()：取消映射；
?????? ·cudaD3D9UnregisterResource()：取消注冊。
?????? 具體如何使用請查詢CUDA編程文檔。
?????? Direct3D部分的知識，請查詢MSDN。
?
5.5. 驅動API
?????? 驅動API是比運行時API更底層的一套接口，運行時API是在驅動API的基礎上封裝而成。驅動API是一種基于句柄、命令式的API：大多數對象都通過不透明的句柄引用。
?????? 以下列了主要的句柄：
| 對象 | 句柄 | 描述 |
| 設備 | CUdevice | 支持CUDA的設備 |
| 上下文 | CUcontext | 大致等同于CPU進程 |
| 模塊 | CUmodule | 大致等同于動態庫 |
| 函數 | CUfunction | 內核 |
| 堆存儲器 | CUdeviceptr | 設備存儲器的指針 |
| CUDA數組 | CUarray | 設備上一維或二維數據的不透明容器，可通過紋理參考讀取 |
| 紋理參考 | CUtexref | 描述如何解釋紋理存儲器數據的對象 |
?????? 運用驅動API開發CUDA程序需要了解：初始化、設備管理、上下文管理、模塊管理、執行控制、存儲器管理、流管理、事件管理、紋理參考管理、OpenGL互操作、Direct3D互操作。
?????? 驅動API文檔地址為：http://docs.nvidia.com/cuda/cuda-driver-api/index.html?。
?
5.5.1. 初始化
?????? 驅動API與運行時API不同，需要在調用任何驅動API函數（不包括初始化函數）之前初始化。初始化函數為cuInit()。
?
5.5.2. 設備管理
?????? 驅動API與運行時API不同，不需要設置設備，而是直接使用得到的設備句柄操作設備。
?????? 設備管理的主要函數有：
?????? ·cuDeviceGetCount()：獲得主機上設備總數；
?????? ·cuDeviceGet()：獲得對應設備句柄；
?????? ·cuDeviceGetProperties()：獲得設備信息。
?????? 具體解釋及其中參數信息請查閱驅動API文檔。
?
5.5.3. 上下文管理
?????? CUDA上下文類似于CPU進程。在驅動程序API中執行的所有資源和操作都封裝在CUDA上下文內在該上下文被銷毀時，系統將自動清除這些資源。除了模塊和紋理參考之類的對象之外，每個上下文都有自己獨特的32位地址空間。
?????? 一個主機線程只能有一個當前設備上下文。每個主機線程都有一個當前上下文堆棧，并為每個上下文維護一個使用計數。
?????? 上下文管理的主要函數有：
?????? ·cuCtxCreate()：創建上下文；
?????? ·cuCtxDestroy()：銷毀上下文；
?????? ·cuCtxPopCurrent()：使當前上下文離開當前主機線程上下文堆棧；
?????? ·cuCtxPushCurrent()：壓入上下文到當前主機線程上下文堆棧；
?????? ·cuCtxAttach()：增加一個上下文計數；
?????? ·cuCtxDetach()：消耗一個上下文計數（當上下文使用計數為0時，自動銷毀上下文）。
?????? 具體使用請查詢CUDA編程文檔及驅動API文檔。
?
5.5.4. 模塊管理
?????? 模塊是可獨立加載的設備代碼和數據包，類似于windows中的DLL。所有符號的名稱（包括函數、全局變量和紋理參考）均在模塊范圍內維護，從而使獨立的第三方編寫的模塊可在相同的CUDA上下文中進行互操作。
?????? 模塊管理的主要函數有：
?????? ·cuModuleLoad()：模塊加載；
?????? ·cuModuleGetFunction()：得到模塊中相應函數。
?????? 具體使用請查詢CUDA編程文檔及驅動API文檔。
?
5.5.5. 執行控制
?????? 執行控制是指，執行和控制設備代碼（內核）。驅動API內核的執行，不同運行時API一樣方便，需要設置額外設置grid、block和參數等，還是使用特定的launch函數。
?????? 執行控制的主要函數有：
?????? ·cuFuncSetCacheConfig()：設置函數對應的cache偏好（是設置cache多還是共享內存多）；
?????? ·cuFuncSetSharedMemConfig()：設置共享內存bank的大小；
?????? ·cuLaunchKernel()：launch函數；
?????? ·cuFuncSetBlockShape()：設置block的函數；
?????? ·cuFuncSetSharedSize()：設置共享內存大小；
?????? ·cuLaunch()：launch函數；
?????? ·cuLaunchGrid()：launch函數；
?????? ·cuParamSetSize()：設置內核函數參數的長度；
?????? ·cuParamSet*()：設置內核函數參數。
?????? cuParam*()系列函數用于指定在下一次調用launch函數來啟動內核時為內核提供的參數。其第二個參數指定參數在參數堆棧中的偏移。這個偏移量必須與參數類型的對齊要求相匹配。
?????? 具體使用請查詢CUDA編程文檔及驅動API文檔。
?
5.5.6. 存儲器管理
?????? 驅動API存儲器管理與運行時API類似，只是API接口不同。
?????? 存儲器管理的主要函數有：
?????? ·cuMemAlloc()：分配線性存儲器；
?????? ·cuMemAllocPitch()：分配線性存儲器；
?????? ·cuMemFree()：釋放線性存儲器；
?????? ·cuArrayCreate()：創建數組；
?????? ·cuArrayDestroy()：銷毀數組；
?????? ·cuMemcpy()：數據拷貝；
?????? ·cuMemcpy2D()：數據拷貝；
?????? ·cuMemcpy3D()：數據拷貝；
?????? ·cuMemcpyHtoD()：數據拷貝，從主機拷貝到設備。
?????? 具體使用請查詢驅動API文檔。
?
5.5.7. 流管理
?????? 驅動API流管理與運行時API類似，只是API接口不同。
?????? 流管理的主要函數有：
?????? ·cuStreamCreate()：流創建；
?????? ·cuStreamDestroy()：流銷毀；
?????? ·cuStreamQuery()：流查詢；
?????? ·cuStreamSynchronize()：同步流；
?????? ·cuCtxSynchronize()：同步上下文。
?????? 具體使用請查詢驅動API文檔。
?
5.5.8. 事件管理
?????? 驅動API事件管理與運行時API類似，只是API接口不同。
?????? 事件管理的主要函數有：
?????? ·cuEventCreate()：事件創建；
?????? ·cuEventDestroy()：事件銷毀；
?????? ·cuEventElapsedTime()：計算兩事件的時間差；
?????? ·cuEventQuery()：查詢事件；
?????? ·cuEventRecord()：記錄事件；
?????? ·cuEventSynchronize()：事件同步。
?????? 具體使用請查詢驅動API文檔。
?
5.5.9. 紋理參考管理
?????? 驅動API紋理參考管理與運行時API類似，只是API接口不同。
?????? 紋理參考管理的主要函數有：
?????? ·cuTexRefCreate()：創建參考紋理；
?????? ·cuTexRefDestroy()：銷毀紋理參考；
?????? ·cuTexRefSetAddress()：綁定紋理參考；
?????? ·cuTexRefSetArray()：綁定紋理參考；
?????? ·…：其他一系列與紋理參考有關的函數。
?????? 具體使用請查詢驅動API文檔。
?
5.5.10. OpenGL互操作
?????? 驅動API必須使用cuGLInit()初始化與OpenGL的互操作性，其他與運行時API類似。
?????? OpenGL互操作的主要參數有：
?????? ·cuGLInit()：初始化OpenGL互操作性；
?????? ·cuGLRegisterBufferObject()：注冊緩沖對象；
?????? ·cuGLUnregisterBufferObject()：取消注冊緩沖對象；
?????? ·cuGLMapBufferObject()：綁定緩沖對象；
?????? ·cuGLUnmapBufferObject()：取消綁定緩沖對象。
?????? 具體使用請查詢驅動API文檔。
?
5.5.11. Direct3D互操作
?????? 驅動API Direct3D互操作性要求在創建CUDA上下文時指定Direct3D設備。通過使用cuD3D9CtxCreate()而非cuCtxCreate()創建CUDA上下文即可實現此目標。其他與運行時API類型。
?????? Direct3D互操作的主要函數有：
?????? ·cuD3D9CtxCreate()：創建與Direct3D互操作的CUDA上下文；
?????? ·cuD3D9RegisterResource()：注冊資源；
?????? ·cuD3D9UnregisterResource()：取消注冊資源；
?????? ·cuD3D9MapResources()：綁定資源；
?????? ·cuD3D9UnmapResources()：取消綁定資源；
?????? ·cuD3D9ResourceGetMappedPointer()：獲取資源映射后的CUDA設備存儲器地址；
?????? ·cuD3D9ResourceGetMappedSize()：獲取大小；
?????? ·cuD3D9ResourceGetMappedPitch()：獲取間隔。
?????? 具體使用請查詢驅動API文檔。
?
5.6. 性能優化
?????? 性能優化主要有：warp中減少控制指令、合理使用共享內存、防止共享內存bank沖突、單個線程中寄存器使用的量、block中線程數、常量存儲器的合理利用、線程對全局存儲器的合理訪問等。
?????? 多處理器是以warp為單位處理線程的，有控制指令時，會執行完所有的控制指令對應的指令后才會繼續執行下面的指令。舉個例子，if/else語句兩個方向的線程在同一個warp中，線程1執行if方向，線程2執行else方向，它們可能的執行順序是這樣的：
?????? ·線程1執行if方向，線程2等待
?????? ·線程1執行if方向完畢等待，線程2執行else方向
?????? ·線程2執行else方向完畢，線程1、2共同執行后面的指令
?????? 共享內存屬于片上緩存比全局存儲器讀寫速度更快。把一部分全局存儲器上的數據放入共享內存中處理可有效提高性能。共享存儲器的訪問速度和寄存器差不多，大約讀寫4B的數據需要兩個時鐘周期。共享存儲器的讀取是以半warp為單位的，當半warp中的多個線程訪問數組元素處于同一個bank時會發生bank沖突。但假如線程1、2訪問bank1中的同一塊4字節數據，其他的線程訪問互不沖突的bank時不會有bank沖突。當半warp中所有的線程都訪問同一個bank中同一塊4字節的數據時也不會發生bank沖突，稱為廣播訪問，此時只訪問一次bank。每個多處理器中的共享存儲器大小是有限的，應按照block的大小分配合適的共享存儲器。Block的大小會影響多處理器每次激活的block數。
?????? 每個多處理器寄存器數量是有些的，而且在每個線程中寄存器是線程私有的。按照每個多處理器激活的線程數，合理分配寄存器。如果每個線程分配太多線程，則每個多處理器同時激活的線程數就會減少，從而影響并行效果。
?????? Block中的線程數（NThread）也會影響每個多處理器同時激活的線程數。每個多處理器有最大同時激活線程數（NMThread），且每個多處理器有最大同時激活block數（NMBlock）。Block中的線程數滿足：NThread >= NMThread / NMBlock會激活在一個多處理器中可激活的所有block。在其他資源可充分利用的情況下，多處理器上同時激活的線程數越多，效率越高。
?????? 常量存儲器也是帶片上緩存的存儲器。充分利用常量存儲器可有效提升性能。
?????? 最新的設備，全局存儲器都帶有片上緩存。可以利用多處理器處理線程的特性合理訪問全局存儲器的數據，可使更多數據命中。
?
5.7. NVCC編譯器
?????? NVCC編譯器會分離源碼中設備代碼和主機代碼，主機代碼交由一般的C/C++編譯器（gcc等）編譯，設備代碼由NVCC編譯。
?????? 具體編譯命令請查閱NVCC文檔：
http://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html
?
5.8. 設備模擬
?????? 對于未裝支持CUDA的設備或者調試時，可使用設備模擬。設備模擬是將設備執行的代碼由主機模擬執行，設備代碼并不是在設備上執行，而是主機上模擬出多個線程執行。設備模擬的結果和實際設備實際的結果可能不同。
?
5.9. 其他
?????? CUDA并不支持windows的默認遠程登入客戶端（mstsc）登入遠程主機執行設備。需要遠程登入主機執行CUDA設備，可使用VNC工具。
?
參考文獻：
《GPGPU編程技術——從GLSL、CUDA到OpenCL》——仇德元
http://zh.wikipedia.org/wiki/CUDA
http://blogs.nvidia.com/2012/09/what-is-cuda-2/
http://docs.nvidia.com/cuda/cuda-getting-started-guide-for-microsoft-windows/index.html
http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html
http://docs.nvidia.com/cuda/cuda-runtime-api/index.html
總結
- 上一篇： CUDA学习2
- 下一篇： c++的:: . ：-的区别