Target CUDA RTL --> CUDA error is: an illegal memory access was encountered

Hi,
C++ sample code like below:

class AbsBase {
public:
virtual int f() = 0;
};

class Derived : public AbsBase {
private:
int *arr = new int[100];
public:
int f() { return arr[0]; }
void fillarray() {
arr[0] = 1234;
}
};

int main() {
AbsBase absBase = new Derived();
#pragma omp target enter data map(to: absBase[0:1])
#pragma omp target parallel for map(to: absBase)
for (int i=0;i<10;i++) {
static_cast<Derived
>(absBase)->f();
}
}

if it is built with the offloading support and executed on POWER8, I observe illegal memory access. Am I programming incorrectly or the possible bug in the
current trunk of Clang?

Virtual functions are not supported.

Best regards,
Alexey Bataev

Not yet ;).

We'll get it with 5.1 so we might actually implement it soon. Till then,
you have to create the object on the device you call the virtual
function.

Do you guys have a timeframe for that feature

Implemented?

Do you guys have a timeframe for that feature
Implemented?

I do not and I don't know anyone who will drive this right now.

As mentioned before, you should be able to "move/copy-create" the
elements on the device in order to use virtual functions.

This doesn’t get compile:

class AbsBase {
virtual int f() = 0;
};
class Derived : public AbsBase {
private:
int a;
public:
int f() { a = 1; return a; }
};

int main() {
#pragma omp target parallel for
for (int i=0;i<10;i++) {
Derived d;
}
}

$ clang++ -g -fopenmp -fopenmp-targets=nvptx64 test2.cpp
nvlink error : Undefined reference to ‘__cxa_pure_virtual’ in ‘/tmp/test2-d564cf.cubin’
nvlink error : Undefined reference to ‘_ZTVN10__cxxabiv117__class_type_infoE’ in ‘/tmp/test2-d564cf.cubin’
nvlink error : Undefined reference to ‘_ZTVN10__cxxabiv120__si_class_type_infoE’ in ‘/tmp/test2-d564cf.cubin’
clang-10: error: nvlink command failed with exit code 255 (use -v to see invocation)

Can you try it with -fno-rtti.

[kitayama1@juronc12 ~]$ clang++ -fno-rtti -g -fopenmp -fopenmp-targets=nvptx64 test2.cpp
nvlink error : Undefined reference to ‘__cxa_pure_virtual’ in ‘/tmp/test2-ed87e2.cubin’
clang-10: error: nvlink command failed with exit code 255 (use -v to see invocation)

Not good, this should work. Could you open a bug on this?

Alexey,

Can we have build time error messages?
Clang does not complain at this moment.

Alexey - then, a warning saying it is going to fail at run time or something if the developer is calling a virtual function in the target region.

This piece of C++ program execution fails at run time.

#include <stdio.h>

class AbsBase {
public:
virtual int f() = 0;
};

class Derived : public AbsBase {
private:
int *arr = new int[100];
public:
int f() { return arr[0]; }
void fillarray() {
arr[0] = 1234;
}
};

int main() {
AbsBase absBase = new Derived();
static_cast<Derived
>(absBase)->fillarray();
#pragma omp target parallel for map(to: absBase[0:1])
for (int i=0;i<10;i++) {
Derived d1(static_cast<Derived>(absBase));
printf(“arr[0] is %d\n”, d1.f());
}
}

And it should fail since you’re mapping non trivially copyable type.

Best regards,
Alexey Bataev

#include <stdio.h>

class Base {
public:
virtual int f() { return 2;}
};
class D : public Base {
public:
int f() { return 1234; }
};

int main() {
D *d = new D();
#pragma omp target parallel for map(to: d[0:1])
for (int i=0;i<10;i++) {
D d1(*d);
printf(“arr[0] is %d\n”, d1.f());
}
}

The above works even though an instance of class D is not a trivially mappable type.

Yes, it works, since you’re actually don’t use mapped data. Instead you create the local one, which can be used directly without any issues. Local objects should work fine,but not the mapped ones.

Best regards,
Alexey Bataev

Today’s Trunk can not build below code:

#include <stdio.h>

class AbsBase {
public:
virtual int f(int&) = 0;
};

class Derived : public AbsBase {
private:
int d_;
public:
int f(int &x) { return d_ = x; }

};

int main() {
#pragma omp target parallel for
for (int i=0;i<10;i++) {
int x=123;
Derived d;
printf("%d\n",d.f(x));
}
}

[kitayama1@juronc12 pcp0151]$ clang++ -g -fopenmp -fopenmp-targets=nvptx64 test2.cpp
nvlink error : Undefined reference to ‘_ZTVN10__cxxabiv117__class_type_infoE’ in ‘/tmp/test2-515e4a.cubin’
nvlink error : Undefined reference to ‘_ZTVN10__cxxabiv120__si_class_type_infoE’ in ‘/tmp/test2-515e4a.cubin’
clang-11: error: nvlink command failed with exit code 255 (use -v to see invocation)

Are we supposed to supply this flag from now on?

This works with -fno-rtti, correct? It should work either way but I am
just curious.

Yes, with -fno-rtti it does.

Johannes,

You weren’t including at all those objects that can’t be mapped trivially for the “move/copy-create” trick? Am I correct?