也許有人注意到了，我在【CUDA教程】二、主存與顯存文章中提到了部分常見的異常。實際上，cuda編程最終Boss則是debug。本文將重點講解cuda中錯誤的成因，作為“報錯詞典”供各位開發者們debug。
本文將盡可能全面地列舉所有異常的可能出現情況，如需快速找到問題原因，請使用ctrl+F的頁內搜索功能，檢索內容為cudaError_t枚舉類型的成員名（如"cudaErrorLaunchOutOfResources"）、錯誤代碼（如"701"）和出錯信息（如"too many resources requested for launch"）。
調試技巧
cuda的樣例代碼使用了如下方法處理異常：
1 cudaError_t cudaStatus; 2
3 // Choose which GPU to run on, change this on a multi-GPU system.
4 cudaStatus = cudaSetDevice(0); 5 if (cudaStatus != cudaSuccess) { 6 fprintf(stderr, "cudaSetDevice failed! Do you have a CUDA-capable GPU installed?"); 7 goto Error; 8 }
事實上，這種做法太過冗長，每次使用cuda函數都要寫六行來處理異常，一個函數的代碼量甚至會提升2~3倍；同時代碼的連貫性也被打破，后期維護起來非常困難。因此我在前面代碼量不大的代碼中使用了這種方法：
assert(cudaSuccess == cudaSetDevice(0));
這種方法解決了上述所有問題，但引入了新的問題，如果發生Assert Failed，我們需要獲得異常的返回值來進行處理，因此需要修改代碼。但如果發生了無法重現（Unable to Reproduce）的錯誤，在多次后續實驗中難以再次觸發，白白浪費了潛在的漏洞修復時機。作為折衷，我們可以定義如下的宏和函數：
1 #define HANDLE_ERROR(err) (HandleError(err, __FILE__, __LINE__))
2
3 static void HandleError(cudaError_t err, const char *file, int line) { 4 if (err != cudaSuccess) { 5 fprintf(stderr, "Error %d: \"%s\" in %s at line %d\n", int(err), cudaGetErrorString(err), file, line); 6 exit(int(err)); 7 } 8 }
實際使用時則只需要：
HANDLE_ERROR(cudaSetDevice(0));
便可將異常無緩沖地打印在標准錯誤流中。
實際編程時請將每一個cuda函數都加上判斷返回值的異常處理體系，否則調試量將會十分顯著地上升！
編譯時錯誤
在討論運行時錯誤（Runtime Error）之前，我們先討論一下編譯時錯誤（Compile Error）的成因。
error C2059: 語法錯誤:“<”
也許是你在新建工程時並沒有建一個CUDA Runtime工程，檢查你是否使用這種方法建立項目：
使用CUDA項目模板建立CUDA工程的方法
還有可能是報錯的cu文件並沒有用CUDA C/C++方法被編譯（但與后綴無關，“.cu”文件后綴名完全可以改為“.cpp”，只要保持使用nvcc編譯器）：
修改項類型，以使用nvcc編譯
還有可能是不規范的include所導致，如cpp文件包含了cu文件：
此時注釋掉第一行即可解決問題
error LNK2005: 函數名 已經在 xxx.cu.obj 中定義
這個錯誤往往會伴隨“fatal error LNK1169: 找到一個或多個多重定義的符號”出現。
在此吐槽一下英某達鏈接器的程序猿——只能鏈接多個lib和一個.cu文件。並建議開發者們摒棄原有的多文件編程方法，函數和類聲明即實現：
傳統多文件編程技巧——頭文件聲明函數，源文件實現函數
cuda正確的多文件編程方法，聲明即實現
運行時錯誤
cuda的Runtime API全部都會帶有返回值，其類型為cudaError_t；而cuFFT API全部帶有cufftResult類型的返回值，curand API則全部帶有curandStatus_t類型返回值。
這里重點介紹Runtime API返回值各錯誤碼的成因（實驗環境：cuda 10.2，sm_61，compute_61）。
cudaSuccess = 0，"no error"
傻孩子，no error就是沒問題，success就是成功，這個函數沒有報錯，繼續運行吧。
不過要注意的是，__global__函數是異步執行的，如果需要與CPU同步，還需要使用cudaDeviceSynchronize()函數實現同步。所以如果調用完核函數后馬上調用cudaGetLastError()，很可能返回cudaSuccess，但核函數運行到某一位置時仍然報錯。建議將cudaDeviceSynchronize()函數放在cudaGetLastError()前（但不要用cudaDeviceSynchronize()直接替換掉cudaGetLastError()，有某些異常如cudaErrorInvalidConfiguration並不會在cudaDeviceSynchronize()中報錯，而是在cudaGetLastError()中被返回）。
cudaErrorInvalidValue = 1，"invalid argument"
如果出現這個問題，大概率是指針問題。請檢查報錯的函數傳遞的參數是不是空指針或野指針，是不是錯把指向host端內存的指針當作指向device端內存的指針（或相反）傳進了API等。此外還有傳入API參數時超過了其范圍，如不正常值的枚舉等。如下述四種情況代碼均將返回這類錯誤代碼：
1.使用空指針或野指針
1 double** pp; 2 cudaError_t ct = cudaMalloc(pp, sizeof(double) * 1024); //野指針傳入API作為參數
3 printf("%s\n", cudaGetErrorString(ct)); //"invalid argument"
2. 搞混host指針與device指針
1 double* p; 2 HANDLE_ERROR(cudaMallocHost(&p, sizeof(double) * 1024)); //申請host端內存的函數
3 cudaError_t ct = cudaFree(p); //釋放device端內存的函數
4 printf("%s\n", cudaGetErrorString(ct)); //"invalid argument"
3. 使用已經釋放的指針
1 double* p; 2 HANDLE_ERROR(cudaMalloc(&p, sizeof(double) * 1024)); 3 HANDLE_ERROR(cudaFree(p)); //釋放已經申請的device端內存
4 cudaError_t ct = cudaFree(p); //再次釋放或修改對應值（如cudaMemset等）
5 printf("%s\n", cudaGetErrorString(ct)); //"invalid argument"
4. 使用錯誤的枚舉值
1 int limit_type = 16; //不存在的cudaLimit類型枚舉值
2 cudaError_t ct = cudaDeviceSetLimit((cudaLimit)limit_type, 100 * 1024); 3 printf("%s\n", cudaGetErrorString(ct)); //"invalid argument"
出現這類問題后，cuda仍可繼續提供服務。
但是下面這幾種情況並不會返回cudaErrorInvalidValue，甚至會返回cudaSuccess：
1.過大的內存申請量（返回cudaErrorMemoryAllocation = 2）
1 double* p; //顯存4GB，申請8TB試試
2 cudaError_t ct = cudaMalloc(&p, sizeof(double) * 1024ull * 1024ull * 1024ull * 1024ull); 3 printf("%s\n", cudaGetErrorString(ct)); //"out of memory"
2. 過大的空間上限（返回cudaSuccess = 0）
1 cudaError_t ct = cudaDeviceSetLimit(cudaLimitMallocHeapSize, 1024ull * 1024ull * 1024ull * 1024ull); 2 printf("%s\n", cudaGetErrorString(ct)); //"no error"
cudaErrorMemoryAllocation = 2，"out of memory"
顧名思義，內存爆掉了。最常見的就是調用cudaMalloc時超過了堆內存的最大限制，當然cudaHostAlloc、cudaMallocHost等都會返回這一問題，核函數內的malloc也有產生這一問題的可能性。此時要使用cudaDeviceSetLimit增大堆內存上限，或檢查是否發生了內存泄漏，並及時Free掉多余的內存。
出現這類問題后，cuda仍可繼續提供服務，僅拒絕分配給用戶所申請的內存空間而已。如下例：
1 HANDLE_ERROR(cudaDeviceSetLimit(cudaLimitMallocHeapSize, 1024 * 1024 * 1024)); 2 int *p, *q; 3 cudaError_t ct = cudaMalloc(&p, sizeof(int) * 200 * 1024 * 1024); 4 printf("%s\n", cudaGetErrorString(ct)); //"no error"
5 ct = cudaMalloc(&q, sizeof(int) * 200 * 1024 * 1024); 6 printf("%s\n", cudaGetErrorString(ct)); //"out of memory"
7 HANDLE_ERROR(cudaFree(p)); 8 ct = cudaMalloc(&q, sizeof(int) * 200 * 1024 * 1024); 9 printf("%s\n", cudaGetErrorString(ct)); //"no error"
但是下面這幾種情況並不會返回cudaErrorMemoryAllocation：
1.棧溢出（返回cudaErrorLaunchFailure = 719）
1 __device__ int func(int n) { 2 if(n <= 1) return 0; 3 double fxxker[256]; //較深的遞歸與較多的局域變量使棧溢出
4 return func(n - 1) + func(n - 2) + 1; 5 } 6 __global__ void fxxk_stack() { 7 printf("%d\n", func(32)); 8 } 9
10 //main
11 fxxk_stack<<<1, 1>>>(); 12 cudaError_t ct = cudaDeviceSynchronize(); 13 printf("%s\n", cudaGetErrorString(ct)); //"unspecified launch failure"
cudaErrorInitializationError = 3，"initialization error"
正如其名，初始化錯誤。由於cuda使用了lazy context initialization，這一錯誤在任何API中都可以返回，但只有第一次被調用時會返回。
出現這類問題后，程序必須終止后重啟才能重新使用cuda服務，畢竟初始化只會進行一次，失敗了就也辦法重來了。
出現這類問題，往往是cuda動態或靜態庫文件被誤刪造成的，重裝cuda幾乎一定可以解決。
cudaErrorCudartUnloading = 4，"driver shutting down"
面壁吧，沒事卸載cuda驅動干什么（狗頭）。
出現這類問題后，程序必須終止后重啟才能重新使用cuda服務，畢竟cuda driver都沒了。
重裝cuda驅動可以解決。
cudaErrorInvalidConfiguration = 9，"invalid configuration argument"
運行時參數傳遞得太大了。比如：
1 subFunc<<<dim3(16, 16), dim3(64, 64)>>> (); 2 HANDLE_ERROR(cudaDeviceSynchronize()); 3 cudaError_t ct = cudaGetLastError(); 4 printf("%s\n", cudaGetErrorString(ct));
通過deviceQuery我們發現，每個block中只有1024個thread，而64 × 64 = 4096 > 1024：
deviceQuery查詢結果
出現這類問題后，cuda仍可繼續提供服務，僅單純拒絕了啟動核函數。
cudaErrorInvalidPitchValue = 12，"invalid pitch argument"
這個錯誤只會出現在cudaMemcpy2D、cudaMemcpy2DFromArray、cudaMemcpy2DToArray、cudaMemcpy3D及其異步形式函數的返回值中（當然也會被cudaGetLastError和cudaPeekAtLastError捕獲到）。
Pitch是通過cudaMallocPitch（申請二維數組）、cudaMalloc3D（申請三維數組）時產生的，用於數據對齊，加速尋址訪問速度。如下例：
1 double* p1; 2 size_t pitch1 = 0; 3 HANDLE_ERROR(cudaMallocPitch(&p1, &pitch1, 1280, 720)); 4 doSome<<<B, T>>> (p1, pitch1); 5
6 double* p2; 7 size_t pitch2 = 0; //未初始化或未改變其0值
8 HANDLE_ERROR(cudaMallocPitch(&p2, &pitch1, 1280, 720)); //【Bug】錯將pitch1的地址傳入
9 printf("%zu\n", pitch1); //1536 = 3×512
10 cudaError_t ct = cudaMemcpy2D(p2, pitch2, p1, pitch1, 1280, 720, cudaMemcpyDeviceToDevice); 11 printf("%s\n", cudaGetErrorString(ct)); //invalid pitch argument 12 //將【Bug】行處“&pitch1”改為“&pitch2”，程序將運行正常
由於cuda最早只支持C語言，因此保留了大量函數式編程的風格，沒有封裝高維數組便是其中一個例子，因此對於Pitch的使用需要開發者額外注意。
出現這類問題后，cuda仍可繼續提供服務，僅拒絕執行了當前被錯誤傳參的cudaMemcpy類函數的執行。
下述情況中並不會返回cudaErrorInvalidPitchValue：
1.錯將pitch的地址傳入pitch參數（返回cudaErrorInvalidValue = 1，CUDA C++直接報編譯錯誤）
1 double* dst_p; 2 size_t dst_pitch; 3 HANDLE_ERROR(cudaMallocPitch(&dst_p, &dst_pitch, 1280, 720)); 4 cudaError_t ct = cudaMemcpy2D(dst_p, &dst_pitch, p, &pitch, 1280, 720, cudaMemcpyDeviceToDevice); 5 printf("%s\n", cudaGetErrorString(ct));
cudaErrorInvalidSymbol = 13，"invalid device symbol"
這個錯誤只會出現在cudaGetSymbolAddress、cudaGetSymbolSize，以及cudaMemcpyFromSymbol、cudaMemcpyToSymbol及其異步形式函數的返回值中（當然也會被cudaGetLastError和cudaPeekAtLastError捕獲到）。
Symbol指定義在全局的__device__或__constant__修飾的顯存變量。cuda 4.1之前尚支持用變量名的字符串來表示Symbol的地址傳入上述API中，但cuda 4.1后廢棄了這一寫法，cuda 5.0后更是將其刪除。目前僅支持用顯存變量本身來指代Symbol，如：
1 __device__ size_t bounds[10]; 2 void SetBounds(size_t* host_bounds) { 3 HANDLE_ERROR(cudaMemcpyToSymbol(bounds, host_bounds, sizeof(size_t) * 10)); 4 }
如下述兩種情況代碼均將返回這類錯誤代碼：
1.使用cuda 5.0后廢棄的以字符串代替Symbol地址的寫法
1 __constant__ double MAXP[10]; 2
3 double maxp[] = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9}; 4 cudaError_t ct = cudaMemcpyToSymbol("MAXP", maxp, sizeof(double) * 10); 5 printf("%s\n", cudaGetErrorString(ct));
2. 錯誤使用了Symbol的地址而不是Symbol本身（盡管使用Symbol本身時，VS會顯示紅色波浪線，但並不影響編譯，使用Symbol本身才是正確寫法！）
1 __constant__ double r; 2
3 double host_r; 4 cudaError_t ct = cudaMemcpyFromSymbol(&host_r, &r, sizeof(double)); 5 //正確寫法：cudaMemcpyFromSymbol(&host_r, r, sizeof(double))
6 printf("%s\n", cudaGetErrorString(ct)); //invalid device symbol
3. 錯誤使用了非Symbol傳入API
1 double MAXP[10]; 2
3 double maxr[] = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9}; 4 cudaError_t ct = cudaMemcpyToSymbol(MAXP, maxr, sizeof(double) * 10); 5 printf("%s\n", cudaGetErrorString(ct)); //invalid device symbol
出現這類問題后，cuda仍可繼續提供服務，僅拒絕了當前報錯的與Symbol相關操作語句的執行。
cudaErrorInvalidMemcpyDirection = 21，"invalid copy direction for memcpy"
傳入API的cudaMemcpyKind類型的值有問題。cudaMemcpyKind是一個枚舉類型，目前只支持五個內存拷貝方向：主存到主存（相當於cstring的memcpy）、主存到顯存（I/O）、顯存到主存（I/O）、顯存到顯存（kernel中的memcpy）、默認（交給cuda自己判斷方向）。因此枚舉值位於0~4范圍內，如果超過這一范圍，則會報這個錯誤。
1 double *hp, *dp; 2 HANDLE_ERROR(cudaMallocHost(&hp, sizeof(double))); 3 HANDLE_ERROR(cudaMalloc(&dp, sizeof(double))); 4 cudaError_t ct = cudaMemcpy(hp, dp, sizeof(double), (cudaMemcpyKind)5); 5 printf("%s\n", cudaGetErrorString(ct)); //invalid copy direction for memcpy
如果正常編程通常不會出現這一問題，往往是一個cudaMemcpyKind類型變量從幾層函數之外層層傳進來的時候，中間某步被錯誤修改其值，出現了這個錯誤。
cudaErrorInsufficientDriver = 35，"CUDA driver version is insufficient for CUDA runtime version"
和cudaErrorCudartUnloading相似，這個錯誤也是因為驅動，也是只發生在第一次調用API的初始化階段。不過這個錯誤的產生原因是電腦上cuda驅動比cuda runtime庫版本要老。
你需要deviceQuery一下查詢cuda driver與runtime的版本，來決定下一步怎么做——要么安裝更老版本的cuda runtime，要么升級cuda driver，兩種方法都能解決這一問題。
cudaErrorDuplicateVariableName = 43，"duplicate global variable looked up by string name"
按照官方文檔，如果出現這個錯誤，說明你聲明了兩個及以上的相同符號名的設備端變量，即使是在不同文件中，如：
1 //1.cu
2 __device__ int m; 3 //2.cu
4 __constant__ double m;
事實上這個錯誤無法被觸發，甚至cudaGetLastError和cudaPeekAtLastError都不會返回這個問題了，盡管在截止至目前最新的文檔（v11.2.2）中也沒有將它廢棄掉。
以筆者的拙見，至少在cuda 5.0開始就應該將它廢棄甚至刪除掉了，因為在編譯和鏈接階段就足以將這個問題檢出，並且也不再使用符號名代替Symbol傳入API的寫法了。
cudaErrorDevicesUnavailable = 46，"all CUDA-capable devices are busy or unavailable"
可能你此前的某一步操作使得顯卡陷入忙碌，無法處理你的請求。但也有另一種可能，cuda安裝存在問題或有誤刪文件。
出現這類問題后，理論上cuda仍可繼續提供服務，但如果cuda安裝不正確，或存在文件誤刪情況，可能以后的API調用都會返回這一錯誤。
cudaErrorIncompatibleDriverContext = 49，"incompatible driver context"
學過操作系統的我們知道，一個計算機系統在處理用戶請求時，往往以上下文（context）作為進程執行的環境與單位——cuda driver也不例外。但如果你使用了廢棄的API創建了上下文，對於某個用戶請求，無法與驅動兼容，則有可能返回這一錯誤。
cudaErrorMissingConfiguration = 52，"__global__ function call is not configured"
很遺憾，有關這一問題的錯誤代碼我並沒有從任何網站中找到，也從未親身經理過。文檔中稱出錯原因是在調用cudaLaunchKernel()函數前沒有調用cudaConfigureCall()函數產生配置，但事實上所有的核函數調用都可以使用三連尖括號<<<>>>配置運行時參數。
拋磚引玉，等待賢士來補充。
cudaErrorLaunchMaxDepthExceeded = 65，"launch would exceed maximum depth of nested launches"
我嚴重懷疑官方文檔寫錯了！
官方文檔給出的出現這個錯誤的原因是核函數在較深的調用位置處使用了顯式同步：
官方文檔對這一錯誤的解釋，以及提供了修改限制的解決方案
事實上親自實驗之后你就會發現，其實這里完全就是在講cudaErrorSyncDepthExceeded而不是本錯誤cudaErrorLaunchMaxDepthExceeded。
真正有可能返回這一錯誤的API，文檔中只提供了一個cudaGetParameterBufferV2()。這是一個底層API，僅僅PTX可以訪問到，用於在核函數內部啟動核函數實現多級並行（用戶可以直接使用三連尖括號<<<>>>實現多級並行）。但事實上cudaGetLastError和cudaPeekAtLastError都不會捕獲到這個錯誤，所以我也不知道實踐中如何觸發這一異常。
出現這類問題后，理論上cuda仍可繼續提供服務，僅拒絕了報錯位置核函數的啟動。
cudaErrorSyncDepthExceeded = 68，"cudaDeviceSynchronize failed because caller's grid depth exceeds cudaLimitDevRuntimeSyncDepth"
錯誤信息說得很明白了，只有設備端代碼中的cudaDeviceSynchronize()函數會返回這一錯誤，原因是設置的可調用cudaDeviceSynchronize的核函數深度太淺了（默認值為2）。如下例，只允許第一層核函數調用同步：
1 __global__ void doSomething2() { 2 int a = 1 + 1; 3 } 4 __global__ void doSomething1() { 5 doSomething2<<<1, 2>>>(); 6 cudaError_t ct = cudaDeviceSynchronize(); 7 printf("%s\n", cudaGetErrorString(ct)); //cudaDeviceSynchronize failed because caller's grid depth exceeds cudaLimitDevRuntimeSyncDepth
8 } 9 __global__ void doSomething0() { 10 doSomething1<<<1, 2>>>(); 11 cudaError_t ct = cudaDeviceSynchronize(); 12 printf("%s\n", cudaGetErrorString(ct)); //no error
13 } 14
15 HANDLE_ERROR(cudaDeviceSetLimit(cudaLimitDevRuntimeSyncDepth, 1)); 16 doSomething0<<<1, 1>>>(); 17 cudaError_t ct = cudaDeviceSynchronize(); 18 printf("%s\n", cudaGetErrorString(ct)); //no error
19 ct = cudaGetLastError(); 20 printf("%s\n", cudaGetErrorString(ct)); //no error
因為太深的核函數需要消耗大量通信資源才能實現同步，所以設置這個限度對實際編程和性能優化還是很有幫助的。
出現這類問題后，cuda仍可繼續提供服務，僅拒絕了在核函數內實現顯式同步。
注意，如果真的需要在核函數內與其他一同啟動的核函數實現同步，可以考慮使用設備端的__syncthreads()函數。
cudaErrorLaunchPendingCountExceeded = 69，"launch failed because launch would exceed cudaLimitDevRuntimePendingLaunchCount"
只有cudaGetLastError和cudaPeekAtLastError會捕獲到這一錯誤（貌似只有設備端代碼內才有可能觸發）。原因是核函數無法啟動，因為超過了設備端等待啟動的核函數個數的最大限制。
出現這個問題，通過cudaDeviceSetLimit調高cudaLimitDevRuntimePendingLaunchCount限制是一個治標不治本的做法。最推薦開發者修改代碼邏輯，優化多級並行、提升核函數利用效率，巧妙運用同步等。
出現這類問題后，cuda仍可繼續提供服務，僅拒絕了啟動報錯位置的核函數。
cudaErrorInvalidDeviceFunction = 98，"invalid device function"
一般出現在動態修改函數緩存配置（如cudaFuncSetCacheConfig）、獲取函數屬性（如cudaFuncGetAttributes）、設置函數屬性（如cudaFuncSetAttribute）、啟動核函數（如cudaLaunchKernel）等對設備端函數的操作中傳參出現了問題。常見的情況有以下三種：
1.錯將設備端變量符號或聲明的變量地址傳入API
1 double* p = &output; 2 cudaFuncAttributes attributes; 3 cudaError_t ct = cudaFuncGetAttributes(&attributes, p); 4 printf("%s\n", cudaGetErrorString(ct)); //invalid device function
2. 需要傳__global__函數符號的API被錯傳__device__函數的符號
1 __device__ void func() { 2 int a = 1 + 1; 3 } 4
5 cudaError_t ct = cudaFuncSetAttribute(func, cudaFuncAttributeMaxDynamicSharedMemorySize, 1024); 6 printf("%s\n", cudaGetErrorString(ct)); //invalid device function
3. architecture太高，請降低sm和compute等級（如sm_70→sm_61，compute_70→compute_61）
VS中修改CUDA C/C++ -&amp;gt; Device -&amp;gt; Code Generation改變architecture
出現這類問題后，cuda仍可繼續提供服務，僅拒絕了對設備端函數的操作。
cudaErrorNoDevice = 100，"no CUDA-capable device is detected"
你需要反思一下自己有沒有NVIDIA的顯卡，買了之后有沒有安裝在自己的電腦上，是不是真的支持CUDA，有沒有老得過氣。
評論區禁止AMD yes。
cudaErrorInvalidDevice = 101，"invalid device ordinal"
如果你報了這個錯誤，而且你沒有高估自己顯卡數量的話，我敢保證你的deviceQuery就沒跑通。
這個錯誤通常出現在cudaDeviceGetAttribute、cudaSetDevice等需要傳遞device編號的API中返回，同樣也會被cudaGetLastError和cudaPeekAtLastError捕獲到。請注意，這個編號和數組下標一樣，是從0開始的。常見的情況往往是高估了自己電腦安裝過的顯卡數量：
1 cudaError_t ct = cudaSetDevice(2); 2 printf("%s\n", cudaGetErrorString(ct)); //invalid device ordinal
出現這類問題后，cuda仍可繼續提供服務，僅拒絕了顯式切換顯卡、獲取設備屬性等操作。
cudaErrorFileNotFound = 301，"file not found"
關於cuda核函數內如何進行文件操作我至今也不知道答案，希望有賢士為我補充。
cudaErrorOperatingSystem = 304，"OS call failed or operation not supported on this OS"
我懷疑問題出在你使用了集成顯卡。建議直接使用英偉達卡，禁用集顯，防止交互時出現問題。
cudaErrorSymbolNotFound = 500，"named symbol not found"
又是一個應該被廢棄的錯誤碼。cuda 5.0開始就已經廢棄根據符號名來確定顯存符號的操作了，所以這個錯誤正常來講不會被觸發。
cudaErrorIllegalAddress = 700，"an illegal memory access was encountered"
通常是數組訪問越界造成的，包括但不限於：
1.共享內存越界
1 __global__ void Kernel(double* ptr) { 2 __shared__ int s[16]; 3
4 unsigned int Idx = blockIdx.x * blockDim.x + threadIdx.x; 5 if(Idx >= 1024) return; 6
7 s[Idx] = ptr[Idx % 16]; 8 } 9
10 double *device_p; 11 HANDLE_ERROR(cudaMalloc(&device_p, sizeof(double) * 16)); 12 Kernel<<<64, 16>>>(device_p); 13 cudaError_t ct = cudaDeviceSynchronize(); 14 printf("%s\n", cudaGetErrorString(ct)); //an illegal memory access was encountered
2. 全局或常量內存越界
1 __constant__ int limits[16]; 2
3 __global__ void Kernel(double* ptr) { 4 unsigned int Idx = blockIdx.x * blockDim.x + threadIdx.x; 5 if(Idx >= 256) return; 6
7 double sum = 0; 8 for(unsigned int i = 0; i < limits[Idx]; ++i) { 9 sum += ptr[i]; 10 } 11 } 12
13 double *device_p; 14 HANDLE_ERROR(cudaMalloc(&device_p, sizeof(double) * 16)); 15 Kernel<<<16, 16>>>(device_p); 16 cudaError_t ct = cudaDeviceSynchronize(); 17 printf("%s\n", cudaGetErrorString(ct)); //an illegal memory access was encountered
3. 在核函數內使用了未初始化的野指針或空指針（在__host__函數內使用未初始化的指針將會報cudaErrorInvalidValue錯誤）
1 __global__ void Kernel(double* ptr) { 2 unsigned int Idx = blockIdx.x * blockDim.x + threadIdx.x; 3 if(Idx >= 256) return; 4
5 double* tmp = nullptr; 6
7 tmp[Idx] = ptr[Idx]; 8 } 9
10 double *device_p; 11 HANDLE_ERROR(cudaMalloc(&device_p, sizeof(double) * 16)); 12 Kernel<<<16, 16>>>(device_p); 13 cudaError_t ct = cudaDeviceSynchronize(); 14 printf("%s\n", cudaGetErrorString(ct)); //an illegal memory access was encountered
出現這類問題后，程序必須終止后重啟才能重新使用cuda服務。
下面這幾種情況並不會返回cudaErrorIllegalAddress：
1.cudaMemcpy時使用大於dst或src指針所申請內存大小的count參數（返回cudaErrorInvalidValue = 1）
1 double *host_p, *device_p; 2 HANDLE_ERROR(cudaMallocHost(&host_p, sizeof(double) * 32)); 3 HANDLE_ERROR(cudaMalloc(&device_p, sizeof(double) * 16)); 4 cudaError_t ct = cudaMemcpy(device_p, host_p, sizeof(double) * 32, cudaMemcpyHostToDevice); 5 printf("%s\n", cudaGetErrorString(ct)); //"invalid argument"
2. 使用野指針、空指針、已經刪除的指針傳入API（返回cudaErrorInvalidValue = 1）
1 double *device_p; 2 HANDLE_ERROR(cudaMalloc(&device_p, sizeof(double) * 32)); 3 HANDLE_ERROR(cudaFree(device_p)); 4 cudaError_t ct = cudaMemset(device_p, 0, sizeof(double) * 32); 5 printf("%s\n", cudaGetErrorString(ct)); //"invalid argument"
cudaErrorLaunchOutOfResources = 701，"too many resources requested for launch"
字面意思是啟動核函數時請求的資源太多，超過了閑置的資源，導致核函數無法啟動。
90%以上的博客認為這個異常只和寄存器有關。確實，開啟--ptxas-options=-v我們可以查看一個核函數使用了多少資源。
開啟ptxas輸出選項
通過deviceQuery，可以查看常量內存（cmem）、共享內存（smem）和寄存器數（register）的大小：
常量內存、共享內存、每核寄存器數的大小
下面的代碼運行時會報cudaErrorLaunchOutOfResources錯誤：
1 __global__ void Kernel(double* ptr1, double* ptr2, double* ptr3, double* ptr4, 2 double* ptr5, double* ptr6, double* ptr7, double* ptr8, 3 double* ptr9, double* ptr10, double* ptr11, double* ptr12, 4 double* ptr13, double* ptr14, double* ptr15, double* ptr16, 5 double* ptr17, double* ptr18, double* ptr19, double* ptr20, 6 double* ptr21, double* ptr22, double* ptr23, double* ptr24, 7 double* ptr25, double* ptr26, double* ptr27, double* ptr28, 8 double* ptr29, double* ptr30, double* ptr31, double* ptr32, 9 size_t tot_thrd) { 10 unsigned int Idx = blockIdx.x * blockDim.x + threadIdx.x; 11 if(Idx >= tot_thrd) return; 12 } 13
14 double *device_p; 15 HANDLE_ERROR(cudaMalloc(&device_p, sizeof(double) * 16)); 16 Kernel<<<16, 1024>>>( 17 device_p, device_p, device_p, device_p, device_p, device_p, device_p, device_p, 18 device_p, device_p, device_p, device_p, device_p, device_p, device_p, device_p, 19 device_p, device_p, device_p, device_p, device_p, device_p, device_p, device_p, 20 device_p, device_p, device_p, device_p, device_p, device_p, device_p, device_p, 21 1024*16); 22 HANDLE_ERROR(cudaDeviceSynchronize()); 23 cudaError_t ct = cudaGetLastError(); 24 printf("%s\n", cudaGetErrorString(ct)); //"too many resources requested for launch"
此時ptxas的輸出如下：
1 1>ptxas info : 64448 bytes gmem, 72 bytes cmem[3] 2 1>ptxas info : Compiling entry function '_Z6KernelPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_y' for 'sm_61'
3 1>ptxas info : Function properties for _Z6KernelPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_S_y 4 1> 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 5 1>ptxas info : Used 70 registers, 584 bytes cmem[0]
這個一長串的函數是什么？我們可以用c++filt工具來查看函數原型：
說明ptxas輸出分析的函數正是核函數Kernel。而每個線程內Kernel占用了70個寄存器，
，因此報出cudaErrorLaunchOutOfResources錯誤。
這時，我們可以優化一下核函數傳參、減少每個block中的thread（如<<<16, 1024>>>改為<<<64, 256>>>）、使用--maxrregcount限制一個線程內核函數最多使用的寄存器數等：
通過--maxrregcount限制每個核函數最多使用的寄存器數
事實上，除此之外還有一些可能，如：
1.使用了過大的設備限制，如堆內存限制過大（每錯，不會返回cudaErrorInvalidValue的情況，雖然會返回cudaSuccess，但會帶來副作用）
1 template <typename T>
2 __global__ void subConstructArray(T* dsts, size_t Len) { 3 size_t Idx = blockIdx.x * blockDim.x + threadIdx.x; 4 if (Idx >= Len) return; 5
6 new (dsts + Idx)T; 7 } 8
9 template <typename T>
10 __global__ void subDestructArray(T* dsts, size_t Len) { 11 size_t Idx = blockIdx.x * blockDim.x + threadIdx.x; 12 if (Idx >= Len) return; 13
14 (dsts + Idx)->~T(); 15 } 16
17 template <typename T>
18 void ConstructArray(T*& dsts, size_t Len) { 19 HANDLE_ERROR(cudaMalloc(&dsts, sizeof(T) * Len)); 20 subConstructArray<<<1, Len>>>
21 (dsts, Len); 22 cudaDeviceSynchronize(); 23 HANDLE_ERROR(cudaGetLastError()); 24 } 25
26 template <typename T>
27 void DestructArray(T* dsts, size_t Len) { 28 subDestructArray<<<1, Len>>>
29 (dsts, Len); 30 cudaDeviceSynchronize(); 31
32 cudaError_t ct = cudaGetLastError(); 33 printf("%s\n", cudaGetErrorString(ct)); //too many resources requested for launch
34 HANDLE_ERROR(cudaFree(dsts)); 35 } 36
37 template<typename T>
38 class tensor { 39 private: 40 T* _elems; 41 int _Ply; 42 int _Height; 43 int _Width; 44 public: 45 __device__ tensor() : _elems(nullptr), _Ply(0), _Height(0), _Width(0) {} 46 __device__ ~tensor() { free(_elems); } 47
48 __device__ void Set_size(int _ply, int _hgt, int _wid) { 49 if (_ply * _hgt * _wid != _Ply * _Height * _Width) { 50 free(_elems); 51 _elems = (T*)malloc(sizeof(T) * _ply * _hgt * _wid); 52 } 53 _Ply = _ply; 54 _Height = _hgt; 55 _Width = _wid; 56 } 57 }; 58
59 size_t Tot_Thrd = 256; 60 tensor<double>* out; 61 tensor<double>* er; 62
63 void Set_Tot_Thrd(size_t _num_of_thrds) { 64 DestructArray(out, Tot_Thrd); 65 DestructArray(er, Tot_Thrd); 66 ConstructArray(out, _num_of_thrds); 67 ConstructArray(er, _num_of_thrds); 68 Tot_Thrd = _num_of_thrds; 69 } 70
71 HANDLE_ERROR(cudaDeviceSetLimit(cudaLimitMallocHeapSize, 2000 * 1024 * 1024)); 72 ConstructArray(out, Tot_Thrd); 73 ConstructArray(er, Tot_Thrd); 74 Set_Tot_Thrd(16);
其原理是占用了太多的顯存空間，導致核函數無法申請足夠的顯存來啟動。
2. 共享內存使用過多
每個核內使用了太多共享內存，和使用了太多寄存器一樣，在此不作贅述。
據說核函數內使用printf也容易報這個錯誤，因此盡量用其它替代方法。
出現這類問題后，cuda仍可繼續提供服務，僅拒絕了啟動報錯位置的核函數。
下面這幾種情況並不會返回cudaErrorLaunchOutOfResources：
1.動態申請過大的內存（返回cudaErrorMemoryAllocation = 2）
前文已經討論過了，在此不再贅述了。注意報錯信息的資源並不包括堆內存資源。
2. 過深的遞歸或過多的遞歸函數內變量（棧溢出，返回cudaErrorLaunchFailure = 719）
后文將要討論，不再贅述。注意報錯信息的資源並不包括棧內存資源。
cudaErrorAssert = 710，"device-side assert triggered"
設備端代碼中的斷言失敗。請檢查__global__或__device__函數內報錯的assert語句是否真的成立。
出現這類問題后，程序必須終止后重啟才能重新使用cuda服務。
cudaErrorLaunchFailure = 719，"unspecified launch failure"
最常見的是棧溢出造成的，而我也沒見過其它情況導致這一問題。
棧溢出可能是遞歸層數太深、甚至無限遞歸：
1 __global__ void fxxk_stack() { 2 int buff[1024]; 3 fxxk_stack<<<16, 256>>>(); 4 } 5
6 fxxk_stack<<<16, 256>>>(); 7 cudaError_t ct = cudaDeviceSynchronize(); 8 printf("%s\n", cudaGetErrorString(ct)); //"unspecified launch failure"
也有可能是開了太多的局域變量，導致較淺的遞歸也爆棧：
1 __global__ void fxxk_stack() { 2 int buff[1024000]; 3 } 4
5 fxxk_stack<<<16, 256>>>(); 6 cudaError_t ct = cudaDeviceSynchronize(); 7 printf("%s\n", cudaGetErrorString(ct)); //"unspecified launch failure"
出現這類問題后，程序必須終止后重啟才能重新使用cuda服務，畢竟顯卡棧並沒有那么魯棒。
有關cudaError的討論就是這些了，如果開發者朋友們遇到了其它問題，或文章中提到的錯誤碼但並不屬於提到的觸發方法之一，歡迎在評論區或私信中給出！