Index: lib/Sema/SemaOpenMP.cpp =================================================================== --- lib/Sema/SemaOpenMP.cpp +++ lib/Sema/SemaOpenMP.cpp @@ -333,25 +333,37 @@ const llvm::function_ref< bool(OMPClauseMappableExprCommon::MappableExprComponentListRef, OpenMPClauseKind)> &Check) { - auto SI = Stack.rbegin(); - auto SE = Stack.rend(); + if (Stack.empty()) + return false; + + if (CurrentRegionOnly) + return checkMappableExprComponentListsForDeclAtLevel(VD, Stack.size() - 1, Check); + + if (Stack.size() > 2) + for (unsigned I = Stack.size() - 2; I > 0; --I) + if (checkMappableExprComponentListsForDeclAtLevel(VD, I, Check)) + return true; + return false; + } - if (SI == SE) + /// 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 (Stack.size() <= Level) return false; - if (CurrentRegionOnly) { - SE = std::next(SI); - } else { - ++SI; - } + auto StartI = Stack.begin(); + std::advance(StartI, Level); - for (; SI != SE; ++SI) { - auto MI = SI->MappedExprComponents.find(VD); - if (MI != SI->MappedExprComponents.end()) - for (auto &L : MI->second.Components) - if (Check(L, MI->second.Kind)) - return true; - } + 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; } @@ -912,9 +924,8 @@ bool IsVariableUsedInMapClause = false; bool IsVariableAssociatedWithSection = false; - DSAStack->checkMappableExprComponentListsForDecl( - D, /*CurrentRegionOnly=*/true, - [&](OMPClauseMappableExprCommon::MappableExprComponentListRef + DSAStack->checkMappableExprComponentListsForDeclAtLevel( + D, Level + 1, [&](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,20 @@ } #endif #endif + +///==========================================================================/// +// RUN: %clang_cc1 -DCK30 -std=c++11 -fopenmp -S -emit-llvm -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda %s -o - 2>&1 | FileCheck -check-prefix=CK30 %s + +#ifdef CK30 + +void target_maps_parallel_integer(int a){ + int ParamToKernel = a; +#pragma omp target map(tofrom: ParamToKernel) + { + ParamToKernel += 1; + } +} + +// CK30: {{.*}}void @__omp_offloading_{{.*}}(i32* dereferenceable(4) %ParamToKernel) #0 { + +#endif