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
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
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
