CUDA统一内存未向内核传输问题排查求助

调试项目时遇到CUDA统一内存未正确传输到设备的问题——将TestClass指针传入内核时触发cudaErrorIllegalAddress错误，但直接传入char**参数则运行正常。按统一内存机制，内核启动时应自动完成数据传输，现排查原因。

出错代码 #

#include "DebuggerDeviceHelper.cuh"
#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <string>
#include <iostream>

class TestClass {
public:
    TestClass(int elems) {
        cudaMallocManaged((void**)&map, elems * sizeof(char*));
    }

    void set_index(std::string str, int index) {
        cudaMallocManaged((void**)&map[index], (str.length() + 1) * sizeof(char));
        memcpy(map[index], str.c_str(), (str.length() + 1) * sizeof(char));
    }

    __device__ void test_print() {
        printf("test\n");
    }

    char** map;
};

__global__ void test_kernel(TestClass* test_class) {
    char** arr = test_class->map;

    for (int i = 0; i < 5; i++) {
        for (int j = 0; j < 6; j++) {
            printf("%c", arr[i][j]);
        }
        printf("\n");
    }
}

void main() {
    TestClass h_test_class(5);
    h_test_class.set_index("Test 1", 0);
    h_test_class.set_index("Test 2", 1);
    h_test_class.set_index("Test 3", 2);
    h_test_class.set_index("Test 4", 3);
    h_test_class.set_index("Test 5", 4);
    TestClass* d_test_class;
    cudaMemcpyToSymbol(d_test_class, &h_test_class, sizeof(TestClass*), 0, cudaMemcpyHostToDevice);
    test_kernel <<<1, 1>>> (d_test_class);
    std::cout << cudaGetErrorName(cudaDeviceSynchronize());
}

正常运行代码 #

__global__ void test_kernel(char** arr) {
    for (int i = 0; i < 5; i++) {
        for (int j = 0; j < 6; j++) {
            printf("%c", arr[i][j]);
        }
        printf("\n");
    }
}

void main() {
    TestClass h_test_class(5);
    h_test_class.set_index("Test 1", 0);
    h_test_class.set_index("Test 2", 1);
    h_test_class.set_index("Test 3", 2);
    h_test_class.set_index("Test 4", 3);
    h_test_class.set_index("Test 5", 4);
    char** arr = h_test_class.map;
    test_kernel <<<1, 1>>> (arr);
    std::cout << cudaGetErrorName(cudaDeviceSynchronize());
}

问题原因分析 #

1. cudaMemcpyToSymbol用法完全错误：该API用于修改全局符号变量的内存，而代码中d_test_class是未初始化的指针，并非全局符号，导致内核收到无效地址，触发非法访问错误。
2. 栈上对象无法被设备访问：h_test_class是主机栈上的普通对象，不属于统一内存范畴。设备无法直接访问主机栈内存，直接传入栈对象地址必然报错。
3. 正常代码的核心逻辑：h_test_class.map是通过cudaMallocManaged分配的统一内存指针，这类指针的地址空间在主机和设备间共享，直接传入内核时，CUDA会自动处理数据的可见性与传输，因此能正常运行。

修复方案 #

方式1：将TestClass对象分配为统一内存 #

把TestClass对象也放到统一内存中，让设备能直接访问：

void main() {
    TestClass* h_test_class;
    // 分配统一内存存储TestClass对象
    cudaMallocManaged((void**)&h_test_class, sizeof(TestClass));
    // 手动调用构造函数（cudaMallocManaged不会自动执行C++构造逻辑）
    new(h_test_class) TestClass(5);
    
    h_test_class->set_index("Test 1", 0);
    h_test_class->set_index("Test 2", 1);
    h_test_class->set_index("Test 3", 2);
    h_test_class->set_index("Test 4", 3);
    h_test_class->set_index("Test 5", 4);
    
    // 直接传入统一内存指针，CUDA自动处理传输
    test_kernel <<<1, 1>>> (h_test_class);
    std::cout << cudaGetErrorName(cudaDeviceSynchronize());
    
    // 手动调用析构函数并释放内存
    h_test_class->~TestClass();
    cudaFree(h_test_class);
}

方式2：显式复制主机对象到设备内存 #

如果不想用统一内存存储TestClass对象，可显式分配设备内存并复制：

void main() {
    TestClass h_test_class(5);
    h_test_class.set_index("Test 1", 0);
    h_test_class.set_index("Test 2", 1);
    h_test_class.set_index("Test 3", 2);
    h_test_class.set_index("Test 4", 3);
    h_test_class.set_index("Test 5", 4);
    
    TestClass* d_test_class;
    cudaMalloc((void**)&d_test_class, sizeof(TestClass));
    // 复制主机对象到设备内存
    cudaMemcpy(d_test_class, &h_test_class, sizeof(TestClass), cudaMemcpyHostToDevice);
    
    test_kernel <<<1, 1>>> (d_test_class);
    std::cout << cudaGetErrorName(cudaDeviceSynchronize());
    
    cudaFree(d_test_class);
}

此方法中，h_test_class.map作为统一内存指针，复制到设备后仍能访问共享地址空间中的数据，因此可正常运行。

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