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

Itaru Kitayama via Openmp-dev openmp-dev at lists.llvm.org
Wed Jan 15 16:58:46 PST 2020


Are we supposed to supply this flag from now on?

On Thu, Jan 16, 2020 at 9:55 AM Alexey Bataev <a.bataev at outlook.com> wrote:

> Compile without RTTI, -fno-rtti
>
> -------------
> Best regards,
> Alexey Bataev
>
> 15.01.2020 2:52 PM, Itaru Kitayama пишет:
>
> 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
>>>>
>>>
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/openmp-dev/attachments/20200116/b2c3d383/attachment-0001.html>


More information about the Openmp-dev mailing list