CUDA教程四、异常处理与编程技巧

也许有人注意到了，我在【CUDA教程】二、主存与显存文章中提到了部分常见的异常。实际上，cuda编程最终Boss则是debug。本文将重点讲解cuda中错误的成因，作为“报错词典”供各位开发者们debug。

本文将尽可能全面地列举所有异常的可能出现情况，如需快速找到问题原因，请使用ctrl+F的页内搜索功能，检索内容为cudaError_t枚举类型的成员名（如"cudaErrorLaunchOutOfResources"）、错误代码（如"701"）和出错信息（如"too many resources requested for launch"）。

---

调试技巧

cuda的样例代码使用了如下方法处理异常：

cudaError_t cudaStatus;

// Choose which GPU to run on, change this on a multi-GPU system.
cudaStatus = cudaSetDevice(0);
if (cudaStatus != cudaSuccess) {
    fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
    goto Error;
}

事实上，这种做法太过冗长，每次使用cuda函数都要写六行来处理异常，一个函数的代码量甚至会提升2~3倍；同时代码的连贯性也被打破，后期维护起来非常困难。因此我在前面代码量不大的代码中使用了这种方法：

assert(cudaSuccess == cudaSetDevice(0));

这种方法解决了上述所有问题，但引入了新的问题，如果发生Assert Failed，我们需要获得异常的返回值来进行处理，因此需要修改代码。但如果发生了无法重现（Unable to Reproduce）的错误，在多次后续实验中难以再次触发，白白浪费了潜在的漏洞修复时机。作为折衷，我们可以定义如下的宏和函数：

#define HANDLE_ERROR(err) (HandleError(err, __FILE__, __LINE__))

static void HandleError(cudaError_t err, const char *file, int line) {
    if (err != cudaSuccess) {
        fprintf(stderr, "Error %d: \"%s\" in %s at line %d\n", int(err), cudaGetErrorString(err), file, line);
        exit(int(err));
    }
}

实际使用时则只需要：

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.使用空指针或野指针

double** pp;
cudaError_t ct = cudaMalloc(pp, sizeof(double) * 1024);        //野指针传入API作为参数
printf("%s\n", cudaGetErrorString(ct));                //"invalid argument"

2. 搞混host指针与device指针

double* p;
HANDLE_ERROR(cudaMallocHost(&p, sizeof(double) * 1024));    //申请host端内存的函数
cudaError_t ct = cudaFree(p);                    //释放device端内存的函数
printf("%s\n", cudaGetErrorString(ct));                //"invalid argument"

3. 使用已经释放的指针

double* p;
HANDLE_ERROR(cudaMalloc(&p, sizeof(double) * 1024));
HANDLE_ERROR(cudaFree(p));                    //释放已经申请的device端内存
cudaError_t ct = cudaFree(p);                    //再次释放或修改对应值（如cudaMemset等）
printf("%s\n", cudaGetErrorString(ct));                //"invalid argument"

4. 使用错误的枚举值

int limit_type = 16;                        //不存在的cudaLimit类型枚举值
cudaError_t ct = cudaDeviceSetLimit((cudaLimit)limit_type, 100 * 1024);
printf("%s\n", cudaGetErrorString(ct));                //"invalid argument"

出现这类问题后，cuda仍可继续提供服务。

但是下面这几种情况并不会返回cudaErrorInvalidValue，甚至会返回cudaSuccess：

1.过大的内存申请量（返回cudaErrorMemoryAllocation = 2）

double* p;                            //显存4GB，申请8TB试试
cudaError_t ct = cudaMalloc(&p, sizeof(double) * 1024ull * 1024ull * 1024ull * 1024ull);
printf("%s\n", cudaGetErrorString(ct));                //"out of memory"

2. 过大的空间上限（返回cudaSuccess = 0）

cudaError_t ct = cudaDeviceSetLimit(cudaLimitMallocHeapSize, 1024ull * 1024ull * 1024ull * 1024ull);
printf("%s\n", cudaGetErrorString(ct));                //"no error"

cudaErrorMemoryAllocation = 2，"out of memory"

顾名思义，内存爆掉了。最常见的就是调用cudaMalloc时超过了堆内存的最大限制，当然cudaHostAlloc、cudaMallocHost等都会返回这一问题，核函数内的malloc也有产生这一问题的可能性。此时要使用cudaDeviceSetLimit增大堆内存上限，或检查是否发生了内存泄漏，并及时Free掉多余的内存。

出现这类问题后，cuda仍可继续提供服务，仅拒绝分配给用户所申请的内存空间而已。如下例：

