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

[Itaru Kitayama via Openmp-dev]

Alexey,
Thanks! It builds without the -fno-rtti.

On Fri, Jan 17, 2020 at 8:19 AM Alexey Bataev <a.bataev at outlook.com> wrote:

> Hi Itaru, update the compiler and try to compile your example without
> -fno-rtti. It should compile fine.
>
> -------------
> Best regards,
> Alexey Bataev
>
> 15.01.2020 3:36 PM, Itaru Kitayama пишет:
>
> Yes, with -fno-rtti it does.
>
> On Thu, Jan 16, 2020 at 10:34 AM Doerfert, Johannes <jdoerfert at anl.gov>
> wrote:
>
>> This works with -fno-rtti, correct? It should work either way but I am
>> just curious.
>>
>> On 01/16, Itaru Kitayama wrote:
>> > 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
>> > >>>
>> > >>
>>
>> --
>>
>> 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/20200117/7392ecaf/attachment-0001.html>