CUDA code=700(cudaErrorIllegalAddress) 报错与排查方法

最近笔者在调试自己写的 CUDA 代码时, 遇到了 code=700(cudaErrorIllegalAddress) 的报错, 在此记录一下排查和解决方法.

报错

报错是由 CUDA API 函数执行时产生的, 由 checkCudaErrors() 函数检测出(CUDA 常用错误检测实现, 如下所示).

template <typename T>
void check(T result, char const *const func, const char *const file,int const line) {if (result) {fprintf(stderr, "CUDA error at %s:%d code=%d(%s) \"%s\" \n", file, line,static_cast<unsigned int>(result), cudaGetErrorName(result), func);exit(EXIT_FAILURE);}
}
#define checkCudaErrors(val) check((val), #val, __FILE__, __LINE__)

代码运行时报错如下所示, 显示是执行 cudaMemGetInfo() 函数时错误.

huanghy@node8:~/CL/src/cuda/build$ ./example 
[sample_cuda] start
[sample_kernel] grid_size:1, block_size:512, shm_size:6144
[sample_kernel] finished
CUDA error at /home/huanghy/CL/src/cuda/sample.cu:53 code=700(cudaErrorIllegalAddress) "cudaMemGetInfo(&freeMem, &totalMem)" 

原因

简单查阅资料可知, code=700(cudaErrorIllegalAddress) 的报错原因是 “an illegal memory access was encountered”, 即"遇到了一个非法的内存访问".

大多数情况下, 该问题产生都与数组越界访问的情况有关, 但值得一提的是, 往往报错的地方并不是问题实际存在的地方, 而由之前的 kernel 代码中的错误访问导致的.
比如, 此处报错是在 API 函数cudaMemGetInfo() 执行时, 也有可能是在自己定义的 kernel 函数执行时, 但可能一直排查当前报错的 kernel 不能解决问题的.

排查

一个很好的排查上述问题, 也是对自己的 CUDA 代码进行内存访问检查的方法是使用 CUDA 的 compute-sanitizer 工具.
该工具功能很多, 其中一个功能就是进行内存检测.

使用如下指令进行内存检查:

compute-sanitizer --launch-timeout=0 --tool=memcheck ./example > opt.txt 2>&1

其中, ./example 为检测的可执行文件. 由于输出可能比较多, 所以这里重定向到文件中. --launch-timeout=0 是将等待 kernel 加载的时间设置为无限, 以避免 compute-sanitizer 出现终止的情况, 如下所示.

========= COMPUTE-SANITIZER
========= Error: No attachable process found. compute-sanitizer timed-out.
========= Default timeout can be adjusted with --launch-timeout. Awaiting target completion.

最终 compute-sanitizer 会输出检测到的内存访问错误, 如下所示:

