An illegal instruction was encountered in Virtual Class Inheritance On Device

[ak-yoshi]

I try to use virtual inheritance in device with referring following site : How to implement device side CUDA virtual functions? - Stack Overflow.

But program was illegally terminated with error command “an illegal instruction was encountered”.
(however, same program normaly works when compiling with “-G” option of nvcc.)

Could anyone tell me the relosution?

Sample code is as follows;

#include"test.h"

__global__
void calcFunction(Base **ptr) {
  const double val0 = (*ptr)->calc();
  printf("val0 : %3.3e\n", val0);
}
int main( int argc, char **argv ) {


  Test test;
  test.malloc();

  calcFunction<<<1, 1>>>(test.ptr);
  CUDA_SAFE_CALL( cudaDeviceSynchronize() );

}

#pragma once
#include <stdio.h>

#define CUDA_SAFE_CALL( func ) \
  do { \
    cudaError_t err = ( func ); \
    if ( err != cudaSuccess ) { \
      fprintf( stderr, "  Error : %s ( error code: %d ) at %s line %d\n", cudaGetErrorString( err ), err, __FILE__, __LINE__ ); \
      exit( err ); \
    } \
  } while ( 0 )

class Base {

public:
  __host__ __device__
  Base() {}

  __host__ __device__
  virtual ~Base() {}

  __host__ __device__
  virtual double calc() = 0;

};

class Derived : public Base {

public:
  __host__ __device__
  Derived() { printf("allocated\n"); }

  __host__ __device__
  double calc() override { return 1.0; };

};

class Test {
public:
  Base **ptr;

  void malloc();

};

#include"test.h"

__global__
void allocateKernel(Base **ptr) {

    (*ptr) = new Derived();

}

void Test::malloc() {

  CUDA_SAFE_CALL( cudaMalloc( (void**)&ptr, sizeof( Base* ) ) );

  allocateKernel<<<1, 1>>>(ptr);
  CUDA_SAFE_CALL( cudaDeviceSynchronize() );

}

[striker159]

Your sample code works fine on my machine.

[Robert_Crovella]

it might be helpful if you share the cuda version you are using as well as the exact compile command lines you are using. identifying whether you are on windows or linux may also be helpful

Furthermore, for a problem such as you describe, if I personally ran into that, I would first test on the latest CUDA (12.1, currently) before I did any more head-scratching.

[ak-yoshi]

Thanks for your reply!
CUDA version I used is 12.0 , and exact compile command is as follows.

nvcc -o main.cu.o -c main.cu -O3 -std=c++20  -arch=sm_70 --expt-relaxed-constexpr  -I/usr/local/openmpi-4.1.4/include -I/usr/local/cuda-12.0/targets/x86_64-linux/include

nvcc -o test.cu.o -c test.cu -O3 -std=c++20  -arch=sm_70 --expt-relaxed-constexpr  -I/usr/local/openmpi-4.1.4/include -I/usr/local/cuda-12.0/targets/x86_64-linux/include

mpicxx -o bin/out main.cu.o test.cu.o  -L/usr/local/cuda-12.0/targets/x86_64-linux/lib -lcudart

[Robert_Crovella]

You have device code scattered in two different modules, calling each other. For that code structure, then, you need to enable relocatable device code with device linking.

In a simplistic fashion, something like this:

nvcc -std=c++17 -o test main.cu test.cu -rdc=true

If you require that all the compile and link steps be broken out separately, then you will need to issue the right commands at each step of the way.

Here is an example.

[ak-yoshi]

Great! Thanks to tell me a solution.
As you said, an illegal istruction error was surpressed by compiliation with “-rdc=true” option.
And compilation & linking using nvcc & mpicxx was successed by using command as follows.

# create object file from source file
nvcc -o main.cu.o -c main.cu -O3 -std=c++20  -arch=sm_70 --expt-relaxed-constexpr  --relocatable-device-code=true -c -I/usr/local/openmpi-4.1.4/include -I/usr/local/cuda-12.0/targets/x86_64-linux/include
nvcc -o test.cu.o -c test.cu -O3 -std=c++20  -arch=sm_70 --expt-relaxed-constexpr  --relocatable-device-code=true -c -I/usr/local/openmpi-4.1.4/include -I/usr/local/cuda-12.0/targets/x86_64-linux/include

# create link file from object file
nvcc -o main.cu.l main.cu.o -dlink -arch=sm_70 -lcudart -lcudadevrt
nvcc -o test.cu.l test.cu.o -dlink -arch=sm_70 -lcudart -lcudadevrt

# create executable file from link file & object file
mpicxx -o bin/out main.cu.l test.cu.l main.cu.o test.cu.o  -L/usr/local/cuda-12.0/targets/x86_64-linux/lib -lcudart -lcudadevrt