HANDLE_ERROR(cudaDeviceSetLimit(cudaLimitMallocHeapSize, 1024 * 1024 * 1024));
int *p, *q;
cudaError_t ct = cudaMalloc(&p, sizeof(int) * 200 * 1024 * 1024);
printf("%s\n", cudaGetErrorString(ct));                //"no error"
ct = cudaMalloc(&q, sizeof(int) * 200 * 1024 * 1024);
printf("%s\n", cudaGetErrorString(ct));                //"out of memory"
HANDLE_ERROR(cudaFree(p));
ct = cudaMalloc(&q, sizeof(int) * 200 * 1024 * 1024);
printf("%s\n", cudaGetErrorString(ct));                //"no error"

但是下面这几种情况并不会返回cudaErrorMemoryAllocation：

1.栈溢出（返回cudaErrorLaunchFailure = 719）

__device__ int func(int n) {
    if(n <= 1) return 0;
    double fxxker[256];                    //较深的递归与较多的局域变量使栈溢出
    return func(n - 1) + func(n - 2) + 1;
}
__global__ void fxxk_stack() {
    printf("%d\n", func(32));
}

//main
fxxk_stack<<<1, 1>>>();
cudaError_t ct = cudaDeviceSynchronize();
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"

运行时参数传递得太大了。比如：

subFunc<<<dim3(16, 16), dim3(64, 64)>>> ();
HANDLE_ERROR(cudaDeviceSynchronize());
cudaError_t ct = cudaGetLastError();
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（申请三维数组）时产生的，用于数据对齐，加速寻址访问速度。如下例：

double* p1;
size_t pitch1 = 0;
HANDLE_ERROR(cudaMallocPitch(&p1, &pitch1, 1280, 720));
doSome<<<B, T>>> (p1, pitch1);

double* p2;
size_t pitch2 = 0;                        //未初始化或未改变其0值
HANDLE_ERROR(cudaMallocPitch(&p2, &pitch1, 1280, 720));        //【Bug】错将pitch1的地址传入
printf("%zu\n", pitch1);                    //1536 = 3×512
cudaError_t ct = cudaMemcpy2D(p2, pitch2, p1, pitch1, 1280, 720, cudaMemcpyDeviceToDevice);
printf("%s\n", cudaGetErrorString(ct));                //invalid pitch argument
//将【Bug】行处“&pitch1”改为“&pitch2”，程序将运行正常

由于cuda最早只支持C语言，因此保留了大量函数式编程的风格，没有封装高维数组便是其中一个例子，因此对于Pitch的使用需要开发者额外注意。

出现这类问题后，cuda仍可继续提供服务，仅拒绝执行了当前被错误传参的cudaMemcpy类函数的执行。

下述情况中并不会返回cudaErrorInvalidPitchValue：

1.错将pitch的地址传入pitch参数（返回cudaErrorInvalidValue = 1，CUDA C++直接报编译错误）

double* dst_p;
size_t dst_pitch;
HANDLE_ERROR(cudaMallocPitch(&dst_p, &dst_pitch, 1280, 720));
cudaError_t ct = cudaMemcpy2D(dst_p, &dst_pitch, p, &pitch, 1280, 720, cudaMemcpyDeviceToDevice);
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，如：

__device__ size_t bounds[10];
void SetBounds(size_t* host_bounds) {
    HANDLE_ERROR(cudaMemcpyToSymbol(bounds, host_bounds, sizeof(size_t) * 10));
}

如下述两种情况代码均将返回这类错误代码：

1.使用cuda 5.0后废弃的以字符串代替Symbol地址的写法

__constant__ double MAXP[10];

double maxp[] = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9};
cudaError_t ct = cudaMemcpyToSymbol("MAXP", maxp, sizeof(double) * 10);
printf("%s\n", cudaGetErrorString(ct));    

2. 错误使用了Symbol的地址而不是Symbol本身（尽管使用Symbol本身时，VS会显示红色波浪线，但并不影响编译，使用Symbol本身才是正确写法！）

__constant__ double r;

double host_r;
cudaError_t ct = cudaMemcpyFromSymbol(&host_r, &r, sizeof(double));
//正确写法：cudaMemcpyFromSymbol(&host_r, r, sizeof(double))
printf("%s\n", cudaGetErrorString(ct));                //invalid device symbol

3. 错误使用了非Symbol传入API

double MAXP[10];

double maxr[] = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9};
cudaError_t ct = cudaMemcpyToSymbol(MAXP, maxr, sizeof(double) * 10);
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范围内，如果超过这一范围，则会报这个错误。

