<div dir="ltr">Alexey,<div>Builds without a single warning, but run time error.</div><div><br></div><div>[kitayama1@juronc15 pcp0151]$ ./a.out<br>Libomptarget --> Loading RTLs...<br>Libomptarget --> Loading library '<a href="http://libomptarget.rtl.ppc64.so">libomptarget.rtl.ppc64.so</a>'...<br>Libomptarget --> Successfully loaded library '<a href="http://libomptarget.rtl.ppc64.so">libomptarget.rtl.ppc64.so</a>'!<br>Libomptarget --> Registering RTL <a href="http://libomptarget.rtl.ppc64.so">libomptarget.rtl.ppc64.so</a> supporting 4 devices!<br>Libomptarget --> Loading library '<a href="http://libomptarget.rtl.x86_64.so">libomptarget.rtl.x86_64.so</a>'...<br>Libomptarget --> Unable to load library '<a href="http://libomptarget.rtl.x86_64.so">libomptarget.rtl.x86_64.so</a>': <a href="http://libomptarget.rtl.x86_64.so">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">libomptarget.rtl.cuda.so</a>'...<br>Target CUDA RTL --> Start initializing CUDA<br>Libomptarget --> Successfully loaded library '<a href="http://libomptarget.rtl.cuda.so">libomptarget.rtl.cuda.so</a>'!<br>Libomptarget --> Registering RTL <a href="http://libomptarget.rtl.cuda.so">libomptarget.rtl.cuda.so</a> supporting 1 devices!<br>Libomptarget --> Loading library '<a href="http://libomptarget.rtl.aarch64.so">libomptarget.rtl.aarch64.so</a>'...<br>Libomptarget --> Unable to load library '<a href="http://libomptarget.rtl.aarch64.so">libomptarget.rtl.aarch64.so</a>': <a href="http://libomptarget.rtl.aarch64.so">libomptarget.rtl.aarch64.so</a>: cannot open shared object file: No such file or directory!<br>Libomptarget --> RTLs loaded!<br>Libomptarget --> Image 0x0000000010001c90 is NOT compatible with RTL <a href="http://libomptarget.rtl.ppc64.so">libomptarget.rtl.ppc64.so</a>!<br>Libomptarget --> Image 0x0000000010001c90 is compatible with RTL <a href="http://libomptarget.rtl.cuda.so">libomptarget.rtl.cuda.so</a>!<br>Libomptarget --> RTL 0x00000100153fd6d0 has index 0!<br>Libomptarget --> Registering image 0x0000000010001c90 with RTL <a href="http://libomptarget.rtl.cuda.so">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 0x0000000010001c90<br>Target CUDA RTL --> CUDA module successfully loaded!<br>Target CUDA RTL --> Entry point 0x0000000000000000 maps to __omp_offloading_30_804572b9_main_l32 (0x00001100003ad4a0)<br>Target CUDA RTL --> Sending global device environment data 4 bytes<br>Libomptarget --> Entry 0: Base=0x0000010015451260, Begin=0x0000010015451268, Size=8, Type=0x20<br>Libomptarget --> Entry 1: Base=0x0000010015451260, Begin=0x0000010015451260, Size=16, Type=0x1000000000001<br>Libomptarget --> Entry 2: Base=0x0000010015451268, Begin=0x0000010015451280, Size=4, Type=0x1000000000011<br>Libomptarget --> Looking up mapping(HstPtrBegin=0x0000010015451268, Size=8)...<br>Libomptarget --> Creating new map entry: HstBase=0x0000010015451260, HstBegin=0x0000010015451268, HstEnd=0x0000010015451270, TgtBegin=0x0000110048600000<br>Libomptarget --> There are 8 bytes allocated at target address 0x0000110048600000 - is new<br>Libomptarget --> Looking up mapping(HstPtrBegin=0x0000010015451260, 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:0x0000010015451260) -> (tgt:0x0000000000000000)<br>Target CUDA RTL --> Error when copying data from host to device. Pointers: host = 0x0000010015451260, 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 0x0000000010001c90 is compatible with RTL 0x00000100153fd6d0!<br>Libomptarget --> Unregistered image 0x0000000010001c90 from RTL 0x00000100153fd6d0!<br>Libomptarget --> Done unregistering images!<br>Libomptarget --> Removing translation table for descriptor 0x00000000100825f8<br>Libomptarget --> Done unregistering library!<br></div></div><br><div class="gmail_quote"><div dir="ltr" class="gmail_attr">On Mon, Jan 20, 2020 at 10:16 AM Alexey Bataev <<a href="mailto:a.bataev@hotmail.com">a.bataev@hotmail.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="auto">
Itaru, try #pragma omp target parallel for is_device_ptr(absBase)<br>
<br>
<div dir="ltr">Best regards,
<div>Alexey Bataev</div>
</div>
<div dir="ltr"><br>
<blockquote type="cite">19 янв. 2020 г., в 20:01, Itaru Kitayama <<a href="mailto:itaru.kitayama@gmail.com" target="_blank">itaru.kitayama@gmail.com</a>> написал(а):<br>
<br>
</blockquote>
</div>
<blockquote type="cite">
<div dir="ltr">
<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_-442678189435418515gmail-m_-8815972695588059338gmail-m_5502872398520314696appendonsend"></div>
<hr style="display:inline-block;width:98%">
<div id="gmail-m_-442678189435418515gmail-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>
</div>
</blockquote>
</div>
</blockquote></div>