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

[Itaru Kitayama via Openmp-dev]

Johannes,
Tested with Trunk on POWER8:

test5.cpp:20:40: error: expected expression containing only member accesses
and/or array sections based on named variables
                #pragma omp target enter data map(to:*this,
this->arr[0:100])
                                                     ^~~~~
test5.cpp:23:40: error: expected expression containing only member accesses
and/or array sections based on named variables
                #pragma omp target enter data map(to:*this,
this->arr[0:100])


On Sat, Jan 18, 2020 at 4:50 AM Doerfert, Johannes <jdoerfert at anl.gov>
wrote:

> 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 HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/openmp-dev/attachments/20200118/12c11f64/attachment.html>