diff options
| -rw-r--r-- | clang/lib/Sema/SemaOpenMP.cpp | 29 | ||||
| -rw-r--r-- | clang/test/OpenMP/target_map_codegen.cpp | 26 |
2 files changed, 52 insertions, 3 deletions
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp index 2b7733d2adb..49da0e49977 100644 --- a/clang/lib/Sema/SemaOpenMP.cpp +++ b/clang/lib/Sema/SemaOpenMP.cpp @@ -412,6 +412,30 @@ public: 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) { + if (isStackEmpty()) + return false; + + auto StartI = Stack.back().first.begin(); + auto EndI = Stack.back().first.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( @@ -994,9 +1018,8 @@ bool Sema::IsOpenMPCapturedByRef(ValueDecl *D, unsigned Level) { 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 diff --git a/clang/test/OpenMP/target_map_codegen.cpp b/clang/test/OpenMP/target_map_codegen.cpp index 4933bf31c2a..69e80bb9505 100644 --- a/clang/test/OpenMP/target_map_codegen.cpp +++ b/clang/test/OpenMP/target_map_codegen.cpp @@ -1056,6 +1056,9 @@ void implicit_maps_template_type_capture (int a){ // CK19: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4] // CK19: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] [i32 32] +// CK19: [[SIZE00n:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4] +// CK19: [[MTYPE00n:@.+]] = private {{.*}}constant [1 x i32] [i32 32] + // CK19: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 400] // CK19: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 33] @@ -1194,6 +1197,28 @@ void explicit_maps_single (int ii){ ++a; } + // Map of a scalar in nested region. + int b = a; + + // Region 00n + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00n]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00n]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32** + // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** + // CK19-DAG: store i32* [[VAR0:%.+]], i32** [[CBP0]] + // CK19-DAG: store i32* [[VAR0]], i32** [[CP0]] + + // CK19: call void [[CALL00n:@.+]](i32* {{[^,]+}}) + #pragma omp target map(alloc:b) + #pragma omp parallel + { + ++b; + } + // Map of an array. int arra[100]; @@ -2388,6 +2413,7 @@ void explicit_maps_single (int ii){ } // CK19: define {{.+}}[[CALL00]] +// CK19: define {{.+}}[[CALL00n]] // CK19: define {{.+}}[[CALL01]] // CK19: define {{.+}}[[CALL02]] // CK19: define {{.+}}[[CALL03]] |
