上一篇我介紹了cuda的基本知識，本篇我將會介紹有關主存和顯存的相關概念和二者的聯系。
__host__，__device__與__global__修飾函數
cuda中引入了三個宏：__host__、__device__與__global__，用於修飾函數，使得函數被定位到不同的位置。
那修飾后的函數有什么作用呢？
__host__函數，其實就是我們平常寫C/C++所定義的運行在CPU中的函數，這個修飾符通常可以不寫，效果是等價的。而__device__函數和__global__函數則是必須運行在GPU的函數，因此必須要顯式聲明在函數前。
我們來看下例：
1 #include <cstring>
2 #include <cstdlib>
3 #include <cassert>
4
5 #include "cuda_runtime.h"
6 #include "device_launch_parameters.h"
7
8 __device__ double triple(double x) { 9 //返回x的三倍
10 return x * 3; 11 } 12
13 __global__ void kern_AddVector(double* c, double const* a, double const* b, size_t n) { 14 //求向量c = a + 3b
15 size_t Idx = blockIdx.x * blockDim.x + threadIdx.x; 16 if(Idx >= n) return; //超過數組大小，直接返回
17
18 c[Idx] = a[Idx] + triple(b[Idx]); //實現向量相加
19 } 20
21 __host__ void addVector(double* c, double const* a, double const* b, size_t n) { 22 //申請顯存內地址
23 double *device_c, *device_a, *device_b; 24 assert(cudaSuccess == cudaMalloc(&device_c, sizeof(double) * n)); 25 assert(cudaSuccess == cudaMalloc(&device_a, sizeof(double) * n)); 26 assert(cudaSuccess == cudaMalloc(&device_b, sizeof(double) * n)); 27 //將數據拷貝到顯存之中
28 assert(cudaSuccess == cudaMemcpy(device_a, a, sizeof(double) * n, cudaMemcpyHostToDevice)); 29 assert(cudaSuccess == cudaMemcpy(device_b, b, sizeof(double) * n, cudaMemcpyHostToDevice)); 30 //執行核函數
31 size_t thread_count = 1024; 32 size_t block_count = (n - 1) / thread_count + 1; 33 kern_AddVector<<<block_count, thread_count>>> (device_c, device_a, device_b, n); 34 cudaDeviceSynchronize(); 35 cudaError_t ct = cudaGetLastError(); 36 assert(cudaSuccess == ct); 37 //將顯存中的數據拷貝到主存中
38 assert(cudaSuccess == cudaMemcpy(c, device_c, sizeof(double) * n, cudaMemcpyDeviceToHost)); 39 //釋放臨時變量
40 assert(cudaSuccess == cudaFree(device_a)); 41 assert(cudaSuccess == cudaFree(device_b)); 42 assert(cudaSuccess == cudaFree(device_c)); 43 } 44
45 #include <cstdio>
46
47 int main() { 48 const size_t N = 10; 49 double a[N] = {0.1, 0.2, -0.3, 0.1, 0.5, -0.2, 0.2, -0.3, 0.4, 0.1}; 50 double b[N] = {0.2, -0.1, -0.1, 0.2, -0.2, 0.2, 0.1, 0.1, 0.1, 0.3}; 51 double c[N]; 52 addVector(c, a, b, N); 53 for(double& e: c) { 54 printf("%lf, ", e); 55 } 56
57 return 0; 58 }
例子中addVector函數就是__host__函數，當然main函數也是__host__函數。
__host__函數可以直接調用__host__函數，但不能直接調用__device__函數；__host__函數可以通過傳遞運行時參數來調用__global__函數，同樣也不能像調用__host__函數那樣直接調用。而能調用__device__函數的只有__global__函數或者__device__函數。
比如如果main函數這樣寫，就會報錯：
1 int main() { 2 //__host__函數直接調用__device__函數
3 double e3 = triple(e); //error: calling a __device__ function("triple") from a __host__ function("main") is not allowed 4
5 //不傳遞運行時參數調用__global__函數
6 kern_AddVector(c, a, b, N); //error: a __global__ function call must be configured
7 }
事實上，我們現在大多數的顯卡都已經支持了sm_50, compute_50及以上的計算能力（我們可以使用上一篇中提到的deviceQuery來獲取自己顯卡的計算能力），而這更加豐富了我們的調用關系——50之前cuda沒有調用棧，所有__device__函數在編譯的時候都是內聯的；但50之后，__device__函數可以通過調用__device__函數實現直接或者間接的遞歸；而__device__和__global__函數也可以繼續通過傳遞運行時參數調用__global__函數，實現二級甚至二級以上的並行。用拓撲結構圖來表示則是：
cuda的函數調用關系
除了函數被宏修飾，變量也可以被修飾。
__device__，__shared__與__constant__修飾變量
__device__、__shared__與__constant__也是cuda的宏，用於修飾變量（別忘了__device__也可以修飾函數）。三種變量都不會被聲明在CPU中，而是在GPU中。
__device__變量即設備端的全局變量，和C/C++的全局變量聲明位置一樣，只能在所有類和函數外聲明。__host__函數無法直接訪問__device__變量，但可以通過cuda運行庫中的cudaMemcpyToSymbol()以及cudaMemcpyFromSymbol()函數傳遞或獲取到它的值。__device__函數和__global__函數可以直接訪問它們，只需要注意不要線程沖突就好。
__shared__變量即塊內共享變量，只能在__device__函數或者__global__函數內被聲明。__shared__變量不能跨過一個線程塊，所以聲明時其所在的__global__函數的運行時變量中的塊數往往是1——當然也可以是更大的值，但某一個塊中的__shared__變量就無法被其他塊所訪問到。變量聲明時不能初始化，但可以對它進行賦值。
__constant__變量即設備端的常量，並不像它的名字那樣一成不變——但至少它在__device__函數和__global__函數中的訪問權限是只讀的，這樣它就可以被放在高速緩存中，極大地提升訪問效率。聲明方法又和C/C++不同：聲明時賦初值是無效的，必須在__host__函數中通過cuda運行庫中的cudaMemcpyToSymbol()函數傳遞給它；當然，__host__函數內部也可以用cudaMemcpyFromSymbol()函數獲取到它的值。
於是我們又可以豐富上圖：
cuda的函數調用關系，以及設備端變量的訪問權限
圖片是1920×1080的，所以想拿去做壁紙也是沒問題的（狗頭）。
而如果變量前面沒有修飾，那就是寄存器變量（就像C/C++里的寄存器變量），如果是在__device__函數或者__global__函數內，那么每個線程分別持有一個該變量，不會共享，對其讀取和修改也只會發生在該線程內。
不過要注意，__device__和__constant__只能聲明在全局變量區域，__shared__變量只能聲明在核函數內部，類的成員變量和其他函數內的局部變量是無法被上述關鍵詞修飾的。
下面代碼是一個例子：
1 #include <cstdio>
2 #include <cassert>
3
4 #include "cuda_runtime.h"
5 #include "device_launch_parameters.h"
6
7 #define N 10
8 __device__ int arr[N]; 9
10 __global__ void print() { 11 size_t Idx = blockIdx.x * blockDim.x + threadIdx.x; 12 if(Idx >= N) return; 13
14 printf("%d\n", arr[Idx]); 15 } 16
17 int main() { 18 int a[N] = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9}; 19 assert(cudaSuccess == cudaMemcpyToSymbol(arr, a, sizeof(a))); 20 print<<<10, 1>>> (); 21 cudaDeviceSynchronize(); 22 assert(cudaSuccess == cudaGetLastError()); 23 }
這里沒有使用前文例子中的thread_count和block_count，因為我們明確知道線程數是遠小於1024的，甚至核函數內也不需要寫大於N則返回的邏輯。
當然，由於是多線程，所以輸出是亂序的。但如果交換核數和線程數，因為一個核內線程是輪轉調度的，所以輸出是順序的。
cudaMalloc、cudaFree、cudaMemset與cudaMemcpy
cuda_runtime.h庫中包含了一些和標准C語言庫中的函數非常相近的__host__函數——注意，他們只能在__host__函數中被調用，__global__函數和__device__函數要調用函數原型。
這些函數原型是：
1 cudaError_t cudaMalloc(void **devPtr, size_t size); 2 cudaError_t cudaFree(void *devPtr); 3 cudaError_t cudaMemset(void *devPtr, int value, size_t count); 4 cudaError_t cudaMemcpy(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind); 5 cudaError_t cudaMemcpyToSymbol(const void *symbol, const void *src, size_t count, size_t offset = 0, enum cudaMemcpyKind kind = cudaMemcpyHostToDevice); 6 cudaError_t cudaMemcpyFromSymbol(void *dst, const void *symbol, size_t count, size_t offset = 0, enum cudaMemcpyKind kind = cudaMemcpyDeviceToHost);
前四個函數，我們可以通過名字找到它們在C/C++里的“近親”：
1 void* malloc(size_t _Size); 2 void free(void* _Block); 3 void* memset(void* _Dst, int _Val, size_t _Size); 4 void* memcpy(void* _Dst, void const* _Src, size_t _Size);
但是這些函數又清一色地返回了cudaError_t這一枚舉類型，所以，我們在申請顯存空間時，寫法應為：
1 double* p; const size_t N = 10; 2 cudaError_t ct = cudaMalloc(&p, sizeof(double) * N); 3 assert(cudaSuccess == ct); 4
5 //對比在內存中的malloc：
6 p = (double*)malloc(sizeof(double) * N);
相信大家也猜到了cudaFree和cudaMemset的用法，事實上我並不喜歡額外創建一個cudaError_t變量，而是直接放在assert中：
1 assert(cudaSuccess == cudaMemset(p, 0, sizeof(double) * N)); 2 assert(cudaSuccess == cudaFree(p)); 3
4 //對比在內存中的memset、free：
5 memset(p, 0, sizeof(double) * N); 6 free(p);
cudaMemcpy函數不同於memcpy，它有第四個參數，是cudaMemcpyKind枚舉類型，其聲明如下：
1 enum __device_builtin__ cudaMemcpyKind 2 { 3 cudaMemcpyHostToHost = 0, /**< Host -> Host */
4 cudaMemcpyHostToDevice = 1, /**< Host -> Device */
5 cudaMemcpyDeviceToHost = 2, /**< Device -> Host */
6 cudaMemcpyDeviceToDevice = 3, /**< Device -> Device */
7 cudaMemcpyDefault = 4 /**< Direction of the transfer is inferred from the pointer values. Requires unified virtual addressing */
8 };
相信注釋也寫得非常清楚了：
- cudaMemcpyHostToHost就是從主機端拷貝到主機端，即此時cudaMemcpy等價於memcpy，不屬於I/O，耗時最短；
- cudaMemcpyHostToDevice則是從主機端傳送到設備端，即源數據在內存中，目標指針指向了一段顯存范圍，屬於I/O，消耗時間較長；
- cudaMemcpyDeviceToHost則是從設備端傳送到主機端，即源數據在顯存中，目標指針指向了一段內存范圍，同樣屬於I/O，消耗時間較長；
- cudaMemcpyDeviceToDevice則是從設備端拷貝到設備端，CPU只給顯卡發送一個信號，不涉及數據交互，因此不屬於I/O，不會消耗太多時間。大多數情況下可以異步執行。
下列代碼則是一些例子：
1 /*
2 * host_a、host_b是經過malloc或new，或者全局、局部變量的數組，包含有N個int 3 * device_a、device_b是經過cudaMalloc的數組，同樣包含有N個int 4 */
5
6 cudaMemcpy(host_a, host_b, sizeof(int) * N, cudaMemcpyHostToHost); //正確
7 cudaMemcpy(device_b, host_b, sizeof(int) * N, cudaMemcpyHostToDevice); //正確
8 cudaMemcpy(host_a, device_a, sizeof(int) * N, cudaMemcpyDeviceToHost); //正確
9 cudaMemcpy(device_b, device_a, sizeof(int) * N, cudaMemcpyDeviceToDevice); //正確
10
11 cudaMemcpy(host_a, device_a, sizeof(int) * N, cudaMemcpyHostToDevice); //錯誤，函數返回一個cudaErrorInvalidValue
12 cudaMemcpy(device_b, device_a, sizeof(int) * N, cudaMemcpyHostToHost); //錯誤，函數返回一個cudaErrorInvalidValue
13 cudaMemcpy(device_b, host_b, sizeof(int) * N, cudaMemcpyDeviceToDevice); //錯誤，函數返回一個cudaErrorInvalidValue
14 cudaMemcpy(host_a, host_b, sizeof(int) * N, cudaMemcpyDeviceToHost); //錯誤，函數返回一個cudaErrorInvalidValue
而cudaMemcpyToSymbol()和cudaMemcpyFromSymbol()兩個函數，前文也提到了，是用來初始化__device__顯存全局變量和__constant__顯存常量的。雖然函數有五個變量，但后兩個變量我們一般只用其初始值，所以寫法通常為：
1 __constant__ int arr[N]; 2
3 __host__ void init() { 4 int a[N] = {9, 8, 7, 6, 5, 4, 3, 2, 1, 0}; 5 int b[N]; 6 assert(cudaSuccess == cudaMemcpyToSymbol(arr, a, sizeof(a))); 7 assert(cudaSuccess == cudaMemcpyFromSymbol(b, arr, sizeof(b))); 8 }
可是正如前文所說的，這些函數都是__host__函數，只能在__host__函數中被調用。那么__global__和__device__函數該如何申請、復制、修改和釋放顯存數據呢？
設備端的malloc、free、memset和memcpy
沒錯，這就是答案——在__global__和__device__函數中使用函數的原型：
1 template<typename T>
2 __global__ void buildList(T** arrs, size_t size, size_t tot_list) { 3 size_t Idx = blockIdx.x * blockDim.x + threadIdx.x; 4 if(Idx >= tot_list) return; 5
6 arrs[Idx] = (T*)malloc(sizeof(T) * size); 7 memset(arrs[Idx], 0, sizeof(T) * size); 8 } 9
10 template<typename T>
11 __global__ void copyList(T** dsts, const T* const* srcs, size_t size, size_t tot_list) { 12 size_t Idx = blockIdx.x * blockDim.x + threadIdx.x; 13 if(Idx >= tot_list) return; 14
15 memcpy(dsts[Idx], srcs[Idx], sizeof(T) * size); 16 } 17
18 template<typename T>
19 __global__ void clearList(T** arrs, size_t tot_list) { 20 size_t Idx = blockIdx.x * blockDim.x + threadIdx.x; 21 if(Idx >= tot_list) return; 22
23 free(arrs[Idx]); 24 }
這三個函數實現了長度為tot_list的指針數組的每一個元素並行申請大小為size的內存並初始化，深拷貝指針數組，以及並行釋放指針數組中的每一個元素的功能。
編程中可能出現的異常
說了半天，也沒有提到cudaError_t具體會返回什么異常。
首先我們看一下枚舉類型cudaError_t的常見值：
- cudaSuccess = 0。這是幾乎所有程序繼續運行下去的基礎，即未發生任何錯誤。
- cudaErrorInvalidValue =1。在初學者身上比較常見意為傳入API函數的值不在合法區間范圍內。通常是一些低級錯誤，比如在初始化常量時沒有使用cudaMemcpyToSymbol而是錯誤使用了cudaMemcpy、或是在cudaMalloc、cudaMemcpy等函數中傳遞了空指針等。
- cudaErrorMemoryAllocation = 2。通常是需要申請內存的函數如cudaMalloc會返回這個錯誤，一般是申請的內存超過了可用顯存大小。
- cudaErrorInitializationError = 3。任何runtime庫中的函數都有可能返回這個異常，但只有可能在第一次調用時返回。因為cuda的初始化方法是lazy context initialization，即直到調用才會初始化，並不會在程序一開始就初始化。
- cudaErrorCudartUnloading = 4。出現這個異常大多都是誤刪了cuda驅動。如果出現這個異常，請自行懺悔。
- cudaErrorInvalidConfiguration = 9。通常是傳遞運行時參數時超過了顯卡的負載，如線程數大於deviceQuery輸出的每個核的最大線程數、核數大於網格中最大核數等等。
- cudaErrorInvalidPitchValue = 12。一般是在使用cudaMemcpy2D、cudaMemcpy3D等拷貝高維數組時，Pitch出現了問題——可能沒有申請Pitch，或者Pitch的地址出錯等等。
- cudaErrorInvalidSymbol = 13。即對顯存全局變量和常量進行相關操作時，符號名稱出錯，或進行了多余的格式轉換。如你想將數組a拷貝給顯存常量arr時，傳遞的第一個參數可以是單純的arr，也可以是加引號的"arr"，如果寫成是轉化過的(void*)arr，就會返回這一錯誤。
- cudaErrorDuplicateVariableName = 43。意為你在定義全局變量時，出現了多個變量重名的情況，可能在同一文件中，也可能在鏈接前的不同文件中。
- cudaErrorNoDevice = 100。你需要檢查你的顯卡是否支持cuda。
- cudaErrorFileNotFound = 301。找不到指定文件。
- cudaErrorSymbolNotFound = 500。找不到符號名。通常是在通過字符串尋找設備符號時出現的，此時需要檢查你的拼寫。
- cudaErrorIllegalAddress = 700。你可能搞錯了傳入API的指針究竟指向了內存空間還是顯存空間，或者在核函數訪問時發生了數組越界等等，產生了非法地址。一旦出現了這個問題，程序就必須終止才能繼續使用cuda。
- cudaErrorLaunchOutOfResources = 701。你可能使用了過多的線程數或寄存器數，可以deviceQuery一下，然后在項目設置中限制一下寄存器的使用。
- cudaErrorAssert = 710。即在__global__或__device__函數中的斷言assert被觸發，在觸發的同時cuda往往也會將具體行數、核坐標、線程坐標的信息打印出來。一旦出現了這個問題，程序就必須終止才能繼續使用cuda。
- cudaErrorHardwareStackError = 714。通常是棧溢出，可能是你在遞歸__global__或__device__函數的層數太多，或函數內局部變量數組開得太大。一旦出現了這個問題，程序就必須終止才能繼續使用cuda。
- cudaErrorLaunchFailure = 719。在執行核函數時發生了內核異常，通常是設備共享內存越界、取消引用無效設備指針等等。一旦出現了這個問題，程序同樣必須終止之后才能繼續使用cuda。
有關錯誤代碼我將在CUDA教程四中詳細介紹。
有關主存、顯存、runtime庫的基本函數以及cudaError_t的介紹就是這些了。