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

[Itaru Kitayama via Openmp-dev]

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/20200116/e5ea3923/attachment-0001.html>