在CUDA kernel中使用C++类对象及成员触发编译报错如何解决

我是C++新手，对CUDA计算工具的使用经验较少。目前我尝试使用CUDA将代码在GPU上并行化，需求要求我可以实现类实例化，或是至少能在kernel中调用类成员。

首先是我的类定义：

//Header
#pragma once
#include<cuda.h>
#include<cuda_runtime.h>
#include<cuda_runtime_api.h>
#include<stdio.h>
#include<iostream>

class Ding
{
private:
    int index;
    double vector_1[100];
    double vector_2[100];
    double prop;

public:
    __host__ __device__ Ding(int);
    __host__ __device__ ~Ding();

    __device__ void calculate_stuff(double, int);
    __device__ double get_prop();
};

//Source
#include "Ding.h"
#include <math.h>

__host__ __device__ Ding::Ding(int ind) {
    index = ind;
    prop = 1;
    for (int ii = 0; ii < 100; ii++) {
        vector_1[ii] = (4 * ii + ind) / (ind + ii + 1);
        vector_2[ii] = (-2.14 * ii + ind) / (2*ind + ii + 1);
    }
}

__host__ __device__ Ding::~Ding() {};

__device__ void Ding::calculate_stuff(double coeff, int N) {
    prop = 1;

    for (int ii = 0; ii < N; ii++) {
        for (int jj = 0; jj < 100; jj++) {
            prop += pow(-1, ii) * vector_1[jj] * vector_2[jj]*coeff;
        }
    }
}

__device__ double Ding::get_prop() {
    return prop;
}

该类仅包含一些无意义计算，仅作为我跑通CUDA相关逻辑的示例使用。

以下是主源文件代码：

#include<cuda.h>
#include<cuda_runtime.h>
#include<cuda_runtime_api.h>
#include "device_launch_parameters.h"
#include<stdio.h>
#include<iostream>
#include "Ding.h"

using namespace std;

__global__ void some_kernel(double *vec_a, Ding* teil, int size) {
    int ii = blockIdx.x * blockDim.x + threadIdx.x;

    if (ii < size) {
        vec_a[ii] += teil[ii].get_prop();
        vec_a[ii] += ii;
    }
}

int main() {
    double* vec_1, * d_vec_1;
    int N = 300;
    double result = 0;
    Ding* teil;
    Ding* d_teil;

    vec_1 = (double*)malloc(N * sizeof(double));
    teil = (Ding*)malloc(N * sizeof(Ding));
    for (int ii = 0; ii < N; ii++) {
        vec_1[ii] = 0;
        teil[ii] = Ding::Ding(ii);
    }

    cudaMalloc(&d_vec_1, N * sizeof(double));
    cudaMalloc(&d_teil, N * sizeof(Ding));

    cudaMemcpy(d_vec_1, vec_1, N * sizeof(double), cudaMemcpyHostToDevice);
    cudaMemcpy(d_teil, teil, N * sizeof(Ding), cudaMemcpyHostToDevice);


    some_kernel <<< 256/N + 1, 256 >>> (d_vec_1, d_teil,  N);

    for (int ii = 0; ii < N; ii++) {
        result += vec_1[ii];
    }
    
    cout << "Old result: " << result << endl;

    cudaMemcpy(vec_1, d_vec_1, N * sizeof(double), cudaMemcpyDeviceToHost);
    
    result = 0;

    for (int ii = 0; ii < N; ii++) {
        result += vec_1[ii];
    }

    cout << "New result: " << result << endl;
}

我在主机端创建了对象数组并拷贝到设备端，设备端仅调用成员函数get_prop获取prop属性值并累加到vec_a向量中。注释掉vec_a[ii] += teil[ii].get_prop();行时代码运行正常，但只要访问类成员就会触发如下编译错误，我暂时无法定位问题原因：

