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

Question

我有一个基于 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), data(nullptr), 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); }
};

但是，当我尝试运行一个 thrust::transform 和一个仿函数（此处简化）时，如

#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,0) in block (0,0,0) ========= Address 0x55965a2ce530 is out of bounds 在应用程序上使用 cuda-memcheck 时。错误报告中的行是矩阵库中的this->data[n_cols * i + j] = other.data[n_cols * i + j];行。就调试而言，我还没有发现任何索引超出范围。 对问题有什么建议吗？在构造中注释掉矩阵对象使代码运行可用。

Answer 1

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

原则上设计是可以的，只要你要么在主机内存中实例化使用，要么在设备内存中实例化使用。 new 运算符可以很好地用于主机和设备代码（尽管设备分配在运行时堆上进行，该堆必须是先验大小并且无法从主机 API 访问），但是 class 实例不可移植。数据分配在主机内存或设备内存中，并且在未分配数据的运行时 space 中不可访问。当您这样做时：

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

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

根据经验，任何你使用推力的东西都应该是 POD type。这意味着它们可以轻松复制到设备上，而不会出现便携性问题。您很可能需要重构代码以消除 new 和 delete 的使用，并将指针传递给将保存 class 数据的内存（您在设备上手动分配的内存）或在托管内存中）为此目的构造函数的版本。或者考虑通过模板专门化静态声明数据数组，这样您的 class 根本不需要动态内存分配。 Eigen 在其 CUDA 支持中使用了这种方法。

虽然这里不适用（因为 thrust 不使用仿函数的引用传递），请注意还有一个非常优雅的 design pattern 用于使用托管内存分配 class es 可以安全地通过引用传递。类以这种方式分配的内存可移植到主机和设备，但由于它们使用的托管内存的开销，它们的最终性能会较差。

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

Thrust: invalid global read when initializing thrust transform functor

c++

cuda

thrust

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

Thrust: invalid __global__ read when initializing thrust transform functor

c++

cuda

thrust

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

Thrust: invalid global read when initializing thrust transform functor