CUDA最简内核触发非法内存访问错误的原因与调试咨询

环境信息 #

• 系统：Ubuntu 22.04
• CUDA版本：12.2.128
• GPU：GeForce RTX 2060

问题代码 #

#include "cuda_runtime.h"

#include <iostream>
#include <stdexcept>
#include <cstdlib>
#include <cmath>

#define CUDACHECK(E) if (E != cudaSuccess) \
    { \
        std::cerr << "Fatal error [line " << __LINE__ << "]: " << cudaGetErrorString(E) << ".\n"; \
        std::abort(); \
    }

#define ASSERT(B) if (!(B)) \
    { \
        std::cerr << "Fatal error: test failed in line " << __LINE__ << ".\n"; \
        std::abort(); \
    }

__global__ void fill(float* x, const std::size_t& size, const float& value)
{
    const auto tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid < size)
    {
        x[tid] = value;
    }
}


int main()
{
    // allocate device data
    const auto size = std::size_t{10000};
    auto* xD = static_cast<float*>(nullptr);
    CUDACHECK(cudaMalloc(&xD, size * sizeof(float)));
    ASSERT(xD != nullptr);
    
    // run fill kernel
    const auto fill_value = static_cast<float>(-1.0);
    fill<<<1,1>>>(xD, size, fill_value);
    CUDACHECK(cudaDeviceSynchronize());
    
    // <--- control doesn't get past here!
    
    // copy to host
    //auto* xH = reinterpret_cast<float*>(std::malloc(size * sizeof(float)));
    //ASSERT(xH != nullptr);
    //CUDACHECK(cudaMemcpy(xH, xD, size * sizeof(float), cudaMemcpyDeviceToHost));
    
    //constexpr auto tol = std::is_same<float, float>::value ? static_cast<float>(1E-5) : static_cast<float>(1E-10);
    //for (auto ii = std::size_t{}; ii < size; ++ii)
    //{
    //    ASSERT(std::fabs(fill_value - xH[ii]) < tol);
    //}
    
    // free
    CUDACHECK(cudaFree(reinterpret_cast<void*>(xD)));
    //std::free(xH);
    
    // tests have been passed
    std::cout << "Tests passed! Hallo, CUDA world!" << std::endl;
    
    return 0;
}

运行输出 #

Fatal error [line 42]: an illegal memory access was encountered.
Aborted (core dumped)

Compute-sanitizer检测结果 #

