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

[Doerfert, Johannes via Openmp-dev]

On 01/16, Itaru Kitayama via Openmp-dev wrote:
> Are we supposed to supply this flag from now on?

As a workaround for now maybe but this should work *with RTTI enabled*,
thus without the flag.

Cheers,
  Johannes

 
> On Thu, Jan 16, 2020 at 9:55 AM Alexey Bataev <a.bataev at outlook.com> wrote:
> 
> > Compile without RTTI, -fno-rtti
> >
> > -------------
> > Best regards,
> > Alexey Bataev
> >
> > 15.01.2020 2:52 PM, Itaru Kitayama пишет:
> >
> > 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
> >>>>
> >>>

> _______________________________________________
> Openmp-dev mailing list
> Openmp-dev at lists.llvm.org
> https://lists.llvm.org/cgi-bin/mailman/listinfo/openmp-dev


-- 

Johannes Doerfert
Researcher

Argonne National Laboratory
Lemont, IL 60439, USA

jdoerfert at anl.gov
-------------- next part --------------
A non-text attachment was scrubbed...
Name: signature.asc
Type: application/pgp-signature
Size: 228 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-dev/attachments/20200116/461924df/attachment-0001.sig>