[OpenMP] Add support for close map modifier in Clang

[gtbercea]

Summary

This patch adds support for the close map modifier in Clang.

This ensures that the new map type is marked and passed to the OpenMP runtime appropriately.

Additional regression tests have been merged from patch D55892 (author @saghir).

[ABataev]

lib/CodeGen/CGOpenMPRuntime.cpp

7695

Why?

[gtbercea]

lib/CodeGen/CGOpenMPRuntime.cpp

7695

If the pointee has been mapped as TO/FROM/etc already no need to map it TO/FROM/etc again.

[gtbercea]

• Improve test.

[ABataev]

LG

[Hahnfeld]

There's already D55892 with a better set of tests, including target enter data / target exit data.

[ABataev]

In D65341#1604411, @Hahnfeld wrote:

There's already D55892 with a better set of tests, including target enter data / target exit data.

Better to merge those two patches into one.

[gtbercea]

In D65341#1604440, @ABataev wrote:

In D65341#1604411, @Hahnfeld wrote:

There's already D55892 with a better set of tests, including target enter data / target exit data.

Better to merge those two patches into one.

How would you like me to proceed? It looks like the other patch has been sitting there approved for many many months.

[kkwli0]

In D65341#1604561, @gtbercea wrote:

In D65341#1604440, @ABataev wrote:

In D65341#1604411, @Hahnfeld wrote:

There's already D55892 with a better set of tests, including target enter data / target exit data.

Better to merge those two patches into one.

How would you like me to proceed? It looks like the other patch has been sitting there approved for many many months.

Yes, the patch has been waiting for the runtime implementation first and then commit.

[gtbercea]

In D65341#1604562, @kkwli0 wrote:

In D65341#1604561, @gtbercea wrote:

In D65341#1604440, @ABataev wrote:

In D65341#1604411, @Hahnfeld wrote:

There's already D55892 with a better set of tests, including target enter data / target exit data.

Better to merge those two patches into one.

How would you like me to proceed? It looks like the other patch has been sitting there approved for many many months.

Yes, the patch has been waiting for the runtime implementation first and then commit.

The runtime implementation is here: D65340

[kkwli0]

In D65341#1604562, @kkwli0 wrote:

In D65341#1604561, @gtbercea wrote:

In D65341#1604440, @ABataev wrote:

In D65341#1604411, @Hahnfeld wrote:

There's already D55892 with a better set of tests, including target enter data / target exit data.

Better to merge those two patches into one.

How would you like me to proceed? It looks like the other patch has been sitting there approved for many many months.

Yes, the patch has been waiting for the runtime implementation first and then commit.

I am not sure if it is easier to rebase D55892 and merge this patch to D55892. I am okay to go either way.

[gtbercea]

• Improve test.
• Add tests from previous patch.

[gtbercea]

@Hahnfeld I have merged the tests from the previous patch as per Alexey's suggestion - with minor changes to make them pass. Let me know if this now addresses your previous comments.

[ABataev]

LG

[kkwli0]

Looks fine to me.

[Hahnfeld]

lib/CodeGen/CGOpenMPRuntime.cpp

7695

I'm not sure if there's a test for this. I can't verify right now, so I trust you that it's covered.

[gtbercea]

lib/CodeGen/CGOpenMPRuntime.cpp

7695

I added the test as per Alexey's previous request.

lib/CodeGen/CGOpenMPRuntime.cpp

enum OpenMPOffloadMappingFlags : uint64_t {
     OMP_MAP_RETURN_PARAM = 0x40,
     /// This flag signals that the reference being passed is a pointer to
     /// private data.
     OMP_MAP_PRIVATE = 0x80,
     /// Pass the element to the device by value.
     OMP_MAP_LITERAL = 0x100,
     /// Implicit map
     OMP_MAP_IMPLICIT = 0x200,
+    /// Close is a hint to the runtime to allocate memory close to
+    /// the target device.
+    OMP_MAP_CLOSE = 0x400,
     /// The 16 MSBs of the flags indicate whether the entry is member of some
     /// struct/class.
     OMP_MAP_MEMBER_OF = 0xffff000000000000,
     LLVM_MARK_AS_BITMASK_ENUM(/* LargestFlag = */ OMP_MAP_MEMBER_OF),
   };
 
   /// Class that associates information with a base pointer to be passed to the
   /// runtime library.
