CUDA學習一：CUDA C簡介
#1、一個程式來了解CUDA C
#include <stdio.h> #include "commom/book.h" //__global__表示函式在裝置而非主機上執行，add函式由編譯裝置程式碼的編譯器編譯 __global__ void add(int a, int b, int *c){ *c = a + b; } int main(void) { int c; int *dev_c; HANDLE_ERROR（ cudaMalloc（（void**）&dev_c, sizeof(int)））; add<<<1,1>>>(2,7,dev_c); HANDLE_ERROR( cudaMemcpy( &c dev_c, sizeof(int), cudaMemcpyDeviceToHost)); printf("2 + 7 = %d\n",c); cudaFree(dev_c); return 0; }
上面的這個小程式就是利用CUDA C實現一個加法運算，其中加法的運算在裝置上執行，結果拷貝到主句的記憶體中並列印輸出。
以上是函式的大致功能，下面重點對程式內部關於裝置/主機程式碼與C語言區別部分進行說明：
##1.1、核函式的呼叫
//__global__表示函式在裝置而非主機上執行，add函式由編譯裝置程式碼的編譯器編譯
__global__ void add(int a, int b, int *c){
*c = a + b;
}
CUDA C為標準的C增加了__ global_ _ 修飾符，這個修飾符告訴編譯器函式應該編譯為在裝置而不是在主機上執行，上面的例子中，add函式將被交給編譯裝置程式碼的編譯器而main函式將被交給主機的編譯器。
add<<<1,1>>>(2,7,dev_c);
add<<<1,1>>>:尖括號表示將一些引數傳遞各執行時系統，這些引數並不是傳遞給裝置程式碼的引數而是告訴執行時如何啟動裝置程式碼。傳遞給裝置程式碼本身的引數是放在圓括號中傳遞的，和標準的函式呼叫相同。
##1.2、傳遞引數
引數傳遞包含的概念：
可以像呼叫C函式那樣將引數傳遞給核函式；
當裝置執行任何有用的操作時，都需要分配記憶體，例如將計算值返回給主機。
###1.2.1、cudaMalloc和malloc
裝置通過cudaMalloc來分配記憶體，第一個引數是一個指標用於儲存新分配記憶體地址的變數，第二個引數是分配記憶體的大小，但是與malloc不同的是分配的這部分記憶體的指標並不是作為cudaMalloc的返回值，並且返回值的型別為void*。
HANDLE_ERROR（）是定義的一個巨集，這個巨集只判斷函式呼叫是否返回了一個錯誤的值，如果是返回錯誤的資訊，退出程式並將退出碼設定為EXIT_FAILURE。
注意：
裝置指標的使用限制如下：
1、可以將cudaMalloc（）分配的指標傳遞給在裝置上執行的函式；
2、可以在裝置的程式碼中使用cudaMalloc（）分配的指標進行記憶體的讀/寫操作；
3、可以將cudaMalloc()分配的指標傳遞給主機上執行的函式；
4、不能在主機的程式碼中使用cudaMalloc（）分配的指標進行記憶體的讀/寫操作。
5、cudaMalloc（）申請的記憶體釋放時需要呼叫cudaFree()
###1.2.2、cudaMemcpy和Memcpy
對裝置記憶體的訪問常見的有兩種方法：（1）在裝置程式碼中直接使用裝置指標（2）呼叫cudaMemcpy（）函式。
主機的指標只能訪問主機程式碼中的記憶體，裝置指標只能訪問裝置程式碼中的記憶體。
主機程式碼可以通過呼叫cudaMemory（）來訪問裝置上的記憶體，這個函式的呼叫類似於標準的C中的memcpy（），只不過多了一個引數來設定裝置記憶體指標究竟是源指標還是目標指標。
如：
HANDLE_ERROR( cudaMemcpy( &c
dev_c,
sizeof(int),
cudaMemcpyDeviceToHost));
中的cudaMemcpyDeviceToHost這個引數告訴我們執行時源指標是一個這杯指標，目的指標是一個主機指標，所表達的意思是訪問裝置中對應記憶體地址中的資料，將其存入到主機的地址中。
#2、裝置屬性的認識
一方面，一個主機可能含有整合以及外掛的獨立GPU，這就形成了多GPU支援CUDA的情況；另一方面，由於CUDA執行時本身並不能保證應用程式選擇最優或者最合適的GPU，這就要求我們能通過CUDA提供的介面來搞清楚有多少支援CUDA的裝置以及每個裝置的具體屬性。
##2.1、查詢裝置
在深入的學習裝置程式設計之前，我們要學會通過某種機制來判斷計算機當前有哪些裝置以及每個裝置支援哪些功能。我們可以通過一個簡單的介面獲得這些資訊。
首選我們希望知道有多少個裝置支援CUDA架構並且這些裝置能夠執行CUDA的核函式，通過呼叫函式cudaGetDeviceCount（）：
int count;
HANDLE_ERROR(cudaGetDeviceCount(&count));
這個過程對每個裝置進行迭代，查詢每個裝置的信心，對於支援CUDA的裝置返回該裝置的屬性資訊，也即是cudaDeviceProp的資料結構。該資料結構包含以下資訊：
/**
* CUDA device properties
*/
struct cudaDeviceProp
{
char name[256]; /**< 標示裝置的ASCII字串(例如，“GeForce GT 750M”) */
size_t totalGlobalMem; /**< 裝置上全域性記憶體的總量，單位為位元組 */
size_t sharedMemPerBlock; /**< 在一個執行緒塊中（block）可以使用的最大共享記憶體的數量，單位為位元組 */
int regsPerBlock; /**< 每個執行緒中可用的32位暫存器的數量 */
int warpSize; /**< 在一個執行緒束（Warp）中包含的執行緒的數量 */
size_t memPitch; /**< 記憶體複製中最大的修正量（pitch），單位為位元組 */
int maxThreadsPerBlock; /**< 一個執行緒塊中可以包含的最大的執行緒數量 */
int maxThreadsDim[3]; /**< 在多維的執行緒塊陣列中每一維可以包含的最大執行緒數量 */
int maxGridSize[3]; /**< 在多維的執行緒格（Grid）中，每一維可以包含的執行緒塊的數量 */
int clockRate; /**< 時鐘頻率 單位KHZ */
size_t totalConstMem; /**< 常量記憶體的總量 */
int major; /**< 裝置計算功能集（compute capability）的主版本號 */
int minor; /**< 裝置計算功能集的次版本號 */
size_t textureAlignment; /**< 裝置的紋理對齊（texture Alignment）要求 */
size_t texturePitchAlignment; /**< Pitch alignment requirement for texture references bound to pitched memory */
int deviceOverlap; /**< 裝置是否可以同時執行一個cudaMemory()呼叫和一個核函式的呼叫 */
int multiProcessorCount; /**< 裝置上多處理器的數量 */
int kernelExecTimeoutEnabled; /**< 裝置上執行的核函式是否存在執行時的限制 */
int integrated; /**< 裝置是否是一個整合的GPU */
int canMapHostMemory; /**< 裝置是否將主機記憶體對映到CUDA的裝置地址空間 */
int computeMode; /**< 裝置的計算模式預設（default），獨佔（Exclusive），禁止（Prohibited） (See ::cudaComputeMode) */
int maxTexture1D; /**< 一維紋理的最大大小 */
int maxTexture1DMipmap; /**< Maximum 1D mipmapped texture size */
int maxTexture1DLinear; /**< Maximum size for 1D textures bound to linear memory */
int maxTexture2D[2]; /**< 二維紋理的最大大小 */
int maxTexture2DMipmap[2]; /**< Maximum 2D mipmapped texture dimensions */
int maxTexture2DLinear[3]; /**< Maximum dimensions (width, height, pitch) for 2D textures bound to pitched memory */
int maxTexture2DGather[2]; /**< Maximum 2D texture dimensions if texture gather operations have to be performed */
int maxTexture3D[3]; /**< 三維紋理的最大大小 */
int maxTextureCubemap; /**< Maximum Cubemap texture dimensions */
int maxTexture1DLayered[2]; /**< Maximum 1D layered texture dimensions */
int maxTexture2DLayered[3]; /**< Maximum 2D layered texture dimensions */
int maxTextureCubemapLayered[2];/**< Maximum Cubemap layered texture dimensions */
int maxSurface1D; /**< Maximum 1D surface size */
int maxSurface2D[2]; /**< Maximum 2D surface dimensions */
int maxSurface3D[3]; /**< Maximum 3D surface dimensions */
int maxSurface1DLayered[2]; /**< Maximum 1D layered surface dimensions */
int maxSurface2DLayered[3]; /**< Maximum 2D layered surface dimensions */
int maxSurfaceCubemap; /**< Maximum Cubemap surface dimensions */
int maxSurfaceCubemapLayered[2];/**< Maximum Cubemap layered surface dimensions */
size_t surfaceAlignment; /**< Alignment requirements for surfaces */
int concurrentKernels; /**< 是否支援在同一個上下文中同時執行多個核函式 */
int ECCEnabled; /**< Device has ECC support enabled */
int pciBusID; /**< PCI bus ID of the device */
int pciDeviceID; /**< PCI device ID of the device */
int pciDomainID; /**< PCI domain ID of the device */
int tccDriver; /**< 1 if device is a Tesla device using TCC driver, 0 otherwise */
int asyncEngineCount; /**< Number of asynchronous engines */
int unifiedAddressing; /**< Device shares a unified address space with the host */
int memoryClockRate; /**< Peak memory clock frequency in kilohertz */
int memoryBusWidth; /**< Global memory bus width in bits */
int l2CacheSize; /**< Size of L2 cache in bytes */
int maxThreadsPerMultiProcessor;/**< Maximum resident threads per multiprocessor */
};
##2.2、裝置屬性的使用
在不同的計算任務到來的時候，我們並不一定需要了解所有的裝置的所有的屬性，雖然我們可以呼叫函式cudaGetDevicePropertise()來查詢。
有一種自動的方式來執行迭代訪問和選擇操作：
1、首先找到我們希望裝置擁有的屬性並將這些屬性填充到一個cudaDeviceProp結構中。
CUDADeviceProp prop;
memset( &prop, 0, sizeof(cudaDeviceProp));
prop.major = 1;
prop.minor = 3;
2、填充完cudaDeviceProp結構後，將其傳遞給cudaChooseDevice()，這樣CUDA執行時會檢查滿足要求的裝置。
3、cudaChooseDevice（）函式將返回一個裝置ID，然後我們將這個裝置的ID傳遞給cudaSetDevice（），隨後所有的裝置操作都將在這個裝置上執行。