错误 MSB3721 命令""C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.4\bin\nvcc.exe" -gencode=arch=compute_52,code=sm_52 --use-local-env -ccbin "C:\Program Files (x86)\Microsoft Visual Studio\2019\Community\VC\Tools\MSVC\14.29.30133\bin\HostX86\x64" -x cu -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.4\include" -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.4\include" -G --keep-dir x64\Debug -maxrregcount=0 --machine 64 -cuda -cudart static -g -D_DEBUG -D_CONSOLE -D_UNICODE -DUNICODE -Xcompiler "/EHsc /W3 /nologo /Od /Fdx64\Debug\vc142.pdb /FS /Zi /RTC1 /MDd " -o x64\Debug\File.cu.obj "C:\Users\ronal\source\repos\Basic_Cuda_Test\Basic_Cuda_Test\File.cu"" 已退出，返回代码为 255。 Basic_Cuda_Test C:\Program Files (x86)\Microsoft Visual Studio\2019\Community\MSBuild\Microsoft\VC\v160\BuildCustomizations\CUDA 11.4.targets 785

从报错信息可知我使用的是CUDA v11.4和Microsoft Visual Studio 2019环境，不访问类成员时计算逻辑运行正常，说明CUDA环境安装配置无误。我最优期望是可以直接在设备端创建类对象，当前阶段能够在设备端正常调用类成员即可满足需求，应该是比较基础的使用问题，期待各位的解答与解决方案。

核心问题修复 #

1. 解决设备函数符号无法解析问题 #

你把Ding类的成员函数实现放在了单独的源文件中，而get_prop是__device__修饰的设备函数，nvcc默认编译模式下不支持跨编译单元的设备函数调用，所以编译报错。两种修复方案二选一：

• 把Ding类的所有成员函数实现全部移动到Ding.h头文件中，编译主文件时就能直接看到函数实现
• 打开Visual Studio项目属性，在「CUDA C/C++ - Common」中开启「Generate Relocatable Device Code」选项，启用CUDA分离编译模式

2. 修正核函数启动参数 #

你写的启动参数<<< 256/N + 1, 256 >>>当N=300时，256/300结果为0，实际仅启动1个block共256个线程，无法覆盖全部300个元素，会导致部分数据没有被计算。修改为：

some_kernel <<< (N + 255) / 256, 256 >>> (d_vec_1, d_teil, N);

这个写法可以保证不论N取任意正整数，都能启动足够的block覆盖所有计算索引。

3. 规范主机端对象初始化 #

你用malloc分配Ding对象数组，malloc仅分配原始内存不会调用构造函数，后续的赋值操作虽然暂时能运行，但属于非规范写法，容易触发未定义行为。修改为用new分配对象数组：

teil = new Ding[N];
vec_1 = new double[N];

分配时会自动为每个元素调用构造函数，符合C++对象初始化规范。

4. 增加核函数错误检查 #

核函数调用是异步执行的，当前没有做执行状态检查，就算核函数运行出错也无法及时感知。可以在核函数调用后增加同步和错误检查逻辑：

some_kernel <<< (N + 255) / 256, 256 >>> (d_vec_1, d_teil, N);
// 检查核函数启动是否出错
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
    cout << "Kernel launch failed: " << cudaGetErrorString(err) << endl;
    return 1;
}
// 等待核函数执行完成
err = cudaDeviceSynchronize();
if (err != cudaSuccess) {
    cout << "Kernel execute failed: " << cudaGetErrorString(err) << endl;
    return 1;
}

设备端创建类对象的方法 #

只要类的构造函数加了__device__修饰符，直接在核函数中定义局部对象即可，也可以用new在设备堆上分配对象：

__global__ void test_create_obj_kernel() {
    // 直接在设备栈上创建对象
    Ding local_obj(threadIdx.x);
    local_obj.calculate_stuff(1.5, 5);
    printf("Thread %d prop: %f\n", threadIdx.x, local_obj.get_prop());

    // 在设备堆上分配对象
    Ding* heap_obj = new Ding(threadIdx.x);
    // 使用完成后手动释放
    delete heap_obj;
}

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