gtbercea created this revision.
When a scalar argument is explicitly mapped, the device kernel needs to be
passed that argument by reference.
This is a partial fix to this bug
<https://bugs.llvm.org//show_bug.cgi?id=31920>.
The full fix requires data sharing support which will land in future patches.
Repository:
rL LLVM
https://reviews.llvm.org/D29905
Files:
lib/Sema/SemaOpenMP.cpp
test/OpenMP/target_map_codegen.cpp
Index: test/OpenMP/target_map_codegen.cpp
===================================================================
--- test/OpenMP/target_map_codegen.cpp
+++ test/OpenMP/target_map_codegen.cpp
@@ -4756,3 +4756,24 @@
}
#endif
#endif
+
+///==========================================================================///
+// RUN: %clang -DCK30 -fopenmp -S -emit-llvm -fopenmp
-fopenmp-targets=nvptx64-nvidia-cuda %s -o - 2>&1 | FileCheck
-check-prefix=CK30 %s
+#ifdef CK30
+#include <stdio.h>
+
+void target_maps_parallel_integer(int a){
+ int ParamToKernel = a;
+ #pragma omp target map(tofrom:ParamToKernel)
+ {
+ printf("Print ParamToKernel = %d\n", ParamToKernel);
+ #pragma omp parallel
+ {
+ printf("Do nothing\n");
+ }
+ }
+}
+
+// CK30: define void
@__omp_offloading_802_{{.*}}target_maps_parallel_integer{{.*}}l{{.*}}(i32*
dereferenceable(4) %ParamToKernel) #0 {
+
+#endif
Index: lib/Sema/SemaOpenMP.cpp
===================================================================
--- lib/Sema/SemaOpenMP.cpp
+++ lib/Sema/SemaOpenMP.cpp
@@ -355,6 +355,27 @@
return false;
}
+ /// Do the check specified in \a Check to all component lists at a given
level
+ /// and return true if any issue is found.
+ bool checkMappableExprComponentListsForDeclAtLevel(
+ ValueDecl *VD, unsigned Level,
+ const llvm::function_ref<
+ bool(OMPClauseMappableExprCommon::MappableExprComponentListRef,
+ OpenMPClauseKind)> &Check) {
+ auto StartI = std::next(Stack.begin());
+ auto EndI = Stack.end();
+ if (std::distance(StartI, EndI) <= (int)Level)
+ return false;
+ std::advance(StartI, Level);
+
+ auto MI = StartI->MappedExprComponents.find(VD);
+ if (MI != StartI->MappedExprComponents.end())
+ for (auto &L : MI->second.Components)
+ if (Check(L, MI->second.Kind))
+ return true;
+ return false;
+ }
+
/// Create a new mappable expression component list associated with a given
/// declaration and initialize it with the provided list of components.
void addMappableExpressionComponents(
@@ -912,9 +933,8 @@
bool IsVariableUsedInMapClause = false;
bool IsVariableAssociatedWithSection = false;
- DSAStack->checkMappableExprComponentListsForDecl(
- D, /*CurrentRegionOnly=*/true,
- [&](OMPClauseMappableExprCommon::MappableExprComponentListRef
+ DSAStack->checkMappableExprComponentListsForDeclAtLevel(
+ D, Level, [&](OMPClauseMappableExprCommon::MappableExprComponentListRef
MapExprComponents,
OpenMPClauseKind WhereFoundClauseKind) {
// Only the map clause information influences how a variable is
Index: test/OpenMP/target_map_codegen.cpp
===================================================================
--- test/OpenMP/target_map_codegen.cpp
+++ test/OpenMP/target_map_codegen.cpp
@@ -4756,3 +4756,24 @@
}
#endif
#endif
+
+///==========================================================================///
+// RUN: %clang -DCK30 -fopenmp -S -emit-llvm -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda %s -o - 2>&1 | FileCheck -check-prefix=CK30 %s
+#ifdef CK30
+#include <stdio.h>
+
+void target_maps_parallel_integer(int a){
+ int ParamToKernel = a;
+ #pragma omp target map(tofrom:ParamToKernel)
+ {
+ printf("Print ParamToKernel = %d\n", ParamToKernel);
+ #pragma omp parallel
+ {
+ printf("Do nothing\n");
+ }
+ }
+}
+
+// CK30: define void @__omp_offloading_802_{{.*}}target_maps_parallel_integer{{.*}}l{{.*}}(i32* dereferenceable(4) %ParamToKernel) #0 {
+
+#endif
Index: lib/Sema/SemaOpenMP.cpp
===================================================================
--- lib/Sema/SemaOpenMP.cpp
+++ lib/Sema/SemaOpenMP.cpp
@@ -355,6 +355,27 @@
return false;
}
+ /// Do the check specified in \a Check to all component lists at a given level
+ /// and return true if any issue is found.
+ bool checkMappableExprComponentListsForDeclAtLevel(
+ ValueDecl *VD, unsigned Level,
+ const llvm::function_ref<
+ bool(OMPClauseMappableExprCommon::MappableExprComponentListRef,
+ OpenMPClauseKind)> &Check) {
+ auto StartI = std::next(Stack.begin());
+ auto EndI = Stack.end();
+ if (std::distance(StartI, EndI) <= (int)Level)
+ return false;
+ std::advance(StartI, Level);
+
+ auto MI = StartI->MappedExprComponents.find(VD);
+ if (MI != StartI->MappedExprComponents.end())
+ for (auto &L : MI->second.Components)
+ if (Check(L, MI->second.Kind))
+ return true;
+ return false;
+ }
+
/// Create a new mappable expression component list associated with a given
/// declaration and initialize it with the provided list of components.
void addMappableExpressionComponents(
@@ -912,9 +933,8 @@
bool IsVariableUsedInMapClause = false;
bool IsVariableAssociatedWithSection = false;
- DSAStack->checkMappableExprComponentListsForDecl(
- D, /*CurrentRegionOnly=*/true,
- [&](OMPClauseMappableExprCommon::MappableExprComponentListRef
+ DSAStack->checkMappableExprComponentListsForDeclAtLevel(
+ D, Level, [&](OMPClauseMappableExprCommon::MappableExprComponentListRef
MapExprComponents,
OpenMPClauseKind WhereFoundClauseKind) {
// Only the map clause information influences how a variable is
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
