6.7. C++ Virtual Functions#

The 5.2 OpenMP Specification clarified restrictions on the use of polymorphic classes and virtual functions when used within target regions. The following example identifies problem cases in which the restrictions are not followed (for Unified Shared Memory, as prescribed by the requires directive).

The first section illustrates the restriction that when mapping an object for the first time, the static and dynamic types must match.

For the first target region the behavior of the implicit map of ar is not specified - its static type (A) doesn’t match its dynamic type (D).
Hence access to the virtual functions is undefined. However, the second target region can access D::vf() since the object to which ap points is not mapped and therefore the restriction does not apply.

The second section illustrates the restriction:

“Invoking a virtual member function of an object on a device other than the device on which the object was constructed results in unspecified behavior, unless the object is accessible and was constructed on the host device.”

An instantiation of a polymorphic class (A) occurs in the target region, and access of its virtual function is incorrectly attempted on the host (another device). However, once the object is deleted on the target device and instantiated on the host, access within the next target region is permitted.

//%compiler: clang
//%cflags: -fopenmp

/*
* @@name:       virtual_functions.1
* @@type:       C++
* @@compilable: yes
* @@linkable:   no
* @@expect:     success
* @@version:	omp_5.2
*/
#include <iostream>
#pragma omp requires unified_shared_memory

#pragma omp begin declare target
class A {  
  public:
   virtual void vf()  { std::cout << "In A\n"; }
};

class D: public A {
  public:
   void vf() override { std::cout << "In D\n"; }
};
#pragma omp end declare target

int main(){

   // Section 1 --------------------------------------------------------
   D d;               // D derives from A, and A::vf() is virtual
   A &ar = d;         // reference to Derived object d

   #pragma omp target // implicit map of ar is illegal here
   {
      ar.vf();        // unspecified whether A::vf() or D::vf() is called
   }
   
   A *ap = &d;        // pointer to derived object d
   #pragma omp target // No need for mapping with Unified Share Memory
   {                  // implicit ap[:0] map is fine 
      ap->vf();       // calls D::vf()
   }

   // Section 2 --------------------------------------------------------
   ap = nullptr;
   #pragma omp target map(ap)
   {
        ap = new A();
   }
   
   ap->vf();     // illegal

   #pragma omp target
   {
      delete ap;
   }
   ap = new A();
   #pragma omp target  // No need for mapping with Unified Share Memory
   {
      ap->vf();  // ok
   }

   return 0;
}