========= COMPUTE-SANITIZER
[sample_cuda] start
[sample_kernel] grid_size:1, block_size:512, shm_size:6144
========= Invalid __global__ write of size 4 bytes
=========     at 0x1190 in sample_kernel(int *, at::GenericPackedTensorAccessor<int, (unsigned long)1, at::RestrictPtrTraits, int>, at::GenericPackedTensorAccessor<int, (unsigned long)1, at::RestrictPtrTraits, int>, at::GenericPackedTensorAccessor<int, (unsigned long)1, at::RestrictPtrTraits, int>, at::GenericPackedTensorAccessor<int, (unsigned long)1, at::RestrictPtrTraits, int>, curandStateXORWOW *, unsigned int, int, unsigned int)
=========     by thread (32,0,0) in block (0,0,0)
=========     Address 0x7f40c00275a4 is out of bounds
=========     and is 23,461 bytes after the nearest allocation at 0x7f40c001fc00 of size 7,680 bytes
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame: [0x305c18]
=========                in /usr/lib/x86_64-linux-gnu/libcuda.so.1
=========     Host Frame: [0x1488c]
=========                in /usr/local/cuda-11.8/lib64/libcudart.so.11.0
=========     Host Frame:cudaLaunchKernel [0x6c318]
=========                in /usr/local/cuda-11.8/lib64/libcudart.so.11.0
=========     Host Frame:cudaError cudaLaunchKernel<char>(char const*, dim3, dim3, void**, unsigned long, CUstream_st*) [0x1f2f7]
=========                in /home/huanghy/CL/src/cuda/build/./example
=========     Host Frame:__device_stub__Z23sample_kernelPiN2at27GenericPackedTensorAccessorIiLm1ENS0_17RestrictPtrTraitsEiEES3_S3_S3_P17curandStateXORWOWjij(int*, at::GenericPackedTensorAccessor<int, 1ul, at::RestrictPtrTraits, int>&, at::GenericPackedTensorAccessor<int, 1ul, at::RestrictPtrTraits, int>&, at::GenericPackedTensorAccessor<int, 1ul, at::RestrictPtrTraits, int>&, at::GenericPackedTensorAccessor<int, 1ul, at::RestrictPtrTraits, int>&, curandStateXORWOW*, unsigned int, int, unsigned int) [0x1aec2]
=========                in /home/huanghy/CL/src/cuda/build/./example
=========     Host Frame:sample_kernel(int*, at::GenericPackedTensorAccessor<int, 1ul, at::RestrictPtrTraits, int>, at::GenericPackedTensorAccessor<int, 1ul, at::RestrictPtrTraits, int>, at::GenericPackedTensorAccessor<int, 1ul, at::RestrictPtrTraits, int>, at::GenericPackedTensorAccessor<int, 1ul, at::RestrictPtrTraits, int>, curandStateXORWOW*, unsigned int, int, unsigned int) [0x1af3a]
=========                in /home/huanghy/CL/src/cuda/build/./example
=========     Host Frame:sample_cuda(std::vector<at::Tensor, std::allocator<at::Tensor> >&, std::vector<CSR, std::allocator<CSR> >&, at::Tensor&, CSR const&, unsigned int, unsigned int, unsigned int, unsigned long long) [0x1a57d]
=========                in /home/huanghy/CL/src/cuda/build/./example
=========     Host Frame:sample(std::vector<at::Tensor, std::allocator<at::Tensor> >&, std::vector<CSR, std::allocator<CSR> >&, at::Tensor&, CSR const&, unsigned int, unsigned int, unsigned int, unsigned long long) [0x18900]
=========                in /home/huanghy/CL/src/cuda/build/./example
=========     Host Frame:main [0x87d0]
=========                in /home/huanghy/CL/src/cuda/build/./example
=========     Host Frame:__libc_start_main [0x21c87]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========     Host Frame:_start [0x804a]
=========                in /home/huanghy/CL/src/cuda/build/./example
========= 
=========  下面有很多重复单不同 threadIdx 和 blockIdx 的报错, 在此省略
========= 
[sample_kernel] finished
========= Program hit cudaErrorLaunchFailure (error 719) due to "unspecified launch failure" on CUDA API call to cudaMemGetInfo.
=========     Saved host backtrace up to driver entry point at error
=========     Host Frame: [0x4545f6]
=========                in /usr/lib/x86_64-linux-gnu/libcuda.so.1
=========     Host Frame:cudaMemGetInfo [0x533ab]
=========                in /usr/local/cuda-11.8/lib64/libcudart.so.11.0
=========     Host Frame:print_device_mem() [0x19796]
=========                in /home/huanghy/CL/src/cuda/build/./example
=========     Host Frame:sample_cuda(std::vector<at::Tensor, std::allocator<at::Tensor> >&, std::vector<CSR, std::allocator<CSR> >&, at::Tensor&, CSR const&, unsigned int, unsigned int, unsigned int, unsigned long long) [0x1a5bb]
=========                in /home/huanghy/CL/src/cuda/build/./example
=========     Host Frame:sample(std::vector<at::Tensor, std::allocator<at::Tensor> >&, std::vector<CSR, std::allocator<CSR> >&, at::Tensor&, CSR const&, unsigned int, unsigned int, unsigned int, unsigned long long) [0x18900]
=========                in /home/huanghy/CL/src/cuda/build/./example
=========     Host Frame:main [0x87d0]
=========                in /home/huanghy/CL/src/cuda/build/./example
=========     Host Frame:__libc_start_main [0x21c87]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========     Host Frame:_start [0x804a]
=========                in /home/huanghy/CL/src/cuda/build/./example
========= 
CUDA error at /home/huanghy/CL/src/cuda/sample.cu:53 code=719(cudaErrorLaunchFailure) "cudaMemGetInfo(&freeMem, &totalMem)" 
========= Target application returned an error
========= ERROR SUMMARY: 34 errors

在输出中, compute-sanitizer 会指明在具体的哪个 kernel 函数中发生了越界访问, 并指明相关的 threadIdx 和 blockIdx 以及内存地址.

以上述输出为例, 可以看到是在 sample_kernel() 函数中 threadIdx 为 (32,0,0) blockIdx (0,0,0) 处出现了 Address 0x7f40c00275a4 is out of bounds 的越界访问问题.

虽然地址信息很难让我们确定具体越界访问的位置, 但是通过该工具的输出, 可以确定到具体的 kernel 函数, 对于问题排查已经有了很大帮助.

参考

• cuda - Unspecified launch failure on Memcpy - Stack Overflow
• Compute-sanitizer not quite a drop-in replacement of cuda-memcheck - CUDA Developer Tools / Compute Sanitizer - NVIDIA Developer Forums