<div dir="ltr">All,<div><br></div><div>The similar code looks like below:</div><div>#include <stdio.h><br>#include <iostream><br><br>class AbsBase {<br>public:<br> virtual int f() = 0;<br> virtual void map_in() = 0;<br> virtual void map_out() = 0;<br>};<br><br>class Derived : public AbsBase {<br>private:<br> int a,b,c;<br> //int *arr = new int[100];<br>public:<br> virtual int f() { return a; }<br> int f2() { return a; }<br> void fillarray() {<br> //arr[0] = 1234;<br> a =56;<br> }<br> virtual void map_in() {<br> #pragma omp target enter data map(to:this[0:1])<br> //#pragma omp target enter data map(to:this->arr[0:100])<br> }<br> virtual void map_out() {<br> #pragma omp target exit data map(from:this[0:1])<br> //#pragma omp target exit data map(from:this->arr[0:1])<br> }<br>};<br><br>int main() {<br> AbsBase *absBase[2];<br> absBase[0] = new Derived();<br> std::cout << "pointer is " << absBase[0] << std::endl;<br> absBase[1] = new Derived();<br> static_cast<Derived*>(absBase[0])->fillarray();<br> static_cast<Derived*>(absBase[0])->map_in();<br> static_cast<Derived*>(absBase[1])->map_in();<br> std::cout << sizeof(Derived) << std::endl;<br> int a1, a2, a3;<br>#pragma omp target parallel for map(to: a1,a2,a3) is_device_ptr(absBase)<br> for (int i=0;i<10;i++) {<br> //Derived d1(*static_cast<Derived*>(absBase));<br> //printf("arr[0] is %d\n", d1.f2());<br> printf("%d\n", static_cast<Derived*>(absBase[0])->f2());<br><br> }<br> //static_cast<Derived*>(absBase)->map_out();<br>}<br></div><div><br></div><div>ends up an illegal memory error at run time. Can you spot immediately what is wrong with this code?</div><div>In the debugging log, mapping in for two objects seem to have succeeded onto the device, in areas</div><div>separately clearly.</div><div><br></div></div><br><div class="gmail_quote"><div dir="ltr" class="gmail_attr">On Tue, Jan 21, 2020 at 5:30 PM Itaru Kitayama <<a href="mailto:itaru.kitayama@gmail.com">itaru.kitayama@gmail.com</a>> wrote:<br></div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex"><div dir="ltr">By mapping this pointer first and then followed by the mapping of the array, the sample code worked as expected.<div>Now I am wondering in the map clause, the order of objects appear in the list respected by the compiler? </div></div><br><div class="gmail_quote"><div dir="ltr" class="gmail_attr">On Mon, Jan 20, 2020 at 10:01 AM Itaru Kitayama <<a href="mailto:itaru.kitayama@gmail.com" target="_blank">itaru.kitayama@gmail.com</a>> wrote:<br></div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex"><div dir="ltr">I have updated the code, as a dynamically allocated array is in practice, much important:<div>#include <stdio.h><br>#include <iostream><br><br>class AbsBase {<br>public:<br> virtual int f() = 0;<br> virtual void map_in() = 0;<br> virtual void map_out() = 0;<br>};<br><br>class Derived : public AbsBase {<br>private:<br> int *arr = new int[100];<br> //int arr[100];<br>public:<br> int f() { return arr[0]; }<br> void fillarray() {<br> arr[0] = 1234;<br> }<br> virtual void map_in() {<br> #pragma omp target enter data map(to:this[0:1], this->arr[0:1])<br> }<br> virtual void map_out() {<br> #pragma omp target exit data map(from:this[0:1], this->arr[0:1])<br> }<br>};<br><br>int main() {<br> AbsBase *absBase = new Derived();<br> static_cast<Derived*>(absBase)->fillarray();<br> static_cast<Derived*>(absBase)->map_in();<br>#pragma omp target parallel for<br> for (int i=0;i<10;i++) {<br> Derived d1(*static_cast<Derived*>(absBase));<br> printf("arr[0] is %d\n", d1.f());<br> }<br> static_cast<Derived*>(absBase)->map_out();<br>}<br></div><div><br></div><div>Here are the relevant errors:</div><div>[...]</div><div>Libomptarget --> Looking up mapping(HstPtrBegin=0x0000010023511268, Size=8)...<br>Libomptarget --> Creating new map entry: HstBase=0x0000010023511260, HstBegin=0x0000010023511268, HstEnd=0x0000010023511270, TgtBegin=0x0000110048600000<br>Libomptarget --> There are 8 bytes allocated at target address 0x0000110048600000 - is new<br>Libomptarget --> Looking up mapping(HstPtrBegin=0x0000010023511260, Size=16)...<br>Libomptarget --> WARNING: Pointer is not mapped but section extends into already mapped data<br>Libomptarget --> Explicit extension of mapping is not allowed.<br>Libomptarget --> Call to getOrAllocTgtPtr returned null pointer (device failure or illegal mapping).<br>Libomptarget --> There are 16 bytes allocated at target address 0x0000000000000000 - is new<br>Libomptarget --> Moving 16 bytes (hst:0x0000010023511260) -> (tgt:0x0000000000000000)<br>Target CUDA RTL --> Error when copying data from host to device. Pointers: host = 0x0000010023511260, device = 0x0000000000000000, size = 16<br>Target CUDA RTL --> CUDA error is: invalid argument<br>Libomptarget --> Copying data to device failed.<br>Libomptarget fatal error 1: failure of target construct while offloading is mandatory<br></div><div>[...]</div></div><br><div class="gmail_quote"><div dir="ltr" class="gmail_attr">On Sun, Jan 19, 2020 at 3:26 PM Itaru Kitayama <<a href="mailto:itaru.kitayama@gmail.com" target="_blank">itaru.kitayama@gmail.com</a>> wrote:<br></div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex"><div dir="ltr">Tom,<div><br></div><div>Below builds, but fails at run time:</div><div><br></div><div>#include <stdio.h><br><br>class AbsBase {<br>public:<br> virtual int f() = 0;<br> virtual void map_in() = 0;<br> virtual void map_out() = 0;<br>};<br><br>class Derived : public AbsBase {<br>private:<br> int *arr = new int[100];<br>public:<br> int f() { return arr[0]; }<br> void fillarray() {<br> arr[0] = 1234;<br> }<br> virtual void map_in() {<br> #pragma omp target enter data map(to:this[0:1], this->arr[0:1])<br> }<br> virtual void map_out() {<br> #pragma omp target exit data map(from:this[0:1], this->arr[0:1])<br> }<br>};<br><br>int main() {<br> AbsBase *absBase = new Derived();<br> static_cast<Derived*>(absBase)->fillarray();<br> static_cast<Derived*>(absBase)->map_in();<br><br>#pragma omp target parallel for<br> for (int i=0;i<10;i++) {<br> Derived d1(*static_cast<Derived*>(absBase));<br> printf("arr[0] is %d\n", d1.f());<br> }<br> absBase->map_out();<br>}<br></div><div><br></div><div>What's wrong with my code?</div><div><br></div><div>[kitayama1@juronc11 pcp0151]$ ./a.out<br>Libomptarget --> Loading RTLs...<br>Libomptarget --> Loading library '<a href="http://libomptarget.rtl.ppc64.so" target="_blank">libomptarget.rtl.ppc64.so</a>'...<br>Libomptarget --> Successfully loaded library '<a href="http://libomptarget.rtl.ppc64.so" target="_blank">libomptarget.rtl.ppc64.so</a>'!<br>Libomptarget --> Registering RTL <a href="http://libomptarget.rtl.ppc64.so" target="_blank">libomptarget.rtl.ppc64.so</a> supporting 4 devices!<br>Libomptarget --> Loading library '<a href="http://libomptarget.rtl.x86_64.so" target="_blank">libomptarget.rtl.x86_64.so</a>'...<br>Libomptarget --> Unable to load library '<a href="http://libomptarget.rtl.x86_64.so" target="_blank">libomptarget.rtl.x86_64.so</a>': <a href="http://libomptarget.rtl.x86_64.so" target="_blank">libomptarget.rtl.x86_64.so</a>: cannot open shared object file: No such file or directory!<br>Libomptarget --> Loading library '<a href="http://libomptarget.rtl.cuda.so" target="_blank">libomptarget.rtl.cuda.so</a>'...<br>Target CUDA RTL --> Start initializing CUDA<br>Libomptarget --> Successfully loaded library '<a href="http://libomptarget.rtl.cuda.so" target="_blank">libomptarget.rtl.cuda.so</a>'!<br>Libomptarget --> Registering RTL <a href="http://libomptarget.rtl.cuda.so" target="_blank">libomptarget.rtl.cuda.so</a> supporting 1 devices!<br>Libomptarget --> Loading library '<a href="http://libomptarget.rtl.aarch64.so" target="_blank">libomptarget.rtl.aarch64.so</a>'...<br>Libomptarget --> Unable to load library '<a href="http://libomptarget.rtl.aarch64.so" target="_blank">libomptarget.rtl.aarch64.so</a>': <a href="http://libomptarget.rtl.aarch64.so" target="_blank">libomptarget.rtl.aarch64.so</a>: cannot open shared object file: No such file or directory!<br>Libomptarget --> RTLs loaded!<br>Libomptarget --> Image 0x0000000010001a90 is NOT compatible with RTL <a href="http://libomptarget.rtl.ppc64.so" target="_blank">libomptarget.rtl.ppc64.so</a>!<br>Libomptarget --> Image 0x0000000010001a90 is compatible with RTL <a href="http://libomptarget.rtl.cuda.so" target="_blank">libomptarget.rtl.cuda.so</a>!<br>Libomptarget --> RTL 0x0000010000a6d6d0 has index 0!<br>Libomptarget --> Registering image 0x0000000010001a90 with RTL <a href="http://libomptarget.rtl.cuda.so" target="_blank">libomptarget.rtl.cuda.so</a>!<br>Libomptarget --> Done registering entries!<br>Libomptarget --> Call to omp_get_num_devices returning 1<br>Libomptarget --> Default TARGET OFFLOAD policy is now mandatory (devices were found)<br>Libomptarget --> Entering data begin region for device -1 with 3 mappings<br>Libomptarget --> Use default device id 0<br>Libomptarget --> Checking whether device 0 is ready.<br>Libomptarget --> Is the device 0 (local ID 0) initialized? 0<br>Target CUDA RTL --> Init requires flags to 1<br>Target CUDA RTL --> Getting device 0<br>Target CUDA RTL --> Max CUDA blocks per grid 2147483647 exceeds the hard team limit 65536, capping at the hard limit<br>Target CUDA RTL --> Using 1024 CUDA threads per block<br>Target CUDA RTL --> Max number of CUDA blocks 65536, threads 1024 & warp size 32<br>Target CUDA RTL --> Default number of teams set according to library's default 128<br>Target CUDA RTL --> Default number of threads set according to library's default 128<br>Libomptarget --> Device 0 is ready to use.<br>Target CUDA RTL --> Load data from image 0x0000000010001a90<br>Target CUDA RTL --> CUDA module successfully loaded!<br>Target CUDA RTL --> Entry point 0x0000000000000000 maps to __omp_offloading_30_809bddc5_main_l31 (0x00001100003a99d0)<br>Target CUDA RTL --> Sending global device environment data 4 bytes<br>Libomptarget --> Entry 0: Base=0x0000010000ac1260, Begin=0x0000010000ac1268, Size=8, Type=0x20<br>Libomptarget --> Entry 1: Base=0x0000010000ac1260, Begin=0x0000010000ac1260, Size=16, Type=0x1000000000001<br>Libomptarget --> Entry 2: Base=0x0000010000ac1268, Begin=0x0000010000ac1280, Size=4, Type=0x1000000000011<br>Libomptarget --> Looking up mapping(HstPtrBegin=0x0000010000ac1268, Size=8)...<br>Libomptarget --> Creating new map entry: HstBase=0x0000010000ac1260, HstBegin=0x0000010000ac1268, HstEnd=0x0000010000ac1270, TgtBegin=0x0000110048600000<br>Libomptarget --> There are 8 bytes allocated at target address 0x0000110048600000 - is new<br>Libomptarget --> Looking up mapping(HstPtrBegin=0x0000010000ac1260, Size=16)...<br>Libomptarget --> WARNING: Pointer is not mapped but section extends into already mapped data<br>Libomptarget --> Explicit extension of mapping is not allowed.<br>Libomptarget --> Call to getOrAllocTgtPtr returned null pointer (device failure or illegal mapping).<br>Libomptarget --> There are 16 bytes allocated at target address 0x0000000000000000 - is new<br>Libomptarget --> Moving 16 bytes (hst:0x0000010000ac1260) -> (tgt:0x0000000000000000)<br>Target CUDA RTL --> Error when copying data from host to device. Pointers: host = 0x0000010000ac1260, device = 0x0000000000000000, size = 16<br>Target CUDA RTL --> CUDA error is: invalid argument<br>Libomptarget --> Copying data to device failed.<br>Libomptarget fatal error 1: failure of target construct while offloading is mandatory<br>Libomptarget --> Unloading target library!<br>Libomptarget --> Image 0x0000000010001a90 is compatible with RTL 0x0000010000a6d6d0!<br>Libomptarget --> Unregistered image 0x0000000010001a90 from RTL 0x0000010000a6d6d0!<br>Libomptarget --> Done unregistering images!<br>Libomptarget --> Removing translation table for descriptor 0x000000001007dff8<br>Libomptarget --> Done unregistering library!<br></div></div><br><div class="gmail_quote"><div dir="ltr" class="gmail_attr">On Sat, Jan 18, 2020 at 11:49 AM Scogland, Tom <<a href="mailto:scogland1@llnl.gov" target="_blank">scogland1@llnl.gov</a>> wrote:<br></div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex">
<div dir="ltr">
<div>That's the map of "*this" talking, which we made possible in a newer standard than is implemented. If you change it over to mapping either this by an array section or a reference to *this it should work. The general pattern is to use that kind of mapping
method to get the correct static type at the point of the map.</div>
<div style="font-family:Consolas,Courier,monospace;font-size:12pt;color:rgb(0,0,0);background-color:rgb(255,255,255)">
<br>
</div>
<div style="font-family:Consolas,Courier,monospace;font-size:12pt;color:rgb(0,0,0);background-color:rgb(255,255,255)">
-Tom<br>
</div>
<div id="gmail-m_9059459022016626655gmail-m_9209131696094346748gmail-m_-8815972695588059338gmail-m_5502872398520314696appendonsend"></div>
<hr style="display:inline-block;width:98%">
<div id="gmail-m_9059459022016626655gmail-m_9209131696094346748gmail-m_-8815972695588059338gmail-m_5502872398520314696divRplyFwdMsg" dir="ltr"><font face="Calibri, sans-serif" style="font-size:11pt" color="#000000"><b>From:</b> Itaru Kitayama <<a href="mailto:itaru.kitayama@gmail.com" target="_blank">itaru.kitayama@gmail.com</a>><br>
<b>Sent:</b> Friday, January 17, 2020 5:42 PM<br>
<b>To:</b> Doerfert, Johannes <<a href="mailto:jdoerfert@anl.gov" target="_blank">jdoerfert@anl.gov</a>><br>
<b>Cc:</b> Alexey Bataev <<a href="mailto:a.bataev@hotmail.com" target="_blank">a.bataev@hotmail.com</a>>; openmp-dev <<a href="mailto:openmp-dev@lists.llvm.org" target="_blank">openmp-dev@lists.llvm.org</a>>; Scogland, Tom <<a href="mailto:scogland1@llnl.gov" target="_blank">scogland1@llnl.gov</a>><br>
<b>Subject:</b> Re: [Openmp-dev] Target CUDA RTL --> CUDA error is: an illegal memory access was encountered</font>
<div> </div>
</div>
<div>
<div dir="ltr">Johannes,
<div>Tested with Trunk on POWER8:</div>
<div><br>
</div>
<div>test5.cpp:20:40: error: expected expression containing only member accesses and/or array sections based on named variables<br>
#pragma omp target enter data map(to:*this, this->arr[0:100])<br>
^~~~~<br>
test5.cpp:23:40: error: expected expression containing only member accesses and/or array sections based on named variables<br>
#pragma omp target enter data map(to:*this, this->arr[0:100])<br>
</div>
<div><br>
</div>
</div>
<br>
<div>
<div dir="ltr">On Sat, Jan 18, 2020 at 4:50 AM Doerfert, Johannes <<a href="mailto:jdoerfert@anl.gov" target="_blank">jdoerfert@anl.gov</a>> wrote:<br>
</div>
<blockquote style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex">
Check the attached email from Tom for a way to make this work. Let me know what you think<br>
<br>
<br>
<br>
---------------------------------------<br>
Johannes Doerfert<br>
Researcher<br>
<br>
Argonne National Laboratory<br>
Lemont, IL 60439, USA<br>
<br>
<a href="mailto:jdoerfert@anl.gov" target="_blank">jdoerfert@anl.gov</a><br>
<br>
________________________________________<br>
From: Alexey Bataev <<a href="mailto:a.bataev@hotmail.com" target="_blank">a.bataev@hotmail.com</a>><br>
Sent: Friday, January 17, 2020 03:33<br>
To: Itaru Kitayama<br>
Cc: Doerfert, Johannes; openmp-dev<br>
Subject: Re: [Openmp-dev] Target CUDA RTL --> CUDA error is: an illegal memory access was encountered<br>
<br>
The compiler does not work the way you want. It is not an issue of OpenMP but of C++ itself.<br>
<br>
Best regards,<br>
Alexey Bataev<br>
<br>
16 янв. 2020 г., в 23:20, Itaru Kitayama <<a href="mailto:itaru.kitayama@gmail.com" target="_blank">itaru.kitayama@gmail.com</a>> написал(а):<br>
<br>
<br>
#include <stdio.h><br>
<br>
class AbsBase {<br>
public:<br>
virtual int f() = 0;<br>
};<br>
<br>
class Derived : public AbsBase {<br>
private:<br>
//int *arr = new int[100];<br>
int arr[100];<br>
public:<br>
int f() { return arr[0]; }<br>
void fillarray() {<br>
arr[0] = 1234;<br>
}<br>
};<br>
<br>
int main() {<br>
AbsBase *absBase = new Derived();<br>
Derived *p = static_cast<Derived*>(absBase);<br>
p->fillarray();<br>
#pragma omp target parallel for map(to: p[0:1])<br>
for (int i=0;i<10;i++) {<br>
//Derived d1(*static_cast<Derived*>(absBase));<br>
Derived d1(*p);<br>
printf("arr[0] is %d\n", d1.f());<br>
}<br>
}<br>
<br>
Above gives me what I wanted to see, but I would like to avoid doing a cast on a pointer<br>
to an abstract base class, but the instantiation is done by the derived class.<br>
<br>
<br>
On Fri, Jan 17, 2020 at 6:07 PM Alexey Bataev <<a href="mailto:a.bataev@hotmail.com" target="_blank">a.bataev@hotmail.com</a><mailto:<a href="mailto:a.bataev@hotmail.com" target="_blank">a.bataev@hotmail.com</a>>> wrote:<br>
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.<br>
<br>
Best regards,<br>
Alexey Bataev<br>
<br>
16 янв. 2020 г., в 22:58, Itaru Kitayama <<a href="mailto:itaru.kitayama@gmail.com" target="_blank">itaru.kitayama@gmail.com</a><mailto:<a href="mailto:itaru.kitayama@gmail.com" target="_blank">itaru.kitayama@gmail.com</a>>> написал(а):<br>
<br>
<br>
If I change class Derived's private data to an array of 100 integers, it executes, but<br>
the first element of the array is reported as 0, instead of 1234. Shouldn't bitwise copy<br>
work when mapping the pointer to AbsBase class like an array of one?<br>
<br>
On Wed, Jan 15, 2020 at 2:32 PM Alexey Bataev <<a href="mailto:a.bataev@hotmail.com" target="_blank">a.bataev@hotmail.com</a><mailto:<a href="mailto:a.bataev@hotmail.com" target="_blank">a.bataev@hotmail.com</a>>> wrote:<br>
And it should fail since you're mapping non trivially copyable type.<br>
<br>
Best regards,<br>
Alexey Bataev<br>
<br>
14 янв. 2020 г., в 19:14, Itaru Kitayama <<a href="mailto:itaru.kitayama@gmail.com" target="_blank">itaru.kitayama@gmail.com</a><mailto:<a href="mailto:itaru.kitayama@gmail.com" target="_blank">itaru.kitayama@gmail.com</a>>> написал(а):<br>
<br>
<br>
This piece of C++ program execution fails at run time.<br>
<br>
#include <stdio.h><br>
<br>
class AbsBase {<br>
public:<br>
virtual int f() = 0;<br>
};<br>
<br>
class Derived : public AbsBase {<br>
private:<br>
int *arr = new int[100];<br>
public:<br>
int f() { return arr[0]; }<br>
void fillarray() {<br>
arr[0] = 1234;<br>
}<br>
};<br>
<br>
int main() {<br>
AbsBase *absBase = new Derived();<br>
static_cast<Derived*>(absBase)->fillarray();<br>
#pragma omp target parallel for map(to: absBase[0:1])<br>
for (int i=0;i<10;i++) {<br>
Derived d1(*static_cast<Derived*>(absBase));<br>
printf("arr[0] is %d\n", d1.f());<br>
}<br>
}<br>
<br>
On Sun, Jan 12, 2020 at 1:45 PM Doerfert, Johannes <<a href="mailto:jdoerfert@anl.gov" target="_blank">jdoerfert@anl.gov</a><mailto:<a href="mailto:jdoerfert@anl.gov" target="_blank">jdoerfert@anl.gov</a>>> wrote:<br>
On 01/12, Itaru Kitayama wrote:<br>
> Do you guys have a timeframe for that feature<br>
> Implemented?<br>
<br>
I do not and I don't know anyone who will drive this right now.<br>
<br>
As mentioned before, you should be able to "move/copy-create" the<br>
elements on the device in order to use virtual functions.<br>
<br>
<br>
> On Sun, Jan 12, 2020 at 12:51 Doerfert, Johannes <<a href="mailto:jdoerfert@anl.gov" target="_blank">jdoerfert@anl.gov</a><mailto:<a href="mailto:jdoerfert@anl.gov" target="_blank">jdoerfert@anl.gov</a>>> wrote:<br>
><br>
> > On 01/11, Alexey Bataev via Openmp-dev wrote:<br>
> > > Virtual functions are not supported.<br>
> ><br>
> > Not yet ;).<br>
> ><br>
> > We'll get it with 5.1 so we might actually implement it soon. Till then,<br>
> > you have to create the object on the device you call the virtual<br>
> > function.<br>
> ><br>
<br>
--<br>
<br>
Johannes Doerfert<br>
Researcher<br>
<br>
Argonne National Laboratory<br>
Lemont, IL 60439, USA<br>
<br>
<a href="mailto:jdoerfert@anl.gov" target="_blank">jdoerfert@anl.gov</a><mailto:<a href="mailto:jdoerfert@anl.gov" target="_blank">jdoerfert@anl.gov</a>><br>
</blockquote>
</div>
</div>
</div>
</blockquote></div>
</blockquote></div>
</blockquote></div>
</blockquote></div>