OpenMPOffloadMappingFlags getMapTypeBits(
     }
     if (AddPtrFlag)
       Bits |= OMP_MAP_PTR_AND_OBJ;
     if (AddIsTargetParamFlag)
       Bits |= OMP_MAP_TARGET_PARAM;
     if (llvm::find(MapModifiers, OMPC_MAP_MODIFIER_always)
         != MapModifiers.end())
       Bits |= OMP_MAP_ALWAYS;
+    if (llvm::find(MapModifiers, OMPC_MAP_MODIFIER_close)
+        != MapModifiers.end())
+      Bits |= OMP_MAP_CLOSE;
     return Bits;
   }
 
   /// Return true if the provided expression is a final array section. A
   /// final array section, is one whose length can't be proved to be one.
   bool isFinalArraySectionExpression(const Expr *E) const {
     const auto *OASE = dyn_cast<OMPArraySectionExpr>(E);
 
for (; I != CE; ++I) {
           // (there is a set of entries for each capture).
           OpenMPOffloadMappingFlags Flags = getMapTypeBits(
               MapType, MapModifiers, IsImplicit,
               !IsExpressionFirstInfo || RequiresReference,
               IsCaptureFirstInfo && !RequiresReference);
 
           if (!IsExpressionFirstInfo) {
             // If we have a PTR_AND_OBJ pair where the OBJ is a pointer as well,
-            // then we reset the TO/FROM/ALWAYS/DELETE flags.
+            // then we reset the TO/FROM/ALWAYS/DELETE/CLOSE flags.
             if (IsPointer)
               Flags &= ~(OMP_MAP_TO | OMP_MAP_FROM | OMP_MAP_ALWAYS |
-                         OMP_MAP_DELETE);
+                         OMP_MAP_DELETE | OMP_MAP_CLOSE);
 
             if (ShouldBeMemberOf) {
               // Set placeholder value MEMBER_OF=FFFF to indicate that the flag
               // should be later updated with the correct value of MEMBER_OF.
               Flags |= OMP_MAP_MEMBER_OF;
               // From now on, all subsequent PTR_AND_OBJ entries should not be
               // marked as MEMBER_OF.
               ShouldBeMemberOf = false;

test/OpenMP/target_data_codegen.cpp

void no_target_devices(int arg) {
   // CK3-NOT: tgt_target_data_begin
   // CK3: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
   // CK3-NOT: tgt_target_data_end
   // CK3: ret
   #pragma omp target data map(to: arg) if(arg) device(4)
   {++arg;}
 }
 #endif
+///==========================================================================///
+//// RUN: %clang_cc1 -DCK4 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s --check-prefix CK4
+#ifdef CK4
+void test_close_modifier(int arg) {
+  // CK4: private unnamed_addr constant [1 x i64] [i64 1059]
+  #pragma omp target data map(close,tofrom: arg)
+  {++arg;}
+}
+#endif
+///==========================================================================///
+//// RUN: %clang_cc1 -DCK5 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s --check-prefix CK5
+#ifdef CK5
+
+struct S1 {
+  int i;
+};
+struct S2 {
+  S1 s;
+  struct S2 *ps;
+};
+
+void test_close_modifier(int arg) {
+  S2 *ps;
+  // CK5: private unnamed_addr constant [5 x i64] [i64 1059, i64 32, {{.*}}, i64 16, i64 1043]
+  #pragma omp target data map(close,tofrom: arg, ps->ps->ps->ps->s)
+  {
+    ++(arg);
+  }
+}
+#endif
 #endif