[Openmp-dev] Target CUDA RTL --> CUDA error is: an illegal memory access was encountered

[Itaru Kitayama via Openmp-dev]

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 at 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)

On Wed, Jan 15, 2020 at 11:31 PM Alexey Bataev <a.bataev at hotmail.com> wrote:

> 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
>
> 14 янв. 2020 г., в 23:17, Itaru Kitayama <itaru.kitayama at gmail.com>
> написал(а):
>
> 
> #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.
>
> On Wed, Jan 15, 2020 at 2:32 PM Alexey Bataev <a.bataev at hotmail.com>
> wrote:
>
>> And it should fail since you're mapping non trivially copyable type.
>>
>> Best regards,
>> Alexey Bataev
>>
>> 14 янв. 2020 г., в 19:14, Itaru Kitayama <itaru.kitayama at gmail.com>
>> написал(а):
>>
>> 
>> 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());
>>         }
>> }
>>
>> On Sun, Jan 12, 2020 at 1:45 PM Doerfert, Johannes <jdoerfert at anl.gov>
>> wrote:
>>
>>> On 01/12, Itaru Kitayama wrote:
>>> > 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.
>>>
>>>
>>> > On Sun, Jan 12, 2020 at 12:51 Doerfert, Johannes <jdoerfert at anl.gov>
>>> wrote:
>>> >
>>> > > On 01/11, Alexey Bataev via Openmp-dev wrote:
>>> > > > Virtual functions are not supported.
>>> > >
>>> > > 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.
>>> > >
>>>
>>> --
>>>
>>> Johannes Doerfert
>>> Researcher
>>>
>>> Argonne National Laboratory
>>> Lemont, IL 60439, USA
>>>
>>> jdoerfert at anl.gov
>>>
>>
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/openmp-dev/attachments/20200116/f75bd249/attachment.html>