无法在CUDA核函数中使用自定义模板类的问题求助

嘿，我看到你遇到的这个CUDA模板类的问题了，简直太熟悉了——早期CUDA版本对C++类的默认生成函数支持有坑，尤其是涉及到主机/设备跨端传递对象的时候！

核心问题分析 #

你遇到的非法内存访问，根源在于两个关键点：

1. 默认拷贝构造函数的设备端缺失 #

你把主机端创建的MyArray对象直接传值给核函数时，设备端需要生成一个类的副本。但CUDA 9.1这类早期版本不会自动为带const成员的类生成同时支持主机和设备的默认拷贝构造函数。这就导致设备端的副本没有正确复制size和data成员，data变成了无效的垃圾指针，自然触发非法内存访问。

而你直接传x.data能正常工作，是因为这个指针本身是有效的设备内存地址，绕过了类对象拷贝的问题。

2. 模板实例化位置错误 #

你在头文件里写template class MyArray<double>;会导致每个包含头文件的编译单元都重复实例化模板，可能引发链接冲突，也会干扰设备端的函数实例化。

具体修复步骤 #

步骤1：手动添加跨端支持的拷贝构造函数 #

在class.h的MyArray类声明中，添加显式的拷贝构造函数，确保主机和设备端都能正确拷贝类对象：

#pragma once
template <typename T>
class MyArray {
public:
    const int size;
    T *data;
    __host__ MyArray(int size);
    // 添加跨端支持的拷贝构造函数
    __device__ __host__ MyArray(const MyArray& other) : size(other.size), data(other.data) {}
    __device__ __host__ T GetValue(int);
    __device__ __host__ void SetValue(T, int);
    __device__ __host__ T& operator()(int);
    ~MyArray();
};
// 移除头文件里的模板实例化

步骤2：调整模板实例化位置 #

把模板实例化语句移到class.cu的末尾，确保模板只在实现文件中实例化一次：

// class.cu 中的实现代码
template <typename T>
__host__ MyArray<T>::MyArray(int size) : size(size) {
    gpuErrchk(cudaMalloc(&data, size * sizeof(T)));
}

template <typename T>
__device__ __host__ T MyArray<T>::GetValue(int i) {
    return data[i];
}

template <typename T>
__device__ __host__ void MyArray<T>::SetValue(T val, int i) {
    data[i] = val;
}

template <typename T>
__device__ __host__ T& MyArray<T>::operator()(int i) {
    return data[i];
}

template <typename T>
MyArray<T>::~MyArray() {
    gpuErrchk(cudaFree(data));
}

// 在实现文件末尾添加模板实例化
template class MyArray<double>;

步骤3：确保成员函数的修饰符正确 #

所有需要在主机和设备端都调用的成员函数，必须加上__device__ __host__修饰符，这一点你在声明里已经做了，但实现时也要保持一致（如上例所示）。

修复后核函数的验证 #

修改完成后，你的核函数里的三种访问方式都能正常运行了：

__global__ void test(MyArray<double> array, double *data, int size) {
    int j = threadIdx.x;
    array.SetValue(1, j); // 现在可以正常运行
    array(j) = 2; // 现在可以正常运行
    array.data[j] = 3; // 现在可以正常运行
    printf("Thread %d: array.data[j] = %f\n", j, array.data[j]);
}

关于无警告的说明 #

CUDA编译器在处理主机/设备跨端的类拷贝时，不会默认检查是否需要设备端的拷贝构造函数，这类运行时内存错误只能通过显式的CUDA错误检查（比如你用的cudaPeekAtLastError和cudaDeviceSynchronize）才能捕获，编译器没法提前预警——这也是早期CUDA C++开发的一个常见“坑”。

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