========= COMPUTE-SANITIZER
========= Invalid __global__ read of size 8 bytes
=========     at 0x20 in fill(float *, const unsigned long &, const float &)
=========     by thread (0,0,0) in block (0,0,0)
=========     Address 0x7ffc902b1f30 is out of bounds
=========     and is 639439150385 bytes after the nearest allocation at 0x7f67aea00000 of size 512 bytes
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame: [0x32e950]
=========                in /lib/x86_64-linux-gnu/libcuda.so.1
=========     Host Frame:libcudart_static_4d8b33a106dceb3c07a56e26de61f2d53bb62a68 [0x1093e]
=========                in /home/nitin/Documents/code/gpu/device/./hw_cuda
=========     Host Frame:cudaLaunchKernel [0x70b4e]
=========                in /home/nitin/Documents/code/gpu/device/./hw_cuda
=========     Host Frame:cudaError cudaLaunchKernel<char>(char const*, dim3, dim3, void**, unsigned long, CUstream_st*) [0xb235]
=========                in /home/nitin/Documents/code/gpu/device/./hw_cuda
=========     Host Frame:__device_stub__Z4fillPfRKmRKf(float*, unsigned long const*, float const*) [0xb094]
=========                in /home/nitin/Documents/code/gpu/device/./hw_cuda
=========     Host Frame:fill(float*, unsigned long const&, float const&) [0xb0f7]
=========                in /home/nitin/Documents/code/gpu/device/./hw_cuda
=========     Host Frame:main [0xadaa]
=========                in /home/nitin/Documents/code/gpu/device/./hw_cuda
=========     Host Frame:../sysdeps/nptl/libc_start_call_main.h:58:__libc_start_call_main [0x29d90]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========     Host Frame:../csu/libc-start.c:379:__libc_start_main [0x29e40]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========     Host Frame:_start [0xab15]
=========                in /home/nitin/Documents/code/gpu/device/./hw_cuda
========= 
========= Program hit cudaErrorLaunchFailure (error 719) due to "unspecified launch failure" on CUDA API call to cudaDeviceSynchronize.
=========     Saved host backtrace up to driver entry point at error
=========     Host Frame: [0x47e786]
=========                in /lib/x86_64-linux-gnu/libcuda.so.1
=========     Host Frame:cudaDeviceSynchronize [0x48a64]
=========                in /home/nitin/Documents/code/gpu/device/./hw_cuda
=========     Host Frame:main [0xadaf]
=========                in /home/nitin/Documents/code/gpu/device/./hw_cuda
=========     Host Frame:../sysdeps/nptl/libc_start_call_main.h:58:__libc_start_call_main [0x29d90]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========     Host Frame:../csu/libc-start.c:379:__libc_start_main [0x29e40]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========     Host Frame:_start [0xab15]
=========                in /home/nitin/Documents/code/gpu/device/./hw_cuda
========= 
Fatal error [line 42]: ========= Program hit cudaErrorLaunchFailure (error 719) due to "unspecified launch failure" on CUDA API call to cudaDeviceSynchronize.
=========     Saved host backtrace up to driver entry point at error
=========     Host Frame: [0x47e786]
=========                in /lib/x86_64-linux-gnu/libcuda.so.1
=========     Host Frame:cudaDeviceSynchronize [0x48a64]
=========                in /home/nitin/Documents/code/gpu/device/./hw_cuda
=========     Host Frame:main [0xadfb]
=========                in /home/nitin/Documents/code/gpu/device/./hw_cuda
=========     Host Frame:../sysdeps/nptl/libc_start_call_main.h:58:__libc_start_call_main [0x29d90]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========     Host Frame:../csu/libc-start.c:379:__libc_start_main [0x29e40]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========     Host Frame:_start [0xab15]
=========                in /home/nitin/Documents/code/gpu/device/./hw_cuda
========= 
unspecified launch failure.
========= Error: process didn't terminate successfully
========= Target application returned an error
========= ERROR SUMMARY: 3 errors

问题解答 #

1. 非法内存访问错误的成因 #

错误核心在于CUDA Kernel不能直接接收主机端的引用类型参数。代码中fill Kernel的参数const std::size_t& size和const float& value是引用类型，这些引用指向的是主机栈上的变量内存地址。而GPU设备无法直接访问主机的栈内存，当Kernel尝试读取size和value时，会去访问主机端的内存地址（从sanitizer输出的0x7ffc902b1f30可以判断是主机栈地址，设备内存地址通常以0x7f开头），导致非法内存访问。

2. 调试定位与解决方法 #

定位步骤

• 查看compute-sanitizer的错误输出：明确指出错误发生在fill Kernel中，是对主机地址的非法读取。
• 区分内存地址类型：0x7ffc开头的地址属于主机栈内存，GPU无法直接访问，说明Kernel尝试访问了主机端的数据。
• 检查Kernel参数：发现参数使用了引用类型，这是CUDA编程的常见误区——Kernel参数传递应避免使用主机端的引用或指针（除非使用统一内存且配置正确）。

解决方法

将Kernel的引用参数改为值传递，修改后的fill Kernel如下：

__global__ void fill(float* x, const std::size_t size, const float value)
{
    const auto tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid < size)
    {
        x[tid] = value;
    }
}

这样，size和value会被直接拷贝到设备端供Kernel使用，避免了对主机内存的非法访问。

内容的提问来源于stack exchange，提问作者Nitin Malapally