[PATCH] D65341: [OpenMP] Add support for close map modifier in Clang

[Gheorghe-Teodor Bercea via Phabricator via cfe-commits]

gtbercea updated this revision to Diff 212011.
gtbercea added a comment.

- Improve test.


Repository:
  rC Clang

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D65341/new/

https://reviews.llvm.org/D65341

Files:
  lib/CodeGen/CGOpenMPRuntime.cpp
  test/OpenMP/target_data_codegen.cpp


Index: test/OpenMP/target_data_codegen.cpp
===================================================================
--- test/OpenMP/target_data_codegen.cpp
+++ test/OpenMP/target_data_codegen.cpp
@@ -283,4 +283,34 @@
   {++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
Index: lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- lib/CodeGen/CGOpenMPRuntime.cpp
+++ lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7087,6 +7087,9 @@
     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,
@@ -7255,6 +7258,9 @@
     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;
   }
 
@@ -7683,10 +7689,10 @@
 
           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


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D65341.212011.patch
Type: text/x-patch
Size: 2730 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20190726/4ce038af/attachment.bin>