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

So, I have been working on a code that relies on a diamond class inheritance structure. Following some articles online (and a very popular FAQ on the dreaded diamond class inheritance), I switched my middle class of the diamond to inherit virtually from the base class. Now, when I construct and pass in the Join class, I encounter illegal memory access.

I would like to understand if what I am doing is fundamentally wrong, so, I ended up creating a simple example to show when the error occurs, and when it doesn't occur.

A simple test kernel:

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);
}

My classes (I can recreate the issue with just 3 classes, I don't even need the full diamond):

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();
  }
};

Some testing code:

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 I get as an output:

what():  an illegal memory access was encountered

What I expect, note this code runs fine if I don't use the keyword virtual for der1 class.

value[1] = 1.000000
size = 20
Blizzard
  • 1,117
  • 2
  • 11
  • 28
  • 1
    https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#virtual-functions – Robert Crovella Oct 02 '20 at 23:34
  • Oh god... It's written in a bit convoluted way, to clarify, what I am trying to do is not supported within CUDA? The part I don't understand that is is it virtual functions that aren't supported in my example or virtual base class (I am guessing the later)? – Blizzard Oct 02 '20 at 23:38
  • 2
    It is not allowed to initialize an object in host code, pass that object to device code, then attempt to call one of its virtual functions. You are doing exactly that. If you create/initialize the object in device code, then call its virtual method in device code, it should work. This is a somewhat common question, I will locate a duplicate eventually. – Robert Crovella Oct 02 '20 at 23:40
  • Seems very limiting. :( Thank you for your answer, that helps a lot! – Blizzard Oct 02 '20 at 23:44

0 Answers0