.. option:: -fopenmp-version=<arg>
+.. option:: -fopenmp-extensions, -fno-openmp-extensions
+
+Enable or disable all Clang extensions for OpenMP directives and clauses. By
+default, they are enabled.
+
.. program:: clang1
.. option:: -fopenmp=<arg>
.. program:: clang
+------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
| 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 |
+| | <https://openmp.llvm.org/docs/openacc/OpenMPExtensions.html#ompx-hold>`_ | | |
++------------------------------+---------------------------------------------------------------------------+--------------------------+--------------------------------------------------------+
/// 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];
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<
"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<
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, 1, "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")
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)
#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"
};
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
HelpText<"Parse OpenMP pragmas and generate parallel code.">;
def fno_openmp : Flag<["-"], "fno-openmp">, Group<f_Group>, Flags<[NoArgumentUnused]>;
def fopenmp_version_EQ : Joined<["-"], "fopenmp-version=">, Group<f_Group>, Flags<[CC1Option, NoArgumentUnused]>;
+defm openmp_extensions: BoolFOption<"openmp-extensions",
+ LangOpts<"OpenMPExtensions">, DefaultTrue,
+ PosFlag<SetTrue, [CC1Option, NoArgumentUnused],
+ "Enable all Clang extensions for OpenMP directives and clauses">,
+ NegFlag<SetFalse, [CC1Option, NoArgumentUnused],
+ "Disable all Clang extensions for OpenMP directives and clauses">>;
def fopenmp_EQ : Joined<["-"], "fopenmp=">, Group<f_Group>;
def fopenmp_use_tls : Flag<["-"], "fopenmp-use-tls">, Group<f_Group>,
Flags<[NoArgumentUnused, HelpHidden]>;
using namespace llvm::omp;
unsigned clang::getOpenMPSimpleClauseType(OpenMPClauseKind Kind, StringRef Str,
- unsigned OpenMPVersion) {
+ const LangOptions &LangOpts) {
switch (Kind) {
case OMPC_default:
return llvm::StringSwitch<unsigned>(Str)
.Case(#Name, static_cast<unsigned>(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;
}
.Case(#Name, static_cast<unsigned>(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;
}
/// 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
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;
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
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=*/true))
+ CmdArgs.push_back("-fno-openmp-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,
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=*/true))
+ CmdArgs.push_back("-fno-openmp-extensions");
}
const SanitizerArgs &Sanitize = TC.getSanitizerArgs();
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))
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;
// 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;
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();
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))
} 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)
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))
NextToken().is(tok::colon)) {
// Parse optional <device modifier> ':'
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 ':'
Preprocessor &PP = P.getPreprocessor();
OpenMPMapModifierKind TypeModifier =
static_cast<OpenMPMapModifierKind>(getOpenMPSimpleClauseType(
- OMPC_map, PP.getSpelling(Tok), P.getLangOpts().OpenMP));
+ OMPC_map, PP.getSpelling(Tok), P.getLangOpts()));
return TypeModifier;
}
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();
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))
Preprocessor &PP = P.getPreprocessor();
OpenMPMapClauseKind MapType =
static_cast<OpenMPMapClauseKind>(getOpenMPSimpleClauseType(
- OMPC_map, PP.getSpelling(Tok), P.getLangOpts().OpenMP));
+ OMPC_map, PP.getSpelling(Tok), P.getLangOpts()));
return MapType;
}
(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.");
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,
// 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;
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.");
Data.ColonLoc = ConsumeToken();
} else if (Kind == OMPC_to || Kind == OMPC_from) {
while (Tok.is(tok::identifier)) {
- auto Modifier =
- static_cast<OpenMPMotionModifierKind>(getOpenMPSimpleClauseType(
- Kind, PP.getSpelling(Tok), getLangOpts().OpenMP));
+ auto Modifier = static_cast<OpenMPMotionModifierKind>(
+ getOpenMPSimpleClauseType(Kind, PP.getSpelling(Tok), getLangOpts()));
if (Modifier == OMPC_MOTION_MODIFIER_unknown)
break;
Data.MotionModifiers.push_back(Modifier);
CXXScopeSpec &MapperIdScopeSpec, DeclarationNameInfo MapperId,
ArrayRef<Expr *> UnresolvedMappers,
OpenMPMapClauseKind MapType = OMPC_MAP_unknown,
+ ArrayRef<OpenMPMapModifierKind> Modifiers = None,
bool IsMapTypeImplicit = false, bool NoDiagnose = false) {
// We only expect mappable expressions in 'to', 'from', and 'map' clauses.
assert((CKind == OMPC_map || CKind == OMPC_to || CKind == OMPC_from) &&
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.
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]
ArrayRef<Expr *> 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.
MappableVarListInfo MVLI(VarList);
checkMappableExpressionList(*this, DSAStack, OMPC_map, MVLI, Locs.StartLoc,
MapperIdScopeSpec, MapperId, UnresolvedMappers,
- MapType, IsMapTypeImplicit, NoDiagnose);
+ MapType, Modifiers, IsMapTypeImplicit,
+ NoDiagnose);
// 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.
--- /dev/null
+// RUN: %clang -c -Xclang -verify=ompx -fopenmp %s
+// RUN: %clang -c -Xclang -verify=ompx -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 -c -Xclang -verify=omp -fopenmp -fno-openmp-extensions %s
+// RUN: %clang -c -Xclang -verify=omp -fopenmp-simd -fno-openmp-extensions %s
+
+// RUN: %clang -c -Xclang -verify=omp -fopenmp \
+// RUN: -fopenmp-extensions -fno-openmp-extensions %s
+// RUN: %clang -c -Xclang -verify=omp -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)
+ ;
+}
return tmain<int, 5>(argc, &argc) + tmain<char, 1>(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 <typename T, int C>
+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 <typename T, int C> 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<int, 5>(argc, &argc) + tmain<char, 1>(argv[0][0], argv[0]);
+}
+
+#endif
#endif
// 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
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)
{
#pragma omp target map(present, alloc: e)
foo();
#endif
+#ifdef OMPX
+ #pragma omp target map(ompx_hold, alloc: e)
+ foo();
+#endif
}
return 0;
// 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)
// 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, 5>(int argc, int *argv) {
// CHECK-NEXT: int i, j, b, c, d, e, x[20];
// CHECK-NEXT: #pragma omp target data map(to: c)
// 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)
// 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, 1>(char argc, char *argv) {
// CHECK-NEXT: char i, j, b, c, d, e, x[20];
// CHECK-NEXT: #pragma omp target data map(to: c)
// 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)
// 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];
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)
--- /dev/null
+// 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
-// 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 -fno-openmp-extensions -ferror-limit 100 -o - %s -Wuninitialized
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp -fopenmp -fno-openmp-extensions -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 -fno-openmp-extensions -ferror-limit 100 -o - %s -Wuninitialized
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp -fopenmp-simd -fno-openmp-extensions -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}}
#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;
}
-// 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 -fno-openmp-extensions -ferror-limit 100 -o - %s -Wuninitialized
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp -fopenmp -fno-openmp-extensions -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 -fno-openmp-extensions -ferror-limit 100 -o - %s -Wuninitialized
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp -fopenmp-simd -fno-openmp-extensions -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) {
#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;
}
--- /dev/null
+// 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
-// 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, -fno-openmp-extensions
+// RUN: %clang_cc1 -verify=expected,ge50,lt51,omp,lt51-omp -fopenmp -fno-openmp-extensions -ferror-limit 300 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt50,lt51,omp,lt51-omp -fopenmp -fno-openmp-extensions -fopenmp-version=40 -ferror-limit 300 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt50,lt51,omp,lt51-omp -fopenmp -fno-openmp-extensions -fopenmp-version=45 -ferror-limit 300 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51,omp,lt51-omp -fopenmp -fno-openmp-extensions -fopenmp-version=50 -ferror-limit 300 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,ge51,omp,ge51-omp -fopenmp -fno-openmp-extensions -fopenmp-version=51 -ferror-limit 300 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -DCCODE -verify -fopenmp -fno-openmp-extensions -ferror-limit 300 -x c %s -Wno-openmp -Wuninitialized
+
+// -fopenmp-simd, -fno-openmp-extensions
+// RUN: %clang_cc1 -verify=expected,ge50,lt51,omp,lt51-omp -fopenmp-simd -fno-openmp-extensions -ferror-limit 300 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt50,lt51,omp,lt51-omp -fopenmp-simd -fno-openmp-extensions -fopenmp-version=40 -ferror-limit 300 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt50,lt51,omp,lt51-omp -fopenmp-simd -fno-openmp-extensions -fopenmp-version=45 -ferror-limit 300 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51,omp,lt51-omp -fopenmp-simd -fno-openmp-extensions -fopenmp-version=50 -ferror-limit 300 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,ge51,omp,ge51-omp -fopenmp-simd -fno-openmp-extensions -fopenmp-version=51 -ferror-limit 300 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -DCCODE -verify -fopenmp-simd -fno-openmp-extensions -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;
{}
#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}}
{}
{}
#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)
{}
{}
#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'}}
#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)
#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();
#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)
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)
-// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp -fopenmp-version=51 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp -fno-openmp-extensions -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp -fno-openmp-extensions -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp -fno-openmp-extensions -fopenmp-version=51 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp-simd -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp-simd -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp-simd -fopenmp-version=51 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp-simd -fno-openmp-extensions -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp-simd -fno-openmp-extensions -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp-simd -fno-openmp-extensions -fopenmp-version=51 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
void foo() {
}
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();
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();
-// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp -fopenmp-version=45 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp -fopenmp-version=50 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp -fopenmp-version=51 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp -fno-openmp-extensions -fopenmp-version=45 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp -fno-openmp-extensions -fopenmp-version=50 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp -fno-openmp-extensions -fopenmp-version=51 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp-simd -fopenmp-version=45 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp-simd -fopenmp-version=50 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp-simd -fopenmp-version=51 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp-simd -fno-openmp-extensions -fopenmp-version=45 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp-simd -fno-openmp-extensions -fopenmp-version=50 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp-simd -fno-openmp-extensions -fopenmp-version=51 %s -Wno-openmp-mapping -Wuninitialized
void foo() {
}
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();
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();
-// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp -fopenmp-version=51 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp -fno-openmp-extensions -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp -fno-openmp-extensions -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp -fno-openmp-extensions -fopenmp-version=51 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp-simd -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp-simd -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp-simd -fopenmp-version=51 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp-simd -fno-openmp-extensions -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp-simd -fno-openmp-extensions -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp-simd -fno-openmp-extensions -fopenmp-version=51 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
void foo() {
}
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();
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();
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();
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();
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();
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();
-// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp -fopenmp-version=51 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp -fno-openmp-extensions -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp -fno-openmp-extensions -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp -fno-openmp-extensions -fopenmp-version=51 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp-simd -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp-simd -fno-openmp-extensions -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
void foo() {
}
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();
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();
-// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp -fopenmp-version=51 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp -fno-openmp-extensions -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp -fno-openmp-extensions -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp -fno-openmp-extensions -fopenmp-version=51 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp-simd -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp-simd -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp-simd -fopenmp-version=51 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp-simd -fno-openmp-extensions -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp-simd -fno-openmp-extensions -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp-simd -fno-openmp-extensions -fopenmp-version=51 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
void foo() {
}
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();
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();
-// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp -fopenmp-version=51 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp -fno-openmp-extensions -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp -fno-openmp-extensions -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp -fno-openmp-extensions -fopenmp-version=51 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp-simd -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp-simd -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp-simd -fopenmp-version=51 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp-simd -fno-openmp-extensions -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp-simd -fno-openmp-extensions -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp-simd -fno-openmp-extensions -fopenmp-version=51 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
void foo() {
}
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();
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();
-// RUN: %clang_cc1 -verify=expected,ge45,ge50,lt51 -fopenmp -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,lt45,lt50,lt51 -fopenmp-version=40 -fopenmp -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,ge45,lt50,lt51 -fopenmp-version=45 -fopenmp -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,ge45,ge50,lt51 -fopenmp-version=50 -fopenmp -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,ge45,ge50,ge51 -fopenmp-version=51 -fopenmp -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
-
-// RUN: %clang_cc1 -verify=expected,ge45,ge50,lt51 -fopenmp-simd -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -DCCODE -verify=expected,ge45,ge50,lt51 -fopenmp -ferror-limit 200 -x c %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge45,ge50,lt51 -fopenmp -fno-openmp-extensions -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt45,lt50,lt51 -fopenmp -fno-openmp-extensions -fopenmp-version=40 -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge45,lt50,lt51 -fopenmp -fno-openmp-extensions -fopenmp-version=45 -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge45,ge50,lt51 -fopenmp -fno-openmp-extensions -fopenmp-version=50 -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge45,ge50,ge51 -fopenmp -fno-openmp-extensions -fopenmp-version=51 -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
+
+// RUN: %clang_cc1 -verify=expected,ge45,ge50,lt51 -fopenmp-simd -fno-openmp-extensions -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -DCCODE -verify=expected,ge45,ge50,lt51 -fopenmp -fno-openmp-extensions -ferror-limit 200 -x c %s -Wno-openmp-mapping -Wuninitialized
#ifdef CCODE
void foo(int arg) {
const int n = 0;
#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)
#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)
design/Overview
+OpenACC Support
+===============
+
+:doc:`OpenACC support <openacc/Overview>` 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 <openacc/OpenMPExtensions>` to make
+the additional functionality also available in OpenMP applications.
+
+.. toctree::
+ :hidden:
+ :maxdepth: 1
+
+ openacc/Overview
LLVM/OpenMP Optimizations
=========================
--- /dev/null
+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 these extensions. The command-line
+option ``-fno-openmp-extensions`` can be specified to disable 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 the associated data is not unmapped until the end of the
+ construct. As usual, the standard OpenMP reference count for the
+ data must also reach zero before the data is unmapped.
+* 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 meaningful 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 <ompx-holdBehavior>` 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
+<ompx-holdExample>`, *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.
--- /dev/null
+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
projects / platform / upstream / llvm.git / commitdiff