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

[Doerfert, Johannes via Openmp-dev]

Check the attached email from Tom for a way to make this work. Let me know what you think



---------------------------------------
Johannes Doerfert
Researcher

Argonne National Laboratory
Lemont, IL 60439, USA

jdoerfert at anl.gov

________________________________________
From: Alexey Bataev <a.bataev at hotmail.com>
Sent: Friday, January 17, 2020 03:33
To: Itaru Kitayama
Cc: Doerfert, Johannes; openmp-dev
Subject: Re: [Openmp-dev] Target CUDA RTL --> CUDA error is: an illegal memory access was encountered

The compiler does not work the way you want. It is not an issue of OpenMP but of C++ itself.

Best regards,
Alexey Bataev

16 янв. 2020 г., в 23:20, Itaru Kitayama <itaru.kitayama at gmail.com> написал(а):


#include <stdio.h>

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

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

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

Above gives me what I wanted to see, but I would like to avoid doing a cast on a pointer
to an abstract base class, but the instantiation is done by the derived class.


On Fri, Jan 17, 2020 at 6:07 PM Alexey Bataev <a.bataev at hotmail.com<mailto:a.bataev at hotmail.com>> wrote:
AbsBase is a pointer to a base class? Then not, compiler is not aware that it is a pointer to the derived class and copies the array as an array of base class, nit derived one.

Best regards,
Alexey Bataev

16 янв. 2020 г., в 22:58, Itaru Kitayama <itaru.kitayama at gmail.com<mailto:itaru.kitayama at gmail.com>> написал(а):


If I change class Derived's private data to an array of 100 integers, it executes, but
the first element of the array is reported as 0, instead of 1234. Shouldn't bitwise copy
work when mapping the pointer to AbsBase class like an array of one?

On Wed, Jan 15, 2020 at 2:32 PM Alexey Bataev <a.bataev at hotmail.com<mailto: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<mailto: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<mailto: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<mailto: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<mailto:jdoerfert at anl.gov>
-------------- next part --------------
An embedded message was scrubbed...
From: "Scogland, Tom" <scogland1 at llnl.gov>
Subject: Re: [Openmp-dev] Target CUDA RTL --> CUDA error is: an illegal memory access was encountered
Date: Fri, 17 Jan 2020 19:39:07 +0000
Size: 38262
URL: <http://lists.llvm.org/pipermail/openmp-dev/attachments/20200117/de58d643/attachment-0001.mht>