推力：初始化推力变换函子时读取的__global__无效

问题描述

我有一个经过测试的基于CRTP的自定义矩阵库，用于动态矩阵：

#include <thrust/device_vector.h>
#include <assert.h>

namespace CML
{
template<class T,class Derived>

    template<class T,class Derived> class Matrix_Base {};
    template <class T>
    class Dynamic_Matrix : public Matrix_Base<T,Dynamic_Matrix<T>> 
    {
    private:
        size_t n_rows,n_cols;

        T* data;

        bool is_allocated;

        __host__ __device__ void allocate_data()
        {
              assert(n_rows > 0 && n_cols > 0);
              data = new T[n_rows * n_cols];
              is_allocated = true;
        }

        __host__ __device__ void deallocate_data() { delete[] data; data = nullptr; is_allocated = false; }

        __host__ __device__ void assign_data(const Dynamic_Matrix &other)
        {
             if (!other.is_allocated){ return; }

             n_rows = other.n_rows;
             n_cols = other.n_cols;

             if (!is_allocated){ allocate_data(); }
        
             printf("Dynamic matrix assign data,is_allocated? %d,is_other_allocated? %d \n",is_allocated,other.is_allocated);
             if (other.n_rows == 0 || other.n_cols == 0)
             {
                  printf("Error: n_rows == 0 or n_cols == 0! \n");
             } 
             for (size_t i = 0; i < n_rows; i++)
             {
                  for (size_t j = 0; j < n_cols; j++)
                  {
                        this->data[n_cols * i + j] = other.data[n_cols * i + j]; //<-- this line gives error
                  }
             }
        }

    public:
        
        __host__ __device__ Dynamic_Matrix() : n_rows(0),n_cols(0),data(nullptr),is_allocated(false) {}

        __host__ __device__ Dynamic_Matrix(const size_t n_rows,const size_t n_cols) :
        n_rows(n_rows),n_cols(n_cols),is_allocated(false)
        {
            allocate_data();
        }

        __host__ __device__ Dynamic_Matrix(const Dynamic_Matrix &other):
        data(nullptr),is_allocated(false)
        {
            assign_data(other);
        }

        __host__ __device__ ~Dynamic_Matrix() { deallocate_data(); }

        __host__ __device__ Dynamic_Matrix& operator=(const Dynamic_Matrix &rhs)
        {
             if (this == &rhs)
             {
                 return *this;
             }
             deallocate_data(); 
             assign_data(rhs);
             return *this;
        }

        __host__ __device__ void resize(const size_t n_rows,const size_t n_cols)
        {
              assert(n_rows > 0 && n_cols > 0);
              *this = Dynamic_Matrix<T>(n_rows,n_cols);
        }
    };

    using MatrixXd = Dynamic_Matrix<double>;
};

在我的项目的某些类中使用，类似于这个简单的类：

#include <thrust/device_vector.h>

class My_Class
{
private:

    CML::MatrixXd mat1;
    CML::MatrixXd mat2;
    CML::MatrixXd mat3;
    CML::MatrixXd mat4;

public:
    __host__ __device__ My_Class() { mat1.resize(3,1); mat2.resize(3,1); mat3.resize(3,1); mat4.resize(3,1); }
};

但是，当我尝试使用functor（在这里简化）运行推力:::转换时，

#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <thrust/iterator/counting_iterator.h>

class myFunctor
{
private: 
    My_Class my_class;
public: 
    __host__ __device__ myFunctor() {}

    __device__ double operator()(const unsigned int n) { double ret = 0; return ret; }
};

int main()
{
    int n_cbs = 10;
    thrust::device_vector<double> cb_costs(n_cbs);

    thrust::counting_iterator<unsigned int> index_iter(0);
    
    thrust::transform(thrust::device,index_iter,index_iter + n_cbs,cb_costs.begin(),myFunctor());
    return 0;
}

这给了我错误Invalid __global__ read of size 8 ========= at 0x00001250 in /../dynamic_matrix.cuh:724:CML::Dynamic_Matrix<double>::assign_data(CML::Dynamic_Matrix<double> const &) ========= by thread (255,0) in block (0,0) ========= Address 0x55965a2ce530 is out of bounds 在应用程序上使用cuda-memcheck时。错误报告中的行是矩阵库中的this->data[n_cols * i + j] = other.data[n_cols * i + j];行。 Sofar作为调试工具带给我，我还没有发现任何索引超出此范围。 有什么问题的建议吗？注释掉构造中的矩阵对象使代码可运行。

解决方法

正如评论中指出的那样，问题在于您的Dynamic_Matrix类在主机和设备之间的可移植性。

原则上，只要您仅实例化并在主机存储器中使用或实例化并在设备存储器中使用，设计就可以。 new运算符可以很好地在主机代码和设备代码中使用（尽管设备分配是在运行时堆上进行的，该堆必须具有先验大小，并且不能从主机API进行访问），但是类实例不是便携式的。数据分配位于主机内存或设备内存中，并且在未分配数据的运行时空间中不可访问。当您这样做时：

thrust::transform(thrust::device,index_iter,index_iter + n_cbs,cb_costs.begin(),myFunctor());

您暗中要求将您的myFunctor实例复制到设备，这是代码中断的地方。数据指针是主机指针，分配给主机new，并且设备代码因无效的指针错误而崩溃。

根据经验，推力使用的任何物体都应为POD type。这意味着它们可以轻而易举地复制到设备上，而不会出现可移植性问题。您极有可能需要重构代码以消除对new和delete的使用，并将指针传递到将保存类数据的内存（您在设备或托管内存中手动分配的数据） ）用于该目的的构造函数版本。另外，也可以考虑通过模板专门化来静态声明数据数组，因此您的类完全不需要动态内存分配。 Eigen在CUDA支持中使用了这种方法。

尽管它不适用于此处（因为推力未对函子使用引用传递），但请注意，还有一个非常优雅的design pattern用于使用托管内存分配可以安全地传递的类参考。以这种方式分配的类可移植到主机和设备，尽管它们的最终性能会很差，这是因为它们使用的托管内存的开销。