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

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

gtbercea created this revision.
gtbercea added reviewers: ABataev, caomhin.
Herald added subscribers: cfe-commits, guansong.
Herald added a reviewer: jdoerfert.
Herald added a project: clang.

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.


Repository:
  rC Clang

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,13 @@
   {++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
 #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;
   }
 
@@ -7686,7 +7692,7 @@
             // then we reset the TO/FROM/ALWAYS/DELETE 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.211969.patch
Type: text/x-patch
Size: 1939 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20190726/00d3b282/attachment.bin>