Index: clang/docs/ClangCommandLineReference.rst =================================================================== --- clang/docs/ClangCommandLineReference.rst +++ clang/docs/ClangCommandLineReference.rst @@ -2039,6 +2039,10 @@ .. option:: -fopenmp-version= +.. option:: -fopenmp-extensions, -fno-openmp-extensions + +Enable all Clang extensions for OpenMP directives and clauses. + .. program:: clang1 .. option:: -fopenmp= .. program:: clang Index: clang/docs/OpenMPSupport.rst =================================================================== --- clang/docs/OpenMPSupport.rst +++ clang/docs/OpenMPSupport.rst @@ -360,3 +360,20 @@ +------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+ | task extension | nowait clause on taskwait | :none:`unclaimed` | | +------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+ + +OpenMP Extensions +================= + +The following table provides a quick overview over various OpenMP +extensions and their implementation status. These extensions are not +currently defined by any standard, so links to associated LLVM +documentation are provided. As these extensions mature, they will be +considered for standardization. Please contact *openmp-dev* at +*lists.llvm.org* to provide feedback. + ++------------------------------+---------------------------------------------------------------------------+--------------------------+--------------------------------------------------------+ +|Category | Feature | Status | Reviews | ++==============================+===========================================================================+==========================+========================================================+ +| device extension | `'ompx_hold' map type modifier | :good:`prototyped` | D106509, D106510 | +| | `_ | | | ++------------------------------+---------------------------------------------------------------------------+--------------------------+--------------------------------------------------------+ Index: clang/include/clang/AST/OpenMPClause.h =================================================================== --- clang/include/clang/AST/OpenMPClause.h +++ clang/include/clang/AST/OpenMPClause.h @@ -5606,7 +5606,8 @@ /// Map-type-modifiers for the 'map' clause. OpenMPMapModifierKind MapTypeModifiers[NumberOfOMPMapClauseModifiers] = { OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown, - OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown}; + OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown, + OMPC_MAP_MODIFIER_unknown}; /// Location of map-type-modifiers for the 'map' clause. SourceLocation MapTypeModifiersLoc[NumberOfOMPMapClauseModifiers]; Index: clang/include/clang/Basic/DiagnosticParseKinds.td =================================================================== --- clang/include/clang/Basic/DiagnosticParseKinds.td +++ clang/include/clang/Basic/DiagnosticParseKinds.td @@ -1303,8 +1303,8 @@ def err_omp_unknown_map_type : Error< "incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'">; def err_omp_unknown_map_type_modifier : Error< - "incorrect map type modifier, expected 'always', 'close', " - "%select{or 'mapper'|'mapper', or 'present'}0">; + "incorrect map type modifier, expected one of: 'always', 'close', 'mapper'" + "%select{|, 'present'}0%select{|, 'ompx_hold'}1">; def err_omp_map_type_missing : Error< "missing map type">; def err_omp_map_type_modifier_missing : Error< Index: clang/include/clang/Basic/DiagnosticSemaKinds.td =================================================================== --- clang/include/clang/Basic/DiagnosticSemaKinds.td +++ clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -10532,6 +10532,8 @@ "variable already marked as mapped in current construct">; def err_omp_invalid_map_type_for_directive : Error< "%select{map type '%1' is not allowed|map type must be specified}0 for '#pragma omp %2'">; +def err_omp_invalid_map_type_modifier_for_directive : Error< + "map type modifier '%0' is not allowed for '#pragma omp %1'">; def err_omp_no_clause_for_directive : Error< "expected at least one %0 clause for '#pragma omp %1'">; def err_omp_threadprivate_in_clause : Error< Index: clang/include/clang/Basic/LangOptions.def =================================================================== --- clang/include/clang/Basic/LangOptions.def +++ clang/include/clang/Basic/LangOptions.def @@ -231,6 +231,7 @@ LANGOPT(CUDA , 1, 0, "CUDA") LANGOPT(HIP , 1, 0, "HIP") LANGOPT(OpenMP , 32, 0, "OpenMP support and version of OpenMP (31, 40 or 45)") +LANGOPT(OpenMPExtensions , 1, 0, "Enable all Clang extensions for OpenMP directives and clauses") LANGOPT(OpenMPSimd , 1, 0, "Use SIMD only OpenMP support.") LANGOPT(OpenMPUseTLS , 1, 0, "Use TLS for threadprivates or runtime calls") LANGOPT(OpenMPIsDevice , 1, 0, "Generate code only for OpenMP target device") Index: clang/include/clang/Basic/OpenMPKinds.h =================================================================== --- clang/include/clang/Basic/OpenMPKinds.h +++ clang/include/clang/Basic/OpenMPKinds.h @@ -14,6 +14,7 @@ #ifndef LLVM_CLANG_BASIC_OPENMPKINDS_H #define LLVM_CLANG_BASIC_OPENMPKINDS_H +#include "clang/Basic/LangOptions.h" #include "llvm/ADT/StringRef.h" #include "llvm/Frontend/OpenMP/OMPConstants.h" @@ -167,7 +168,7 @@ }; unsigned getOpenMPSimpleClauseType(OpenMPClauseKind Kind, llvm::StringRef Str, - unsigned OpenMPVersion); + const LangOptions &LangOpts); const char *getOpenMPSimpleClauseTypeName(OpenMPClauseKind Kind, unsigned Type); /// Checks if the specified directive is a directive with an associated Index: clang/include/clang/Basic/OpenMPKinds.def =================================================================== --- clang/include/clang/Basic/OpenMPKinds.def +++ clang/include/clang/Basic/OpenMPKinds.def @@ -123,6 +123,8 @@ OPENMP_MAP_MODIFIER_KIND(close) OPENMP_MAP_MODIFIER_KIND(mapper) OPENMP_MAP_MODIFIER_KIND(present) +// This is an OpenMP extension for the sake of OpenACC support. +OPENMP_MAP_MODIFIER_KIND(ompx_hold) // Modifiers for 'to' or 'from' clause. OPENMP_MOTION_MODIFIER_KIND(mapper) Index: clang/include/clang/Driver/Options.td =================================================================== --- clang/include/clang/Driver/Options.td +++ clang/include/clang/Driver/Options.td @@ -2379,6 +2379,12 @@ HelpText<"Parse OpenMP pragmas and generate parallel code.">; def fno_openmp : Flag<["-"], "fno-openmp">, Group, Flags<[NoArgumentUnused]>; def fopenmp_version_EQ : Joined<["-"], "fopenmp-version=">, Group, Flags<[CC1Option, NoArgumentUnused]>; +defm openmp_extensions: BoolFOption<"openmp-extensions", + LangOpts<"OpenMPExtensions">, DefaultFalse, + PosFlag, + NegFlag>; def fopenmp_EQ : Joined<["-"], "fopenmp=">, Group; def fopenmp_use_tls : Flag<["-"], "fopenmp-use-tls">, Group, Flags<[NoArgumentUnused, HelpHidden]>; Index: clang/lib/Basic/OpenMPKinds.cpp =================================================================== --- clang/lib/Basic/OpenMPKinds.cpp +++ clang/lib/Basic/OpenMPKinds.cpp @@ -21,7 +21,7 @@ using namespace llvm::omp; unsigned clang::getOpenMPSimpleClauseType(OpenMPClauseKind Kind, StringRef Str, - unsigned OpenMPVersion) { + const LangOptions &LangOpts) { switch (Kind) { case OMPC_default: return llvm::StringSwitch(Str) @@ -59,7 +59,9 @@ .Case(#Name, static_cast(OMPC_MAP_MODIFIER_##Name)) #include "clang/Basic/OpenMPKinds.def" .Default(OMPC_MAP_unknown); - if (OpenMPVersion < 51 && Type == OMPC_MAP_MODIFIER_present) + if (LangOpts.OpenMP < 51 && Type == OMPC_MAP_MODIFIER_present) + return OMPC_MAP_MODIFIER_unknown; + if (!LangOpts.OpenMPExtensions && Type == OMPC_MAP_MODIFIER_ompx_hold) return OMPC_MAP_MODIFIER_unknown; return Type; } @@ -70,7 +72,7 @@ .Case(#Name, static_cast(OMPC_MOTION_MODIFIER_##Name)) #include "clang/Basic/OpenMPKinds.def" .Default(OMPC_MOTION_MODIFIER_unknown); - if (OpenMPVersion < 51 && Type == OMPC_MOTION_MODIFIER_present) + if (LangOpts.OpenMP < 51 && Type == OMPC_MOTION_MODIFIER_present) return OMPC_MOTION_MODIFIER_unknown; return Type; } Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp =================================================================== --- clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -7269,6 +7269,14 @@ /// 0x800 is reserved for compatibility with XLC. /// Produce a runtime error if the data is not already allocated. OMP_MAP_PRESENT = 0x1000, + // Increment and decrement a separate reference counter so that the data + // cannot be unmapped within the associated region. Thus, this flag is + // intended to be used on 'target' and 'target data' directives because they + // are inherently structured. It is not intended to be used on 'target + // enter data' and 'target exit data' directives because they are inherently + // dynamic. + // This is an OpenMP extension for the sake of OpenACC support. + OMP_MAP_OMPX_HOLD = 0x2000, /// Signal that the runtime library should use args as an array of /// descriptor_dim pointers and use args_size as dims. Used when we have /// non-contiguous list items in target update directive @@ -7570,6 +7578,9 @@ llvm::find(MotionModifiers, OMPC_MOTION_MODIFIER_present) != MotionModifiers.end()) Bits |= OMP_MAP_PRESENT; + if (llvm::find(MapModifiers, OMPC_MAP_MODIFIER_ompx_hold) != + MapModifiers.end()) + Bits |= OMP_MAP_OMPX_HOLD; if (IsNonContiguous) Bits |= OMP_MAP_NON_CONTIG; return Bits; @@ -8923,6 +8934,20 @@ CombinedInfo.Types.back() |= OMP_MAP_PRESENT; // Remove TARGET_PARAM flag from the first element (*CurTypes.begin()) &= ~OMP_MAP_TARGET_PARAM; + // If any element has the ompx_hold modifier, then make sure the runtime + // uses the hold reference count for the struct as a whole so that it won't + // be unmapped by an extra dynamic reference count decrement. Add it to all + // elements as well so the runtime knows which reference count to check + // when determining whether it's time for device-to-host transfers of + // individual elements. + if (CurTypes.end() != + llvm::find_if(CurTypes, [](OpenMPOffloadMappingFlags Type) { + return Type & OMP_MAP_OMPX_HOLD; + })) { + CombinedInfo.Types.back() |= OMP_MAP_OMPX_HOLD; + for (auto &M : CurTypes) + M |= OMP_MAP_OMPX_HOLD; + } // All other current entries will be MEMBER_OF the combined entry // (except for PTR_AND_OBJ entries which do not have a placeholder value Index: clang/lib/Driver/ToolChains/Clang.cpp =================================================================== --- clang/lib/Driver/ToolChains/Clang.cpp +++ clang/lib/Driver/ToolChains/Clang.cpp @@ -5771,6 +5771,9 @@ options::OPT_fno_openmp_simd); Args.AddAllArgs(CmdArgs, options::OPT_fopenmp_enable_irbuilder); Args.AddAllArgs(CmdArgs, options::OPT_fopenmp_version_EQ); + if (Args.hasFlag(options::OPT_fopenmp_extensions, + options::OPT_fno_openmp_extensions, /*Default=*/false)) + CmdArgs.push_back("-fopenmp-extensions"); Args.AddAllArgs(CmdArgs, options::OPT_fopenmp_cuda_number_of_sm_EQ); Args.AddAllArgs(CmdArgs, options::OPT_fopenmp_cuda_blocks_per_sm_EQ); Args.AddAllArgs(CmdArgs, @@ -5806,6 +5809,9 @@ Args.AddLastArg(CmdArgs, options::OPT_fopenmp_simd, options::OPT_fno_openmp_simd); Args.AddAllArgs(CmdArgs, options::OPT_fopenmp_version_EQ); + if (Args.hasFlag(options::OPT_fopenmp_extensions, + options::OPT_fno_openmp_extensions, /*Default=*/false)) + CmdArgs.push_back("-fopenmp-extensions"); } const SanitizerArgs &Sanitize = TC.getSanitizerArgs(); Index: clang/lib/Parse/ParseOpenMP.cpp =================================================================== --- clang/lib/Parse/ParseOpenMP.cpp +++ clang/lib/Parse/ParseOpenMP.cpp @@ -1651,7 +1651,7 @@ unsigned Type = getOpenMPSimpleClauseType( Kind, Tok.isAnnotation() ? "" : P.getPreprocessor().getSpelling(Tok), - P.getLangOpts().OpenMP); + P.getLangOpts()); SourceLocation TypeLoc = Tok.getLocation(); if (Tok.isNot(tok::r_paren) && Tok.isNot(tok::comma) && Tok.isNot(tok::annot_pragma_openmp_end)) @@ -3310,8 +3310,7 @@ Arg[Modifier2] = OMPC_SCHEDULE_MODIFIER_unknown; Arg[ScheduleKind] = OMPC_SCHEDULE_unknown; unsigned KindModifier = getOpenMPSimpleClauseType( - Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok), - getLangOpts().OpenMP); + Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok), getLangOpts()); if (KindModifier > OMPC_SCHEDULE_unknown) { // Parse 'modifier' Arg[Modifier1] = KindModifier; @@ -3323,8 +3322,7 @@ // Parse ',' 'modifier' ConsumeAnyToken(); KindModifier = getOpenMPSimpleClauseType( - Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok), - getLangOpts().OpenMP); + Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok), getLangOpts()); Arg[Modifier2] = KindModifier > OMPC_SCHEDULE_unknown ? KindModifier : (unsigned)OMPC_SCHEDULE_unknown; @@ -3339,8 +3337,7 @@ else Diag(Tok, diag::warn_pragma_expected_colon) << "schedule modifier"; KindModifier = getOpenMPSimpleClauseType( - Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok), - getLangOpts().OpenMP); + Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok), getLangOpts()); } Arg[ScheduleKind] = KindModifier; KLoc[ScheduleKind] = Tok.getLocation(); @@ -3354,8 +3351,7 @@ DelimLoc = ConsumeAnyToken(); } else if (Kind == OMPC_dist_schedule) { Arg.push_back(getOpenMPSimpleClauseType( - Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok), - getLangOpts().OpenMP)); + Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok), getLangOpts())); KLoc.push_back(Tok.getLocation()); if (Tok.isNot(tok::r_paren) && Tok.isNot(tok::comma) && Tok.isNot(tok::annot_pragma_openmp_end)) @@ -3365,8 +3361,7 @@ } else if (Kind == OMPC_defaultmap) { // Get a defaultmap modifier unsigned Modifier = getOpenMPSimpleClauseType( - Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok), - getLangOpts().OpenMP); + Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok), getLangOpts()); // Set defaultmap modifier to unknown if it is either scalar, aggregate, or // pointer if (Modifier < OMPC_DEFAULTMAP_MODIFIER_unknown) @@ -3384,8 +3379,7 @@ Diag(Tok, diag::warn_pragma_expected_colon) << "defaultmap modifier"; // Get a defaultmap kind Arg.push_back(getOpenMPSimpleClauseType( - Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok), - getLangOpts().OpenMP)); + Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok), getLangOpts())); KLoc.push_back(Tok.getLocation()); if (Tok.isNot(tok::r_paren) && Tok.isNot(tok::comma) && Tok.isNot(tok::annot_pragma_openmp_end)) @@ -3400,8 +3394,7 @@ NextToken().is(tok::colon)) { // Parse optional ':' Arg.push_back(getOpenMPSimpleClauseType( - Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok), - getLangOpts().OpenMP)); + Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok), getLangOpts())); KLoc.push_back(Tok.getLocation()); ConsumeAnyToken(); // Parse ':' @@ -3512,7 +3505,7 @@ Preprocessor &PP = P.getPreprocessor(); OpenMPMapModifierKind TypeModifier = static_cast(getOpenMPSimpleClauseType( - OMPC_map, PP.getSpelling(Tok), P.getLangOpts().OpenMP)); + OMPC_map, PP.getSpelling(Tok), P.getLangOpts())); return TypeModifier; } @@ -3554,7 +3547,8 @@ OpenMPMapModifierKind TypeModifier = isMapModifier(*this); if (TypeModifier == OMPC_MAP_MODIFIER_always || TypeModifier == OMPC_MAP_MODIFIER_close || - TypeModifier == OMPC_MAP_MODIFIER_present) { + TypeModifier == OMPC_MAP_MODIFIER_present || + TypeModifier == OMPC_MAP_MODIFIER_ompx_hold) { Data.MapTypeModifiers.push_back(TypeModifier); Data.MapTypeModifiersLoc.push_back(Tok.getLocation()); ConsumeToken(); @@ -3577,7 +3571,8 @@ if (PP.LookAhead(0).is(tok::colon)) return false; Diag(Tok, diag::err_omp_unknown_map_type_modifier) - << (getLangOpts().OpenMP >= 51 ? 1 : 0); + << (getLangOpts().OpenMP >= 51 ? 1 : 0) + << getLangOpts().OpenMPExtensions; ConsumeToken(); } if (getCurToken().is(tok::comma)) @@ -3596,7 +3591,7 @@ Preprocessor &PP = P.getPreprocessor(); OpenMPMapClauseKind MapType = static_cast(getOpenMPSimpleClauseType( - OMPC_map, PP.getSpelling(Tok), P.getLangOpts().OpenMP)); + OMPC_map, PP.getSpelling(Tok), P.getLangOpts())); return MapType; } @@ -3749,8 +3744,8 @@ (Tok.is(tok::identifier) || Tok.is(tok::kw_default)) && NextToken().is(tok::comma)) { // Parse optional reduction modifier. - Data.ExtraModifier = getOpenMPSimpleClauseType(Kind, PP.getSpelling(Tok), - getLangOpts().OpenMP); + Data.ExtraModifier = + getOpenMPSimpleClauseType(Kind, PP.getSpelling(Tok), getLangOpts()); Data.ExtraModifierLoc = Tok.getLocation(); ConsumeToken(); assert(Tok.is(tok::comma) && "Expected comma."); @@ -3796,7 +3791,7 @@ ColonProtectionRAIIObject ColonRAII(*this); Data.ExtraModifier = getOpenMPSimpleClauseType( Kind, Tok.is(tok::identifier) ? PP.getSpelling(Tok) : "", - getLangOpts().OpenMP); + getLangOpts()); Data.ExtraModifierLoc = Tok.getLocation(); if (Data.ExtraModifier == OMPC_DEPEND_unknown) { SkipUntil(tok::colon, tok::r_paren, tok::annot_pragma_openmp_end, @@ -3821,8 +3816,8 @@ // Try to parse modifier if any. Data.ExtraModifier = OMPC_LINEAR_val; if (Tok.is(tok::identifier) && PP.LookAhead(0).is(tok::l_paren)) { - Data.ExtraModifier = getOpenMPSimpleClauseType(Kind, PP.getSpelling(Tok), - getLangOpts().OpenMP); + Data.ExtraModifier = + getOpenMPSimpleClauseType(Kind, PP.getSpelling(Tok), getLangOpts()); Data.ExtraModifierLoc = ConsumeToken(); LinearT.consumeOpen(); NeedRParenForLinear = true; @@ -3835,8 +3830,8 @@ if ((getLangOpts().OpenMP >= 50 && !isOpenMPDistributeDirective(DKind) && !isOpenMPTaskLoopDirective(DKind)) && Tok.is(tok::identifier) && PP.LookAhead(0).is(tok::colon)) { - Data.ExtraModifier = getOpenMPSimpleClauseType(Kind, PP.getSpelling(Tok), - getLangOpts().OpenMP); + Data.ExtraModifier = + getOpenMPSimpleClauseType(Kind, PP.getSpelling(Tok), getLangOpts()); Data.ExtraModifierLoc = Tok.getLocation(); ConsumeToken(); assert(Tok.is(tok::colon) && "Expected colon."); @@ -3879,9 +3874,8 @@ Data.ColonLoc = ConsumeToken(); } else if (Kind == OMPC_to || Kind == OMPC_from) { while (Tok.is(tok::identifier)) { - auto Modifier = - static_cast(getOpenMPSimpleClauseType( - Kind, PP.getSpelling(Tok), getLangOpts().OpenMP)); + auto Modifier = static_cast( + getOpenMPSimpleClauseType(Kind, PP.getSpelling(Tok), getLangOpts())); if (Modifier == OMPC_MOTION_MODIFIER_unknown) break; Data.MotionModifiers.push_back(Modifier); Index: clang/lib/Sema/SemaOpenMP.cpp =================================================================== --- clang/lib/Sema/SemaOpenMP.cpp +++ clang/lib/Sema/SemaOpenMP.cpp @@ -19318,6 +19318,7 @@ CXXScopeSpec &MapperIdScopeSpec, DeclarationNameInfo MapperId, ArrayRef UnresolvedMappers, OpenMPMapClauseKind MapType = OMPC_MAP_unknown, + ArrayRef Modifiers = None, bool IsMapTypeImplicit = false) { // We only expect mappable expressions in 'to', 'from', and 'map' clauses. assert((CKind == OMPC_map || CKind == OMPC_to || CKind == OMPC_from) && @@ -19339,6 +19340,10 @@ bool UpdateUMIt = false; Expr *UnresolvedMapper = nullptr; + bool HasHoldModifier = + Modifiers.end() != std::find(Modifiers.begin(), Modifiers.end(), + OMPC_MAP_MODIFIER_ompx_hold); + // Keep track of the mappable components and base declarations in this clause. // Each entry in the list is going to have a list of components associated. We // record each set of the components so that we can build the clause later on. @@ -19537,6 +19542,21 @@ continue; } + // The 'ompx_hold' modifier is specifically intended to be used on a + // 'target' or 'target data' directive to prevent data from being unmapped + // during the associated statement. It is not permitted on a 'target + // enter data' or 'target exit data' directive, which have no associated + // statement. + if ((DKind == OMPD_target_enter_data || DKind == OMPD_target_exit_data) && + HasHoldModifier) { + SemaRef.Diag(StartLoc, + diag::err_omp_invalid_map_type_modifier_for_directive) + << getOpenMPSimpleClauseTypeName(OMPC_map, + OMPC_MAP_MODIFIER_ompx_hold) + << getOpenMPDirectiveName(DKind); + continue; + } + // target, target data // OpenMP 5.0 [2.12.2, Restrictions, p. 163] // OpenMP 5.0 [2.12.5, Restrictions, p. 174] @@ -19611,7 +19631,8 @@ const OMPVarListLocTy &Locs, ArrayRef UnresolvedMappers) { OpenMPMapModifierKind Modifiers[] = { OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown, - OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown}; + OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown, + OMPC_MAP_MODIFIER_unknown}; SourceLocation ModifiersLoc[NumberOfOMPMapClauseModifiers]; // Process map-type-modifiers, flag errors for duplicate modifiers. @@ -19632,7 +19653,7 @@ MappableVarListInfo MVLI(VarList); checkMappableExpressionList(*this, DSAStack, OMPC_map, MVLI, Locs.StartLoc, MapperIdScopeSpec, MapperId, UnresolvedMappers, - MapType, IsMapTypeImplicit); + MapType, Modifiers, IsMapTypeImplicit); // We need to produce a map clause even if we don't have variables so that // other diagnostics related with non-existing map clauses are accurate. Index: clang/test/OpenMP/driver-fopenmp-extensions.c =================================================================== --- /dev/null +++ clang/test/OpenMP/driver-fopenmp-extensions.c @@ -0,0 +1,26 @@ +// RUN: %clang -Xclang -verify=omp -c -fopenmp %s +// RUN: %clang -Xclang -verify=omp -c -fopenmp-simd %s + +// RUN: %clang -c -Xclang -verify=ompx -fopenmp -fopenmp-extensions %s +// RUN: %clang -c -Xclang -verify=ompx -fopenmp-simd -fopenmp-extensions %s + +// RUN: %clang -Xclang -verify=omp -c -fopenmp -fno-openmp-extensions %s +// RUN: %clang -Xclang -verify=omp -c -fopenmp-simd -fno-openmp-extensions %s + +// RUN: %clang -Xclang -verify=omp -c -fopenmp \ +// RUN: -fopenmp-extensions -fno-openmp-extensions %s +// RUN: %clang -Xclang -verify=omp -c -fopenmp-simd \ +// RUN: -fopenmp-extensions -fno-openmp-extensions %s + +// RUN: %clang -c -Xclang -verify=ompx -fopenmp \ +// RUN: -fno-openmp-extensions -fopenmp-extensions %s +// RUN: %clang -c -Xclang -verify=ompx -fopenmp-simd \ +// RUN: -fno-openmp-extensions -fopenmp-extensions %s + +void foo() { + int x; + // ompx-no-diagnostics + // omp-error@+1 {{incorrect map type modifier}} + #pragma omp target map(ompx_hold, alloc: x) + ; +} Index: clang/test/OpenMP/target_ast_print.cpp =================================================================== --- clang/test/OpenMP/target_ast_print.cpp +++ clang/test/OpenMP/target_ast_print.cpp @@ -1125,4 +1125,76 @@ return tmain(argc, &argc) + tmain(argv[0][0], argv[0]); } #endif // OMP51 + +#ifdef OMPX + +// RUN: %clang_cc1 -DOMPX -verify -fopenmp -fopenmp-extensions -ast-print %s | FileCheck %s --check-prefix=OMPX +// RUN: %clang_cc1 -DOMPX -fopenmp -fopenmp-extensions -x c++ -std=c++11 -emit-pch -o %t %s +// RUN: %clang_cc1 -DOMPX -fopenmp -fopenmp-extensions -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s --check-prefix=OMPX + +// RUN: %clang_cc1 -DOMPX -verify -fopenmp-simd -fopenmp-extensions -ast-print %s | FileCheck %s --check-prefix=OMPX +// RUN: %clang_cc1 -DOMPX -fopenmp-simd -fopenmp-extensions -x c++ -std=c++11 -emit-pch -o %t %s +// RUN: %clang_cc1 -DOMPX -fopenmp-simd -fopenmp-extensions -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s --check-prefix=OMPX + +void foo() {} + +template +T tmain(T argc, T *argv) { + T i, ompx_hold; +#pragma omp target map(ompx_hold,alloc: i) + foo(); +#pragma omp target map(ompx_hold from: i) + foo(); +#pragma omp target map(ompx_hold) + {ompx_hold++;} +#pragma omp target map(ompx_hold,i) + {ompx_hold++;i++;} + return 0; +} + +// OMPX: template T tmain(T argc, T *argv) { +// OMPX-NEXT: T i, ompx_hold; +// OMPX-NEXT: #pragma omp target map(ompx_hold,alloc: i) +// OMPX-NEXT: foo() +// OMPX-NEXT: #pragma omp target map(ompx_hold,from: i) +// OMPX-NEXT: foo() +// OMPX-NEXT: #pragma omp target map(tofrom: ompx_hold) +// OMPX-NEXT: { +// OMPX-NEXT: ompx_hold++; +// OMPX-NEXT: } +// OMPX-NEXT: #pragma omp target map(tofrom: ompx_hold,i) +// OMPX-NEXT: { +// OMPX-NEXT: ompx_hold++; +// OMPX-NEXT: i++; +// OMPX-NEXT: } + +// OMPX-LABEL: int main(int argc, char **argv) { +// OMPX-NEXT: int i, ompx_hold; +// OMPX-NEXT: #pragma omp target map(ompx_hold,alloc: i) +// OMPX-NEXT: foo(); +// OMPX-NEXT: #pragma omp target map(ompx_hold,from: i) +// OMPX-NEXT: foo(); +// OMPX-NEXT: #pragma omp target map(tofrom: ompx_hold) +// OMPX-NEXT: { +// OMPX-NEXT: ompx_hold++; +// OMPX-NEXT: } +// OMPX-NEXT: #pragma omp target map(tofrom: ompx_hold,i) +// OMPX-NEXT: { +// OMPX-NEXT: ompx_hold++; +// OMPX-NEXT: i++; +// OMPX-NEXT: } +int main (int argc, char **argv) { + int i, ompx_hold; + #pragma omp target map(ompx_hold,alloc: i) + foo(); + #pragma omp target map(ompx_hold from: i) + foo(); + #pragma omp target map(ompx_hold) + {ompx_hold++;} + #pragma omp target map(ompx_hold,i) + {ompx_hold++;i++;} + return tmain(argc, &argc) + tmain(argv[0][0], argv[0]); +} + +#endif #endif Index: clang/test/OpenMP/target_data_ast_print.cpp =================================================================== --- clang/test/OpenMP/target_data_ast_print.cpp +++ clang/test/OpenMP/target_data_ast_print.cpp @@ -6,13 +6,13 @@ // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s -// RUN: %clang_cc1 -DOMP51 -verify -fopenmp -fopenmp-version=51 -ast-print %s | FileCheck -check-prefixes=CHECK,OMP51 %s -// RUN: %clang_cc1 -DOMP51 -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -emit-pch -o %t %s -// RUN: %clang_cc1 -DOMP51 -fopenmp -fopenmp-version=51 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck -check-prefixes=CHECK,OMP51 %s +// RUN: %clang_cc1 -DOMP51 -DOMPX -verify -fopenmp -fopenmp-version=51 -fopenmp-extensions -ast-print %s | FileCheck -check-prefixes=CHECK,OMP51,OMPX %s +// RUN: %clang_cc1 -DOMP51 -DOMPX -fopenmp -fopenmp-version=51 -fopenmp-extensions -x c++ -std=c++11 -emit-pch -o %t %s +// RUN: %clang_cc1 -DOMP51 -DOMPX -fopenmp -fopenmp-version=51 -fopenmp-extensions -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck -check-prefixes=CHECK,OMP51,OMPX %s -// RUN: %clang_cc1 -DOMP51 -verify -fopenmp-simd -fopenmp-version=51 -ast-print %s | FileCheck -check-prefixes=CHECK,OMP51 %s -// RUN: %clang_cc1 -DOMP51 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -emit-pch -o %t %s -// RUN: %clang_cc1 -DOMP51 -fopenmp-simd -fopenmp-version=51 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck -check-prefixes=CHECK,OMP51 %s +// RUN: %clang_cc1 -DOMP51 -DOMPX -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-extensions -ast-print %s | FileCheck -check-prefixes=CHECK,OMP51,OMPX %s +// RUN: %clang_cc1 -DOMP51 -DOMPX -fopenmp-simd -fopenmp-version=51 -fopenmp-extensions -x c++ -std=c++11 -emit-pch -o %t %s +// RUN: %clang_cc1 -DOMP51 -DOMPX -fopenmp-simd -fopenmp-version=51 -fopenmp-extensions -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck -check-prefixes=CHECK,OMP51,OMPX %s // expected-no-diagnostics #ifndef HEADER @@ -56,6 +56,11 @@ foo(); #endif +#ifdef OMPX +#pragma omp target data map(ompx_hold,alloc: e) + foo(); +#endif + // nesting a target region #pragma omp target data map(e) { @@ -67,6 +72,10 @@ #pragma omp target map(present, alloc: e) foo(); #endif +#ifdef OMPX + #pragma omp target map(ompx_hold, alloc: e) + foo(); +#endif } return 0; @@ -94,6 +103,8 @@ // CHECK-NEXT: foo(); // OMP51-NEXT: #pragma omp target data map(present,alloc: e) // OMP51-NEXT: foo(); +// OMPX-NEXT: #pragma omp target data map(ompx_hold,alloc: e) +// OMPX-NEXT: foo(); // CHECK-NEXT: #pragma omp target data map(tofrom: e) // CHECK-NEXT: { // CHECK-NEXT: #pragma omp target map(always,alloc: e) @@ -102,6 +113,8 @@ // CHECK-NEXT: foo(); // OMP51-NEXT: #pragma omp target map(present,alloc: e) // OMP51-NEXT: foo(); +// OMPX-NEXT: #pragma omp target map(ompx_hold,alloc: e) +// OMPX-NEXT: foo(); // CHECK: template<> int tmain(int argc, int *argv) { // CHECK-NEXT: int i, j, b, c, d, e, x[20]; // CHECK-NEXT: #pragma omp target data map(to: c) @@ -124,6 +137,8 @@ // CHECK-NEXT: foo(); // OMP51-NEXT: #pragma omp target data map(present,alloc: e) // OMP51-NEXT: foo(); +// OMPX-NEXT: #pragma omp target data map(ompx_hold,alloc: e) +// OMPX-NEXT: foo(); // CHECK-NEXT: #pragma omp target data map(tofrom: e) // CHECK-NEXT: { // CHECK-NEXT: #pragma omp target map(always,alloc: e) @@ -132,6 +147,8 @@ // CHECK-NEXT: foo(); // OMP51-NEXT: #pragma omp target map(present,alloc: e) // OMP51-NEXT: foo(); +// OMPX-NEXT: #pragma omp target map(ompx_hold,alloc: e) +// OMPX-NEXT: foo(); // CHECK: template<> char tmain(char argc, char *argv) { // CHECK-NEXT: char i, j, b, c, d, e, x[20]; // CHECK-NEXT: #pragma omp target data map(to: c) @@ -154,6 +171,8 @@ // CHECK-NEXT: foo(); // OMP51-NEXT: #pragma omp target data map(present,alloc: e) // OMP51-NEXT: foo(); +// OMPX-NEXT: #pragma omp target data map(ompx_hold,alloc: e) +// OMPX-NEXT: foo(); // CHECK-NEXT: #pragma omp target data map(tofrom: e) // CHECK-NEXT: { // CHECK-NEXT: #pragma omp target map(always,alloc: e) @@ -162,6 +181,8 @@ // CHECK-NEXT: foo(); // OMP51-NEXT: #pragma omp target map(present,alloc: e) // OMP51-NEXT: foo(); +// OMPX-NEXT: #pragma omp target map(ompx_hold,alloc: e) +// OMPX-NEXT: foo(); int main (int argc, char **argv) { int b = argc, c, d, e, f, g, x[20]; @@ -221,6 +242,13 @@ foo(); #endif +// OMPX-NEXT: #pragma omp target data map(ompx_hold,alloc: e) +// OMPX-NEXT: foo(); +#ifdef OMPX +#pragma omp target data map(ompx_hold,alloc: e) + foo(); +#endif + // nesting a target region #pragma omp target data map(e) // CHECK-NEXT: #pragma omp target data map(tofrom: e) Index: clang/test/OpenMP/target_data_map_codegen_hold.cpp =================================================================== --- /dev/null +++ clang/test/OpenMP/target_data_map_codegen_hold.cpp @@ -0,0 +1,608 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals --global-value-regex ".offload_maptypes.*" ".offload_sizes.*" --global-hex-value-regex ".offload_maptypes.*" +// expected-no-diagnostics +#ifndef HEADER +#define HEADER + +// powerpc64le-ibm-linux-gnu + +// RUN: %clang_cc1 -verify -fopenmp -fopenmp-extensions \ +// RUN: -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ \ +// RUN: -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | \ +// RUN: FileCheck %s --check-prefixes=CHECK-PPC64LE +// RUN: %clang_cc1 -fopenmp -fopenmp-extensions \ +// RUN: -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 \ +// RUN: -triple powerpc64le-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -fopenmp-extensions \ +// RUN: -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ \ +// RUN: -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t \ +// RUN: -verify %s -emit-llvm -o - | \ +// RUN: FileCheck %s --check-prefixes=CHECK-PPC64LE + +// i386-pc-linux-gnu + +// RUN: %clang_cc1 -verify -fopenmp -fopenmp-extensions \ +// RUN: -fopenmp-targets=i386-pc-linux-gnu -x c++ \ +// RUN: -triple i386-unknown-unknown -emit-llvm %s -o - | \ +// RUN: FileCheck %s --check-prefixes=CHECK-I386 +// RUN: %clang_cc1 -fopenmp -fopenmp-extensions \ +// RUN: -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 \ +// RUN: -triple i386-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -fopenmp-extensions \ +// RUN: -fopenmp-targets=i386-pc-linux-gnu -x c++ \ +// RUN: -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s \ +// RUN: -emit-llvm -o - | \ +// RUN: FileCheck %s --check-prefixes=CHECK-I386 + +struct S1 { + int i; +}; +struct S2 { + S1 s; + struct S2 *ps; +}; + +// Map flags used in @.offload_maptypes* below: +// +// TO = 0x1 +// FROM = 0x2 +// ALWAYS = 0x4 +// PTR_AND_OBJ = 0x10 +// CLOSE = 0x400 +// OMPX_HOLD = 0x2000 +// MEMBER_OF_1 = 0x1000000000000 +// MEMBER_OF_7 = 0x7000000000000 + +//. +// CHECK-PPC64LE: @.offload_sizes = private unnamed_addr constant [1 x i64] [i64 20] +// CHECK-PPC64LE: @.offload_maptypes = private unnamed_addr constant [1 x i64] [i64 [[#0x2001]]] +// CHECK-PPC64LE: @.offload_sizes.1 = private unnamed_addr constant [1 x i64] [i64 20] +// CHECK-PPC64LE: @.offload_maptypes.2 = private unnamed_addr constant [1 x i64] [i64 [[#0x2405]]] +// CHECK-PPC64LE: @.offload_sizes.3 = private unnamed_addr constant [1 x i64] [i64 4] +// CHECK-PPC64LE: @.offload_maptypes.4 = private unnamed_addr constant [1 x i64] [i64 [[#0x2003]]] +// CHECK-PPC64LE: @.offload_maptypes.5 = private unnamed_addr constant [11 x i64] [i64 [[#0x2000]], i64 [[#0x1000000002003]], i64 [[#0x1000000002010]], i64 [[#0x2010]], i64 [[#0x2013]], i64 [[#0x3]], i64 [[#0x2000]], i64 [[#0x7000000002003]], i64 [[#0x7000000002010]], i64 [[#0x2010]], i64 [[#0x2013]]] +//. +// CHECK-I386: @.offload_sizes = private unnamed_addr constant [1 x i64] [i64 20] +// CHECK-I386: @.offload_maptypes = private unnamed_addr constant [1 x i64] [i64 [[#0x2001]]] +// CHECK-I386: @.offload_sizes.1 = private unnamed_addr constant [1 x i64] [i64 20] +// CHECK-I386: @.offload_maptypes.2 = private unnamed_addr constant [1 x i64] [i64 [[#0x2405]]] +// CHECK-I386: @.offload_sizes.3 = private unnamed_addr constant [1 x i64] [i64 4] +// CHECK-I386: @.offload_maptypes.4 = private unnamed_addr constant [1 x i64] [i64 [[#0x2003]]] +// CHECK-I386: @.offload_maptypes.5 = private unnamed_addr constant [11 x i64] [i64 [[#0x2000]], i64 [[#0x1000000002003]], i64 [[#0x1000000002010]], i64 [[#0x2010]], i64 [[#0x2013]], i64 [[#0x3]], i64 [[#0x2000]], i64 [[#0x7000000002003]], i64 [[#0x7000000002010]], i64 [[#0x2010]], i64 [[#0x2013]]] +//. +// CHECK-PPC64LE-LABEL: @_Z3fooi( +// CHECK-PPC64LE-NEXT: entry: +// CHECK-PPC64LE-NEXT: [[ARG_ADDR:%.*]] = alloca i32, align 4 +// CHECK-PPC64LE-NEXT: [[LB:%.*]] = alloca [5 x float], align 4 +// CHECK-PPC64LE-NEXT: [[PS1:%.*]] = alloca %struct.S2*, align 8 +// CHECK-PPC64LE-NEXT: [[PS2:%.*]] = alloca %struct.S2*, align 8 +// CHECK-PPC64LE-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x i8*], align 8 +// CHECK-PPC64LE-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x i8*], align 8 +// CHECK-PPC64LE-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x i8*], align 8 +// CHECK-PPC64LE-NEXT: [[DOTOFFLOAD_BASEPTRS1:%.*]] = alloca [1 x i8*], align 8 +// CHECK-PPC64LE-NEXT: [[DOTOFFLOAD_PTRS2:%.*]] = alloca [1 x i8*], align 8 +// CHECK-PPC64LE-NEXT: [[DOTOFFLOAD_MAPPERS3:%.*]] = alloca [1 x i8*], align 8 +// CHECK-PPC64LE-NEXT: [[DOTOFFLOAD_BASEPTRS5:%.*]] = alloca [1 x i8*], align 8 +// CHECK-PPC64LE-NEXT: [[DOTOFFLOAD_PTRS6:%.*]] = alloca [1 x i8*], align 8 +// CHECK-PPC64LE-NEXT: [[DOTOFFLOAD_MAPPERS7:%.*]] = alloca [1 x i8*], align 8 +// CHECK-PPC64LE-NEXT: [[DOTOFFLOAD_BASEPTRS29:%.*]] = alloca [11 x i8*], align 8 +// CHECK-PPC64LE-NEXT: [[DOTOFFLOAD_PTRS30:%.*]] = alloca [11 x i8*], align 8 +// CHECK-PPC64LE-NEXT: [[DOTOFFLOAD_MAPPERS31:%.*]] = alloca [11 x i8*], align 8 +// CHECK-PPC64LE-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [11 x i64], align 8 +// CHECK-PPC64LE-NEXT: store i32 [[ARG:%.*]], i32* [[ARG_ADDR]], align 4 +// CHECK-PPC64LE-NEXT: [[TMP0:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK-PPC64LE-NEXT: [[TMP1:%.*]] = bitcast i8** [[TMP0]] to [5 x float]** +// CHECK-PPC64LE-NEXT: store [5 x float]* [[LB]], [5 x float]** [[TMP1]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK-PPC64LE-NEXT: [[TMP3:%.*]] = bitcast i8** [[TMP2]] to [5 x float]** +// CHECK-PPC64LE-NEXT: store [5 x float]* [[LB]], [5 x float]** [[TMP3]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0 +// CHECK-PPC64LE-NEXT: store i8* null, i8** [[TMP4]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK-PPC64LE-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK-PPC64LE-NEXT: call void @__tgt_target_data_begin_mapper(%struct.ident_t* @[[GLOB1:[0-9]+]], i64 -1, i32 1, i8** [[TMP5]], i8** [[TMP6]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null) +// CHECK-PPC64LE-NEXT: [[TMP7:%.*]] = load i32, i32* [[ARG_ADDR]], align 4 +// CHECK-PPC64LE-NEXT: [[INC:%.*]] = add nsw i32 [[TMP7]], 1 +// CHECK-PPC64LE-NEXT: store i32 [[INC]], i32* [[ARG_ADDR]], align 4 +// CHECK-PPC64LE-NEXT: [[TMP8:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK-PPC64LE-NEXT: [[TMP9:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK-PPC64LE-NEXT: call void @__tgt_target_data_end_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i32 1, i8** [[TMP8]], i8** [[TMP9]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null) +// CHECK-PPC64LE-NEXT: [[TMP10:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS1]], i32 0, i32 0 +// CHECK-PPC64LE-NEXT: [[TMP11:%.*]] = bitcast i8** [[TMP10]] to [5 x float]** +// CHECK-PPC64LE-NEXT: store [5 x float]* [[LB]], [5 x float]** [[TMP11]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP12:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS2]], i32 0, i32 0 +// CHECK-PPC64LE-NEXT: [[TMP13:%.*]] = bitcast i8** [[TMP12]] to [5 x float]** +// CHECK-PPC64LE-NEXT: store [5 x float]* [[LB]], [5 x float]** [[TMP13]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP14:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_MAPPERS3]], i64 0, i64 0 +// CHECK-PPC64LE-NEXT: store i8* null, i8** [[TMP14]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP15:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS1]], i32 0, i32 0 +// CHECK-PPC64LE-NEXT: [[TMP16:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS2]], i32 0, i32 0 +// CHECK-PPC64LE-NEXT: call void @__tgt_target_data_begin_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i32 1, i8** [[TMP15]], i8** [[TMP16]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes.1, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.2, i32 0, i32 0), i8** null, i8** null) +// CHECK-PPC64LE-NEXT: [[TMP17:%.*]] = load i32, i32* [[ARG_ADDR]], align 4 +// CHECK-PPC64LE-NEXT: [[INC4:%.*]] = add nsw i32 [[TMP17]], 1 +// CHECK-PPC64LE-NEXT: store i32 [[INC4]], i32* [[ARG_ADDR]], align 4 +// CHECK-PPC64LE-NEXT: [[TMP18:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS1]], i32 0, i32 0 +// CHECK-PPC64LE-NEXT: [[TMP19:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS2]], i32 0, i32 0 +// CHECK-PPC64LE-NEXT: call void @__tgt_target_data_end_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i32 1, i8** [[TMP18]], i8** [[TMP19]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes.1, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.2, i32 0, i32 0), i8** null, i8** null) +// CHECK-PPC64LE-NEXT: [[TMP20:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS5]], i32 0, i32 0 +// CHECK-PPC64LE-NEXT: [[TMP21:%.*]] = bitcast i8** [[TMP20]] to i32** +// CHECK-PPC64LE-NEXT: store i32* [[ARG_ADDR]], i32** [[TMP21]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP22:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS6]], i32 0, i32 0 +// CHECK-PPC64LE-NEXT: [[TMP23:%.*]] = bitcast i8** [[TMP22]] to i32** +// CHECK-PPC64LE-NEXT: store i32* [[ARG_ADDR]], i32** [[TMP23]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP24:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_MAPPERS7]], i64 0, i64 0 +// CHECK-PPC64LE-NEXT: store i8* null, i8** [[TMP24]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP25:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS5]], i32 0, i32 0 +// CHECK-PPC64LE-NEXT: [[TMP26:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS6]], i32 0, i32 0 +// CHECK-PPC64LE-NEXT: call void @__tgt_target_data_begin_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i32 1, i8** [[TMP25]], i8** [[TMP26]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes.3, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.4, i32 0, i32 0), i8** null, i8** null) +// CHECK-PPC64LE-NEXT: [[TMP27:%.*]] = load i32, i32* [[ARG_ADDR]], align 4 +// CHECK-PPC64LE-NEXT: [[INC8:%.*]] = add nsw i32 [[TMP27]], 1 +// CHECK-PPC64LE-NEXT: store i32 [[INC8]], i32* [[ARG_ADDR]], align 4 +// CHECK-PPC64LE-NEXT: [[TMP28:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS5]], i32 0, i32 0 +// CHECK-PPC64LE-NEXT: [[TMP29:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS6]], i32 0, i32 0 +// CHECK-PPC64LE-NEXT: call void @__tgt_target_data_end_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i32 1, i8** [[TMP28]], i8** [[TMP29]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes.3, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.4, i32 0, i32 0), i8** null, i8** null) +// CHECK-PPC64LE-NEXT: [[TMP30:%.*]] = load %struct.S2*, %struct.S2** [[PS1]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP31:%.*]] = load %struct.S2*, %struct.S2** [[PS1]], align 8 +// CHECK-PPC64LE-NEXT: [[S:%.*]] = getelementptr inbounds [[STRUCT_S2:%.*]], %struct.S2* [[TMP31]], i32 0, i32 0 +// CHECK-PPC64LE-NEXT: [[TMP32:%.*]] = load %struct.S2*, %struct.S2** [[PS1]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP33:%.*]] = load %struct.S2*, %struct.S2** [[PS1]], align 8 +// CHECK-PPC64LE-NEXT: [[PS:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP33]], i32 0, i32 1 +// CHECK-PPC64LE-NEXT: [[TMP34:%.*]] = load %struct.S2*, %struct.S2** [[PS1]], align 8 +// CHECK-PPC64LE-NEXT: [[PS9:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP34]], i32 0, i32 1 +// CHECK-PPC64LE-NEXT: [[TMP35:%.*]] = load %struct.S2*, %struct.S2** [[PS9]], align 8 +// CHECK-PPC64LE-NEXT: [[PS10:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP35]], i32 0, i32 1 +// CHECK-PPC64LE-NEXT: [[TMP36:%.*]] = load %struct.S2*, %struct.S2** [[PS1]], align 8 +// CHECK-PPC64LE-NEXT: [[PS11:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP36]], i32 0, i32 1 +// CHECK-PPC64LE-NEXT: [[TMP37:%.*]] = load %struct.S2*, %struct.S2** [[PS11]], align 8 +// CHECK-PPC64LE-NEXT: [[PS12:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP37]], i32 0, i32 1 +// CHECK-PPC64LE-NEXT: [[TMP38:%.*]] = load %struct.S2*, %struct.S2** [[PS12]], align 8 +// CHECK-PPC64LE-NEXT: [[PS13:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP38]], i32 0, i32 1 +// CHECK-PPC64LE-NEXT: [[TMP39:%.*]] = load %struct.S2*, %struct.S2** [[PS1]], align 8 +// CHECK-PPC64LE-NEXT: [[PS14:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP39]], i32 0, i32 1 +// CHECK-PPC64LE-NEXT: [[TMP40:%.*]] = load %struct.S2*, %struct.S2** [[PS14]], align 8 +// CHECK-PPC64LE-NEXT: [[PS15:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP40]], i32 0, i32 1 +// CHECK-PPC64LE-NEXT: [[TMP41:%.*]] = load %struct.S2*, %struct.S2** [[PS15]], align 8 +// CHECK-PPC64LE-NEXT: [[PS16:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP41]], i32 0, i32 1 +// CHECK-PPC64LE-NEXT: [[TMP42:%.*]] = load %struct.S2*, %struct.S2** [[PS16]], align 8 +// CHECK-PPC64LE-NEXT: [[S17:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP42]], i32 0, i32 0 +// CHECK-PPC64LE-NEXT: [[TMP43:%.*]] = getelementptr %struct.S2*, %struct.S2** [[PS]], i32 1 +// CHECK-PPC64LE-NEXT: [[TMP44:%.*]] = bitcast %struct.S1* [[S]] to i8* +// CHECK-PPC64LE-NEXT: [[TMP45:%.*]] = bitcast %struct.S2** [[TMP43]] to i8* +// CHECK-PPC64LE-NEXT: [[TMP46:%.*]] = ptrtoint i8* [[TMP45]] to i64 +// CHECK-PPC64LE-NEXT: [[TMP47:%.*]] = ptrtoint i8* [[TMP44]] to i64 +// CHECK-PPC64LE-NEXT: [[TMP48:%.*]] = sub i64 [[TMP46]], [[TMP47]] +// CHECK-PPC64LE-NEXT: [[TMP49:%.*]] = sdiv exact i64 [[TMP48]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) +// CHECK-PPC64LE-NEXT: [[TMP50:%.*]] = load %struct.S2*, %struct.S2** [[PS2]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP51:%.*]] = load %struct.S2*, %struct.S2** [[PS2]], align 8 +// CHECK-PPC64LE-NEXT: [[S18:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP51]], i32 0, i32 0 +// CHECK-PPC64LE-NEXT: [[TMP52:%.*]] = load %struct.S2*, %struct.S2** [[PS2]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP53:%.*]] = load %struct.S2*, %struct.S2** [[PS2]], align 8 +// CHECK-PPC64LE-NEXT: [[PS19:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP53]], i32 0, i32 1 +// CHECK-PPC64LE-NEXT: [[TMP54:%.*]] = load %struct.S2*, %struct.S2** [[PS2]], align 8 +// CHECK-PPC64LE-NEXT: [[PS20:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP54]], i32 0, i32 1 +// CHECK-PPC64LE-NEXT: [[TMP55:%.*]] = load %struct.S2*, %struct.S2** [[PS20]], align 8 +// CHECK-PPC64LE-NEXT: [[PS21:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP55]], i32 0, i32 1 +// CHECK-PPC64LE-NEXT: [[TMP56:%.*]] = load %struct.S2*, %struct.S2** [[PS2]], align 8 +// CHECK-PPC64LE-NEXT: [[PS22:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP56]], i32 0, i32 1 +// CHECK-PPC64LE-NEXT: [[TMP57:%.*]] = load %struct.S2*, %struct.S2** [[PS22]], align 8 +// CHECK-PPC64LE-NEXT: [[PS23:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP57]], i32 0, i32 1 +// CHECK-PPC64LE-NEXT: [[TMP58:%.*]] = load %struct.S2*, %struct.S2** [[PS23]], align 8 +// CHECK-PPC64LE-NEXT: [[PS24:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP58]], i32 0, i32 1 +// CHECK-PPC64LE-NEXT: [[TMP59:%.*]] = load %struct.S2*, %struct.S2** [[PS2]], align 8 +// CHECK-PPC64LE-NEXT: [[PS25:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP59]], i32 0, i32 1 +// CHECK-PPC64LE-NEXT: [[TMP60:%.*]] = load %struct.S2*, %struct.S2** [[PS25]], align 8 +// CHECK-PPC64LE-NEXT: [[PS26:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP60]], i32 0, i32 1 +// CHECK-PPC64LE-NEXT: [[TMP61:%.*]] = load %struct.S2*, %struct.S2** [[PS26]], align 8 +// CHECK-PPC64LE-NEXT: [[PS27:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP61]], i32 0, i32 1 +// CHECK-PPC64LE-NEXT: [[TMP62:%.*]] = load %struct.S2*, %struct.S2** [[PS27]], align 8 +// CHECK-PPC64LE-NEXT: [[S28:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP62]], i32 0, i32 0 +// CHECK-PPC64LE-NEXT: [[TMP63:%.*]] = getelementptr %struct.S2*, %struct.S2** [[PS19]], i32 1 +// CHECK-PPC64LE-NEXT: [[TMP64:%.*]] = bitcast %struct.S1* [[S18]] to i8* +// CHECK-PPC64LE-NEXT: [[TMP65:%.*]] = bitcast %struct.S2** [[TMP63]] to i8* +// CHECK-PPC64LE-NEXT: [[TMP66:%.*]] = ptrtoint i8* [[TMP65]] to i64 +// CHECK-PPC64LE-NEXT: [[TMP67:%.*]] = ptrtoint i8* [[TMP64]] to i64 +// CHECK-PPC64LE-NEXT: [[TMP68:%.*]] = sub i64 [[TMP66]], [[TMP67]] +// CHECK-PPC64LE-NEXT: [[TMP69:%.*]] = sdiv exact i64 [[TMP68]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) +// CHECK-PPC64LE-NEXT: [[TMP70:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 0 +// CHECK-PPC64LE-NEXT: [[TMP71:%.*]] = bitcast i8** [[TMP70]] to %struct.S2** +// CHECK-PPC64LE-NEXT: store %struct.S2* [[TMP30]], %struct.S2** [[TMP71]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP72:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 0 +// CHECK-PPC64LE-NEXT: [[TMP73:%.*]] = bitcast i8** [[TMP72]] to %struct.S1** +// CHECK-PPC64LE-NEXT: store %struct.S1* [[S]], %struct.S1** [[TMP73]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP74:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0 +// CHECK-PPC64LE-NEXT: store i64 [[TMP49]], i64* [[TMP74]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP75:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i64 0, i64 0 +// CHECK-PPC64LE-NEXT: store i8* null, i8** [[TMP75]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP76:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 1 +// CHECK-PPC64LE-NEXT: [[TMP77:%.*]] = bitcast i8** [[TMP76]] to %struct.S2** +// CHECK-PPC64LE-NEXT: store %struct.S2* [[TMP30]], %struct.S2** [[TMP77]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP78:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 1 +// CHECK-PPC64LE-NEXT: [[TMP79:%.*]] = bitcast i8** [[TMP78]] to %struct.S1** +// CHECK-PPC64LE-NEXT: store %struct.S1* [[S]], %struct.S1** [[TMP79]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP80:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 1 +// CHECK-PPC64LE-NEXT: store i64 4, i64* [[TMP80]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP81:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i64 0, i64 1 +// CHECK-PPC64LE-NEXT: store i8* null, i8** [[TMP81]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP82:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 2 +// CHECK-PPC64LE-NEXT: [[TMP83:%.*]] = bitcast i8** [[TMP82]] to %struct.S2*** +// CHECK-PPC64LE-NEXT: store %struct.S2** [[PS]], %struct.S2*** [[TMP83]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP84:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 2 +// CHECK-PPC64LE-NEXT: [[TMP85:%.*]] = bitcast i8** [[TMP84]] to %struct.S2*** +// CHECK-PPC64LE-NEXT: store %struct.S2** [[PS10]], %struct.S2*** [[TMP85]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP86:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 2 +// CHECK-PPC64LE-NEXT: store i64 8, i64* [[TMP86]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP87:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i64 0, i64 2 +// CHECK-PPC64LE-NEXT: store i8* null, i8** [[TMP87]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP88:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 3 +// CHECK-PPC64LE-NEXT: [[TMP89:%.*]] = bitcast i8** [[TMP88]] to %struct.S2*** +// CHECK-PPC64LE-NEXT: store %struct.S2** [[PS10]], %struct.S2*** [[TMP89]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP90:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 3 +// CHECK-PPC64LE-NEXT: [[TMP91:%.*]] = bitcast i8** [[TMP90]] to %struct.S2*** +// CHECK-PPC64LE-NEXT: store %struct.S2** [[PS13]], %struct.S2*** [[TMP91]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP92:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 3 +// CHECK-PPC64LE-NEXT: store i64 8, i64* [[TMP92]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP93:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i64 0, i64 3 +// CHECK-PPC64LE-NEXT: store i8* null, i8** [[TMP93]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP94:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 4 +// CHECK-PPC64LE-NEXT: [[TMP95:%.*]] = bitcast i8** [[TMP94]] to %struct.S2*** +// CHECK-PPC64LE-NEXT: store %struct.S2** [[PS13]], %struct.S2*** [[TMP95]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP96:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 4 +// CHECK-PPC64LE-NEXT: [[TMP97:%.*]] = bitcast i8** [[TMP96]] to %struct.S1** +// CHECK-PPC64LE-NEXT: store %struct.S1* [[S17]], %struct.S1** [[TMP97]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP98:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 4 +// CHECK-PPC64LE-NEXT: store i64 4, i64* [[TMP98]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP99:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i64 0, i64 4 +// CHECK-PPC64LE-NEXT: store i8* null, i8** [[TMP99]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP100:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 5 +// CHECK-PPC64LE-NEXT: [[TMP101:%.*]] = bitcast i8** [[TMP100]] to i32** +// CHECK-PPC64LE-NEXT: store i32* [[ARG_ADDR]], i32** [[TMP101]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP102:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 5 +// CHECK-PPC64LE-NEXT: [[TMP103:%.*]] = bitcast i8** [[TMP102]] to i32** +// CHECK-PPC64LE-NEXT: store i32* [[ARG_ADDR]], i32** [[TMP103]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP104:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 5 +// CHECK-PPC64LE-NEXT: store i64 4, i64* [[TMP104]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP105:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i64 0, i64 5 +// CHECK-PPC64LE-NEXT: store i8* null, i8** [[TMP105]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP106:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 6 +// CHECK-PPC64LE-NEXT: [[TMP107:%.*]] = bitcast i8** [[TMP106]] to %struct.S2** +// CHECK-PPC64LE-NEXT: store %struct.S2* [[TMP50]], %struct.S2** [[TMP107]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP108:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 6 +// CHECK-PPC64LE-NEXT: [[TMP109:%.*]] = bitcast i8** [[TMP108]] to %struct.S1** +// CHECK-PPC64LE-NEXT: store %struct.S1* [[S18]], %struct.S1** [[TMP109]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP110:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 6 +// CHECK-PPC64LE-NEXT: store i64 [[TMP69]], i64* [[TMP110]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP111:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i64 0, i64 6 +// CHECK-PPC64LE-NEXT: store i8* null, i8** [[TMP111]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP112:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 7 +// CHECK-PPC64LE-NEXT: [[TMP113:%.*]] = bitcast i8** [[TMP112]] to %struct.S2** +// CHECK-PPC64LE-NEXT: store %struct.S2* [[TMP50]], %struct.S2** [[TMP113]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP114:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 7 +// CHECK-PPC64LE-NEXT: [[TMP115:%.*]] = bitcast i8** [[TMP114]] to %struct.S1** +// CHECK-PPC64LE-NEXT: store %struct.S1* [[S18]], %struct.S1** [[TMP115]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP116:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 7 +// CHECK-PPC64LE-NEXT: store i64 4, i64* [[TMP116]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP117:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i64 0, i64 7 +// CHECK-PPC64LE-NEXT: store i8* null, i8** [[TMP117]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP118:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 8 +// CHECK-PPC64LE-NEXT: [[TMP119:%.*]] = bitcast i8** [[TMP118]] to %struct.S2*** +// CHECK-PPC64LE-NEXT: store %struct.S2** [[PS19]], %struct.S2*** [[TMP119]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP120:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 8 +// CHECK-PPC64LE-NEXT: [[TMP121:%.*]] = bitcast i8** [[TMP120]] to %struct.S2*** +// CHECK-PPC64LE-NEXT: store %struct.S2** [[PS21]], %struct.S2*** [[TMP121]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP122:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 8 +// CHECK-PPC64LE-NEXT: store i64 8, i64* [[TMP122]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP123:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i64 0, i64 8 +// CHECK-PPC64LE-NEXT: store i8* null, i8** [[TMP123]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP124:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 9 +// CHECK-PPC64LE-NEXT: [[TMP125:%.*]] = bitcast i8** [[TMP124]] to %struct.S2*** +// CHECK-PPC64LE-NEXT: store %struct.S2** [[PS21]], %struct.S2*** [[TMP125]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP126:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 9 +// CHECK-PPC64LE-NEXT: [[TMP127:%.*]] = bitcast i8** [[TMP126]] to %struct.S2*** +// CHECK-PPC64LE-NEXT: store %struct.S2** [[PS24]], %struct.S2*** [[TMP127]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP128:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 9 +// CHECK-PPC64LE-NEXT: store i64 8, i64* [[TMP128]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP129:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i64 0, i64 9 +// CHECK-PPC64LE-NEXT: store i8* null, i8** [[TMP129]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP130:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 10 +// CHECK-PPC64LE-NEXT: [[TMP131:%.*]] = bitcast i8** [[TMP130]] to %struct.S2*** +// CHECK-PPC64LE-NEXT: store %struct.S2** [[PS24]], %struct.S2*** [[TMP131]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP132:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 10 +// CHECK-PPC64LE-NEXT: [[TMP133:%.*]] = bitcast i8** [[TMP132]] to %struct.S1** +// CHECK-PPC64LE-NEXT: store %struct.S1* [[S28]], %struct.S1** [[TMP133]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP134:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 10 +// CHECK-PPC64LE-NEXT: store i64 4, i64* [[TMP134]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP135:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i64 0, i64 10 +// CHECK-PPC64LE-NEXT: store i8* null, i8** [[TMP135]], align 8 +// CHECK-PPC64LE-NEXT: [[TMP136:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 0 +// CHECK-PPC64LE-NEXT: [[TMP137:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 0 +// CHECK-PPC64LE-NEXT: [[TMP138:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0 +// CHECK-PPC64LE-NEXT: call void @__tgt_target_data_begin_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i32 11, i8** [[TMP136]], i8** [[TMP137]], i64* [[TMP138]], i64* getelementptr inbounds ([11 x i64], [11 x i64]* @.offload_maptypes.5, i32 0, i32 0), i8** null, i8** null) +// CHECK-PPC64LE-NEXT: [[TMP139:%.*]] = load i32, i32* [[ARG_ADDR]], align 4 +// CHECK-PPC64LE-NEXT: [[INC32:%.*]] = add nsw i32 [[TMP139]], 1 +// CHECK-PPC64LE-NEXT: store i32 [[INC32]], i32* [[ARG_ADDR]], align 4 +// CHECK-PPC64LE-NEXT: [[TMP140:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 0 +// CHECK-PPC64LE-NEXT: [[TMP141:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 0 +// CHECK-PPC64LE-NEXT: [[TMP142:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0 +// CHECK-PPC64LE-NEXT: call void @__tgt_target_data_end_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i32 11, i8** [[TMP140]], i8** [[TMP141]], i64* [[TMP142]], i64* getelementptr inbounds ([11 x i64], [11 x i64]* @.offload_maptypes.5, i32 0, i32 0), i8** null, i8** null) +// CHECK-PPC64LE-NEXT: ret void +// +// CHECK-I386-LABEL: @_Z3fooi( +// CHECK-I386-NEXT: entry: +// CHECK-I386-NEXT: [[ARG_ADDR:%.*]] = alloca i32, align 4 +// CHECK-I386-NEXT: [[LB:%.*]] = alloca [5 x float], align 4 +// CHECK-I386-NEXT: [[PS1:%.*]] = alloca %struct.S2*, align 4 +// CHECK-I386-NEXT: [[PS2:%.*]] = alloca %struct.S2*, align 4 +// CHECK-I386-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x i8*], align 4 +// CHECK-I386-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x i8*], align 4 +// CHECK-I386-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x i8*], align 4 +// CHECK-I386-NEXT: [[DOTOFFLOAD_BASEPTRS1:%.*]] = alloca [1 x i8*], align 4 +// CHECK-I386-NEXT: [[DOTOFFLOAD_PTRS2:%.*]] = alloca [1 x i8*], align 4 +// CHECK-I386-NEXT: [[DOTOFFLOAD_MAPPERS3:%.*]] = alloca [1 x i8*], align 4 +// CHECK-I386-NEXT: [[DOTOFFLOAD_BASEPTRS5:%.*]] = alloca [1 x i8*], align 4 +// CHECK-I386-NEXT: [[DOTOFFLOAD_PTRS6:%.*]] = alloca [1 x i8*], align 4 +// CHECK-I386-NEXT: [[DOTOFFLOAD_MAPPERS7:%.*]] = alloca [1 x i8*], align 4 +// CHECK-I386-NEXT: [[DOTOFFLOAD_BASEPTRS29:%.*]] = alloca [11 x i8*], align 4 +// CHECK-I386-NEXT: [[DOTOFFLOAD_PTRS30:%.*]] = alloca [11 x i8*], align 4 +// CHECK-I386-NEXT: [[DOTOFFLOAD_MAPPERS31:%.*]] = alloca [11 x i8*], align 4 +// CHECK-I386-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [11 x i64], align 4 +// CHECK-I386-NEXT: store i32 [[ARG:%.*]], i32* [[ARG_ADDR]], align 4 +// CHECK-I386-NEXT: [[TMP0:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK-I386-NEXT: [[TMP1:%.*]] = bitcast i8** [[TMP0]] to [5 x float]** +// CHECK-I386-NEXT: store [5 x float]* [[LB]], [5 x float]** [[TMP1]], align 4 +// CHECK-I386-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK-I386-NEXT: [[TMP3:%.*]] = bitcast i8** [[TMP2]] to [5 x float]** +// CHECK-I386-NEXT: store [5 x float]* [[LB]], [5 x float]** [[TMP3]], align 4 +// CHECK-I386-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0 +// CHECK-I386-NEXT: store i8* null, i8** [[TMP4]], align 4 +// CHECK-I386-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK-I386-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK-I386-NEXT: call void @__tgt_target_data_begin_mapper(%struct.ident_t* @[[GLOB1:[0-9]+]], i64 -1, i32 1, i8** [[TMP5]], i8** [[TMP6]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null) +// CHECK-I386-NEXT: [[TMP7:%.*]] = load i32, i32* [[ARG_ADDR]], align 4 +// CHECK-I386-NEXT: [[INC:%.*]] = add nsw i32 [[TMP7]], 1 +// CHECK-I386-NEXT: store i32 [[INC]], i32* [[ARG_ADDR]], align 4 +// CHECK-I386-NEXT: [[TMP8:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK-I386-NEXT: [[TMP9:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK-I386-NEXT: call void @__tgt_target_data_end_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i32 1, i8** [[TMP8]], i8** [[TMP9]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null) +// CHECK-I386-NEXT: [[TMP10:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS1]], i32 0, i32 0 +// CHECK-I386-NEXT: [[TMP11:%.*]] = bitcast i8** [[TMP10]] to [5 x float]** +// CHECK-I386-NEXT: store [5 x float]* [[LB]], [5 x float]** [[TMP11]], align 4 +// CHECK-I386-NEXT: [[TMP12:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS2]], i32 0, i32 0 +// CHECK-I386-NEXT: [[TMP13:%.*]] = bitcast i8** [[TMP12]] to [5 x float]** +// CHECK-I386-NEXT: store [5 x float]* [[LB]], [5 x float]** [[TMP13]], align 4 +// CHECK-I386-NEXT: [[TMP14:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_MAPPERS3]], i32 0, i32 0 +// CHECK-I386-NEXT: store i8* null, i8** [[TMP14]], align 4 +// CHECK-I386-NEXT: [[TMP15:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS1]], i32 0, i32 0 +// CHECK-I386-NEXT: [[TMP16:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS2]], i32 0, i32 0 +// CHECK-I386-NEXT: call void @__tgt_target_data_begin_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i32 1, i8** [[TMP15]], i8** [[TMP16]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes.1, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.2, i32 0, i32 0), i8** null, i8** null) +// CHECK-I386-NEXT: [[TMP17:%.*]] = load i32, i32* [[ARG_ADDR]], align 4 +// CHECK-I386-NEXT: [[INC4:%.*]] = add nsw i32 [[TMP17]], 1 +// CHECK-I386-NEXT: store i32 [[INC4]], i32* [[ARG_ADDR]], align 4 +// CHECK-I386-NEXT: [[TMP18:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS1]], i32 0, i32 0 +// CHECK-I386-NEXT: [[TMP19:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS2]], i32 0, i32 0 +// CHECK-I386-NEXT: call void @__tgt_target_data_end_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i32 1, i8** [[TMP18]], i8** [[TMP19]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes.1, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.2, i32 0, i32 0), i8** null, i8** null) +// CHECK-I386-NEXT: [[TMP20:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS5]], i32 0, i32 0 +// CHECK-I386-NEXT: [[TMP21:%.*]] = bitcast i8** [[TMP20]] to i32** +// CHECK-I386-NEXT: store i32* [[ARG_ADDR]], i32** [[TMP21]], align 4 +// CHECK-I386-NEXT: [[TMP22:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS6]], i32 0, i32 0 +// CHECK-I386-NEXT: [[TMP23:%.*]] = bitcast i8** [[TMP22]] to i32** +// CHECK-I386-NEXT: store i32* [[ARG_ADDR]], i32** [[TMP23]], align 4 +// CHECK-I386-NEXT: [[TMP24:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_MAPPERS7]], i32 0, i32 0 +// CHECK-I386-NEXT: store i8* null, i8** [[TMP24]], align 4 +// CHECK-I386-NEXT: [[TMP25:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS5]], i32 0, i32 0 +// CHECK-I386-NEXT: [[TMP26:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS6]], i32 0, i32 0 +// CHECK-I386-NEXT: call void @__tgt_target_data_begin_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i32 1, i8** [[TMP25]], i8** [[TMP26]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes.3, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.4, i32 0, i32 0), i8** null, i8** null) +// CHECK-I386-NEXT: [[TMP27:%.*]] = load i32, i32* [[ARG_ADDR]], align 4 +// CHECK-I386-NEXT: [[INC8:%.*]] = add nsw i32 [[TMP27]], 1 +// CHECK-I386-NEXT: store i32 [[INC8]], i32* [[ARG_ADDR]], align 4 +// CHECK-I386-NEXT: [[TMP28:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS5]], i32 0, i32 0 +// CHECK-I386-NEXT: [[TMP29:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS6]], i32 0, i32 0 +// CHECK-I386-NEXT: call void @__tgt_target_data_end_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i32 1, i8** [[TMP28]], i8** [[TMP29]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes.3, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.4, i32 0, i32 0), i8** null, i8** null) +// CHECK-I386-NEXT: [[TMP30:%.*]] = load %struct.S2*, %struct.S2** [[PS1]], align 4 +// CHECK-I386-NEXT: [[TMP31:%.*]] = load %struct.S2*, %struct.S2** [[PS1]], align 4 +// CHECK-I386-NEXT: [[S:%.*]] = getelementptr inbounds [[STRUCT_S2:%.*]], %struct.S2* [[TMP31]], i32 0, i32 0 +// CHECK-I386-NEXT: [[TMP32:%.*]] = load %struct.S2*, %struct.S2** [[PS1]], align 4 +// CHECK-I386-NEXT: [[TMP33:%.*]] = load %struct.S2*, %struct.S2** [[PS1]], align 4 +// CHECK-I386-NEXT: [[PS:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP33]], i32 0, i32 1 +// CHECK-I386-NEXT: [[TMP34:%.*]] = load %struct.S2*, %struct.S2** [[PS1]], align 4 +// CHECK-I386-NEXT: [[PS9:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP34]], i32 0, i32 1 +// CHECK-I386-NEXT: [[TMP35:%.*]] = load %struct.S2*, %struct.S2** [[PS9]], align 4 +// CHECK-I386-NEXT: [[PS10:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP35]], i32 0, i32 1 +// CHECK-I386-NEXT: [[TMP36:%.*]] = load %struct.S2*, %struct.S2** [[PS1]], align 4 +// CHECK-I386-NEXT: [[PS11:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP36]], i32 0, i32 1 +// CHECK-I386-NEXT: [[TMP37:%.*]] = load %struct.S2*, %struct.S2** [[PS11]], align 4 +// CHECK-I386-NEXT: [[PS12:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP37]], i32 0, i32 1 +// CHECK-I386-NEXT: [[TMP38:%.*]] = load %struct.S2*, %struct.S2** [[PS12]], align 4 +// CHECK-I386-NEXT: [[PS13:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP38]], i32 0, i32 1 +// CHECK-I386-NEXT: [[TMP39:%.*]] = load %struct.S2*, %struct.S2** [[PS1]], align 4 +// CHECK-I386-NEXT: [[PS14:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP39]], i32 0, i32 1 +// CHECK-I386-NEXT: [[TMP40:%.*]] = load %struct.S2*, %struct.S2** [[PS14]], align 4 +// CHECK-I386-NEXT: [[PS15:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP40]], i32 0, i32 1 +// CHECK-I386-NEXT: [[TMP41:%.*]] = load %struct.S2*, %struct.S2** [[PS15]], align 4 +// CHECK-I386-NEXT: [[PS16:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP41]], i32 0, i32 1 +// CHECK-I386-NEXT: [[TMP42:%.*]] = load %struct.S2*, %struct.S2** [[PS16]], align 4 +// CHECK-I386-NEXT: [[S17:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP42]], i32 0, i32 0 +// CHECK-I386-NEXT: [[TMP43:%.*]] = getelementptr %struct.S2*, %struct.S2** [[PS]], i32 1 +// CHECK-I386-NEXT: [[TMP44:%.*]] = bitcast %struct.S1* [[S]] to i8* +// CHECK-I386-NEXT: [[TMP45:%.*]] = bitcast %struct.S2** [[TMP43]] to i8* +// CHECK-I386-NEXT: [[TMP46:%.*]] = ptrtoint i8* [[TMP45]] to i64 +// CHECK-I386-NEXT: [[TMP47:%.*]] = ptrtoint i8* [[TMP44]] to i64 +// CHECK-I386-NEXT: [[TMP48:%.*]] = sub i64 [[TMP46]], [[TMP47]] +// CHECK-I386-NEXT: [[TMP49:%.*]] = sdiv exact i64 [[TMP48]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) +// CHECK-I386-NEXT: [[TMP50:%.*]] = load %struct.S2*, %struct.S2** [[PS2]], align 4 +// CHECK-I386-NEXT: [[TMP51:%.*]] = load %struct.S2*, %struct.S2** [[PS2]], align 4 +// CHECK-I386-NEXT: [[S18:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP51]], i32 0, i32 0 +// CHECK-I386-NEXT: [[TMP52:%.*]] = load %struct.S2*, %struct.S2** [[PS2]], align 4 +// CHECK-I386-NEXT: [[TMP53:%.*]] = load %struct.S2*, %struct.S2** [[PS2]], align 4 +// CHECK-I386-NEXT: [[PS19:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP53]], i32 0, i32 1 +// CHECK-I386-NEXT: [[TMP54:%.*]] = load %struct.S2*, %struct.S2** [[PS2]], align 4 +// CHECK-I386-NEXT: [[PS20:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP54]], i32 0, i32 1 +// CHECK-I386-NEXT: [[TMP55:%.*]] = load %struct.S2*, %struct.S2** [[PS20]], align 4 +// CHECK-I386-NEXT: [[PS21:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP55]], i32 0, i32 1 +// CHECK-I386-NEXT: [[TMP56:%.*]] = load %struct.S2*, %struct.S2** [[PS2]], align 4 +// CHECK-I386-NEXT: [[PS22:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP56]], i32 0, i32 1 +// CHECK-I386-NEXT: [[TMP57:%.*]] = load %struct.S2*, %struct.S2** [[PS22]], align 4 +// CHECK-I386-NEXT: [[PS23:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP57]], i32 0, i32 1 +// CHECK-I386-NEXT: [[TMP58:%.*]] = load %struct.S2*, %struct.S2** [[PS23]], align 4 +// CHECK-I386-NEXT: [[PS24:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP58]], i32 0, i32 1 +// CHECK-I386-NEXT: [[TMP59:%.*]] = load %struct.S2*, %struct.S2** [[PS2]], align 4 +// CHECK-I386-NEXT: [[PS25:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP59]], i32 0, i32 1 +// CHECK-I386-NEXT: [[TMP60:%.*]] = load %struct.S2*, %struct.S2** [[PS25]], align 4 +// CHECK-I386-NEXT: [[PS26:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP60]], i32 0, i32 1 +// CHECK-I386-NEXT: [[TMP61:%.*]] = load %struct.S2*, %struct.S2** [[PS26]], align 4 +// CHECK-I386-NEXT: [[PS27:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP61]], i32 0, i32 1 +// CHECK-I386-NEXT: [[TMP62:%.*]] = load %struct.S2*, %struct.S2** [[PS27]], align 4 +// CHECK-I386-NEXT: [[S28:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP62]], i32 0, i32 0 +// CHECK-I386-NEXT: [[TMP63:%.*]] = getelementptr %struct.S2*, %struct.S2** [[PS19]], i32 1 +// CHECK-I386-NEXT: [[TMP64:%.*]] = bitcast %struct.S1* [[S18]] to i8* +// CHECK-I386-NEXT: [[TMP65:%.*]] = bitcast %struct.S2** [[TMP63]] to i8* +// CHECK-I386-NEXT: [[TMP66:%.*]] = ptrtoint i8* [[TMP65]] to i64 +// CHECK-I386-NEXT: [[TMP67:%.*]] = ptrtoint i8* [[TMP64]] to i64 +// CHECK-I386-NEXT: [[TMP68:%.*]] = sub i64 [[TMP66]], [[TMP67]] +// CHECK-I386-NEXT: [[TMP69:%.*]] = sdiv exact i64 [[TMP68]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) +// CHECK-I386-NEXT: [[TMP70:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 0 +// CHECK-I386-NEXT: [[TMP71:%.*]] = bitcast i8** [[TMP70]] to %struct.S2** +// CHECK-I386-NEXT: store %struct.S2* [[TMP30]], %struct.S2** [[TMP71]], align 4 +// CHECK-I386-NEXT: [[TMP72:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 0 +// CHECK-I386-NEXT: [[TMP73:%.*]] = bitcast i8** [[TMP72]] to %struct.S1** +// CHECK-I386-NEXT: store %struct.S1* [[S]], %struct.S1** [[TMP73]], align 4 +// CHECK-I386-NEXT: [[TMP74:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0 +// CHECK-I386-NEXT: store i64 [[TMP49]], i64* [[TMP74]], align 4 +// CHECK-I386-NEXT: [[TMP75:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i32 0, i32 0 +// CHECK-I386-NEXT: store i8* null, i8** [[TMP75]], align 4 +// CHECK-I386-NEXT: [[TMP76:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 1 +// CHECK-I386-NEXT: [[TMP77:%.*]] = bitcast i8** [[TMP76]] to %struct.S2** +// CHECK-I386-NEXT: store %struct.S2* [[TMP30]], %struct.S2** [[TMP77]], align 4 +// CHECK-I386-NEXT: [[TMP78:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 1 +// CHECK-I386-NEXT: [[TMP79:%.*]] = bitcast i8** [[TMP78]] to %struct.S1** +// CHECK-I386-NEXT: store %struct.S1* [[S]], %struct.S1** [[TMP79]], align 4 +// CHECK-I386-NEXT: [[TMP80:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 1 +// CHECK-I386-NEXT: store i64 4, i64* [[TMP80]], align 4 +// CHECK-I386-NEXT: [[TMP81:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i32 0, i32 1 +// CHECK-I386-NEXT: store i8* null, i8** [[TMP81]], align 4 +// CHECK-I386-NEXT: [[TMP82:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 2 +// CHECK-I386-NEXT: [[TMP83:%.*]] = bitcast i8** [[TMP82]] to %struct.S2*** +// CHECK-I386-NEXT: store %struct.S2** [[PS]], %struct.S2*** [[TMP83]], align 4 +// CHECK-I386-NEXT: [[TMP84:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 2 +// CHECK-I386-NEXT: [[TMP85:%.*]] = bitcast i8** [[TMP84]] to %struct.S2*** +// CHECK-I386-NEXT: store %struct.S2** [[PS10]], %struct.S2*** [[TMP85]], align 4 +// CHECK-I386-NEXT: [[TMP86:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 2 +// CHECK-I386-NEXT: store i64 4, i64* [[TMP86]], align 4 +// CHECK-I386-NEXT: [[TMP87:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i32 0, i32 2 +// CHECK-I386-NEXT: store i8* null, i8** [[TMP87]], align 4 +// CHECK-I386-NEXT: [[TMP88:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 3 +// CHECK-I386-NEXT: [[TMP89:%.*]] = bitcast i8** [[TMP88]] to %struct.S2*** +// CHECK-I386-NEXT: store %struct.S2** [[PS10]], %struct.S2*** [[TMP89]], align 4 +// CHECK-I386-NEXT: [[TMP90:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 3 +// CHECK-I386-NEXT: [[TMP91:%.*]] = bitcast i8** [[TMP90]] to %struct.S2*** +// CHECK-I386-NEXT: store %struct.S2** [[PS13]], %struct.S2*** [[TMP91]], align 4 +// CHECK-I386-NEXT: [[TMP92:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 3 +// CHECK-I386-NEXT: store i64 4, i64* [[TMP92]], align 4 +// CHECK-I386-NEXT: [[TMP93:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i32 0, i32 3 +// CHECK-I386-NEXT: store i8* null, i8** [[TMP93]], align 4 +// CHECK-I386-NEXT: [[TMP94:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 4 +// CHECK-I386-NEXT: [[TMP95:%.*]] = bitcast i8** [[TMP94]] to %struct.S2*** +// CHECK-I386-NEXT: store %struct.S2** [[PS13]], %struct.S2*** [[TMP95]], align 4 +// CHECK-I386-NEXT: [[TMP96:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 4 +// CHECK-I386-NEXT: [[TMP97:%.*]] = bitcast i8** [[TMP96]] to %struct.S1** +// CHECK-I386-NEXT: store %struct.S1* [[S17]], %struct.S1** [[TMP97]], align 4 +// CHECK-I386-NEXT: [[TMP98:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 4 +// CHECK-I386-NEXT: store i64 4, i64* [[TMP98]], align 4 +// CHECK-I386-NEXT: [[TMP99:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i32 0, i32 4 +// CHECK-I386-NEXT: store i8* null, i8** [[TMP99]], align 4 +// CHECK-I386-NEXT: [[TMP100:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 5 +// CHECK-I386-NEXT: [[TMP101:%.*]] = bitcast i8** [[TMP100]] to i32** +// CHECK-I386-NEXT: store i32* [[ARG_ADDR]], i32** [[TMP101]], align 4 +// CHECK-I386-NEXT: [[TMP102:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 5 +// CHECK-I386-NEXT: [[TMP103:%.*]] = bitcast i8** [[TMP102]] to i32** +// CHECK-I386-NEXT: store i32* [[ARG_ADDR]], i32** [[TMP103]], align 4 +// CHECK-I386-NEXT: [[TMP104:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 5 +// CHECK-I386-NEXT: store i64 4, i64* [[TMP104]], align 4 +// CHECK-I386-NEXT: [[TMP105:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i32 0, i32 5 +// CHECK-I386-NEXT: store i8* null, i8** [[TMP105]], align 4 +// CHECK-I386-NEXT: [[TMP106:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 6 +// CHECK-I386-NEXT: [[TMP107:%.*]] = bitcast i8** [[TMP106]] to %struct.S2** +// CHECK-I386-NEXT: store %struct.S2* [[TMP50]], %struct.S2** [[TMP107]], align 4 +// CHECK-I386-NEXT: [[TMP108:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 6 +// CHECK-I386-NEXT: [[TMP109:%.*]] = bitcast i8** [[TMP108]] to %struct.S1** +// CHECK-I386-NEXT: store %struct.S1* [[S18]], %struct.S1** [[TMP109]], align 4 +// CHECK-I386-NEXT: [[TMP110:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 6 +// CHECK-I386-NEXT: store i64 [[TMP69]], i64* [[TMP110]], align 4 +// CHECK-I386-NEXT: [[TMP111:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i32 0, i32 6 +// CHECK-I386-NEXT: store i8* null, i8** [[TMP111]], align 4 +// CHECK-I386-NEXT: [[TMP112:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 7 +// CHECK-I386-NEXT: [[TMP113:%.*]] = bitcast i8** [[TMP112]] to %struct.S2** +// CHECK-I386-NEXT: store %struct.S2* [[TMP50]], %struct.S2** [[TMP113]], align 4 +// CHECK-I386-NEXT: [[TMP114:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 7 +// CHECK-I386-NEXT: [[TMP115:%.*]] = bitcast i8** [[TMP114]] to %struct.S1** +// CHECK-I386-NEXT: store %struct.S1* [[S18]], %struct.S1** [[TMP115]], align 4 +// CHECK-I386-NEXT: [[TMP116:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 7 +// CHECK-I386-NEXT: store i64 4, i64* [[TMP116]], align 4 +// CHECK-I386-NEXT: [[TMP117:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i32 0, i32 7 +// CHECK-I386-NEXT: store i8* null, i8** [[TMP117]], align 4 +// CHECK-I386-NEXT: [[TMP118:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 8 +// CHECK-I386-NEXT: [[TMP119:%.*]] = bitcast i8** [[TMP118]] to %struct.S2*** +// CHECK-I386-NEXT: store %struct.S2** [[PS19]], %struct.S2*** [[TMP119]], align 4 +// CHECK-I386-NEXT: [[TMP120:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 8 +// CHECK-I386-NEXT: [[TMP121:%.*]] = bitcast i8** [[TMP120]] to %struct.S2*** +// CHECK-I386-NEXT: store %struct.S2** [[PS21]], %struct.S2*** [[TMP121]], align 4 +// CHECK-I386-NEXT: [[TMP122:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 8 +// CHECK-I386-NEXT: store i64 4, i64* [[TMP122]], align 4 +// CHECK-I386-NEXT: [[TMP123:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i32 0, i32 8 +// CHECK-I386-NEXT: store i8* null, i8** [[TMP123]], align 4 +// CHECK-I386-NEXT: [[TMP124:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 9 +// CHECK-I386-NEXT: [[TMP125:%.*]] = bitcast i8** [[TMP124]] to %struct.S2*** +// CHECK-I386-NEXT: store %struct.S2** [[PS21]], %struct.S2*** [[TMP125]], align 4 +// CHECK-I386-NEXT: [[TMP126:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 9 +// CHECK-I386-NEXT: [[TMP127:%.*]] = bitcast i8** [[TMP126]] to %struct.S2*** +// CHECK-I386-NEXT: store %struct.S2** [[PS24]], %struct.S2*** [[TMP127]], align 4 +// CHECK-I386-NEXT: [[TMP128:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 9 +// CHECK-I386-NEXT: store i64 4, i64* [[TMP128]], align 4 +// CHECK-I386-NEXT: [[TMP129:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i32 0, i32 9 +// CHECK-I386-NEXT: store i8* null, i8** [[TMP129]], align 4 +// CHECK-I386-NEXT: [[TMP130:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 10 +// CHECK-I386-NEXT: [[TMP131:%.*]] = bitcast i8** [[TMP130]] to %struct.S2*** +// CHECK-I386-NEXT: store %struct.S2** [[PS24]], %struct.S2*** [[TMP131]], align 4 +// CHECK-I386-NEXT: [[TMP132:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 10 +// CHECK-I386-NEXT: [[TMP133:%.*]] = bitcast i8** [[TMP132]] to %struct.S1** +// CHECK-I386-NEXT: store %struct.S1* [[S28]], %struct.S1** [[TMP133]], align 4 +// CHECK-I386-NEXT: [[TMP134:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 10 +// CHECK-I386-NEXT: store i64 4, i64* [[TMP134]], align 4 +// CHECK-I386-NEXT: [[TMP135:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i32 0, i32 10 +// CHECK-I386-NEXT: store i8* null, i8** [[TMP135]], align 4 +// CHECK-I386-NEXT: [[TMP136:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 0 +// CHECK-I386-NEXT: [[TMP137:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 0 +// CHECK-I386-NEXT: [[TMP138:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0 +// CHECK-I386-NEXT: call void @__tgt_target_data_begin_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i32 11, i8** [[TMP136]], i8** [[TMP137]], i64* [[TMP138]], i64* getelementptr inbounds ([11 x i64], [11 x i64]* @.offload_maptypes.5, i32 0, i32 0), i8** null, i8** null) +// CHECK-I386-NEXT: [[TMP139:%.*]] = load i32, i32* [[ARG_ADDR]], align 4 +// CHECK-I386-NEXT: [[INC32:%.*]] = add nsw i32 [[TMP139]], 1 +// CHECK-I386-NEXT: store i32 [[INC32]], i32* [[ARG_ADDR]], align 4 +// CHECK-I386-NEXT: [[TMP140:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 0 +// CHECK-I386-NEXT: [[TMP141:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 0 +// CHECK-I386-NEXT: [[TMP142:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0 +// CHECK-I386-NEXT: call void @__tgt_target_data_end_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i32 11, i8** [[TMP140]], i8** [[TMP141]], i64* [[TMP142]], i64* getelementptr inbounds ([11 x i64], [11 x i64]* @.offload_maptypes.5, i32 0, i32 0), i8** null, i8** null) +// CHECK-I386-NEXT: ret void +// +void foo(int arg) { + float lb[5]; + S2 *ps1; + S2 *ps2; + + #pragma omp target data map(ompx_hold, to: lb) + {++arg;} + + #pragma omp target data map(always close ompx_hold, to: lb) + {++arg;} + + #pragma omp target data map(ompx_hold, tofrom : arg) + {++arg;} + + // Make sure the struct picks up ompx_hold even if another element of the + // struct doesn't have ompx_hold. + #pragma omp target data map(tofrom : ps1->s, arg) \ + map(ompx_hold, tofrom : ps1->ps->ps->ps->s, ps2->s) \ + map(tofrom : ps2->ps->ps->ps->s) + { + ++(arg); + } +} + +#endif Index: clang/test/OpenMP/target_enter_data_map_messages.c =================================================================== --- clang/test/OpenMP/target_enter_data_map_messages.c +++ clang/test/OpenMP/target_enter_data_map_messages.c @@ -1,8 +1,14 @@ -// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp -ferror-limit 100 -o - %s -Wuninitialized -// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp -ferror-limit 100 -o - -x c++ %s -Wuninitialized +// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp -fopenmp -ferror-limit 100 -o - %s -Wuninitialized +// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp -fopenmp -ferror-limit 100 -o - -x c++ %s -Wuninitialized -// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp-simd -ferror-limit 100 -o - %s -Wuninitialized -// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp-simd -ferror-limit 100 -o - -x c++ %s -Wuninitialized +// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp -fopenmp-simd -ferror-limit 100 -o - %s -Wuninitialized +// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp -fopenmp-simd -ferror-limit 100 -o - -x c++ %s -Wuninitialized + +// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,ompx -fopenmp -fopenmp-extensions -ferror-limit 100 -o - %s -Wuninitialized +// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,ompx -fopenmp -fopenmp-extensions -ferror-limit 100 -o - -x c++ %s -Wuninitialized + +// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,ompx -fopenmp-simd -fopenmp-extensions -ferror-limit 100 -o - %s -Wuninitialized +// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,ompx -fopenmp-simd -fopenmp-extensions -ferror-limit 100 -o - -x c++ %s -Wuninitialized void xxx(int argc) { int map; // expected-note {{initialize the variable 'map' to silence this warning}} @@ -25,5 +31,12 @@ #pragma omp target enter data map(release: r) // expected-error {{map type 'release' is not allowed for '#pragma omp target enter data'}} #pragma omp target enter data map(delete: r) // expected-error {{map type 'delete' is not allowed for '#pragma omp target enter data'}} + // omp-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} + // ompx-error@+1 {{map type modifier 'ompx_hold' is not allowed for '#pragma omp target enter data'}} + #pragma omp target enter data map(ompx_hold, alloc: r) + // omp-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} + // ompx-error@+1 {{map type modifier 'ompx_hold' is not allowed for '#pragma omp target enter data'}} + #pragma omp target enter data map(ompx_hold, to: r) + return 0; } Index: clang/test/OpenMP/target_exit_data_map_messages.c =================================================================== --- clang/test/OpenMP/target_exit_data_map_messages.c +++ clang/test/OpenMP/target_exit_data_map_messages.c @@ -1,8 +1,14 @@ -// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp -ferror-limit 100 -o - %s -Wuninitialized -// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp -ferror-limit 100 -o - -x c++ %s -Wuninitialized +// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp -fopenmp -ferror-limit 100 -o - %s -Wuninitialized +// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp -fopenmp -ferror-limit 100 -o - -x c++ %s -Wuninitialized -// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp-simd -ferror-limit 100 -o - %s -Wuninitialized -// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp-simd -ferror-limit 100 -o - -x c++ %s -Wuninitialized +// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp -fopenmp-simd -ferror-limit 100 -o - %s -Wuninitialized +// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp -fopenmp-simd -ferror-limit 100 -o - -x c++ %s -Wuninitialized + +// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,ompx -fopenmp -fopenmp-extensions -ferror-limit 100 -o - %s -Wuninitialized +// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,ompx -fopenmp -fopenmp-extensions -ferror-limit 100 -o - -x c++ %s -Wuninitialized + +// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,ompx -fopenmp-simd -fopenmp-extensions -ferror-limit 100 -o - %s -Wuninitialized +// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,ompx -fopenmp-simd -fopenmp-extensions -ferror-limit 100 -o - -x c++ %s -Wuninitialized int main(int argc, char **argv) { @@ -18,5 +24,15 @@ #pragma omp target exit data map(always, alloc: r) // expected-error {{map type 'alloc' is not allowed for '#pragma omp target exit data'}} #pragma omp target exit data map(to: r) // expected-error {{map type 'to' is not allowed for '#pragma omp target exit data'}} + // omp-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} + // ompx-error@+1 {{map type modifier 'ompx_hold' is not allowed for '#pragma omp target exit data'}} + #pragma omp target exit data map(ompx_hold, from: r) + // omp-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} + // ompx-error@+1 {{map type modifier 'ompx_hold' is not allowed for '#pragma omp target exit data'}} + #pragma omp target exit data map(ompx_hold, release: r) + // omp-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} + // ompx-error@+1 {{map type modifier 'ompx_hold' is not allowed for '#pragma omp target exit data'}} + #pragma omp target exit data map(ompx_hold, delete: r) + return 0; } Index: clang/test/OpenMP/target_map_codegen_hold.cpp =================================================================== --- /dev/null +++ clang/test/OpenMP/target_map_codegen_hold.cpp @@ -0,0 +1,928 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" --global-value-regex ".offload_maptypes.*" ".offload_sizes.*" --global-hex-value-regex ".offload_maptypes.*" +// expected-no-diagnostics +#ifndef HEADER +#define HEADER + +//-------------------------------------------------- +// With -DUSE. +//-------------------------------------------------- + +// powerpc64le-ibm-linux-gnu + +// RUN: %clang_cc1 -DUSE -verify -fopenmp -fopenmp-extensions \ +// RUN: -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ \ +// RUN: -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | \ +// RUN: FileCheck %s --check-prefixes=CHECK-USE-PPC64LE +// RUN: %clang_cc1 -DUSE -fopenmp -fopenmp-extensions \ +// RUN: -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 \ +// RUN: -triple powerpc64le-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -DUSE -fopenmp -fopenmp-extensions \ +// RUN: -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ \ +// RUN: -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t \ +// RUN: -verify %s -emit-llvm -o - | \ +// RUN: FileCheck %s --check-prefixes=CHECK-USE-PPC64LE + +// i386-pc-linux-gnu + +// RUN: %clang_cc1 -DUSE -verify -fopenmp -fopenmp-extensions \ +// RUN: -fopenmp-targets=i386-pc-linux-gnu -x c++ \ +// RUN: -triple i386-unknown-unknown -emit-llvm %s -o - | \ +// RUN: FileCheck %s --check-prefixes=CHECK-USE-I386 +// RUN: %clang_cc1 -DUSE -fopenmp -fopenmp-extensions \ +// RUN: -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 \ +// RUN: -triple i386-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -DUSE -fopenmp -fopenmp-extensions \ +// RUN: -fopenmp-targets=i386-pc-linux-gnu -x c++ \ +// RUN: -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s \ +// RUN: -emit-llvm -o - | \ +// RUN: FileCheck %s --check-prefixes=CHECK-USE-I386 + +//-------------------------------------------------- +// Without -DUSE. +//-------------------------------------------------- + +// powerpc64le-ibm-linux-gnu + +// RUN: %clang_cc1 -verify -fopenmp -fopenmp-extensions \ +// RUN: -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ \ +// RUN: -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | \ +// RUN: FileCheck %s --check-prefixes=CHECK-NOUSE-PPC64LE +// RUN: %clang_cc1 -fopenmp -fopenmp-extensions \ +// RUN: -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 \ +// RUN: -triple powerpc64le-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -fopenmp-extensions \ +// RUN: -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ \ +// RUN: -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t \ +// RUN: -verify %s -emit-llvm -o - | \ +// RUN: FileCheck %s --check-prefixes=CHECK-NOUSE-PPC64LE + +// i386-pc-linux-gnu + +// RUN: %clang_cc1 -verify -fopenmp -fopenmp-extensions \ +// RUN: -fopenmp-targets=i386-pc-linux-gnu -x c++ \ +// RUN: -triple i386-unknown-unknown -emit-llvm %s -o - | \ +// RUN: FileCheck %s --check-prefixes=CHECK-NOUSE-I386 +// RUN: %clang_cc1 -fopenmp -fopenmp-extensions \ +// RUN: -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 \ +// RUN: -triple i386-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -fopenmp-extensions \ +// RUN: -fopenmp-targets=i386-pc-linux-gnu -x c++ \ +// RUN: -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s \ +// RUN: -emit-llvm -o - | \ +// RUN: FileCheck %s --check-prefixes=CHECK-NOUSE-I386 + +// Map flags used in @.offload_maptypes* below: +// +// TO = 0x1 +// FROM = 0x2 +// ALWAYS = 0x4 +// TARGET_PARAM = 0x20 +// CLOSE = 0x400 +// OMPX_HOLD = 0x2000 +// MEMBER_OF_1 = 0x1000000000000 +// MEMBER_OF_5 = 0x5000000000000 + +//. +// CHECK-USE-PPC64LE: @.offload_maptypes = private unnamed_addr constant [7 x i64] [i64 [[#0x2020]], i64 [[#0x1000000002003]], i64 [[#0x1000000002003]], i64 [[#0x2023]], i64 [[#0x2020]], i64 [[#0x5000000002003]], i64 [[#0x5000000002003]]] +// CHECK-USE-PPC64LE: @.offload_sizes = private unnamed_addr constant [1 x i64] [i64 4] +// CHECK-USE-PPC64LE: @.offload_maptypes.1 = private unnamed_addr constant [1 x i64] [i64 [[#0x2427]]] +// CHECK-USE-PPC64LE: @.offload_maptypes.2 = private unnamed_addr constant [3 x i64] [i64 [[#0x2020]], i64 [[#0x1000000002003]], i64 [[#0x1000000002003]]] +//. +// CHECK-USE-I386: @.offload_maptypes = private unnamed_addr constant [7 x i64] [i64 [[#0x2020]], i64 [[#0x1000000002003]], i64 [[#0x1000000002003]], i64 [[#0x2023]], i64 [[#0x2020]], i64 [[#0x5000000002003]], i64 [[#0x5000000002003]]] +// CHECK-USE-I386: @.offload_sizes = private unnamed_addr constant [1 x i64] [i64 4] +// CHECK-USE-I386: @.offload_maptypes.1 = private unnamed_addr constant [1 x i64] [i64 [[#0x2427]]] +// CHECK-USE-I386: @.offload_maptypes.2 = private unnamed_addr constant [3 x i64] [i64 [[#0x2020]], i64 [[#0x1000000002003]], i64 [[#0x1000000002003]]] +//. +// CHECK-NOUSE-PPC64LE: @.offload_maptypes = private unnamed_addr constant [7 x i64] [i64 [[#0x2000]], i64 [[#0x1000000002003]], i64 [[#0x1000000002003]], i64 [[#0x2003]], i64 [[#0x2000]], i64 [[#0x5000000002003]], i64 [[#0x5000000002003]]] +// CHECK-NOUSE-PPC64LE: @.offload_sizes = private unnamed_addr constant [1 x i64] [i64 4] +// CHECK-NOUSE-PPC64LE: @.offload_maptypes.1 = private unnamed_addr constant [1 x i64] [i64 [[#0x2407]]] +// CHECK-NOUSE-PPC64LE: @.offload_maptypes.2 = private unnamed_addr constant [3 x i64] [i64 [[#0x2000]], i64 [[#0x1000000002003]], i64 [[#0x1000000002003]]] +//. +// CHECK-NOUSE-I386: @.offload_maptypes = private unnamed_addr constant [7 x i64] [i64 [[#0x2000]], i64 [[#0x1000000002003]], i64 [[#0x1000000002003]], i64 [[#0x2003]], i64 [[#0x2000]], i64 [[#0x5000000002003]], i64 [[#0x5000000002003]]] +// CHECK-NOUSE-I386: @.offload_sizes = private unnamed_addr constant [1 x i64] [i64 4] +// CHECK-NOUSE-I386: @.offload_maptypes.1 = private unnamed_addr constant [1 x i64] [i64 [[#0x2407]]] +// CHECK-NOUSE-I386: @.offload_maptypes.2 = private unnamed_addr constant [3 x i64] [i64 [[#0x2000]], i64 [[#0x1000000002003]], i64 [[#0x1000000002003]]] +//. +struct ST { + int i; + int j; + void test_present_members(); +}; + +// CHECK-USE-PPC64LE-LABEL: @_Z20explicit_maps_singlei( +// CHECK-USE-PPC64LE-NEXT: entry: +// CHECK-USE-PPC64LE-NEXT: [[II_ADDR:%.*]] = alloca i32, align 4 +// CHECK-USE-PPC64LE-NEXT: [[A:%.*]] = alloca i32, align 4 +// CHECK-USE-PPC64LE-NEXT: [[ST1:%.*]] = alloca [[STRUCT_ST:%.*]], align 4 +// CHECK-USE-PPC64LE-NEXT: [[ST2:%.*]] = alloca [[STRUCT_ST]], align 4 +// CHECK-USE-PPC64LE-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [7 x i8*], align 8 +// CHECK-USE-PPC64LE-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [7 x i8*], align 8 +// CHECK-USE-PPC64LE-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [7 x i8*], align 8 +// CHECK-USE-PPC64LE-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [7 x i64], align 8 +// CHECK-USE-PPC64LE-NEXT: [[DOTOFFLOAD_BASEPTRS3:%.*]] = alloca [1 x i8*], align 8 +// CHECK-USE-PPC64LE-NEXT: [[DOTOFFLOAD_PTRS4:%.*]] = alloca [1 x i8*], align 8 +// CHECK-USE-PPC64LE-NEXT: [[DOTOFFLOAD_MAPPERS5:%.*]] = alloca [1 x i8*], align 8 +// CHECK-USE-PPC64LE-NEXT: store i32 [[II:%.*]], i32* [[II_ADDR]], align 4 +// CHECK-USE-PPC64LE-NEXT: [[TMP0:%.*]] = load i32, i32* [[II_ADDR]], align 4 +// CHECK-USE-PPC64LE-NEXT: store i32 [[TMP0]], i32* [[A]], align 4 +// CHECK-USE-PPC64LE-NEXT: [[I:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[ST1]], i32 0, i32 0 +// CHECK-USE-PPC64LE-NEXT: [[J:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[ST1]], i32 0, i32 1 +// CHECK-USE-PPC64LE-NEXT: [[TMP1:%.*]] = getelementptr i32, i32* [[J]], i32 1 +// CHECK-USE-PPC64LE-NEXT: [[TMP2:%.*]] = bitcast i32* [[I]] to i8* +// CHECK-USE-PPC64LE-NEXT: [[TMP3:%.*]] = bitcast i32* [[TMP1]] to i8* +// CHECK-USE-PPC64LE-NEXT: [[TMP4:%.*]] = ptrtoint i8* [[TMP3]] to i64 +// CHECK-USE-PPC64LE-NEXT: [[TMP5:%.*]] = ptrtoint i8* [[TMP2]] to i64 +// CHECK-USE-PPC64LE-NEXT: [[TMP6:%.*]] = sub i64 [[TMP4]], [[TMP5]] +// CHECK-USE-PPC64LE-NEXT: [[TMP7:%.*]] = sdiv exact i64 [[TMP6]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) +// CHECK-USE-PPC64LE-NEXT: [[I1:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[ST2]], i32 0, i32 0 +// CHECK-USE-PPC64LE-NEXT: [[J2:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[ST2]], i32 0, i32 1 +// CHECK-USE-PPC64LE-NEXT: [[TMP8:%.*]] = getelementptr i32, i32* [[J2]], i32 1 +// CHECK-USE-PPC64LE-NEXT: [[TMP9:%.*]] = bitcast i32* [[I1]] to i8* +// CHECK-USE-PPC64LE-NEXT: [[TMP10:%.*]] = bitcast i32* [[TMP8]] to i8* +// CHECK-USE-PPC64LE-NEXT: [[TMP11:%.*]] = ptrtoint i8* [[TMP10]] to i64 +// CHECK-USE-PPC64LE-NEXT: [[TMP12:%.*]] = ptrtoint i8* [[TMP9]] to i64 +// CHECK-USE-PPC64LE-NEXT: [[TMP13:%.*]] = sub i64 [[TMP11]], [[TMP12]] +// CHECK-USE-PPC64LE-NEXT: [[TMP14:%.*]] = sdiv exact i64 [[TMP13]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) +// CHECK-USE-PPC64LE-NEXT: [[TMP15:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK-USE-PPC64LE-NEXT: [[TMP16:%.*]] = bitcast i8** [[TMP15]] to %struct.ST** +// CHECK-USE-PPC64LE-NEXT: store %struct.ST* [[ST1]], %struct.ST** [[TMP16]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP17:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK-USE-PPC64LE-NEXT: [[TMP18:%.*]] = bitcast i8** [[TMP17]] to i32** +// CHECK-USE-PPC64LE-NEXT: store i32* [[I]], i32** [[TMP18]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP19:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0 +// CHECK-USE-PPC64LE-NEXT: store i64 [[TMP7]], i64* [[TMP19]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP20:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0 +// CHECK-USE-PPC64LE-NEXT: store i8* null, i8** [[TMP20]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP21:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 +// CHECK-USE-PPC64LE-NEXT: [[TMP22:%.*]] = bitcast i8** [[TMP21]] to %struct.ST** +// CHECK-USE-PPC64LE-NEXT: store %struct.ST* [[ST1]], %struct.ST** [[TMP22]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP23:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 1 +// CHECK-USE-PPC64LE-NEXT: [[TMP24:%.*]] = bitcast i8** [[TMP23]] to i32** +// CHECK-USE-PPC64LE-NEXT: store i32* [[I]], i32** [[TMP24]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP25:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 1 +// CHECK-USE-PPC64LE-NEXT: store i64 4, i64* [[TMP25]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP26:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 +// CHECK-USE-PPC64LE-NEXT: store i8* null, i8** [[TMP26]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP27:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2 +// CHECK-USE-PPC64LE-NEXT: [[TMP28:%.*]] = bitcast i8** [[TMP27]] to %struct.ST** +// CHECK-USE-PPC64LE-NEXT: store %struct.ST* [[ST1]], %struct.ST** [[TMP28]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP29:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 2 +// CHECK-USE-PPC64LE-NEXT: [[TMP30:%.*]] = bitcast i8** [[TMP29]] to i32** +// CHECK-USE-PPC64LE-NEXT: store i32* [[J]], i32** [[TMP30]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP31:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 2 +// CHECK-USE-PPC64LE-NEXT: store i64 4, i64* [[TMP31]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP32:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2 +// CHECK-USE-PPC64LE-NEXT: store i8* null, i8** [[TMP32]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP33:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 3 +// CHECK-USE-PPC64LE-NEXT: [[TMP34:%.*]] = bitcast i8** [[TMP33]] to i32** +// CHECK-USE-PPC64LE-NEXT: store i32* [[A]], i32** [[TMP34]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP35:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 3 +// CHECK-USE-PPC64LE-NEXT: [[TMP36:%.*]] = bitcast i8** [[TMP35]] to i32** +// CHECK-USE-PPC64LE-NEXT: store i32* [[A]], i32** [[TMP36]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP37:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 3 +// CHECK-USE-PPC64LE-NEXT: store i64 4, i64* [[TMP37]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP38:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 3 +// CHECK-USE-PPC64LE-NEXT: store i8* null, i8** [[TMP38]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP39:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 4 +// CHECK-USE-PPC64LE-NEXT: [[TMP40:%.*]] = bitcast i8** [[TMP39]] to %struct.ST** +// CHECK-USE-PPC64LE-NEXT: store %struct.ST* [[ST2]], %struct.ST** [[TMP40]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP41:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 4 +// CHECK-USE-PPC64LE-NEXT: [[TMP42:%.*]] = bitcast i8** [[TMP41]] to i32** +// CHECK-USE-PPC64LE-NEXT: store i32* [[I1]], i32** [[TMP42]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP43:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 4 +// CHECK-USE-PPC64LE-NEXT: store i64 [[TMP14]], i64* [[TMP43]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP44:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 4 +// CHECK-USE-PPC64LE-NEXT: store i8* null, i8** [[TMP44]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP45:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 5 +// CHECK-USE-PPC64LE-NEXT: [[TMP46:%.*]] = bitcast i8** [[TMP45]] to %struct.ST** +// CHECK-USE-PPC64LE-NEXT: store %struct.ST* [[ST2]], %struct.ST** [[TMP46]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP47:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 5 +// CHECK-USE-PPC64LE-NEXT: [[TMP48:%.*]] = bitcast i8** [[TMP47]] to i32** +// CHECK-USE-PPC64LE-NEXT: store i32* [[I1]], i32** [[TMP48]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP49:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 5 +// CHECK-USE-PPC64LE-NEXT: store i64 4, i64* [[TMP49]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP50:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 5 +// CHECK-USE-PPC64LE-NEXT: store i8* null, i8** [[TMP50]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP51:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 6 +// CHECK-USE-PPC64LE-NEXT: [[TMP52:%.*]] = bitcast i8** [[TMP51]] to %struct.ST** +// CHECK-USE-PPC64LE-NEXT: store %struct.ST* [[ST2]], %struct.ST** [[TMP52]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP53:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 6 +// CHECK-USE-PPC64LE-NEXT: [[TMP54:%.*]] = bitcast i8** [[TMP53]] to i32** +// CHECK-USE-PPC64LE-NEXT: store i32* [[J2]], i32** [[TMP54]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP55:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 6 +// CHECK-USE-PPC64LE-NEXT: store i64 4, i64* [[TMP55]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP56:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 6 +// CHECK-USE-PPC64LE-NEXT: store i8* null, i8** [[TMP56]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP57:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK-USE-PPC64LE-NEXT: [[TMP58:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK-USE-PPC64LE-NEXT: [[TMP59:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0 +// CHECK-USE-PPC64LE-NEXT: [[TMP60:%.*]] = call i32 @__tgt_target_mapper(%struct.ident_t* @[[GLOB1:[0-9]+]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z20explicit_maps_singlei_l654.region_id, i32 7, i8** [[TMP57]], i8** [[TMP58]], i64* [[TMP59]], i64* getelementptr inbounds ([7 x i64], [7 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null) +// CHECK-USE-PPC64LE-NEXT: [[TMP61:%.*]] = icmp ne i32 [[TMP60]], 0 +// CHECK-USE-PPC64LE-NEXT: br i1 [[TMP61]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] +// CHECK-USE-PPC64LE: omp_offload.failed: +// CHECK-USE-PPC64LE-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z20explicit_maps_singlei_l654(%struct.ST* [[ST1]], i32* [[A]], %struct.ST* [[ST2]]) #[[ATTR2:[0-9]+]] +// CHECK-USE-PPC64LE-NEXT: br label [[OMP_OFFLOAD_CONT]] +// CHECK-USE-PPC64LE: omp_offload.cont: +// CHECK-USE-PPC64LE-NEXT: [[TMP62:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0 +// CHECK-USE-PPC64LE-NEXT: [[TMP63:%.*]] = bitcast i8** [[TMP62]] to i32** +// CHECK-USE-PPC64LE-NEXT: store i32* [[A]], i32** [[TMP63]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP64:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS4]], i32 0, i32 0 +// CHECK-USE-PPC64LE-NEXT: [[TMP65:%.*]] = bitcast i8** [[TMP64]] to i32** +// CHECK-USE-PPC64LE-NEXT: store i32* [[A]], i32** [[TMP65]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP66:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_MAPPERS5]], i64 0, i64 0 +// CHECK-USE-PPC64LE-NEXT: store i8* null, i8** [[TMP66]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP67:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0 +// CHECK-USE-PPC64LE-NEXT: [[TMP68:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS4]], i32 0, i32 0 +// CHECK-USE-PPC64LE-NEXT: [[TMP69:%.*]] = call i32 @__tgt_target_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z20explicit_maps_singlei_l668.region_id, i32 1, i8** [[TMP67]], i8** [[TMP68]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.1, i32 0, i32 0), i8** null, i8** null) +// CHECK-USE-PPC64LE-NEXT: [[TMP70:%.*]] = icmp ne i32 [[TMP69]], 0 +// CHECK-USE-PPC64LE-NEXT: br i1 [[TMP70]], label [[OMP_OFFLOAD_FAILED6:%.*]], label [[OMP_OFFLOAD_CONT7:%.*]] +// CHECK-USE-PPC64LE: omp_offload.failed6: +// CHECK-USE-PPC64LE-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z20explicit_maps_singlei_l668(i32* [[A]]) #[[ATTR2]] +// CHECK-USE-PPC64LE-NEXT: br label [[OMP_OFFLOAD_CONT7]] +// CHECK-USE-PPC64LE: omp_offload.cont7: +// CHECK-USE-PPC64LE-NEXT: ret void +// +// CHECK-USE-I386-LABEL: @_Z20explicit_maps_singlei( +// CHECK-USE-I386-NEXT: entry: +// CHECK-USE-I386-NEXT: [[II_ADDR:%.*]] = alloca i32, align 4 +// CHECK-USE-I386-NEXT: [[A:%.*]] = alloca i32, align 4 +// CHECK-USE-I386-NEXT: [[ST1:%.*]] = alloca [[STRUCT_ST:%.*]], align 4 +// CHECK-USE-I386-NEXT: [[ST2:%.*]] = alloca [[STRUCT_ST]], align 4 +// CHECK-USE-I386-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [7 x i8*], align 4 +// CHECK-USE-I386-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [7 x i8*], align 4 +// CHECK-USE-I386-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [7 x i8*], align 4 +// CHECK-USE-I386-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [7 x i64], align 4 +// CHECK-USE-I386-NEXT: [[DOTOFFLOAD_BASEPTRS3:%.*]] = alloca [1 x i8*], align 4 +// CHECK-USE-I386-NEXT: [[DOTOFFLOAD_PTRS4:%.*]] = alloca [1 x i8*], align 4 +// CHECK-USE-I386-NEXT: [[DOTOFFLOAD_MAPPERS5:%.*]] = alloca [1 x i8*], align 4 +// CHECK-USE-I386-NEXT: store i32 [[II:%.*]], i32* [[II_ADDR]], align 4 +// CHECK-USE-I386-NEXT: [[TMP0:%.*]] = load i32, i32* [[II_ADDR]], align 4 +// CHECK-USE-I386-NEXT: store i32 [[TMP0]], i32* [[A]], align 4 +// CHECK-USE-I386-NEXT: [[I:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[ST1]], i32 0, i32 0 +// CHECK-USE-I386-NEXT: [[J:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[ST1]], i32 0, i32 1 +// CHECK-USE-I386-NEXT: [[TMP1:%.*]] = getelementptr i32, i32* [[J]], i32 1 +// CHECK-USE-I386-NEXT: [[TMP2:%.*]] = bitcast i32* [[I]] to i8* +// CHECK-USE-I386-NEXT: [[TMP3:%.*]] = bitcast i32* [[TMP1]] to i8* +// CHECK-USE-I386-NEXT: [[TMP4:%.*]] = ptrtoint i8* [[TMP3]] to i64 +// CHECK-USE-I386-NEXT: [[TMP5:%.*]] = ptrtoint i8* [[TMP2]] to i64 +// CHECK-USE-I386-NEXT: [[TMP6:%.*]] = sub i64 [[TMP4]], [[TMP5]] +// CHECK-USE-I386-NEXT: [[TMP7:%.*]] = sdiv exact i64 [[TMP6]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) +// CHECK-USE-I386-NEXT: [[I1:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[ST2]], i32 0, i32 0 +// CHECK-USE-I386-NEXT: [[J2:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[ST2]], i32 0, i32 1 +// CHECK-USE-I386-NEXT: [[TMP8:%.*]] = getelementptr i32, i32* [[J2]], i32 1 +// CHECK-USE-I386-NEXT: [[TMP9:%.*]] = bitcast i32* [[I1]] to i8* +// CHECK-USE-I386-NEXT: [[TMP10:%.*]] = bitcast i32* [[TMP8]] to i8* +// CHECK-USE-I386-NEXT: [[TMP11:%.*]] = ptrtoint i8* [[TMP10]] to i64 +// CHECK-USE-I386-NEXT: [[TMP12:%.*]] = ptrtoint i8* [[TMP9]] to i64 +// CHECK-USE-I386-NEXT: [[TMP13:%.*]] = sub i64 [[TMP11]], [[TMP12]] +// CHECK-USE-I386-NEXT: [[TMP14:%.*]] = sdiv exact i64 [[TMP13]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) +// CHECK-USE-I386-NEXT: [[TMP15:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK-USE-I386-NEXT: [[TMP16:%.*]] = bitcast i8** [[TMP15]] to %struct.ST** +// CHECK-USE-I386-NEXT: store %struct.ST* [[ST1]], %struct.ST** [[TMP16]], align 4 +// CHECK-USE-I386-NEXT: [[TMP17:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK-USE-I386-NEXT: [[TMP18:%.*]] = bitcast i8** [[TMP17]] to i32** +// CHECK-USE-I386-NEXT: store i32* [[I]], i32** [[TMP18]], align 4 +// CHECK-USE-I386-NEXT: [[TMP19:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0 +// CHECK-USE-I386-NEXT: store i64 [[TMP7]], i64* [[TMP19]], align 4 +// CHECK-USE-I386-NEXT: [[TMP20:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0 +// CHECK-USE-I386-NEXT: store i8* null, i8** [[TMP20]], align 4 +// CHECK-USE-I386-NEXT: [[TMP21:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 +// CHECK-USE-I386-NEXT: [[TMP22:%.*]] = bitcast i8** [[TMP21]] to %struct.ST** +// CHECK-USE-I386-NEXT: store %struct.ST* [[ST1]], %struct.ST** [[TMP22]], align 4 +// CHECK-USE-I386-NEXT: [[TMP23:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 1 +// CHECK-USE-I386-NEXT: [[TMP24:%.*]] = bitcast i8** [[TMP23]] to i32** +// CHECK-USE-I386-NEXT: store i32* [[I]], i32** [[TMP24]], align 4 +// CHECK-USE-I386-NEXT: [[TMP25:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 1 +// CHECK-USE-I386-NEXT: store i64 4, i64* [[TMP25]], align 4 +// CHECK-USE-I386-NEXT: [[TMP26:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 1 +// CHECK-USE-I386-NEXT: store i8* null, i8** [[TMP26]], align 4 +// CHECK-USE-I386-NEXT: [[TMP27:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2 +// CHECK-USE-I386-NEXT: [[TMP28:%.*]] = bitcast i8** [[TMP27]] to %struct.ST** +// CHECK-USE-I386-NEXT: store %struct.ST* [[ST1]], %struct.ST** [[TMP28]], align 4 +// CHECK-USE-I386-NEXT: [[TMP29:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 2 +// CHECK-USE-I386-NEXT: [[TMP30:%.*]] = bitcast i8** [[TMP29]] to i32** +// CHECK-USE-I386-NEXT: store i32* [[J]], i32** [[TMP30]], align 4 +// CHECK-USE-I386-NEXT: [[TMP31:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 2 +// CHECK-USE-I386-NEXT: store i64 4, i64* [[TMP31]], align 4 +// CHECK-USE-I386-NEXT: [[TMP32:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 2 +// CHECK-USE-I386-NEXT: store i8* null, i8** [[TMP32]], align 4 +// CHECK-USE-I386-NEXT: [[TMP33:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 3 +// CHECK-USE-I386-NEXT: [[TMP34:%.*]] = bitcast i8** [[TMP33]] to i32** +// CHECK-USE-I386-NEXT: store i32* [[A]], i32** [[TMP34]], align 4 +// CHECK-USE-I386-NEXT: [[TMP35:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 3 +// CHECK-USE-I386-NEXT: [[TMP36:%.*]] = bitcast i8** [[TMP35]] to i32** +// CHECK-USE-I386-NEXT: store i32* [[A]], i32** [[TMP36]], align 4 +// CHECK-USE-I386-NEXT: [[TMP37:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 3 +// CHECK-USE-I386-NEXT: store i64 4, i64* [[TMP37]], align 4 +// CHECK-USE-I386-NEXT: [[TMP38:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 3 +// CHECK-USE-I386-NEXT: store i8* null, i8** [[TMP38]], align 4 +// CHECK-USE-I386-NEXT: [[TMP39:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 4 +// CHECK-USE-I386-NEXT: [[TMP40:%.*]] = bitcast i8** [[TMP39]] to %struct.ST** +// CHECK-USE-I386-NEXT: store %struct.ST* [[ST2]], %struct.ST** [[TMP40]], align 4 +// CHECK-USE-I386-NEXT: [[TMP41:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 4 +// CHECK-USE-I386-NEXT: [[TMP42:%.*]] = bitcast i8** [[TMP41]] to i32** +// CHECK-USE-I386-NEXT: store i32* [[I1]], i32** [[TMP42]], align 4 +// CHECK-USE-I386-NEXT: [[TMP43:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 4 +// CHECK-USE-I386-NEXT: store i64 [[TMP14]], i64* [[TMP43]], align 4 +// CHECK-USE-I386-NEXT: [[TMP44:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 4 +// CHECK-USE-I386-NEXT: store i8* null, i8** [[TMP44]], align 4 +// CHECK-USE-I386-NEXT: [[TMP45:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 5 +// CHECK-USE-I386-NEXT: [[TMP46:%.*]] = bitcast i8** [[TMP45]] to %struct.ST** +// CHECK-USE-I386-NEXT: store %struct.ST* [[ST2]], %struct.ST** [[TMP46]], align 4 +// CHECK-USE-I386-NEXT: [[TMP47:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 5 +// CHECK-USE-I386-NEXT: [[TMP48:%.*]] = bitcast i8** [[TMP47]] to i32** +// CHECK-USE-I386-NEXT: store i32* [[I1]], i32** [[TMP48]], align 4 +// CHECK-USE-I386-NEXT: [[TMP49:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 5 +// CHECK-USE-I386-NEXT: store i64 4, i64* [[TMP49]], align 4 +// CHECK-USE-I386-NEXT: [[TMP50:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 5 +// CHECK-USE-I386-NEXT: store i8* null, i8** [[TMP50]], align 4 +// CHECK-USE-I386-NEXT: [[TMP51:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 6 +// CHECK-USE-I386-NEXT: [[TMP52:%.*]] = bitcast i8** [[TMP51]] to %struct.ST** +// CHECK-USE-I386-NEXT: store %struct.ST* [[ST2]], %struct.ST** [[TMP52]], align 4 +// CHECK-USE-I386-NEXT: [[TMP53:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 6 +// CHECK-USE-I386-NEXT: [[TMP54:%.*]] = bitcast i8** [[TMP53]] to i32** +// CHECK-USE-I386-NEXT: store i32* [[J2]], i32** [[TMP54]], align 4 +// CHECK-USE-I386-NEXT: [[TMP55:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 6 +// CHECK-USE-I386-NEXT: store i64 4, i64* [[TMP55]], align 4 +// CHECK-USE-I386-NEXT: [[TMP56:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 6 +// CHECK-USE-I386-NEXT: store i8* null, i8** [[TMP56]], align 4 +// CHECK-USE-I386-NEXT: [[TMP57:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK-USE-I386-NEXT: [[TMP58:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK-USE-I386-NEXT: [[TMP59:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0 +// CHECK-USE-I386-NEXT: [[TMP60:%.*]] = call i32 @__tgt_target_mapper(%struct.ident_t* @[[GLOB1:[0-9]+]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z20explicit_maps_singlei_l654.region_id, i32 7, i8** [[TMP57]], i8** [[TMP58]], i64* [[TMP59]], i64* getelementptr inbounds ([7 x i64], [7 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null) +// CHECK-USE-I386-NEXT: [[TMP61:%.*]] = icmp ne i32 [[TMP60]], 0 +// CHECK-USE-I386-NEXT: br i1 [[TMP61]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] +// CHECK-USE-I386: omp_offload.failed: +// CHECK-USE-I386-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z20explicit_maps_singlei_l654(%struct.ST* [[ST1]], i32* [[A]], %struct.ST* [[ST2]]) #[[ATTR2:[0-9]+]] +// CHECK-USE-I386-NEXT: br label [[OMP_OFFLOAD_CONT]] +// CHECK-USE-I386: omp_offload.cont: +// CHECK-USE-I386-NEXT: [[TMP62:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0 +// CHECK-USE-I386-NEXT: [[TMP63:%.*]] = bitcast i8** [[TMP62]] to i32** +// CHECK-USE-I386-NEXT: store i32* [[A]], i32** [[TMP63]], align 4 +// CHECK-USE-I386-NEXT: [[TMP64:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS4]], i32 0, i32 0 +// CHECK-USE-I386-NEXT: [[TMP65:%.*]] = bitcast i8** [[TMP64]] to i32** +// CHECK-USE-I386-NEXT: store i32* [[A]], i32** [[TMP65]], align 4 +// CHECK-USE-I386-NEXT: [[TMP66:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_MAPPERS5]], i32 0, i32 0 +// CHECK-USE-I386-NEXT: store i8* null, i8** [[TMP66]], align 4 +// CHECK-USE-I386-NEXT: [[TMP67:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0 +// CHECK-USE-I386-NEXT: [[TMP68:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS4]], i32 0, i32 0 +// CHECK-USE-I386-NEXT: [[TMP69:%.*]] = call i32 @__tgt_target_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z20explicit_maps_singlei_l668.region_id, i32 1, i8** [[TMP67]], i8** [[TMP68]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.1, i32 0, i32 0), i8** null, i8** null) +// CHECK-USE-I386-NEXT: [[TMP70:%.*]] = icmp ne i32 [[TMP69]], 0 +// CHECK-USE-I386-NEXT: br i1 [[TMP70]], label [[OMP_OFFLOAD_FAILED6:%.*]], label [[OMP_OFFLOAD_CONT7:%.*]] +// CHECK-USE-I386: omp_offload.failed6: +// CHECK-USE-I386-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z20explicit_maps_singlei_l668(i32* [[A]]) #[[ATTR2]] +// CHECK-USE-I386-NEXT: br label [[OMP_OFFLOAD_CONT7]] +// CHECK-USE-I386: omp_offload.cont7: +// CHECK-USE-I386-NEXT: ret void +// +// CHECK-NOUSE-PPC64LE-LABEL: @_Z20explicit_maps_singlei( +// CHECK-NOUSE-PPC64LE-NEXT: entry: +// CHECK-NOUSE-PPC64LE-NEXT: [[II_ADDR:%.*]] = alloca i32, align 4 +// CHECK-NOUSE-PPC64LE-NEXT: [[A:%.*]] = alloca i32, align 4 +// CHECK-NOUSE-PPC64LE-NEXT: [[ST1:%.*]] = alloca [[STRUCT_ST:%.*]], align 4 +// CHECK-NOUSE-PPC64LE-NEXT: [[ST2:%.*]] = alloca [[STRUCT_ST]], align 4 +// CHECK-NOUSE-PPC64LE-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [7 x i8*], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [7 x i8*], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [7 x i8*], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [7 x i64], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[DOTOFFLOAD_BASEPTRS3:%.*]] = alloca [1 x i8*], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[DOTOFFLOAD_PTRS4:%.*]] = alloca [1 x i8*], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[DOTOFFLOAD_MAPPERS5:%.*]] = alloca [1 x i8*], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: store i32 [[II:%.*]], i32* [[II_ADDR]], align 4 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP0:%.*]] = load i32, i32* [[II_ADDR]], align 4 +// CHECK-NOUSE-PPC64LE-NEXT: store i32 [[TMP0]], i32* [[A]], align 4 +// CHECK-NOUSE-PPC64LE-NEXT: [[I:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[ST1]], i32 0, i32 0 +// CHECK-NOUSE-PPC64LE-NEXT: [[J:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[ST1]], i32 0, i32 1 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP1:%.*]] = getelementptr i32, i32* [[J]], i32 1 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP2:%.*]] = bitcast i32* [[I]] to i8* +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP3:%.*]] = bitcast i32* [[TMP1]] to i8* +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP4:%.*]] = ptrtoint i8* [[TMP3]] to i64 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP5:%.*]] = ptrtoint i8* [[TMP2]] to i64 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP6:%.*]] = sub i64 [[TMP4]], [[TMP5]] +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP7:%.*]] = sdiv exact i64 [[TMP6]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) +// CHECK-NOUSE-PPC64LE-NEXT: [[I1:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[ST2]], i32 0, i32 0 +// CHECK-NOUSE-PPC64LE-NEXT: [[J2:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[ST2]], i32 0, i32 1 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP8:%.*]] = getelementptr i32, i32* [[J2]], i32 1 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP9:%.*]] = bitcast i32* [[I1]] to i8* +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP10:%.*]] = bitcast i32* [[TMP8]] to i8* +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP11:%.*]] = ptrtoint i8* [[TMP10]] to i64 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP12:%.*]] = ptrtoint i8* [[TMP9]] to i64 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP13:%.*]] = sub i64 [[TMP11]], [[TMP12]] +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP14:%.*]] = sdiv exact i64 [[TMP13]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP15:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP16:%.*]] = bitcast i8** [[TMP15]] to %struct.ST** +// CHECK-NOUSE-PPC64LE-NEXT: store %struct.ST* [[ST1]], %struct.ST** [[TMP16]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP17:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP18:%.*]] = bitcast i8** [[TMP17]] to i32** +// CHECK-NOUSE-PPC64LE-NEXT: store i32* [[I]], i32** [[TMP18]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP19:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0 +// CHECK-NOUSE-PPC64LE-NEXT: store i64 [[TMP7]], i64* [[TMP19]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP20:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0 +// CHECK-NOUSE-PPC64LE-NEXT: store i8* null, i8** [[TMP20]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP21:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP22:%.*]] = bitcast i8** [[TMP21]] to %struct.ST** +// CHECK-NOUSE-PPC64LE-NEXT: store %struct.ST* [[ST1]], %struct.ST** [[TMP22]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP23:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 1 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP24:%.*]] = bitcast i8** [[TMP23]] to i32** +// CHECK-NOUSE-PPC64LE-NEXT: store i32* [[I]], i32** [[TMP24]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP25:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 1 +// CHECK-NOUSE-PPC64LE-NEXT: store i64 4, i64* [[TMP25]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP26:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 +// CHECK-NOUSE-PPC64LE-NEXT: store i8* null, i8** [[TMP26]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP27:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP28:%.*]] = bitcast i8** [[TMP27]] to %struct.ST** +// CHECK-NOUSE-PPC64LE-NEXT: store %struct.ST* [[ST1]], %struct.ST** [[TMP28]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP29:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 2 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP30:%.*]] = bitcast i8** [[TMP29]] to i32** +// CHECK-NOUSE-PPC64LE-NEXT: store i32* [[J]], i32** [[TMP30]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP31:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 2 +// CHECK-NOUSE-PPC64LE-NEXT: store i64 4, i64* [[TMP31]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP32:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2 +// CHECK-NOUSE-PPC64LE-NEXT: store i8* null, i8** [[TMP32]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP33:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 3 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP34:%.*]] = bitcast i8** [[TMP33]] to i32** +// CHECK-NOUSE-PPC64LE-NEXT: store i32* [[A]], i32** [[TMP34]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP35:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 3 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP36:%.*]] = bitcast i8** [[TMP35]] to i32** +// CHECK-NOUSE-PPC64LE-NEXT: store i32* [[A]], i32** [[TMP36]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP37:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 3 +// CHECK-NOUSE-PPC64LE-NEXT: store i64 4, i64* [[TMP37]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP38:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 3 +// CHECK-NOUSE-PPC64LE-NEXT: store i8* null, i8** [[TMP38]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP39:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 4 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP40:%.*]] = bitcast i8** [[TMP39]] to %struct.ST** +// CHECK-NOUSE-PPC64LE-NEXT: store %struct.ST* [[ST2]], %struct.ST** [[TMP40]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP41:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 4 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP42:%.*]] = bitcast i8** [[TMP41]] to i32** +// CHECK-NOUSE-PPC64LE-NEXT: store i32* [[I1]], i32** [[TMP42]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP43:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 4 +// CHECK-NOUSE-PPC64LE-NEXT: store i64 [[TMP14]], i64* [[TMP43]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP44:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 4 +// CHECK-NOUSE-PPC64LE-NEXT: store i8* null, i8** [[TMP44]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP45:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 5 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP46:%.*]] = bitcast i8** [[TMP45]] to %struct.ST** +// CHECK-NOUSE-PPC64LE-NEXT: store %struct.ST* [[ST2]], %struct.ST** [[TMP46]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP47:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 5 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP48:%.*]] = bitcast i8** [[TMP47]] to i32** +// CHECK-NOUSE-PPC64LE-NEXT: store i32* [[I1]], i32** [[TMP48]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP49:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 5 +// CHECK-NOUSE-PPC64LE-NEXT: store i64 4, i64* [[TMP49]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP50:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 5 +// CHECK-NOUSE-PPC64LE-NEXT: store i8* null, i8** [[TMP50]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP51:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 6 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP52:%.*]] = bitcast i8** [[TMP51]] to %struct.ST** +// CHECK-NOUSE-PPC64LE-NEXT: store %struct.ST* [[ST2]], %struct.ST** [[TMP52]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP53:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 6 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP54:%.*]] = bitcast i8** [[TMP53]] to i32** +// CHECK-NOUSE-PPC64LE-NEXT: store i32* [[J2]], i32** [[TMP54]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP55:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 6 +// CHECK-NOUSE-PPC64LE-NEXT: store i64 4, i64* [[TMP55]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP56:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 6 +// CHECK-NOUSE-PPC64LE-NEXT: store i8* null, i8** [[TMP56]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP57:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP58:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP59:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP60:%.*]] = call i32 @__tgt_target_mapper(%struct.ident_t* @[[GLOB1:[0-9]+]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z20explicit_maps_singlei_l654.region_id, i32 7, i8** [[TMP57]], i8** [[TMP58]], i64* [[TMP59]], i64* getelementptr inbounds ([7 x i64], [7 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null) +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP61:%.*]] = icmp ne i32 [[TMP60]], 0 +// CHECK-NOUSE-PPC64LE-NEXT: br i1 [[TMP61]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] +// CHECK-NOUSE-PPC64LE: omp_offload.failed: +// CHECK-NOUSE-PPC64LE-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z20explicit_maps_singlei_l654() #[[ATTR2:[0-9]+]] +// CHECK-NOUSE-PPC64LE-NEXT: br label [[OMP_OFFLOAD_CONT]] +// CHECK-NOUSE-PPC64LE: omp_offload.cont: +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP62:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP63:%.*]] = bitcast i8** [[TMP62]] to i32** +// CHECK-NOUSE-PPC64LE-NEXT: store i32* [[A]], i32** [[TMP63]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP64:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS4]], i32 0, i32 0 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP65:%.*]] = bitcast i8** [[TMP64]] to i32** +// CHECK-NOUSE-PPC64LE-NEXT: store i32* [[A]], i32** [[TMP65]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP66:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_MAPPERS5]], i64 0, i64 0 +// CHECK-NOUSE-PPC64LE-NEXT: store i8* null, i8** [[TMP66]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP67:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP68:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS4]], i32 0, i32 0 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP69:%.*]] = call i32 @__tgt_target_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z20explicit_maps_singlei_l668.region_id, i32 1, i8** [[TMP67]], i8** [[TMP68]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.1, i32 0, i32 0), i8** null, i8** null) +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP70:%.*]] = icmp ne i32 [[TMP69]], 0 +// CHECK-NOUSE-PPC64LE-NEXT: br i1 [[TMP70]], label [[OMP_OFFLOAD_FAILED6:%.*]], label [[OMP_OFFLOAD_CONT7:%.*]] +// CHECK-NOUSE-PPC64LE: omp_offload.failed6: +// CHECK-NOUSE-PPC64LE-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z20explicit_maps_singlei_l668() #[[ATTR2]] +// CHECK-NOUSE-PPC64LE-NEXT: br label [[OMP_OFFLOAD_CONT7]] +// CHECK-NOUSE-PPC64LE: omp_offload.cont7: +// CHECK-NOUSE-PPC64LE-NEXT: ret void +// +// CHECK-NOUSE-I386-LABEL: @_Z20explicit_maps_singlei( +// CHECK-NOUSE-I386-NEXT: entry: +// CHECK-NOUSE-I386-NEXT: [[II_ADDR:%.*]] = alloca i32, align 4 +// CHECK-NOUSE-I386-NEXT: [[A:%.*]] = alloca i32, align 4 +// CHECK-NOUSE-I386-NEXT: [[ST1:%.*]] = alloca [[STRUCT_ST:%.*]], align 4 +// CHECK-NOUSE-I386-NEXT: [[ST2:%.*]] = alloca [[STRUCT_ST]], align 4 +// CHECK-NOUSE-I386-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [7 x i8*], align 4 +// CHECK-NOUSE-I386-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [7 x i8*], align 4 +// CHECK-NOUSE-I386-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [7 x i8*], align 4 +// CHECK-NOUSE-I386-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [7 x i64], align 4 +// CHECK-NOUSE-I386-NEXT: [[DOTOFFLOAD_BASEPTRS3:%.*]] = alloca [1 x i8*], align 4 +// CHECK-NOUSE-I386-NEXT: [[DOTOFFLOAD_PTRS4:%.*]] = alloca [1 x i8*], align 4 +// CHECK-NOUSE-I386-NEXT: [[DOTOFFLOAD_MAPPERS5:%.*]] = alloca [1 x i8*], align 4 +// CHECK-NOUSE-I386-NEXT: store i32 [[II:%.*]], i32* [[II_ADDR]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP0:%.*]] = load i32, i32* [[II_ADDR]], align 4 +// CHECK-NOUSE-I386-NEXT: store i32 [[TMP0]], i32* [[A]], align 4 +// CHECK-NOUSE-I386-NEXT: [[I:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[ST1]], i32 0, i32 0 +// CHECK-NOUSE-I386-NEXT: [[J:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[ST1]], i32 0, i32 1 +// CHECK-NOUSE-I386-NEXT: [[TMP1:%.*]] = getelementptr i32, i32* [[J]], i32 1 +// CHECK-NOUSE-I386-NEXT: [[TMP2:%.*]] = bitcast i32* [[I]] to i8* +// CHECK-NOUSE-I386-NEXT: [[TMP3:%.*]] = bitcast i32* [[TMP1]] to i8* +// CHECK-NOUSE-I386-NEXT: [[TMP4:%.*]] = ptrtoint i8* [[TMP3]] to i64 +// CHECK-NOUSE-I386-NEXT: [[TMP5:%.*]] = ptrtoint i8* [[TMP2]] to i64 +// CHECK-NOUSE-I386-NEXT: [[TMP6:%.*]] = sub i64 [[TMP4]], [[TMP5]] +// CHECK-NOUSE-I386-NEXT: [[TMP7:%.*]] = sdiv exact i64 [[TMP6]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) +// CHECK-NOUSE-I386-NEXT: [[I1:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[ST2]], i32 0, i32 0 +// CHECK-NOUSE-I386-NEXT: [[J2:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[ST2]], i32 0, i32 1 +// CHECK-NOUSE-I386-NEXT: [[TMP8:%.*]] = getelementptr i32, i32* [[J2]], i32 1 +// CHECK-NOUSE-I386-NEXT: [[TMP9:%.*]] = bitcast i32* [[I1]] to i8* +// CHECK-NOUSE-I386-NEXT: [[TMP10:%.*]] = bitcast i32* [[TMP8]] to i8* +// CHECK-NOUSE-I386-NEXT: [[TMP11:%.*]] = ptrtoint i8* [[TMP10]] to i64 +// CHECK-NOUSE-I386-NEXT: [[TMP12:%.*]] = ptrtoint i8* [[TMP9]] to i64 +// CHECK-NOUSE-I386-NEXT: [[TMP13:%.*]] = sub i64 [[TMP11]], [[TMP12]] +// CHECK-NOUSE-I386-NEXT: [[TMP14:%.*]] = sdiv exact i64 [[TMP13]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) +// CHECK-NOUSE-I386-NEXT: [[TMP15:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK-NOUSE-I386-NEXT: [[TMP16:%.*]] = bitcast i8** [[TMP15]] to %struct.ST** +// CHECK-NOUSE-I386-NEXT: store %struct.ST* [[ST1]], %struct.ST** [[TMP16]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP17:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK-NOUSE-I386-NEXT: [[TMP18:%.*]] = bitcast i8** [[TMP17]] to i32** +// CHECK-NOUSE-I386-NEXT: store i32* [[I]], i32** [[TMP18]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP19:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0 +// CHECK-NOUSE-I386-NEXT: store i64 [[TMP7]], i64* [[TMP19]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP20:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0 +// CHECK-NOUSE-I386-NEXT: store i8* null, i8** [[TMP20]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP21:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 +// CHECK-NOUSE-I386-NEXT: [[TMP22:%.*]] = bitcast i8** [[TMP21]] to %struct.ST** +// CHECK-NOUSE-I386-NEXT: store %struct.ST* [[ST1]], %struct.ST** [[TMP22]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP23:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 1 +// CHECK-NOUSE-I386-NEXT: [[TMP24:%.*]] = bitcast i8** [[TMP23]] to i32** +// CHECK-NOUSE-I386-NEXT: store i32* [[I]], i32** [[TMP24]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP25:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 1 +// CHECK-NOUSE-I386-NEXT: store i64 4, i64* [[TMP25]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP26:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 1 +// CHECK-NOUSE-I386-NEXT: store i8* null, i8** [[TMP26]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP27:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2 +// CHECK-NOUSE-I386-NEXT: [[TMP28:%.*]] = bitcast i8** [[TMP27]] to %struct.ST** +// CHECK-NOUSE-I386-NEXT: store %struct.ST* [[ST1]], %struct.ST** [[TMP28]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP29:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 2 +// CHECK-NOUSE-I386-NEXT: [[TMP30:%.*]] = bitcast i8** [[TMP29]] to i32** +// CHECK-NOUSE-I386-NEXT: store i32* [[J]], i32** [[TMP30]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP31:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 2 +// CHECK-NOUSE-I386-NEXT: store i64 4, i64* [[TMP31]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP32:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 2 +// CHECK-NOUSE-I386-NEXT: store i8* null, i8** [[TMP32]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP33:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 3 +// CHECK-NOUSE-I386-NEXT: [[TMP34:%.*]] = bitcast i8** [[TMP33]] to i32** +// CHECK-NOUSE-I386-NEXT: store i32* [[A]], i32** [[TMP34]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP35:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 3 +// CHECK-NOUSE-I386-NEXT: [[TMP36:%.*]] = bitcast i8** [[TMP35]] to i32** +// CHECK-NOUSE-I386-NEXT: store i32* [[A]], i32** [[TMP36]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP37:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 3 +// CHECK-NOUSE-I386-NEXT: store i64 4, i64* [[TMP37]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP38:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 3 +// CHECK-NOUSE-I386-NEXT: store i8* null, i8** [[TMP38]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP39:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 4 +// CHECK-NOUSE-I386-NEXT: [[TMP40:%.*]] = bitcast i8** [[TMP39]] to %struct.ST** +// CHECK-NOUSE-I386-NEXT: store %struct.ST* [[ST2]], %struct.ST** [[TMP40]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP41:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 4 +// CHECK-NOUSE-I386-NEXT: [[TMP42:%.*]] = bitcast i8** [[TMP41]] to i32** +// CHECK-NOUSE-I386-NEXT: store i32* [[I1]], i32** [[TMP42]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP43:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 4 +// CHECK-NOUSE-I386-NEXT: store i64 [[TMP14]], i64* [[TMP43]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP44:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 4 +// CHECK-NOUSE-I386-NEXT: store i8* null, i8** [[TMP44]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP45:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 5 +// CHECK-NOUSE-I386-NEXT: [[TMP46:%.*]] = bitcast i8** [[TMP45]] to %struct.ST** +// CHECK-NOUSE-I386-NEXT: store %struct.ST* [[ST2]], %struct.ST** [[TMP46]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP47:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 5 +// CHECK-NOUSE-I386-NEXT: [[TMP48:%.*]] = bitcast i8** [[TMP47]] to i32** +// CHECK-NOUSE-I386-NEXT: store i32* [[I1]], i32** [[TMP48]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP49:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 5 +// CHECK-NOUSE-I386-NEXT: store i64 4, i64* [[TMP49]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP50:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 5 +// CHECK-NOUSE-I386-NEXT: store i8* null, i8** [[TMP50]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP51:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 6 +// CHECK-NOUSE-I386-NEXT: [[TMP52:%.*]] = bitcast i8** [[TMP51]] to %struct.ST** +// CHECK-NOUSE-I386-NEXT: store %struct.ST* [[ST2]], %struct.ST** [[TMP52]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP53:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 6 +// CHECK-NOUSE-I386-NEXT: [[TMP54:%.*]] = bitcast i8** [[TMP53]] to i32** +// CHECK-NOUSE-I386-NEXT: store i32* [[J2]], i32** [[TMP54]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP55:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 6 +// CHECK-NOUSE-I386-NEXT: store i64 4, i64* [[TMP55]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP56:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 6 +// CHECK-NOUSE-I386-NEXT: store i8* null, i8** [[TMP56]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP57:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK-NOUSE-I386-NEXT: [[TMP58:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK-NOUSE-I386-NEXT: [[TMP59:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0 +// CHECK-NOUSE-I386-NEXT: [[TMP60:%.*]] = call i32 @__tgt_target_mapper(%struct.ident_t* @[[GLOB1:[0-9]+]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z20explicit_maps_singlei_l654.region_id, i32 7, i8** [[TMP57]], i8** [[TMP58]], i64* [[TMP59]], i64* getelementptr inbounds ([7 x i64], [7 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null) +// CHECK-NOUSE-I386-NEXT: [[TMP61:%.*]] = icmp ne i32 [[TMP60]], 0 +// CHECK-NOUSE-I386-NEXT: br i1 [[TMP61]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] +// CHECK-NOUSE-I386: omp_offload.failed: +// CHECK-NOUSE-I386-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z20explicit_maps_singlei_l654() #[[ATTR2:[0-9]+]] +// CHECK-NOUSE-I386-NEXT: br label [[OMP_OFFLOAD_CONT]] +// CHECK-NOUSE-I386: omp_offload.cont: +// CHECK-NOUSE-I386-NEXT: [[TMP62:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0 +// CHECK-NOUSE-I386-NEXT: [[TMP63:%.*]] = bitcast i8** [[TMP62]] to i32** +// CHECK-NOUSE-I386-NEXT: store i32* [[A]], i32** [[TMP63]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP64:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS4]], i32 0, i32 0 +// CHECK-NOUSE-I386-NEXT: [[TMP65:%.*]] = bitcast i8** [[TMP64]] to i32** +// CHECK-NOUSE-I386-NEXT: store i32* [[A]], i32** [[TMP65]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP66:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_MAPPERS5]], i32 0, i32 0 +// CHECK-NOUSE-I386-NEXT: store i8* null, i8** [[TMP66]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP67:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0 +// CHECK-NOUSE-I386-NEXT: [[TMP68:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS4]], i32 0, i32 0 +// CHECK-NOUSE-I386-NEXT: [[TMP69:%.*]] = call i32 @__tgt_target_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z20explicit_maps_singlei_l668.region_id, i32 1, i8** [[TMP67]], i8** [[TMP68]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.1, i32 0, i32 0), i8** null, i8** null) +// CHECK-NOUSE-I386-NEXT: [[TMP70:%.*]] = icmp ne i32 [[TMP69]], 0 +// CHECK-NOUSE-I386-NEXT: br i1 [[TMP70]], label [[OMP_OFFLOAD_FAILED6:%.*]], label [[OMP_OFFLOAD_CONT7:%.*]] +// CHECK-NOUSE-I386: omp_offload.failed6: +// CHECK-NOUSE-I386-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z20explicit_maps_singlei_l668() #[[ATTR2]] +// CHECK-NOUSE-I386-NEXT: br label [[OMP_OFFLOAD_CONT7]] +// CHECK-NOUSE-I386: omp_offload.cont7: +// CHECK-NOUSE-I386-NEXT: ret void +// +void explicit_maps_single(int ii) { + + // Map of a scalar. + int a = ii; + + struct ST st1; + struct ST st2; + + // Make sure the struct picks up ompx_hold even if another element of the + // struct doesn't have ompx_hold. + #pragma omp target map(tofrom : st1.i) \ + map(ompx_hold, tofrom : a, st1.j, st2.i) \ + map(tofrom : st2.j) + { +#ifdef USE + st1.i++; + a++; + st1.j++; + st2.i++; + st2.j++; +#endif + } + + // Always Close Hold. + #pragma omp target map(always close ompx_hold tofrom: a) + { +#ifdef USE + a++; +#endif + } +} + +// CHECK-USE-PPC64LE-LABEL: @_ZN2ST20test_present_membersEv( +// CHECK-USE-PPC64LE-NEXT: entry: +// CHECK-USE-PPC64LE-NEXT: [[THIS_ADDR:%.*]] = alloca %struct.ST*, align 8 +// CHECK-USE-PPC64LE-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x i8*], align 8 +// CHECK-USE-PPC64LE-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x i8*], align 8 +// CHECK-USE-PPC64LE-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x i8*], align 8 +// CHECK-USE-PPC64LE-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [3 x i64], align 8 +// CHECK-USE-PPC64LE-NEXT: store %struct.ST* [[THIS:%.*]], %struct.ST** [[THIS_ADDR]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[THIS1:%.*]] = load %struct.ST*, %struct.ST** [[THIS_ADDR]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[I:%.*]] = getelementptr inbounds [[STRUCT_ST:%.*]], %struct.ST* [[THIS1]], i32 0, i32 0 +// CHECK-USE-PPC64LE-NEXT: [[J:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[THIS1]], i32 0, i32 1 +// CHECK-USE-PPC64LE-NEXT: [[TMP0:%.*]] = getelementptr i32, i32* [[J]], i32 1 +// CHECK-USE-PPC64LE-NEXT: [[TMP1:%.*]] = bitcast i32* [[I]] to i8* +// CHECK-USE-PPC64LE-NEXT: [[TMP2:%.*]] = bitcast i32* [[TMP0]] to i8* +// CHECK-USE-PPC64LE-NEXT: [[TMP3:%.*]] = ptrtoint i8* [[TMP2]] to i64 +// CHECK-USE-PPC64LE-NEXT: [[TMP4:%.*]] = ptrtoint i8* [[TMP1]] to i64 +// CHECK-USE-PPC64LE-NEXT: [[TMP5:%.*]] = sub i64 [[TMP3]], [[TMP4]] +// CHECK-USE-PPC64LE-NEXT: [[TMP6:%.*]] = sdiv exact i64 [[TMP5]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) +// CHECK-USE-PPC64LE-NEXT: [[TMP7:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK-USE-PPC64LE-NEXT: [[TMP8:%.*]] = bitcast i8** [[TMP7]] to %struct.ST** +// CHECK-USE-PPC64LE-NEXT: store %struct.ST* [[THIS1]], %struct.ST** [[TMP8]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP9:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK-USE-PPC64LE-NEXT: [[TMP10:%.*]] = bitcast i8** [[TMP9]] to i32** +// CHECK-USE-PPC64LE-NEXT: store i32* [[I]], i32** [[TMP10]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP11:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0 +// CHECK-USE-PPC64LE-NEXT: store i64 [[TMP6]], i64* [[TMP11]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP12:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0 +// CHECK-USE-PPC64LE-NEXT: store i8* null, i8** [[TMP12]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP13:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 +// CHECK-USE-PPC64LE-NEXT: [[TMP14:%.*]] = bitcast i8** [[TMP13]] to %struct.ST** +// CHECK-USE-PPC64LE-NEXT: store %struct.ST* [[THIS1]], %struct.ST** [[TMP14]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP15:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 1 +// CHECK-USE-PPC64LE-NEXT: [[TMP16:%.*]] = bitcast i8** [[TMP15]] to i32** +// CHECK-USE-PPC64LE-NEXT: store i32* [[I]], i32** [[TMP16]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP17:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 1 +// CHECK-USE-PPC64LE-NEXT: store i64 4, i64* [[TMP17]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP18:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 +// CHECK-USE-PPC64LE-NEXT: store i8* null, i8** [[TMP18]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP19:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2 +// CHECK-USE-PPC64LE-NEXT: [[TMP20:%.*]] = bitcast i8** [[TMP19]] to %struct.ST** +// CHECK-USE-PPC64LE-NEXT: store %struct.ST* [[THIS1]], %struct.ST** [[TMP20]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP21:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 2 +// CHECK-USE-PPC64LE-NEXT: [[TMP22:%.*]] = bitcast i8** [[TMP21]] to i32** +// CHECK-USE-PPC64LE-NEXT: store i32* [[J]], i32** [[TMP22]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP23:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 2 +// CHECK-USE-PPC64LE-NEXT: store i64 4, i64* [[TMP23]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP24:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2 +// CHECK-USE-PPC64LE-NEXT: store i8* null, i8** [[TMP24]], align 8 +// CHECK-USE-PPC64LE-NEXT: [[TMP25:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK-USE-PPC64LE-NEXT: [[TMP26:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK-USE-PPC64LE-NEXT: [[TMP27:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0 +// CHECK-USE-PPC64LE-NEXT: [[TMP28:%.*]] = call i32 @__tgt_target_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2ST20test_present_membersEv_l919.region_id, i32 3, i8** [[TMP25]], i8** [[TMP26]], i64* [[TMP27]], i64* getelementptr inbounds ([3 x i64], [3 x i64]* @.offload_maptypes.2, i32 0, i32 0), i8** null, i8** null) +// CHECK-USE-PPC64LE-NEXT: [[TMP29:%.*]] = icmp ne i32 [[TMP28]], 0 +// CHECK-USE-PPC64LE-NEXT: br i1 [[TMP29]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] +// CHECK-USE-PPC64LE: omp_offload.failed: +// CHECK-USE-PPC64LE-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2ST20test_present_membersEv_l919(%struct.ST* [[THIS1]]) #[[ATTR2]] +// CHECK-USE-PPC64LE-NEXT: br label [[OMP_OFFLOAD_CONT]] +// CHECK-USE-PPC64LE: omp_offload.cont: +// CHECK-USE-PPC64LE-NEXT: ret void +// +// CHECK-USE-I386-LABEL: @_ZN2ST20test_present_membersEv( +// CHECK-USE-I386-NEXT: entry: +// CHECK-USE-I386-NEXT: [[THIS_ADDR:%.*]] = alloca %struct.ST*, align 4 +// CHECK-USE-I386-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x i8*], align 4 +// CHECK-USE-I386-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x i8*], align 4 +// CHECK-USE-I386-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x i8*], align 4 +// CHECK-USE-I386-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [3 x i64], align 4 +// CHECK-USE-I386-NEXT: store %struct.ST* [[THIS:%.*]], %struct.ST** [[THIS_ADDR]], align 4 +// CHECK-USE-I386-NEXT: [[THIS1:%.*]] = load %struct.ST*, %struct.ST** [[THIS_ADDR]], align 4 +// CHECK-USE-I386-NEXT: [[I:%.*]] = getelementptr inbounds [[STRUCT_ST:%.*]], %struct.ST* [[THIS1]], i32 0, i32 0 +// CHECK-USE-I386-NEXT: [[J:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[THIS1]], i32 0, i32 1 +// CHECK-USE-I386-NEXT: [[TMP0:%.*]] = getelementptr i32, i32* [[J]], i32 1 +// CHECK-USE-I386-NEXT: [[TMP1:%.*]] = bitcast i32* [[I]] to i8* +// CHECK-USE-I386-NEXT: [[TMP2:%.*]] = bitcast i32* [[TMP0]] to i8* +// CHECK-USE-I386-NEXT: [[TMP3:%.*]] = ptrtoint i8* [[TMP2]] to i64 +// CHECK-USE-I386-NEXT: [[TMP4:%.*]] = ptrtoint i8* [[TMP1]] to i64 +// CHECK-USE-I386-NEXT: [[TMP5:%.*]] = sub i64 [[TMP3]], [[TMP4]] +// CHECK-USE-I386-NEXT: [[TMP6:%.*]] = sdiv exact i64 [[TMP5]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) +// CHECK-USE-I386-NEXT: [[TMP7:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK-USE-I386-NEXT: [[TMP8:%.*]] = bitcast i8** [[TMP7]] to %struct.ST** +// CHECK-USE-I386-NEXT: store %struct.ST* [[THIS1]], %struct.ST** [[TMP8]], align 4 +// CHECK-USE-I386-NEXT: [[TMP9:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK-USE-I386-NEXT: [[TMP10:%.*]] = bitcast i8** [[TMP9]] to i32** +// CHECK-USE-I386-NEXT: store i32* [[I]], i32** [[TMP10]], align 4 +// CHECK-USE-I386-NEXT: [[TMP11:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0 +// CHECK-USE-I386-NEXT: store i64 [[TMP6]], i64* [[TMP11]], align 4 +// CHECK-USE-I386-NEXT: [[TMP12:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0 +// CHECK-USE-I386-NEXT: store i8* null, i8** [[TMP12]], align 4 +// CHECK-USE-I386-NEXT: [[TMP13:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 +// CHECK-USE-I386-NEXT: [[TMP14:%.*]] = bitcast i8** [[TMP13]] to %struct.ST** +// CHECK-USE-I386-NEXT: store %struct.ST* [[THIS1]], %struct.ST** [[TMP14]], align 4 +// CHECK-USE-I386-NEXT: [[TMP15:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 1 +// CHECK-USE-I386-NEXT: [[TMP16:%.*]] = bitcast i8** [[TMP15]] to i32** +// CHECK-USE-I386-NEXT: store i32* [[I]], i32** [[TMP16]], align 4 +// CHECK-USE-I386-NEXT: [[TMP17:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 1 +// CHECK-USE-I386-NEXT: store i64 4, i64* [[TMP17]], align 4 +// CHECK-USE-I386-NEXT: [[TMP18:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 1 +// CHECK-USE-I386-NEXT: store i8* null, i8** [[TMP18]], align 4 +// CHECK-USE-I386-NEXT: [[TMP19:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2 +// CHECK-USE-I386-NEXT: [[TMP20:%.*]] = bitcast i8** [[TMP19]] to %struct.ST** +// CHECK-USE-I386-NEXT: store %struct.ST* [[THIS1]], %struct.ST** [[TMP20]], align 4 +// CHECK-USE-I386-NEXT: [[TMP21:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 2 +// CHECK-USE-I386-NEXT: [[TMP22:%.*]] = bitcast i8** [[TMP21]] to i32** +// CHECK-USE-I386-NEXT: store i32* [[J]], i32** [[TMP22]], align 4 +// CHECK-USE-I386-NEXT: [[TMP23:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 2 +// CHECK-USE-I386-NEXT: store i64 4, i64* [[TMP23]], align 4 +// CHECK-USE-I386-NEXT: [[TMP24:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 2 +// CHECK-USE-I386-NEXT: store i8* null, i8** [[TMP24]], align 4 +// CHECK-USE-I386-NEXT: [[TMP25:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK-USE-I386-NEXT: [[TMP26:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK-USE-I386-NEXT: [[TMP27:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0 +// CHECK-USE-I386-NEXT: [[TMP28:%.*]] = call i32 @__tgt_target_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2ST20test_present_membersEv_l919.region_id, i32 3, i8** [[TMP25]], i8** [[TMP26]], i64* [[TMP27]], i64* getelementptr inbounds ([3 x i64], [3 x i64]* @.offload_maptypes.2, i32 0, i32 0), i8** null, i8** null) +// CHECK-USE-I386-NEXT: [[TMP29:%.*]] = icmp ne i32 [[TMP28]], 0 +// CHECK-USE-I386-NEXT: br i1 [[TMP29]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] +// CHECK-USE-I386: omp_offload.failed: +// CHECK-USE-I386-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2ST20test_present_membersEv_l919(%struct.ST* [[THIS1]]) #[[ATTR2]] +// CHECK-USE-I386-NEXT: br label [[OMP_OFFLOAD_CONT]] +// CHECK-USE-I386: omp_offload.cont: +// CHECK-USE-I386-NEXT: ret void +// +// CHECK-NOUSE-PPC64LE-LABEL: @_ZN2ST20test_present_membersEv( +// CHECK-NOUSE-PPC64LE-NEXT: entry: +// CHECK-NOUSE-PPC64LE-NEXT: [[THIS_ADDR:%.*]] = alloca %struct.ST*, align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x i8*], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x i8*], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x i8*], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [3 x i64], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: store %struct.ST* [[THIS:%.*]], %struct.ST** [[THIS_ADDR]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[THIS1:%.*]] = load %struct.ST*, %struct.ST** [[THIS_ADDR]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[I:%.*]] = getelementptr inbounds [[STRUCT_ST:%.*]], %struct.ST* [[THIS1]], i32 0, i32 0 +// CHECK-NOUSE-PPC64LE-NEXT: [[J:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[THIS1]], i32 0, i32 1 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP0:%.*]] = getelementptr i32, i32* [[J]], i32 1 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP1:%.*]] = bitcast i32* [[I]] to i8* +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP2:%.*]] = bitcast i32* [[TMP0]] to i8* +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP3:%.*]] = ptrtoint i8* [[TMP2]] to i64 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP4:%.*]] = ptrtoint i8* [[TMP1]] to i64 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP5:%.*]] = sub i64 [[TMP3]], [[TMP4]] +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP6:%.*]] = sdiv exact i64 [[TMP5]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP7:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP8:%.*]] = bitcast i8** [[TMP7]] to %struct.ST** +// CHECK-NOUSE-PPC64LE-NEXT: store %struct.ST* [[THIS1]], %struct.ST** [[TMP8]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP9:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP10:%.*]] = bitcast i8** [[TMP9]] to i32** +// CHECK-NOUSE-PPC64LE-NEXT: store i32* [[I]], i32** [[TMP10]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP11:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0 +// CHECK-NOUSE-PPC64LE-NEXT: store i64 [[TMP6]], i64* [[TMP11]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP12:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0 +// CHECK-NOUSE-PPC64LE-NEXT: store i8* null, i8** [[TMP12]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP13:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP14:%.*]] = bitcast i8** [[TMP13]] to %struct.ST** +// CHECK-NOUSE-PPC64LE-NEXT: store %struct.ST* [[THIS1]], %struct.ST** [[TMP14]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP15:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 1 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP16:%.*]] = bitcast i8** [[TMP15]] to i32** +// CHECK-NOUSE-PPC64LE-NEXT: store i32* [[I]], i32** [[TMP16]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP17:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 1 +// CHECK-NOUSE-PPC64LE-NEXT: store i64 4, i64* [[TMP17]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP18:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 +// CHECK-NOUSE-PPC64LE-NEXT: store i8* null, i8** [[TMP18]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP19:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP20:%.*]] = bitcast i8** [[TMP19]] to %struct.ST** +// CHECK-NOUSE-PPC64LE-NEXT: store %struct.ST* [[THIS1]], %struct.ST** [[TMP20]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP21:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 2 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP22:%.*]] = bitcast i8** [[TMP21]] to i32** +// CHECK-NOUSE-PPC64LE-NEXT: store i32* [[J]], i32** [[TMP22]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP23:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 2 +// CHECK-NOUSE-PPC64LE-NEXT: store i64 4, i64* [[TMP23]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP24:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2 +// CHECK-NOUSE-PPC64LE-NEXT: store i8* null, i8** [[TMP24]], align 8 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP25:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP26:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP27:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0 +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP28:%.*]] = call i32 @__tgt_target_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2ST20test_present_membersEv_l919.region_id, i32 3, i8** [[TMP25]], i8** [[TMP26]], i64* [[TMP27]], i64* getelementptr inbounds ([3 x i64], [3 x i64]* @.offload_maptypes.2, i32 0, i32 0), i8** null, i8** null) +// CHECK-NOUSE-PPC64LE-NEXT: [[TMP29:%.*]] = icmp ne i32 [[TMP28]], 0 +// CHECK-NOUSE-PPC64LE-NEXT: br i1 [[TMP29]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] +// CHECK-NOUSE-PPC64LE: omp_offload.failed: +// CHECK-NOUSE-PPC64LE-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2ST20test_present_membersEv_l919() #[[ATTR2]] +// CHECK-NOUSE-PPC64LE-NEXT: br label [[OMP_OFFLOAD_CONT]] +// CHECK-NOUSE-PPC64LE: omp_offload.cont: +// CHECK-NOUSE-PPC64LE-NEXT: ret void +// +// CHECK-NOUSE-I386-LABEL: @_ZN2ST20test_present_membersEv( +// CHECK-NOUSE-I386-NEXT: entry: +// CHECK-NOUSE-I386-NEXT: [[THIS_ADDR:%.*]] = alloca %struct.ST*, align 4 +// CHECK-NOUSE-I386-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x i8*], align 4 +// CHECK-NOUSE-I386-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x i8*], align 4 +// CHECK-NOUSE-I386-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x i8*], align 4 +// CHECK-NOUSE-I386-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [3 x i64], align 4 +// CHECK-NOUSE-I386-NEXT: store %struct.ST* [[THIS:%.*]], %struct.ST** [[THIS_ADDR]], align 4 +// CHECK-NOUSE-I386-NEXT: [[THIS1:%.*]] = load %struct.ST*, %struct.ST** [[THIS_ADDR]], align 4 +// CHECK-NOUSE-I386-NEXT: [[I:%.*]] = getelementptr inbounds [[STRUCT_ST:%.*]], %struct.ST* [[THIS1]], i32 0, i32 0 +// CHECK-NOUSE-I386-NEXT: [[J:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[THIS1]], i32 0, i32 1 +// CHECK-NOUSE-I386-NEXT: [[TMP0:%.*]] = getelementptr i32, i32* [[J]], i32 1 +// CHECK-NOUSE-I386-NEXT: [[TMP1:%.*]] = bitcast i32* [[I]] to i8* +// CHECK-NOUSE-I386-NEXT: [[TMP2:%.*]] = bitcast i32* [[TMP0]] to i8* +// CHECK-NOUSE-I386-NEXT: [[TMP3:%.*]] = ptrtoint i8* [[TMP2]] to i64 +// CHECK-NOUSE-I386-NEXT: [[TMP4:%.*]] = ptrtoint i8* [[TMP1]] to i64 +// CHECK-NOUSE-I386-NEXT: [[TMP5:%.*]] = sub i64 [[TMP3]], [[TMP4]] +// CHECK-NOUSE-I386-NEXT: [[TMP6:%.*]] = sdiv exact i64 [[TMP5]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) +// CHECK-NOUSE-I386-NEXT: [[TMP7:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK-NOUSE-I386-NEXT: [[TMP8:%.*]] = bitcast i8** [[TMP7]] to %struct.ST** +// CHECK-NOUSE-I386-NEXT: store %struct.ST* [[THIS1]], %struct.ST** [[TMP8]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP9:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK-NOUSE-I386-NEXT: [[TMP10:%.*]] = bitcast i8** [[TMP9]] to i32** +// CHECK-NOUSE-I386-NEXT: store i32* [[I]], i32** [[TMP10]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP11:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0 +// CHECK-NOUSE-I386-NEXT: store i64 [[TMP6]], i64* [[TMP11]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP12:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0 +// CHECK-NOUSE-I386-NEXT: store i8* null, i8** [[TMP12]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP13:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 +// CHECK-NOUSE-I386-NEXT: [[TMP14:%.*]] = bitcast i8** [[TMP13]] to %struct.ST** +// CHECK-NOUSE-I386-NEXT: store %struct.ST* [[THIS1]], %struct.ST** [[TMP14]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP15:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 1 +// CHECK-NOUSE-I386-NEXT: [[TMP16:%.*]] = bitcast i8** [[TMP15]] to i32** +// CHECK-NOUSE-I386-NEXT: store i32* [[I]], i32** [[TMP16]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP17:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 1 +// CHECK-NOUSE-I386-NEXT: store i64 4, i64* [[TMP17]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP18:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 1 +// CHECK-NOUSE-I386-NEXT: store i8* null, i8** [[TMP18]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP19:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2 +// CHECK-NOUSE-I386-NEXT: [[TMP20:%.*]] = bitcast i8** [[TMP19]] to %struct.ST** +// CHECK-NOUSE-I386-NEXT: store %struct.ST* [[THIS1]], %struct.ST** [[TMP20]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP21:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 2 +// CHECK-NOUSE-I386-NEXT: [[TMP22:%.*]] = bitcast i8** [[TMP21]] to i32** +// CHECK-NOUSE-I386-NEXT: store i32* [[J]], i32** [[TMP22]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP23:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 2 +// CHECK-NOUSE-I386-NEXT: store i64 4, i64* [[TMP23]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP24:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 2 +// CHECK-NOUSE-I386-NEXT: store i8* null, i8** [[TMP24]], align 4 +// CHECK-NOUSE-I386-NEXT: [[TMP25:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK-NOUSE-I386-NEXT: [[TMP26:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK-NOUSE-I386-NEXT: [[TMP27:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0 +// CHECK-NOUSE-I386-NEXT: [[TMP28:%.*]] = call i32 @__tgt_target_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2ST20test_present_membersEv_l919.region_id, i32 3, i8** [[TMP25]], i8** [[TMP26]], i64* [[TMP27]], i64* getelementptr inbounds ([3 x i64], [3 x i64]* @.offload_maptypes.2, i32 0, i32 0), i8** null, i8** null) +// CHECK-NOUSE-I386-NEXT: [[TMP29:%.*]] = icmp ne i32 [[TMP28]], 0 +// CHECK-NOUSE-I386-NEXT: br i1 [[TMP29]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] +// CHECK-NOUSE-I386: omp_offload.failed: +// CHECK-NOUSE-I386-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2ST20test_present_membersEv_l919() #[[ATTR2]] +// CHECK-NOUSE-I386-NEXT: br label [[OMP_OFFLOAD_CONT]] +// CHECK-NOUSE-I386: omp_offload.cont: +// CHECK-NOUSE-I386-NEXT: ret void +// +void ST::test_present_members() { + // Make sure the struct picks up ompx_hold even if another element of the + // struct doesn't have ompx_hold. + #pragma omp target map(tofrom : i) map(ompx_hold, tofrom : j) + { +#ifdef USE + i++; + j++; +#endif + } +} + +#endif Index: clang/test/OpenMP/target_map_messages.cpp =================================================================== --- clang/test/OpenMP/target_map_messages.cpp +++ clang/test/OpenMP/target_map_messages.cpp @@ -1,16 +1,36 @@ -// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized -// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp -fopenmp-version=40 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized -// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp -fopenmp-version=45 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized -// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp -fopenmp-version=50 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized -// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp -fopenmp-version=51 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized -// RUN: %clang_cc1 -DCCODE -verify -fopenmp -ferror-limit 200 -x c %s -Wno-openmp -Wuninitialized - -// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp-simd -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized -// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp-simd -fopenmp-version=40 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized -// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp-simd -fopenmp-version=45 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized -// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp-simd -fopenmp-version=50 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized -// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp-simd -fopenmp-version=51 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized -// RUN: %clang_cc1 -DCCODE -verify -fopenmp-simd -ferror-limit 200 -x c %s -Wno-openmp-mapping -Wuninitialized +// -fopenmp +// RUN: %clang_cc1 -verify=expected,ge50,lt51,omp,lt51-omp -fopenmp -ferror-limit 300 %s -Wno-openmp-target -Wuninitialized +// RUN: %clang_cc1 -verify=expected,lt50,lt51,omp,lt51-omp -fopenmp -fopenmp-version=40 -ferror-limit 300 %s -Wno-openmp-target -Wuninitialized +// RUN: %clang_cc1 -verify=expected,lt50,lt51,omp,lt51-omp -fopenmp -fopenmp-version=45 -ferror-limit 300 %s -Wno-openmp-target -Wuninitialized +// RUN: %clang_cc1 -verify=expected,ge50,lt51,omp,lt51-omp -fopenmp -fopenmp-version=50 -ferror-limit 300 %s -Wno-openmp-target -Wuninitialized +// RUN: %clang_cc1 -verify=expected,ge50,ge51,omp,ge51-omp -fopenmp -fopenmp-version=51 -ferror-limit 300 %s -Wno-openmp-target -Wuninitialized +// RUN: %clang_cc1 -DCCODE -verify -fopenmp -ferror-limit 300 -x c %s -Wno-openmp -Wuninitialized + +// -fopenmp-simd +// RUN: %clang_cc1 -verify=expected,ge50,lt51,omp,lt51-omp -fopenmp-simd -ferror-limit 300 %s -Wno-openmp-target -Wuninitialized +// RUN: %clang_cc1 -verify=expected,lt50,lt51,omp,lt51-omp -fopenmp-simd -fopenmp-version=40 -ferror-limit 300 %s -Wno-openmp-target -Wuninitialized +// RUN: %clang_cc1 -verify=expected,lt50,lt51,omp,lt51-omp -fopenmp-simd -fopenmp-version=45 -ferror-limit 300 %s -Wno-openmp-target -Wuninitialized +// RUN: %clang_cc1 -verify=expected,ge50,lt51,omp,lt51-omp -fopenmp-simd -fopenmp-version=50 -ferror-limit 300 %s -Wno-openmp-target -Wuninitialized +// RUN: %clang_cc1 -verify=expected,ge50,ge51,omp,ge51-omp -fopenmp-simd -fopenmp-version=51 -ferror-limit 300 %s -Wno-openmp-target -Wuninitialized +// RUN: %clang_cc1 -DCCODE -verify -fopenmp-simd -ferror-limit 300 -x c %s -Wno-openmp-mapping -Wuninitialized + +// -fopenmp -fopenmp-extensions +// RUN: %clang_cc1 -verify=expected,ge50,lt51,ompx,lt51-ompx -fopenmp -fopenmp-extensions -ferror-limit 300 %s -Wno-openmp-target -Wuninitialized +// RUN: %clang_cc1 -verify=expected,lt50,lt51,ompx,lt51-ompx -fopenmp -fopenmp-extensions -fopenmp-version=40 -ferror-limit 300 %s -Wno-openmp-target -Wuninitialized +// RUN: %clang_cc1 -verify=expected,lt50,lt51,ompx,lt51-ompx -fopenmp -fopenmp-extensions -fopenmp-version=45 -ferror-limit 300 %s -Wno-openmp-target -Wuninitialized +// RUN: %clang_cc1 -verify=expected,ge50,lt51,ompx,lt51-ompx -fopenmp -fopenmp-extensions -fopenmp-version=50 -ferror-limit 300 %s -Wno-openmp-target -Wuninitialized +// RUN: %clang_cc1 -verify=expected,ge50,ge51,ompx,ge51-ompx -fopenmp -fopenmp-extensions -fopenmp-version=51 -ferror-limit 300 %s -Wno-openmp-target -Wuninitialized +// RUN: %clang_cc1 -DCCODE -verify -fopenmp -fopenmp-extensions -ferror-limit 300 -x c %s -Wno-openmp -Wuninitialized + +// -fopenmp-simd -fopenmp-extensions +// RUN: %clang_cc1 -verify=expected,ge50,lt51,ompx,lt51-ompx -fopenmp-simd -fopenmp-extensions -ferror-limit 300 %s -Wno-openmp-target -Wuninitialized +// RUN: %clang_cc1 -verify=expected,lt50,lt51,ompx,lt51-ompx -fopenmp-simd -fopenmp-extensions -fopenmp-version=40 -ferror-limit 300 %s -Wno-openmp-target -Wuninitialized +// RUN: %clang_cc1 -verify=expected,lt50,lt51,ompx,lt51-ompx -fopenmp-simd -fopenmp-extensions -fopenmp-version=45 -ferror-limit 300 %s -Wno-openmp-target -Wuninitialized +// RUN: %clang_cc1 -verify=expected,ge50,lt51,ompx,lt51-ompx -fopenmp-simd -fopenmp-extensions -fopenmp-version=50 -ferror-limit 300 %s -Wno-openmp-target -Wuninitialized +// RUN: %clang_cc1 -verify=expected,ge50,ge51,ompx,ge51-ompx -fopenmp-simd -fopenmp-extensions -fopenmp-version=51 -ferror-limit 300 %s -Wno-openmp-target -Wuninitialized +// RUN: %clang_cc1 -DCCODE -verify -fopenmp-simd -fopenmp-extensions -ferror-limit 300 -x c %s -Wno-openmp-mapping -Wuninitialized + +// Check #ifdef CCODE void foo(int arg) { const int n = 0; @@ -118,38 +138,72 @@ {} #pragma omp target map(close) // expected-error {{use of undeclared identifier 'close'}} {} - // lt51-error@+1 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} + // lt51-error@+1 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} #pragma omp target map(present, tofrom: c,f) {} - // lt51-error@+1 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} + // lt51-error@+1 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} #pragma omp target map(present, tofrom: c[1:2],f) {} - // lt51-error@+1 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} + // lt51-error@+1 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} #pragma omp target map(present, tofrom: c,f[1:2]) {} // expected-error@+2 {{section length is unspecified and cannot be inferred because subscripted value is not an array}} - // lt51-error@+1 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} + // lt51-error@+1 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} #pragma omp target map(present, tofrom: c[:],f) {} // expected-error@+2 {{section length is unspecified and cannot be inferred because subscripted value is not an array}} - // lt51-error@+1 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} + // lt51-error@+1 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} #pragma omp target map(present, tofrom: c,f[:]) {} // expected-error@+1 {{use of undeclared identifier 'present'}} #pragma omp target map(present) {} + // ge51-omp-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}} + // lt51-omp-error@+1 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} + #pragma omp target map(ompx_hold, tofrom: c,f) + {} + // ge51-omp-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}} + // lt51-omp-error@+1 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} + #pragma omp target map(ompx_hold, tofrom: c[1:2],f) + {} + // ge51-omp-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}} + // lt51-omp-error@+1 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} + #pragma omp target map(ompx_hold, tofrom: c,f[1:2]) + {} + // expected-error@+3 {{section length is unspecified and cannot be inferred because subscripted value is not an array}} + // ge51-omp-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}} + // lt51-omp-error@+1 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} + #pragma omp target map(ompx_hold, tofrom: c[:],f) + {} + // expected-error@+3 {{section length is unspecified and cannot be inferred because subscripted value is not an array}} + // ge51-omp-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}} + // lt51-omp-error@+1 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} + #pragma omp target map(ompx_hold, tofrom: c,f[:]) + {} + // expected-error@+1 {{use of undeclared identifier 'ompx_hold'}} + #pragma omp target map(ompx_hold) + {} #pragma omp target map(close, close, tofrom: a) // expected-error {{same map type modifier has been specified more than once}} {} #pragma omp target map(always, close, always, close, tofrom: a) // expected-error 2 {{same map type modifier has been specified more than once}} {} // ge51-error@+2 {{same map type modifier has been specified more than once}} - // lt51-error@+1 2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} + // lt51-error@+1 2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} #pragma omp target map(present, present, tofrom: a) {} - // expected-error@+3 2 {{same map type modifier has been specified more than once}} - // ge51-error@+2 1 {{same map type modifier has been specified more than once}} - // lt51-error@+1 2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} - #pragma omp target map(always, close, present, always, close, present, tofrom: a) + // ompx-error@+3 {{same map type modifier has been specified more than once}} + // ge51-omp-error@+2 2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}} + // lt51-omp-error@+1 2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} + #pragma omp target map(ompx_hold, ompx_hold, tofrom: a) + {} + // expected-error@+7 2 {{same map type modifier has been specified more than once}} + // ge51-error@+6 {{same map type modifier has been specified more than once}} + // lt51-ompx-error@+5 2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'ompx_hold'}} + // lt51-omp-error@+4 2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} + // ompx-error@+3 {{same map type modifier has been specified more than once}} + // ge51-omp-error@+2 2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}} + // lt51-omp-error@+1 2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} + #pragma omp target map(always, close, present, ompx_hold, always, close, present, ompx_hold, tofrom: a) {} #pragma omp target map( , tofrom: a) // expected-error {{missing map type modifier}} {} @@ -157,14 +211,14 @@ {} #pragma omp target map( , , : a) // expected-error {{missing map type modifier}} expected-error {{missing map type modifier}} expected-error {{missing map type}} {} - // ge51-error@+3 2 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} - // lt51-error@+2 2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} + // ge51-error@+3 2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}} + // lt51-error@+2 2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} // expected-error@+1 {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}} #pragma omp target map( d, f, bf: a) {} // expected-error@+4 {{missing map type modifier}} - // ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} - // lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} + // ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}} + // lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} // expected-error@+1 {{missing map type}} #pragma omp target map( , f, : a) {} @@ -172,13 +226,13 @@ {} #pragma omp target map(always close bf: a) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}} {} - // ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} - // lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} + // ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}} + // lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} // expected-error@+1 {{missing map type}} #pragma omp target map(always tofrom close: a) {} - // ge51-error@+2 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} - // lt51-error@+1 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} + // ge51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}} + // lt51-error@+1 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} #pragma omp target map(tofrom from: a) {} #pragma omp target map(close bf: a) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}} @@ -600,8 +654,8 @@ #pragma omp target data map(always, tofrom: x) #pragma omp target data map(always: x) // expected-error {{missing map type}} -// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} -// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} +// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}} +// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} // expected-error@+1 {{missing map type}} #pragma omp target data map(tofrom, always: x) #pragma omp target data map(always, tofrom: always, tofrom, x) @@ -610,24 +664,24 @@ #pragma omp target data map(close, tofrom: x) #pragma omp target data map(close: x) // expected-error {{missing map type}} -// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} -// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} +// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}} +// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} // expected-error@+1 {{missing map type}} #pragma omp target data map(tofrom, close: x) #pragma omp target data map(close, tofrom: close, tofrom, x) foo(); -// lt51-error@+1 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} +// lt51-error@+1 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} #pragma omp target data map(present, tofrom: x) // ge51-error@+2 {{missing map type}} // lt51-error@+1 {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}} #pragma omp target data map(present: x) -// ge51-error@+4 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} -// lt51-error@+3 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} +// ge51-error@+4 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}} +// lt51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} // ge51-error@+2 {{missing map type}} // lt51-error@+1 {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}} #pragma omp target data map(tofrom, present: x) -// lt51-error@+1 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} +// lt51-error@+1 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} #pragma omp target data map(present, tofrom: present, tofrom, x) foo(); @@ -735,8 +789,8 @@ #pragma omp target data map(always, tofrom: x) #pragma omp target data map(always: x) // expected-error {{missing map type}} -// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} -// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} +// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}} +// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} // expected-error@+1 {{missing map type}} #pragma omp target data map(tofrom, always: x) #pragma omp target data map(always, tofrom: always, tofrom, x) @@ -744,18 +798,18 @@ foo(); #pragma omp target data map(close, tofrom: x) #pragma omp target data map(close: x) // expected-error {{missing map type}} -// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} -// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} +// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}} +// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} // expected-error@+1 {{missing map type}} #pragma omp target data map(tofrom, close: x) foo(); -// lt51-error@+1 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} +// lt51-error@+1 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} #pragma omp target data map(present, tofrom: x) // ge51-error@+2 {{missing map type}} // lt51-error@+1 {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}} #pragma omp target data map(present: x) -// ge51-error@+4 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} -// lt51-error@+3 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} +// ge51-error@+4 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}} +// lt51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} // ge51-error@+2 {{missing map type}} // lt51-error@+1 {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}} #pragma omp target data map(tofrom, present: x) Index: clang/test/OpenMP/target_parallel_for_map_messages.cpp =================================================================== --- clang/test/OpenMP/target_parallel_for_map_messages.cpp +++ clang/test/OpenMP/target_parallel_for_map_messages.cpp @@ -182,8 +182,8 @@ for (i = 0; i < argc; ++i) foo(); #pragma omp target parallel for map(always: x) // expected-error {{missing map type}} for (i = 0; i < argc; ++i) foo(); -// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} -// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} +// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}} +// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} // expected-error@+1 {{missing map type}} #pragma omp target parallel for map(tofrom, always: x) for (i = 0; i < argc; ++i) foo(); @@ -300,8 +300,8 @@ for (i = 0; i < argc; ++i) foo(); #pragma omp target parallel for map(always: x) // expected-error {{missing map type}} for (i = 0; i < argc; ++i) foo(); -// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} -// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} +// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}} +// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} // expected-error@+1 {{missing map type}} #pragma omp target parallel for map(tofrom, always: x) for (i = 0; i < argc; ++i) foo(); Index: clang/test/OpenMP/target_parallel_for_simd_map_messages.cpp =================================================================== --- clang/test/OpenMP/target_parallel_for_simd_map_messages.cpp +++ clang/test/OpenMP/target_parallel_for_simd_map_messages.cpp @@ -182,8 +182,8 @@ for (i = 0; i < argc; ++i) foo(); #pragma omp target parallel for simd map(always: x) // expected-error {{missing map type}} for (i = 0; i < argc; ++i) foo(); -// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} -// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} +// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}} +// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} // expected-error@+1 {{missing map type}} #pragma omp target parallel for simd map(tofrom, always: x) for (i = 0; i < argc; ++i) foo(); @@ -300,8 +300,8 @@ for (i = 0; i < argc; ++i) foo(); #pragma omp target parallel for simd map(always: x) // expected-error {{missing map type}} for (i = 0; i < argc; ++i) foo(); -// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} -// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} +// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}} +// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} // expected-error@+1 {{missing map type}} #pragma omp target parallel for simd map(tofrom, always: x) for (i = 0; i < argc; ++i) foo(); Index: clang/test/OpenMP/target_parallel_map_messages.cpp =================================================================== --- clang/test/OpenMP/target_parallel_map_messages.cpp +++ clang/test/OpenMP/target_parallel_map_messages.cpp @@ -181,8 +181,8 @@ foo(); #pragma omp target parallel map(always: x) // expected-error {{missing map type}} foo(); -// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} -// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} +// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}} +// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} // expected-error@+1 {{missing map type}} #pragma omp target parallel map(tofrom, always: x) foo(); @@ -296,8 +296,8 @@ foo(); #pragma omp target parallel map(always: x) // expected-error {{missing map type}} foo(); -// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} -// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} +// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}} +// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} // expected-error@+1 {{missing map type}} #pragma omp target parallel map(tofrom, always: x) foo(); Index: clang/test/OpenMP/target_simd_map_messages.cpp =================================================================== --- clang/test/OpenMP/target_simd_map_messages.cpp +++ clang/test/OpenMP/target_simd_map_messages.cpp @@ -176,8 +176,8 @@ for (i = 0; i < argc; ++i) foo(); #pragma omp target simd map(always: x) // expected-error {{missing map type}} for (i = 0; i < argc; ++i) foo(); -// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} -// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} +// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}} +// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} // expected-error@+1 {{missing map type}} #pragma omp target simd map(tofrom, always: x) for (i = 0; i < argc; ++i) foo(); @@ -293,8 +293,8 @@ for (i = 0; i < argc; ++i) foo(); #pragma omp target simd map(always: x) // expected-error {{missing map type}} for (i = 0; i < argc; ++i) foo(); -// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} -// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} +// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}} +// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} // expected-error@+1 {{missing map type}} #pragma omp target simd map(tofrom, always: x) for (i = 0; i < argc; ++i) foo(); Index: clang/test/OpenMP/target_teams_distribute_map_messages.cpp =================================================================== --- clang/test/OpenMP/target_teams_distribute_map_messages.cpp +++ clang/test/OpenMP/target_teams_distribute_map_messages.cpp @@ -182,8 +182,8 @@ for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute map(always: x) // expected-error {{missing map type}} for (i = 0; i < argc; ++i) foo(); -// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} -// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} +// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}} +// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} // expected-error@+1 {{missing map type}} #pragma omp target teams distribute map(tofrom, always: x) for (i = 0; i < argc; ++i) foo(); @@ -300,8 +300,8 @@ for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute map(always: x) // expected-error {{missing map type}} for (i = 0; i < argc; ++i) foo(); -// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} -// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} +// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}} +// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} // expected-error@+1 {{missing map type}} #pragma omp target teams distribute map(tofrom, always: x) for (i = 0; i < argc; ++i) foo(); Index: clang/test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp =================================================================== --- clang/test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp +++ clang/test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp @@ -180,8 +180,8 @@ for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute parallel for map(always: x) // expected-error {{missing map type}} for (i = 0; i < argc; ++i) foo(); -// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} -// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} +// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}} +// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} // expected-error@+1 {{missing map type}} #pragma omp target teams distribute parallel for map(tofrom, always: x) for (i = 0; i < argc; ++i) foo(); @@ -302,8 +302,8 @@ for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute parallel for map(always: x) // expected-error {{missing map type}} for (i = 0; i < argc; ++i) foo(); -// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} -// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} +// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}} +// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} // expected-error@+1 {{missing map type}} #pragma omp target teams distribute parallel for map(tofrom, always: x) for (i = 0; i < argc; ++i) foo(); Index: clang/test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp =================================================================== --- clang/test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp +++ clang/test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp @@ -182,8 +182,8 @@ for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute parallel for simd map(always: x) // expected-error {{missing map type}} for (i = 0; i < argc; ++i) foo(); -// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} -// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} +// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}} +// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} // expected-error@+1 {{missing map type}} #pragma omp target teams distribute parallel for simd map(tofrom, always: x) for (i = 0; i < argc; ++i) foo(); @@ -299,8 +299,8 @@ for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute parallel for simd map(always: x) // expected-error {{missing map type}} for (i = 0; i < argc; ++i) foo(); -// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} -// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} +// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}} +// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} // expected-error@+1 {{missing map type}} #pragma omp target teams distribute parallel for simd map(tofrom, always: x) for (i = 0; i < argc; ++i) foo(); Index: clang/test/OpenMP/target_teams_distribute_simd_map_messages.cpp =================================================================== --- clang/test/OpenMP/target_teams_distribute_simd_map_messages.cpp +++ clang/test/OpenMP/target_teams_distribute_simd_map_messages.cpp @@ -182,8 +182,8 @@ for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute simd map(always: x) // expected-error {{missing map type}} for (i = 0; i < argc; ++i) foo(); -// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} -// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} +// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}} +// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} // expected-error@+1 {{missing map type}} #pragma omp target teams distribute simd map(tofrom, always: x) for (i = 0; i < argc; ++i) foo(); @@ -300,8 +300,8 @@ for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute simd map(always: x) // expected-error {{missing map type}} for (i = 0; i < argc; ++i) foo(); -// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} -// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} +// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}} +// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} // expected-error@+1 {{missing map type}} #pragma omp target teams distribute simd map(tofrom, always: x) for (i = 0; i < argc; ++i) foo(); Index: clang/test/OpenMP/target_teams_map_messages.cpp =================================================================== --- clang/test/OpenMP/target_teams_map_messages.cpp +++ clang/test/OpenMP/target_teams_map_messages.cpp @@ -479,8 +479,8 @@ #pragma omp target data map(always, tofrom: x) #pragma omp target data map(always: x) // expected-error {{missing map type}} -// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} -// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} +// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}} +// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} // expected-error@+1 {{missing map type}} #pragma omp target data map(tofrom, always: x) #pragma omp target data map(always, tofrom: always, tofrom, x) @@ -562,8 +562,8 @@ #pragma omp target data map(always, tofrom: x) #pragma omp target data map(always: x) // expected-error {{missing map type}} -// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} -// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} +// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}} +// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}} // expected-error@+1 {{missing map type}} #pragma omp target data map(tofrom, always: x) #pragma omp target data map(always, tofrom: always, tofrom, x) Index: openmp/docs/index.rst =================================================================== --- openmp/docs/index.rst +++ openmp/docs/index.rst @@ -28,6 +28,20 @@ design/Overview +OpenACC Support +=============== + +:doc:`OpenACC support ` is under development for +both Flang and Clang. For this purpose, LLVM's OpenMP runtimes are +being extended to serve as OpenACC runtimes. In some cases, Clang +supports :doc:`OpenMP extensions ` to make +the additional functionality also available in OpenMP applications. + +.. toctree:: + :hidden: + :maxdepth: 1 + + openacc/Overview LLVM/OpenMP Optimizations ========================= Index: openmp/docs/openacc/OpenMPExtensions.rst =================================================================== --- /dev/null +++ openmp/docs/openacc/OpenMPExtensions.rst @@ -0,0 +1,138 @@ +OpenMP Extensions for OpenACC +============================= + +OpenACC provides some functionality that OpenMP does not. In some +cases, Clang supports OpenMP extensions to provide similar +functionality, taking advantage of the runtime implementation already +required for OpenACC. This section documents those extensions. + +By default, Clang recognizes only standard OpenMP. The Clang +command-line option ``-fopenmp-extensions`` is required to enable all +OpenMP extensions, including those described in this section. + +.. _ompx-motivation: + +Motivation +---------- + +There are multiple benefits to exposing OpenACC functionality as LLVM +OpenMP extensions: + +* OpenMP applications can take advantage of the additional + functionality. +* As LLVM's implementation of these extensions matures, it can serve + as a basis for including these extensions in the OpenMP standard. +* Source-to-source translation from certain OpenACC features to OpenMP + is otherwise impossible. +* Runtime tests can be written in terms of OpenMP instead of OpenACC + or low-level runtime calls. +* More generally, there is a clean separation of concerns between + OpenACC and OpenMP development in LLVM. That is, LLVM's OpenMP + developers can discuss, modify, and debug LLVM's extended OpenMP + implementation and test suite without directly considering OpenACC's + language and execution model, which are handled by LLVM's OpenACC + developers. + +.. _ompx-hold: + +``ompx_hold`` Map Type Modifier +------------------------------- + +.. _ompx-holdExample: + +Example +^^^^^^^ + +.. code-block:: c++ + + #pragma omp target data map(ompx_hold, tofrom: x) // holds onto mapping of x throughout region + { + foo(); // might have map(delete: x) + #pragma omp target map(present, alloc: x) // x is guaranteed to be present + printf("%d\n", x); + } + +The ``ompx_hold`` map type modifier above specifies that the ``target +data`` directive holds onto the mapping for ``x`` throughout the +associated region regardless of any ``target exit data`` directives +executed during the call to ``foo``. Thus, the presence assertion for +``x`` at the enclosed ``target`` construct cannot fail. + +.. _ompx-holdBehavior: + +Behavior +^^^^^^^^ + +* Stated more generally, the ``ompx_hold`` map type modifier specifies + that, regardless of changes in the value of the standard OpenMP + reference count, the specified data is not unmapped until the end of + the construct. If ``ompx_hold`` is specified for the same data on + lexically or dynamically enclosed constructs, there is no additional + effect as the data mapping is already held throughout their regions. +* The ``ompx_hold`` map type modifier is permitted to appear only on + ``target`` constructs (and associated combined constructs) and + ``target data`` constructs. It is not permitted to appear on + ``target enter data`` or ``target exit data`` directives because + there is no associated statement, so it is not meangingful to hold + onto a mapping until the end of the directive. +* The runtime reports an error if ``omp_target_disassociate_ptr`` is + called for a mapping for which the ``ompx_hold`` map type modifier + is in effect. +* Like the ``present`` map type modifier, the ``ompx_hold`` map type + modifier applies to an entire struct if it's specified for any + member of that struct even if other ``map`` clauses on the same + directive specify other members without the ``ompx_hold`` map type + modifier. +* ``ompx_hold`` support is not yet provided for ``defaultmap``. + +Implementation +^^^^^^^^^^^^^^ + +* LLVM uses the term *dynamic reference count* for the standard OpenMP + reference count for host/device data mappings. +* The ``ompx_hold`` map type modifier selects an alternate reference + count, called the *hold reference count*. +* A mapping is removed only once both its reference counts reach zero. +* Because ``ompx_hold`` can appear only constructs, increments and + decrements of the hold reference count are guaranteed to be + balanced, so it is impossible to decrement it below zero. +* The dynamic reference count is used wherever ``ompx_hold`` is not + specified (and possibly cannot be specified). Decrementing the + dynamic reference count has no effect if it is already zero. +* The runtime determines that the ``ompx_hold`` map type modifier is + *in effect* (see :ref:`Behavior ` above) when the + hold reference count is greater than zero. + +Relationship with OpenACC +^^^^^^^^^^^^^^^^^^^^^^^^^ + +OpenACC specifies two reference counts for tracking host/device data +mappings. Which reference count is used to implement an OpenACC +directive is determined by the nature of that directive, either +dynamic or structured: + +* The *dynamic reference count* is always used for ``enter data`` and + ``exit data`` directives and corresponding OpenACC routines. +* The *structured reference count* is always used for ``data`` and + compute constructs, which are similar to OpenMP's ``target data`` + and ``target`` constructs. + +Contrast with OpenMP, where the dynamic reference count is always used +unless the application developer specifies an alternate behavior via +our map type modifier extension. We chose the name *hold* for that +map type modifier because, as demonstrated in the above :ref:`example +`, *hold* concisely identifies the desired behavior +from the application developer's perspective without referencing the +implementation of that behavior. + +The hold reference count is otherwise modeled after OpenACC's +structured reference count. For example, calling ``acc_unmap_data``, +which is similar to ``omp_target_disassociate_ptr``, is an error when +the structured reference count is not zero. + +While Flang and Clang obviously must implement the syntax and +semantics for selecting OpenACC reference counts differently than for +selecting OpenMP reference counts, the implementation is the same at +the runtime level. That is, OpenACC's dynamic reference count is +OpenMP's dynamic reference count, and OpenACC's structured reference +count is our OpenMP hold reference count extension. Index: openmp/docs/openacc/Overview.rst =================================================================== --- /dev/null +++ openmp/docs/openacc/Overview.rst @@ -0,0 +1,13 @@ +OpenACC Support +=============== + +OpenACC support is under development for both Flang and Clang. For +this purpose, LLVM's OpenMP runtimes are being extended to serve as +OpenACC runtimes. + +.. toctree:: + :glob: + :hidden: + :maxdepth: 1 + + OpenMPExtensions