double *hp, *dp;
HANDLE_ERROR(cudaMallocHost(&hp, sizeof(double)));
HANDLE_ERROR(cudaMalloc(&dp, sizeof(double)));
cudaError_t ct = cudaMemcpy(hp, dp, sizeof(double), (cudaMemcpyKind)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.cu
__device__ int m;
//2.cu
__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）。如下例，只允许第一层核函数调用同步：

__global__ void doSomething2() {
    int a = 1 + 1;
}
__global__ void doSomething1() {
    doSomething2<<<1, 2>>>();
    cudaError_t ct = cudaDeviceSynchronize();
    printf("%s\n", cudaGetErrorString(ct));        //cudaDeviceSynchronize failed because caller's grid depth exceeds cudaLimitDevRuntimeSyncDepth
}
__global__ void doSomething0() {
    doSomething1<<<1, 2>>>();
    cudaError_t ct = cudaDeviceSynchronize();
    printf("%s\n", cudaGetErrorString(ct));        //no error
}

HANDLE_ERROR(cudaDeviceSetLimit(cudaLimitDevRuntimeSyncDepth, 1));
doSomething0<<<1, 1>>>();
cudaError_t ct = cudaDeviceSynchronize();
printf("%s\n", cudaGetErrorString(ct));            //no error
ct = cudaGetLastError();
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

double* p = &output;
cudaFuncAttributes attributes;
cudaError_t ct = cudaFuncGetAttributes(&attributes, p);
printf("%s\n", cudaGetErrorString(ct));            //invalid device function

2. 需要传__global__函数符号的API被错传__device__函数的符号

__device__ void func() {
    int a = 1 + 1;
}

cudaError_t ct = cudaFuncSetAttribute(func, cudaFuncAttributeMaxDynamicSharedMemorySize, 1024);
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;amp;gt; Device -&amp;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开始的。常见的情况往往是高估了自己电脑安装过的显卡数量：

cudaError_t ct = cudaSetDevice(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.共享内存越界

__global__ void Kernel(double* ptr) {
    __shared__ int s[16];

    unsigned int Idx = blockIdx.x * blockDim.x + threadIdx.x;
    if(Idx >= 1024) return;

    s[Idx] = ptr[Idx % 16];
}

double *device_p;
HANDLE_ERROR(cudaMalloc(&device_p, sizeof(double) * 16));
Kernel<<<64, 16>>>(device_p);
cudaError_t ct = cudaDeviceSynchronize();
printf("%s\n", cudaGetErrorString(ct));            //an illegal memory access was encountered

2. 全局或常量内存越界

__constant__ int limits[16];

__global__ void Kernel(double* ptr) {
    unsigned int Idx = blockIdx.x * blockDim.x + threadIdx.x;
    if(Idx >= 256) return;

    double sum = 0;
    for(unsigned int i = 0; i < limits[Idx]; ++i) {
        sum += ptr[i];
    }
}

double *device_p;
HANDLE_ERROR(cudaMalloc(&device_p, sizeof(double) * 16));
Kernel<<<16, 16>>>(device_p);
cudaError_t ct = cudaDeviceSynchronize();
printf("%s\n", cudaGetErrorString(ct));            //an illegal memory access was encountered

3. 在核函数内使用了未初始化的野指针或空指针（在__host__函数内使用未初始化的指针将会报cudaErrorInvalidValue错误）

__global__ void Kernel(double* ptr) {
    unsigned int Idx = blockIdx.x * blockDim.x + threadIdx.x;
    if(Idx >= 256) return;

    double* tmp = nullptr;

    tmp[Idx] = ptr[Idx];
}

double *device_p;
HANDLE_ERROR(cudaMalloc(&device_p, sizeof(double) * 16));
Kernel<<<16, 16>>>(device_p);
cudaError_t ct = cudaDeviceSynchronize();
printf("%s\n", cudaGetErrorString(ct));            //an illegal memory access was encountered

出现这类问题后，程序必须终止后重启才能重新使用cuda服务。

下面这几种情况并不会返回cudaErrorIllegalAddress：

1.cudaMemcpy时使用大于dst或src指针所申请内存大小的count参数（返回cudaErrorInvalidValue = 1）

double *host_p, *device_p;
HANDLE_ERROR(cudaMallocHost(&host_p, sizeof(double) * 32));
HANDLE_ERROR(cudaMalloc(&device_p, sizeof(double) * 16));
cudaError_t ct = cudaMemcpy(device_p, host_p, sizeof(double) * 32, cudaMemcpyHostToDevice);
printf("%s\n", cudaGetErrorString(ct));            //"invalid argument"

2. 使用野指针、空指针、已经删除的指针传入API（返回cudaErrorInvalidValue = 1）

double *device_p;
HANDLE_ERROR(cudaMalloc(&device_p, sizeof(double) * 32));
HANDLE_ERROR(cudaFree(device_p));
cudaError_t ct = cudaMemset(device_p, 0, sizeof(double) * 32);
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错误：

__global__ void Kernel(double* ptr1, double* ptr2, double* ptr3, double* ptr4,
                       double* ptr5, double* ptr6, double* ptr7, double* ptr8,
                       double* ptr9, double* ptr10, double* ptr11, double* ptr12,
                       double* ptr13, double* ptr14, double* ptr15, double* ptr16,
                       double* ptr17, double* ptr18, double* ptr19, double* ptr20,
                       double* ptr21, double* ptr22, double* ptr23, double* ptr24,
                       double* ptr25, double* ptr26, double* ptr27, double* ptr28,
                       double* ptr29, double* ptr30, double* ptr31, double* ptr32,
                       size_t tot_thrd) {
    unsigned int Idx = blockIdx.x * blockDim.x + threadIdx.x;
    if(Idx >= tot_thrd) return;
}

double *device_p;
HANDLE_ERROR(cudaMalloc(&device_p, sizeof(double) * 16));
Kernel<<<16, 1024>>>(
    device_p, device_p, device_p, device_p, device_p, device_p, device_p, device_p,
    device_p, device_p, device_p, device_p, device_p, device_p, device_p, device_p,
    device_p, device_p, device_p, device_p, device_p, device_p, device_p, device_p,
    device_p, device_p, device_p, device_p, device_p, device_p, device_p, device_p,
    1024*16);
HANDLE_ERROR(cudaDeviceSynchronize());
cudaError_t ct = cudaGetLastError();
printf("%s\n", cudaGetErrorString(ct));            //"too many resources requested for launch"

此时ptxas的输出如下：

1>ptxas info    : 64448 bytes gmem, 72 bytes cmem[3]
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'
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
1>    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
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，但会带来副作用）

template <typename T>
__global__ void subConstructArray(T* dsts, size_t Len) {
    size_t Idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (Idx >= Len) return;

    new (dsts + Idx)T;
}

template <typename T>
__global__ void subDestructArray(T* dsts, size_t Len) {
    size_t Idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (Idx >= Len) return;

    (dsts + Idx)->~T();
}

template <typename T>
void ConstructArray(T*& dsts, size_t Len) {
    HANDLE_ERROR(cudaMalloc(&dsts, sizeof(T) * Len));
    subConstructArray<<<1, Len>>>
        (dsts, Len);
    cudaDeviceSynchronize();
    HANDLE_ERROR(cudaGetLastError());
}

template <typename T>
void DestructArray(T* dsts, size_t Len) {
    subDestructArray<<<1, Len>>>
        (dsts, Len);
    cudaDeviceSynchronize();
    
    cudaError_t ct = cudaGetLastError();
    printf("%s\n", cudaGetErrorString(ct));            //too many resources requested for launch
    HANDLE_ERROR(cudaFree(dsts));
}

template<typename T>
class tensor {
    private:
        T* _elems;
        int _Ply;
        int _Height;
        int _Width;
    public:
        __device__ tensor() : _elems(nullptr), _Ply(0), _Height(0), _Width(0) {}
        __device__ ~tensor() { free(_elems); }

        __device__ void Set_size(int _ply, int _hgt, int _wid) {
            if (_ply * _hgt * _wid != _Ply * _Height * _Width) {
                free(_elems);
                _elems = (T*)malloc(sizeof(T) * _ply * _hgt * _wid);
            }
            _Ply = _ply;
            _Height = _hgt;
            _Width = _wid;
        }
};

size_t Tot_Thrd = 256;
tensor<double>* out;
tensor<double>* er;

void Set_Tot_Thrd(size_t _num_of_thrds) {
    DestructArray(out, Tot_Thrd);
    DestructArray(er, Tot_Thrd);
    ConstructArray(out, _num_of_thrds);
    ConstructArray(er, _num_of_thrds);
    Tot_Thrd = _num_of_thrds;
}

HANDLE_ERROR(cudaDeviceSetLimit(cudaLimitMallocHeapSize, 2000 * 1024 * 1024));
ConstructArray(out, Tot_Thrd);
ConstructArray(er, Tot_Thrd);
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"

最常见的是栈溢出造成的，而我也没见过其它情况导致这一问题。

栈溢出可能是递归层数太深、甚至无限递归：

__global__ void fxxk_stack() {
    int buff[1024];
    fxxk_stack<<<16, 256>>>();
}

fxxk_stack<<<16, 256>>>();
cudaError_t ct = cudaDeviceSynchronize();
printf("%s\n", cudaGetErrorString(ct));                //"unspecified launch failure"

也有可能是开了太多的局域变量，导致较浅的递归也爆栈：

__global__ void fxxk_stack() {
    int buff[1024000];
}

fxxk_stack<<<16, 256>>>();
cudaError_t ct = cudaDeviceSynchronize();
printf("%s\n", cudaGetErrorString(ct));                //"unspecified launch failure"

出现这类问题后，程序必须终止后重启才能重新使用cuda服务，毕竟显卡栈并没有那么鲁棒。

---

有关cudaError的讨论就是这些了，如果开发者朋友们遇到了其它问题，或文章中提到的错误码但并不属于提到的触发方法之一，欢迎在评论区或私信中给出！