CUDA C++内核调用类内嵌套对象方法返回全零问题排查

你的问题本质是主机与设备地址空间隔离导致的指针非法访问：

• 你在主机侧构造Ding对象时，构造函数调用的是主机侧的malloc，stuff_vector指针指向的是主机内存地址
• 执行cudaMemcpy(d_teil, this, sizeof(Ding), cudaMemcpyHostToDevice)时，仅对Ding对象做了浅拷贝，拷贝到设备侧的d_teil内部的stuff_vector指针仍然指向原来的主机内存地址
• 核函数运行在设备侧，访问主机地址属于非法内存访问，直接触发未定义行为，所以返回结果全为0，你没有加CUDA错误检查所以看不到显式报错。

1. CUDA编程模型中主机和设备拥有完全独立的虚拟地址空间，任何一侧的原生指针都不能直接在另一侧访问，否则会触发非法内存访问错误
2. 对包含指针成员的类/结构体做跨设备的内存拷贝时，仅拷贝结构体本身是无效的，需要手动对指针指向的内存做深拷贝，同时将设备侧的指针值更新为设备内存的有效地址
3. CUDA核函数的运行错误是异步的，如果不主动调用cudaGetLastError()、cudaDeviceSynchronize()做错误检查，程序会静默失败不会抛出异常。

你可以按以下步骤修改calculate_stuff方法的逻辑：

1. 新增设备侧stuff_vector的内存分配、数据拷贝逻辑
2. 浅拷贝Ding对象到设备后，手动更新设备侧Ding对象的stuff_vector指针为设备内存地址
3. 增加CUDA错误检查逻辑便于后续调试
4. 补充cudaFree释放申请的设备内存避免泄漏

以下是修改后的calculate_stuff示例代码：

__host__ void Ding::calculate_stuff(double coeff) {
    prop = 1;
    double* d_result;
    Ding* d_teil;
    test_klasse* d_stuff_vector; // 新增设备侧stuff_vector指针

    // 分配设备内存
    cudaMalloc(&d_teil, sizeof(Ding));
    cudaMalloc(&d_result, size * sizeof(double));
    cudaMalloc(&d_stuff_vector, size * sizeof(test_klasse));

    // 拷贝stuff_vector的实际数据到设备
    cudaMemcpy(d_stuff_vector, stuff_vector, size * sizeof(test_klasse), cudaMemcpyHostToDevice);
    // 先浅拷贝Ding对象到设备
    cudaMemcpy(d_teil, this, sizeof(Ding), cudaMemcpyHostToDevice);
    // 手动修改设备侧Ding对象的stuff_vector指针为设备地址
    // 注意：如果stuff_vector是private成员，需要给Ding类新增__host__方法修改该成员，或者临时改为public
    cudaMemcpy(&(d_teil->stuff_vector), &d_stuff_vector, sizeof(test_klasse*), cudaMemcpyHostToDevice);

    // 启动核函数
    calculation <<< size / 512 + 1, 512 >>> (d_teil, d_result, size, coeff);
    // 检查核函数启动错误
    cudaError_t err = cudaGetLastError();
    if (err != cudaSuccess) {
        std::cerr << "Kernel launch failed: " << cudaGetErrorString(err) << std::endl;
        return;
    }
    // 等待核函数执行完成，检查执行错误
    err = cudaDeviceSynchronize();
    if (err != cudaSuccess) {
        std::cerr << "Kernel execute failed: " << cudaGetErrorString(err) << std::endl;
        return;
    }

    // 拷贝结果回主机
    cudaMemcpy(result_vector, d_result, size * sizeof(double), cudaMemcpyDeviceToHost);

    for (int ii = 0; ii < size; ii++) {
        std::cout << result_vector[ii] << std::endl;
        result += result_vector[ii];
    }

    // 释放设备内存
    cudaFree(d_teil);
    cudaFree(d_result);
    cudaFree(d_stuff_vector);
}

另外你原代码中申请的d_size变量完全没有使用，可以直接删除。

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