CUDA使用虚拟类时非法的内存访问

问题描述

                         Base
                         /  \
                        /    \
                       /      \
                    Der1      Der2
                       \      /
                        \    /
                         \  /
                         Join

因此，我一直在研究依赖于菱形类继承结构的代码。在网上阅读了一些文章（以及关于可怕的钻石类继承的非常受欢迎的常见问题解答）之后，我将钻石的中产阶级转换为实际上从基类继承。现在，当我构造并传递Join类时，遇到了非法的内存访问。

我想了解我所做的事情是否根本上是错误的，因此，我最终创建了一个简单的示例来显示何时发生错误以及何时不发生错误。

一个简单的测试内核：

template<typename join_t>
__global__ void kernel(join_t monster) {
  float val   = monster.get_value_at(1);
  int der1_size = monster.get_total_size();

  printf("value[1] = %f\n",val);
  printf("size = %i\n",der1_size);
}

我的班级（我只需3个班级就可以重新创建问题，我什至不需要完整的菱形）：

struct base {
  base() {}
  __host__ __device__
  virtual int get_total_size() const = 0;

  void set_base_size(int const& s) { base_size = s; }
  protected:
    int base_size;
};

struct der1 : public virtual base {

  der1() : base() {}

  float* ptr1;
  int size1;

  __host__ __device__ 
  float get_value_at(int const& i) const {
    return ptr1[i];
  }

  __host__ __device__ 
  int get_size() const { return size1; }

  __host__ __device__
  int get_total_size() const override {
    return base::base_size + get_size();
  }
};

struct join : public der1/*,public der2 */ {
  join() : base(),der1() /*,der2() */ {}

  __host__ __device__
  int get_total_size() const override {
    return der1::get_total_size();
  }
};

一些测试代码：

template<typename vector_struct_t>
auto set_smart(vector_struct_t& v) {

  join my_container;
  int base_size = 10;
  
  my_container.ptr1 = thrust::raw_pointer_cast(v.data());

  my_container.set_base_size(base_size);
  my_container.size1 = v.size();

  return my_container;
}

int
main(int argc,char** argv)
{
  cudaError_t status = cudaSuccess;

  // let's use thrust vector<type_t> for initial arrays
  thrust::host_vector<value_t>   h_vector(10);
  for (index_t i = 0; i < 10; i++)
    h_vector[i] = i;

  thrust::device_vector<value_t>  d_vector = h_vector;

  auto my_container = set_smart(d_vector);

  // Device Output
  status = cudaDeviceSynchronize();
  if(cudaSuccess != status) return EXIT_FAILURE;
  kernel<<<1,1>>>(my_container);
  if(cudaSuccess != status) return EXIT_FAILURE;

  return 0;
}

我得到的输出是什么

what():  an illegal memory access was encountered

我期望的是，请注意，如果我不为der1类使用关键字virtual，则此代码可以正常运行。

value[1] = 1.000000
size = 20

解决方法

暂无找到可以解决该程序问题的有效方法，小编努力寻找整理中！

如果你已经找到好的解决方法，欢迎将解决方案带上本链接一起发送给小编。

小编邮箱:dio#foxmail.com (将#修改为@）