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

Doerfert, Johannes via Openmp-dev openmp-dev at lists.llvm.org
Wed Jan 15 17:34:05 PST 2020


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 --------------
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/572efa03/attachment.sig>


More information about the Openmp-dev mailing list