<div><div dir="auto">Tom,</div></div><div dir="auto">How do I an array of pointers get attached or</div><div dir="auto">translated? The code (NEST) I’m trying to</div><div dir="auto">offload needs to handle an array of pointers.</div><div><br><div class="gmail_quote"><div dir="ltr" class="gmail_attr">On Mon, Feb 10, 2020 at 17:03 Tom Scogland <<a href="mailto:scogland1@llnl.gov">scogland1@llnl.gov</a>> wrote:<br></div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex"><br>
You have a couple of problems. First it's passing absBase through<br>
is_device_ptr despite the fact its a host array. That clause needs to<br>
be removed so the mapping mechanism can do its job. That's a segfault<br>
right off the bat. Second, it's an array of pointers, those pointers<br>
are host pointers, so even if you fix the first issue it's still a<br>
segfault unless those get attached or otherwise translated. Third, the<br>
objects are constructed on the host with host vtables, and we don't have<br>
support for that yet, so that's another segfault. The static casts also<br>
don't do anything. If the overrides were marked final they might, but<br>
as it is they don't.<br>
<br>
This version should be rather closer to working, but I'm not 100% sure<br>
the construction of d1 will succeed, if it does it should all work.<br>
<br>
```<br>
#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 *ab0;<br>
AbsBase *ab1;<br>
ab0 = new Derived();<br>
std::cout << "pointer is " << ab0 << std::endl;<br>
ab1 = new Derived();<br>
ab0->fillarray();<br>
ab0->map_in();<br>
ab1->map_in();<br>
std::cout << sizeof(Derived) << std::endl;<br>
int a1, a2, a3;<br>
#pragma omp target parallel for<br>
for (int i=0;i<10;i++) {<br>
Derived d1(*static_cast<Derived*>(ab0));<br>
//printf("arr[0] is %d\n", d1.f2());<br>
printf("%d\n", d1.f2());<br>
<br>
}<br>
//static_cast<Derived*>(absBase)->map_out();<br>
}<br>
```<br>
<br>
-Tom<br>
<br>
Itaru Kitayama <<a href="mailto:itaru.kitayama@gmail.com" target="_blank">itaru.kitayama@gmail.com</a>> writes:<br>
<br>
> All,<br>
><br>
> The similar code looks like below:<br>
> #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>
><br>
> ends up an illegal memory error at run time. Can you spot immediately what<br>
> is wrong with this code?<br>
> In the debugging log, mapping in for two objects seem to have succeeded<br>
> onto the device, in areas<br>
> separately clearly.<br>
><br>
><br>
> On Tue, Jan 21, 2020 at 5:30 PM Itaru Kitayama <<a href="mailto:itaru.kitayama@gmail.com" target="_blank">itaru.kitayama@gmail.com</a>><br>
> wrote:<br>
><br>
>> By mapping this pointer first and then followed by the mapping of the<br>
>> array, the sample code worked as expected.<br>
>> Now I am wondering in the map clause, the order of objects appear in the<br>
>> list respected by the compiler?<br>
>><br>
>> 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>><br>
>> wrote:<br>
>><br>
>>> I have updated the code, as a dynamically allocated array is in practice,<br>
>>> much important:<br>
>>> #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>
>>><br>
>>> Here are the relevant errors:<br>
>>> [...]<br>
>>> Libomptarget --> Looking up mapping(HstPtrBegin=0x0000010023511268,<br>
>>> Size=8)...<br>
>>> Libomptarget --> Creating new map entry: HstBase=0x0000010023511260,<br>
>>> HstBegin=0x0000010023511268, HstEnd=0x0000010023511270,<br>
>>> TgtBegin=0x0000110048600000<br>
>>> Libomptarget --> There are 8 bytes allocated at target address<br>
>>> 0x0000110048600000 - is new<br>
>>> Libomptarget --> Looking up mapping(HstPtrBegin=0x0000010023511260,<br>
>>> Size=16)...<br>
>>> Libomptarget --> WARNING: Pointer is not mapped but section extends into<br>
>>> already mapped data<br>
>>> Libomptarget --> Explicit extension of mapping is not allowed.<br>
>>> Libomptarget --> Call to getOrAllocTgtPtr returned null pointer (device<br>
>>> failure or illegal mapping).<br>
>>> Libomptarget --> There are 16 bytes allocated at target address<br>
>>> 0x0000000000000000 - is new<br>
>>> Libomptarget --> Moving 16 bytes (hst:0x0000010023511260) -><br>
>>> (tgt:0x0000000000000000)<br>
>>> Target CUDA RTL --> Error when copying data from host to device.<br>
>>> 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<br>
>>> is mandatory<br>
>>> [...]<br>
>>><br>
>>> 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>><br>
>>> wrote:<br>
>>><br>
>>>> Tom,<br>
>>>><br>
>>>> Below builds, but fails at run time:<br>
>>>><br>
>>>> #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>
>>>><br>
>>>> What's wrong with my code?<br>
>>>><br>
>>>> [kitayama1@juronc11 pcp0151]$ ./a.out<br>
>>>> Libomptarget --> Loading RTLs...<br>
>>>> Libomptarget --> Loading library '<a href="http://libomptarget.rtl.ppc64.so" rel="noreferrer" target="_blank">libomptarget.rtl.ppc64.so</a>'...<br>
>>>> Libomptarget --> Successfully loaded library '<a href="http://libomptarget.rtl.ppc64.so" rel="noreferrer" target="_blank">libomptarget.rtl.ppc64.so</a><br>
>>>> '!<br>
>>>> Libomptarget --> Registering RTL <a href="http://libomptarget.rtl.ppc64.so" rel="noreferrer" target="_blank">libomptarget.rtl.ppc64.so</a> supporting 4<br>
>>>> devices!<br>
>>>> Libomptarget --> Loading library '<a href="http://libomptarget.rtl.x86_64.so" rel="noreferrer" target="_blank">libomptarget.rtl.x86_64.so</a>'...<br>
>>>> Libomptarget --> Unable to load library '<a href="http://libomptarget.rtl.x86_64.so" rel="noreferrer" target="_blank">libomptarget.rtl.x86_64.so</a>':<br>
>>>> <a href="http://libomptarget.rtl.x86_64.so" rel="noreferrer" target="_blank">libomptarget.rtl.x86_64.so</a>: cannot open shared object file: No such<br>
>>>> file or directory!<br>
>>>> Libomptarget --> Loading library '<a href="http://libomptarget.rtl.cuda.so" rel="noreferrer" 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" rel="noreferrer" target="_blank">libomptarget.rtl.cuda.so</a>'!<br>
>>>> Libomptarget --> Registering RTL <a href="http://libomptarget.rtl.cuda.so" rel="noreferrer" target="_blank">libomptarget.rtl.cuda.so</a> supporting 1<br>
>>>> devices!<br>
>>>> Libomptarget --> Loading library '<a href="http://libomptarget.rtl.aarch64.so" rel="noreferrer" target="_blank">libomptarget.rtl.aarch64.so</a>'...<br>
>>>> Libomptarget --> Unable to load library '<a href="http://libomptarget.rtl.aarch64.so" rel="noreferrer" target="_blank">libomptarget.rtl.aarch64.so</a>':<br>
>>>> <a href="http://libomptarget.rtl.aarch64.so" rel="noreferrer" target="_blank">libomptarget.rtl.aarch64.so</a>: cannot open shared object file: No such<br>
>>>> file or directory!<br>
>>>> Libomptarget --> RTLs loaded!<br>
>>>> Libomptarget --> Image 0x0000000010001a90 is NOT compatible with RTL<br>
>>>> <a href="http://libomptarget.rtl.ppc64.so" rel="noreferrer" target="_blank">libomptarget.rtl.ppc64.so</a>!<br>
>>>> Libomptarget --> Image 0x0000000010001a90 is compatible with RTL<br>
>>>> <a href="http://libomptarget.rtl.cuda.so" rel="noreferrer" target="_blank">libomptarget.rtl.cuda.so</a>!<br>
>>>> Libomptarget --> RTL 0x0000010000a6d6d0 has index 0!<br>
>>>> Libomptarget --> Registering image 0x0000000010001a90 with RTL<br>
>>>> <a href="http://libomptarget.rtl.cuda.so" rel="noreferrer" 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<br>
>>>> 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<br>
>>>> 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<br>
>>>> size 32<br>
>>>> Target CUDA RTL --> Default number of teams set according to library's<br>
>>>> default 128<br>
>>>> Target CUDA RTL --> Default number of threads set according to library's<br>
>>>> 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<br>
>>>> __omp_offloading_30_809bddc5_main_l31 (0x00001100003a99d0)<br>
>>>> Target CUDA RTL --> Sending global device environment data 4 bytes<br>
>>>> Libomptarget --> Entry 0: Base=0x0000010000ac1260,<br>
>>>> Begin=0x0000010000ac1268, Size=8, Type=0x20<br>
>>>> Libomptarget --> Entry 1: Base=0x0000010000ac1260,<br>
>>>> Begin=0x0000010000ac1260, Size=16, Type=0x1000000000001<br>
>>>> Libomptarget --> Entry 2: Base=0x0000010000ac1268,<br>
>>>> Begin=0x0000010000ac1280, Size=4, Type=0x1000000000011<br>
>>>> Libomptarget --> Looking up mapping(HstPtrBegin=0x0000010000ac1268,<br>
>>>> Size=8)...<br>
>>>> Libomptarget --> Creating new map entry: HstBase=0x0000010000ac1260,<br>
>>>> HstBegin=0x0000010000ac1268, HstEnd=0x0000010000ac1270,<br>
>>>> TgtBegin=0x0000110048600000<br>
>>>> Libomptarget --> There are 8 bytes allocated at target address<br>
>>>> 0x0000110048600000 - is new<br>
>>>> Libomptarget --> Looking up mapping(HstPtrBegin=0x0000010000ac1260,<br>
>>>> Size=16)...<br>
>>>> Libomptarget --> WARNING: Pointer is not mapped but section extends into<br>
>>>> already mapped data<br>
>>>> Libomptarget --> Explicit extension of mapping is not allowed.<br>
>>>> Libomptarget --> Call to getOrAllocTgtPtr returned null pointer (device<br>
>>>> failure or illegal mapping).<br>
>>>> Libomptarget --> There are 16 bytes allocated at target address<br>
>>>> 0x0000000000000000 - is new<br>
>>>> Libomptarget --> Moving 16 bytes (hst:0x0000010000ac1260) -><br>
>>>> (tgt:0x0000000000000000)<br>
>>>> Target CUDA RTL --> Error when copying data from host to device.<br>
>>>> 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<br>
>>>> is mandatory<br>
>>>> Libomptarget --> Unloading target library!<br>
>>>> Libomptarget --> Image 0x0000000010001a90 is compatible with RTL<br>
>>>> 0x0000010000a6d6d0!<br>
>>>> Libomptarget --> Unregistered image 0x0000000010001a90 from RTL<br>
>>>> 0x0000010000a6d6d0!<br>
>>>> Libomptarget --> Done unregistering images!<br>
>>>> Libomptarget --> Removing translation table for descriptor<br>
>>>> 0x000000001007dff8<br>
>>>> Libomptarget --> Done unregistering library!<br>
>>>><br>
>>>> On Sat, Jan 18, 2020 at 11:49 AM Scogland, Tom <<a href="mailto:scogland1@llnl.gov" target="_blank">scogland1@llnl.gov</a>><br>
>>>> wrote:<br>
>>>><br>
>>>>> That's the map of "*this" talking, which we made possible in a newer<br>
>>>>> standard than is implemented. If you change it over to mapping either this<br>
>>>>> by an array section or a reference to *this it should work. The general<br>
>>>>> pattern is to use that kind of mapping method to get the correct static<br>
>>>>> type at the point of the map.<br>
>>>>><br>
>>>>> -Tom<br>
>>>>> ------------------------------<br>
>>>>> *From:* Itaru Kitayama <<a href="mailto:itaru.kitayama@gmail.com" target="_blank">itaru.kitayama@gmail.com</a>><br>
>>>>> *Sent:* Friday, January 17, 2020 5:42 PM<br>
>>>>> *To:* Doerfert, Johannes <<a href="mailto:jdoerfert@anl.gov" target="_blank">jdoerfert@anl.gov</a>><br>
>>>>> *Cc:* Alexey Bataev <<a href="mailto:a.bataev@hotmail.com" target="_blank">a.bataev@hotmail.com</a>>; openmp-dev <<br>
>>>>> <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>
>>>>> *Subject:* Re: [Openmp-dev] Target CUDA RTL --> CUDA error is: an<br>
>>>>> illegal memory access was encountered<br>
>>>>><br>
>>>>> Johannes,<br>
>>>>> Tested with Trunk on POWER8:<br>
>>>>><br>
>>>>> test5.cpp:20:40: error: expected expression containing only member<br>
>>>>> accesses and/or array sections based on named variables<br>
>>>>> #pragma omp target enter data map(to:*this,<br>
>>>>> this->arr[0:100])<br>
>>>>> ^~~~~<br>
>>>>> test5.cpp:23:40: error: expected expression containing only member<br>
>>>>> accesses and/or array sections based on named variables<br>
>>>>> #pragma omp target enter data map(to:*this,<br>
>>>>> this->arr[0:100])<br>
>>>>><br>
>>>>><br>
>>>>> On Sat, Jan 18, 2020 at 4:50 AM Doerfert, Johannes <<a href="mailto:jdoerfert@anl.gov" target="_blank">jdoerfert@anl.gov</a>><br>
>>>>> wrote:<br>
>>>>><br>
>>>>> Check the attached email from Tom for a way to make this work. Let me<br>
>>>>> 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<br>
>>>>> memory access was encountered<br>
>>>>><br>
>>>>> The compiler does not work the way you want. It is not an issue of<br>
>>>>> 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>
>>>>> <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<br>
>>>>> cast on a pointer<br>
>>>>> to an abstract base class, but the instantiation is done by the derived<br>
>>>>> 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><br>
>>>>> <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<br>
>>>>> that it is a pointer to the derived class and copies the array as an array<br>
>>>>> 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><br>
>>>>> <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,<br>
>>>>> it executes, but<br>
>>>>> the first element of the array is reported as 0, instead of 1234.<br>
>>>>> 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><br>
>>>>> <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><br>
>>>>> <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><br>
>>>>> <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><br>
>>>>> <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<br>
>>>>> 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>
>>>>><br>
>>>>><br>
</blockquote></div></div>