[OpenMP][OpenACC] Implement `ompx_hold` map type modifier extension in Clang (1/2)

This patch implements Clang support for an original OpenMP extension
we have developed to support OpenACC: the `ompx_hold` map type
modifier.  The next patch in this series, D106510, implements OpenMP
runtime support.

Consider the following example:

```
 #pragma omp target data map(ompx_hold, tofrom: x) // holds onto mapping of x
 {
   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.  (As usual, the
standard OpenMP reference count for `x` must also reach zero before
the data is unmapped.)

Justification for inclusion in Clang and LLVM's OpenMP runtime:

* The `ompx_hold` modifier supports OpenACC functionality (structured
  reference count) that cannot be achieved in standard OpenMP, as of
  5.1.
* The runtime implementation for `ompx_hold` (next patch) will thus be
  used by Flang's OpenACC support.
* The Clang implementation for `ompx_hold` (this patch) as well as the
  runtime implementation are required for the Clang OpenACC support
  being developed as part of the ECP Clacc project, which translates
  OpenACC to OpenMP at the directive AST level.  These patches are the
  first step in upstreaming OpenACC functionality from Clacc.
* The Clang implementation for `ompx_hold` is also used by the tests
  in the runtime implementation.  That syntactic support makes the
  tests more readable than low-level runtime calls can.  Moreover,
  upstream Flang and Clang do not yet support OpenACC syntax
  sufficiently for writing the tests.
* More generally, the Clang implementation enables 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 can be
  handled by LLVM's OpenACC developers.
* OpenMP users might find the `ompx_hold` modifier useful, as in the
  above example.

See new documentation introduced by this patch in `openmp/docs` for
more detail on the functionality of this extension and its
relationship with OpenACC.  For example, it explains how the runtime
must support two reference counts, as specified by OpenACC.

Clang recognizes `ompx_hold` unless `-fno-openmp-extensions`, a new
command-line option introduced by this patch, is specified.

Reviewed By: ABataev, jdoerfert, protze.joachim, grokos

Differential Revision: https://reviews.llvm.org/D106509
This commit is contained in:
Joel E. Denny 2021-08-31 15:17:07 -04:00
parent dc37f5374c
commit 83ddfa0d22
34 changed files with 2169 additions and 174 deletions

View File

@ -2039,6 +2039,11 @@ Emit OpenMP code only for SIMD-based constructs.
.. 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

View File

@ -360,3 +360,20 @@ want to help with the implementation.
+------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
| 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>`_ | | |
+------------------------------+---------------------------------------------------------------------------+--------------------------+--------------------------------------------------------+

View File

@ -5606,7 +5606,8 @@ private:
/// 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];

View File

@ -1303,8 +1303,8 @@ def err_omp_decl_in_declare_simd_variant : Error<
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<

View File

@ -10534,6 +10534,8 @@ def err_omp_map_shared_storage : 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<

View File

@ -231,6 +231,7 @@ LANGOPT(HalfArgsAndReturns, 1, 0, "half args and returns")
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")

View File

@ -123,6 +123,8 @@ OPENMP_MAP_MODIFIER_KIND(always)
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)

View File

@ -14,6 +14,7 @@
#ifndef LLVM_CLANG_BASIC_OPENMPKINDS_H
#define LLVM_CLANG_BASIC_OPENMPKINDS_H
#include "clang/Basic/LangOptions.h"
#include "llvm/ADT/StringRef.h"
#include "llvm/Frontend/OpenMP/OMPConstants.h"
@ -167,7 +168,7 @@ enum OpenMPReductionClauseModifier {
};
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

View File

@ -2379,6 +2379,12 @@ def fopenmp : Flag<["-"], "fopenmp">, Group<f_Group>, Flags<[CC1Option, NoArgume
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]>;

View File

@ -21,7 +21,7 @@ using namespace clang;
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)
@ -59,7 +59,9 @@ unsigned clang::getOpenMPSimpleClauseType(OpenMPClauseKind Kind, StringRef 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;
}
@ -70,7 +72,7 @@ unsigned clang::getOpenMPSimpleClauseType(OpenMPClauseKind Kind, StringRef Str,
.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;
}

View File

@ -7269,6 +7269,14 @@ public:
/// 0x800 is reserved for compatibility with XLC.
/// Produce a runtime error if the data is not already allocated.
OMP_MAP_PRESENT = 0x1000,
// Increment and decrement a separate reference counter so that the data
// cannot be unmapped within the associated region. Thus, this flag is
// intended to be used on 'target' and 'target data' directives because they
// are inherently structured. It is not intended to be used on 'target
// enter data' and 'target exit data' directives because they are inherently
// dynamic.
// This is an OpenMP extension for the sake of OpenACC support.
OMP_MAP_OMPX_HOLD = 0x2000,
/// Signal that the runtime library should use args as an array of
/// descriptor_dim pointers and use args_size as dims. Used when we have
/// non-contiguous list items in target update directive
@ -7570,6 +7578,9 @@ private:
llvm::find(MotionModifiers, OMPC_MOTION_MODIFIER_present) !=
MotionModifiers.end())
Bits |= OMP_MAP_PRESENT;
if (llvm::find(MapModifiers, OMPC_MAP_MODIFIER_ompx_hold) !=
MapModifiers.end())
Bits |= OMP_MAP_OMPX_HOLD;
if (IsNonContiguous)
Bits |= OMP_MAP_NON_CONTIG;
return Bits;
@ -8923,6 +8934,20 @@ public:
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

View File

@ -5765,6 +5765,9 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
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,
@ -5800,6 +5803,9 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
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();

View File

@ -1651,7 +1651,7 @@ parseOpenMPSimpleClause(Parser &P, OpenMPClauseKind Kind) {
unsigned Type = getOpenMPSimpleClauseType(
Kind, Tok.isAnnotation() ? "" : P.getPreprocessor().getSpelling(Tok),
P.getLangOpts().OpenMP);
P.getLangOpts());
SourceLocation TypeLoc = Tok.getLocation();
if (Tok.isNot(tok::r_paren) && Tok.isNot(tok::comma) &&
Tok.isNot(tok::annot_pragma_openmp_end))
@ -3310,8 +3310,7 @@ OMPClause *Parser::ParseOpenMPSingleExprWithArgClause(OpenMPDirectiveKind DKind,
Arg[Modifier2] = OMPC_SCHEDULE_MODIFIER_unknown;
Arg[ScheduleKind] = OMPC_SCHEDULE_unknown;
unsigned KindModifier = getOpenMPSimpleClauseType(
Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok),
getLangOpts().OpenMP);
Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok), getLangOpts());
if (KindModifier > OMPC_SCHEDULE_unknown) {
// Parse 'modifier'
Arg[Modifier1] = KindModifier;
@ -3323,8 +3322,7 @@ OMPClause *Parser::ParseOpenMPSingleExprWithArgClause(OpenMPDirectiveKind DKind,
// Parse ',' 'modifier'
ConsumeAnyToken();
KindModifier = getOpenMPSimpleClauseType(
Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok),
getLangOpts().OpenMP);
Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok), getLangOpts());
Arg[Modifier2] = KindModifier > OMPC_SCHEDULE_unknown
? KindModifier
: (unsigned)OMPC_SCHEDULE_unknown;
@ -3339,8 +3337,7 @@ OMPClause *Parser::ParseOpenMPSingleExprWithArgClause(OpenMPDirectiveKind DKind,
else
Diag(Tok, diag::warn_pragma_expected_colon) << "schedule modifier";
KindModifier = getOpenMPSimpleClauseType(
Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok),
getLangOpts().OpenMP);
Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok), getLangOpts());
}
Arg[ScheduleKind] = KindModifier;
KLoc[ScheduleKind] = Tok.getLocation();
@ -3354,8 +3351,7 @@ OMPClause *Parser::ParseOpenMPSingleExprWithArgClause(OpenMPDirectiveKind DKind,
DelimLoc = ConsumeAnyToken();
} else if (Kind == OMPC_dist_schedule) {
Arg.push_back(getOpenMPSimpleClauseType(
Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok),
getLangOpts().OpenMP));
Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok), getLangOpts()));
KLoc.push_back(Tok.getLocation());
if (Tok.isNot(tok::r_paren) && Tok.isNot(tok::comma) &&
Tok.isNot(tok::annot_pragma_openmp_end))
@ -3365,8 +3361,7 @@ OMPClause *Parser::ParseOpenMPSingleExprWithArgClause(OpenMPDirectiveKind DKind,
} else if (Kind == OMPC_defaultmap) {
// Get a defaultmap modifier
unsigned Modifier = getOpenMPSimpleClauseType(
Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok),
getLangOpts().OpenMP);
Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok), getLangOpts());
// Set defaultmap modifier to unknown if it is either scalar, aggregate, or
// pointer
if (Modifier < OMPC_DEFAULTMAP_MODIFIER_unknown)
@ -3384,8 +3379,7 @@ OMPClause *Parser::ParseOpenMPSingleExprWithArgClause(OpenMPDirectiveKind DKind,
Diag(Tok, diag::warn_pragma_expected_colon) << "defaultmap modifier";
// Get a defaultmap kind
Arg.push_back(getOpenMPSimpleClauseType(
Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok),
getLangOpts().OpenMP));
Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok), getLangOpts()));
KLoc.push_back(Tok.getLocation());
if (Tok.isNot(tok::r_paren) && Tok.isNot(tok::comma) &&
Tok.isNot(tok::annot_pragma_openmp_end))
@ -3400,8 +3394,7 @@ OMPClause *Parser::ParseOpenMPSingleExprWithArgClause(OpenMPDirectiveKind DKind,
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 ':'
@ -3512,7 +3505,7 @@ static OpenMPMapModifierKind isMapModifier(Parser &P) {
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;
}
@ -3554,7 +3547,8 @@ bool Parser::parseMapTypeModifiers(OpenMPVarListDataTy &Data) {
OpenMPMapModifierKind TypeModifier = isMapModifier(*this);
if (TypeModifier == OMPC_MAP_MODIFIER_always ||
TypeModifier == OMPC_MAP_MODIFIER_close ||
TypeModifier == OMPC_MAP_MODIFIER_present) {
TypeModifier == OMPC_MAP_MODIFIER_present ||
TypeModifier == OMPC_MAP_MODIFIER_ompx_hold) {
Data.MapTypeModifiers.push_back(TypeModifier);
Data.MapTypeModifiersLoc.push_back(Tok.getLocation());
ConsumeToken();
@ -3577,7 +3571,8 @@ bool Parser::parseMapTypeModifiers(OpenMPVarListDataTy &Data) {
if (PP.LookAhead(0).is(tok::colon))
return false;
Diag(Tok, diag::err_omp_unknown_map_type_modifier)
<< (getLangOpts().OpenMP >= 51 ? 1 : 0);
<< (getLangOpts().OpenMP >= 51 ? 1 : 0)
<< getLangOpts().OpenMPExtensions;
ConsumeToken();
}
if (getCurToken().is(tok::comma))
@ -3596,7 +3591,7 @@ static OpenMPMapClauseKind isMapType(Parser &P) {
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;
}
@ -3749,8 +3744,8 @@ bool Parser::ParseOpenMPVarList(OpenMPDirectiveKind DKind,
(Tok.is(tok::identifier) || Tok.is(tok::kw_default)) &&
NextToken().is(tok::comma)) {
// Parse optional reduction modifier.
Data.ExtraModifier = getOpenMPSimpleClauseType(Kind, PP.getSpelling(Tok),
getLangOpts().OpenMP);
Data.ExtraModifier =
getOpenMPSimpleClauseType(Kind, PP.getSpelling(Tok), getLangOpts());
Data.ExtraModifierLoc = Tok.getLocation();
ConsumeToken();
assert(Tok.is(tok::comma) && "Expected comma.");
@ -3796,7 +3791,7 @@ bool Parser::ParseOpenMPVarList(OpenMPDirectiveKind DKind,
ColonProtectionRAIIObject ColonRAII(*this);
Data.ExtraModifier = getOpenMPSimpleClauseType(
Kind, Tok.is(tok::identifier) ? PP.getSpelling(Tok) : "",
getLangOpts().OpenMP);
getLangOpts());
Data.ExtraModifierLoc = Tok.getLocation();
if (Data.ExtraModifier == OMPC_DEPEND_unknown) {
SkipUntil(tok::colon, tok::r_paren, tok::annot_pragma_openmp_end,
@ -3821,8 +3816,8 @@ bool Parser::ParseOpenMPVarList(OpenMPDirectiveKind DKind,
// Try to parse modifier if any.
Data.ExtraModifier = OMPC_LINEAR_val;
if (Tok.is(tok::identifier) && PP.LookAhead(0).is(tok::l_paren)) {
Data.ExtraModifier = getOpenMPSimpleClauseType(Kind, PP.getSpelling(Tok),
getLangOpts().OpenMP);
Data.ExtraModifier =
getOpenMPSimpleClauseType(Kind, PP.getSpelling(Tok), getLangOpts());
Data.ExtraModifierLoc = ConsumeToken();
LinearT.consumeOpen();
NeedRParenForLinear = true;
@ -3835,8 +3830,8 @@ bool Parser::ParseOpenMPVarList(OpenMPDirectiveKind DKind,
if ((getLangOpts().OpenMP >= 50 && !isOpenMPDistributeDirective(DKind) &&
!isOpenMPTaskLoopDirective(DKind)) &&
Tok.is(tok::identifier) && PP.LookAhead(0).is(tok::colon)) {
Data.ExtraModifier = getOpenMPSimpleClauseType(Kind, PP.getSpelling(Tok),
getLangOpts().OpenMP);
Data.ExtraModifier =
getOpenMPSimpleClauseType(Kind, PP.getSpelling(Tok), getLangOpts());
Data.ExtraModifierLoc = Tok.getLocation();
ConsumeToken();
assert(Tok.is(tok::colon) && "Expected colon.");
@ -3879,9 +3874,8 @@ bool Parser::ParseOpenMPVarList(OpenMPDirectiveKind DKind,
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);

View File

@ -19348,6 +19348,7 @@ static void checkMappableExpressionList(
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) &&
@ -19369,6 +19370,10 @@ static void checkMappableExpressionList(
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.
@ -19569,6 +19574,21 @@ static void checkMappableExpressionList(
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]
@ -19644,7 +19664,8 @@ OMPClause *Sema::ActOnOpenMPMapClause(
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.
@ -19665,7 +19686,8 @@ OMPClause *Sema::ActOnOpenMPMapClause(
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.

View File

@ -0,0 +1,26 @@
// 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)
;
}

View File

@ -1125,4 +1125,76 @@ int main (int argc, char **argv) {
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

View File

@ -6,13 +6,13 @@
// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
// RUN: %clang_cc1 -DOMP51 -verify -fopenmp -fopenmp-version=51 -ast-print %s | FileCheck -check-prefixes=CHECK,OMP51 %s
// RUN: %clang_cc1 -DOMP51 -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -emit-pch -o %t %s
// RUN: %clang_cc1 -DOMP51 -fopenmp -fopenmp-version=51 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck -check-prefixes=CHECK,OMP51 %s
// RUN: %clang_cc1 -DOMP51 -DOMPX -verify -fopenmp -fopenmp-version=51 -fopenmp-extensions -ast-print %s | FileCheck -check-prefixes=CHECK,OMP51,OMPX %s
// RUN: %clang_cc1 -DOMP51 -DOMPX -fopenmp -fopenmp-version=51 -fopenmp-extensions -x c++ -std=c++11 -emit-pch -o %t %s
// RUN: %clang_cc1 -DOMP51 -DOMPX -fopenmp -fopenmp-version=51 -fopenmp-extensions -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck -check-prefixes=CHECK,OMP51,OMPX %s
// RUN: %clang_cc1 -DOMP51 -verify -fopenmp-simd -fopenmp-version=51 -ast-print %s | FileCheck -check-prefixes=CHECK,OMP51 %s
// RUN: %clang_cc1 -DOMP51 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -emit-pch -o %t %s
// RUN: %clang_cc1 -DOMP51 -fopenmp-simd -fopenmp-version=51 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck -check-prefixes=CHECK,OMP51 %s
// RUN: %clang_cc1 -DOMP51 -DOMPX -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-extensions -ast-print %s | FileCheck -check-prefixes=CHECK,OMP51,OMPX %s
// RUN: %clang_cc1 -DOMP51 -DOMPX -fopenmp-simd -fopenmp-version=51 -fopenmp-extensions -x c++ -std=c++11 -emit-pch -o %t %s
// RUN: %clang_cc1 -DOMP51 -DOMPX -fopenmp-simd -fopenmp-version=51 -fopenmp-extensions -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck -check-prefixes=CHECK,OMP51,OMPX %s
// expected-no-diagnostics
#ifndef HEADER
@ -56,6 +56,11 @@ T tmain(T argc, T *argv) {
foo();
#endif
#ifdef OMPX
#pragma omp target data map(ompx_hold,alloc: e)
foo();
#endif
// nesting a target region
#pragma omp target data map(e)
{
@ -67,6 +72,10 @@ T tmain(T argc, T *argv) {
#pragma omp target map(present, alloc: e)
foo();
#endif
#ifdef OMPX
#pragma omp target map(ompx_hold, alloc: e)
foo();
#endif
}
return 0;
@ -94,6 +103,8 @@ T tmain(T argc, T *argv) {
// CHECK-NEXT: foo();
// OMP51-NEXT: #pragma omp target data map(present,alloc: e)
// OMP51-NEXT: foo();
// OMPX-NEXT: #pragma omp target data map(ompx_hold,alloc: e)
// OMPX-NEXT: foo();
// CHECK-NEXT: #pragma omp target data map(tofrom: e)
// CHECK-NEXT: {
// CHECK-NEXT: #pragma omp target map(always,alloc: e)
@ -102,6 +113,8 @@ T tmain(T argc, T *argv) {
// 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)
@ -124,6 +137,8 @@ T tmain(T argc, T *argv) {
// CHECK-NEXT: foo();
// OMP51-NEXT: #pragma omp target data map(present,alloc: e)
// OMP51-NEXT: foo();
// OMPX-NEXT: #pragma omp target data map(ompx_hold,alloc: e)
// OMPX-NEXT: foo();
// CHECK-NEXT: #pragma omp target data map(tofrom: e)
// CHECK-NEXT: {
// CHECK-NEXT: #pragma omp target map(always,alloc: e)
@ -132,6 +147,8 @@ T tmain(T argc, T *argv) {
// 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)
@ -154,6 +171,8 @@ T tmain(T argc, T *argv) {
// CHECK-NEXT: foo();
// OMP51-NEXT: #pragma omp target data map(present,alloc: e)
// OMP51-NEXT: foo();
// OMPX-NEXT: #pragma omp target data map(ompx_hold,alloc: e)
// OMPX-NEXT: foo();
// CHECK-NEXT: #pragma omp target data map(tofrom: e)
// CHECK-NEXT: {
// CHECK-NEXT: #pragma omp target map(always,alloc: e)
@ -162,6 +181,8 @@ T tmain(T argc, T *argv) {
// CHECK-NEXT: foo();
// OMP51-NEXT: #pragma omp target map(present,alloc: e)
// OMP51-NEXT: foo();
// OMPX-NEXT: #pragma omp target map(ompx_hold,alloc: e)
// OMPX-NEXT: foo();
int main (int argc, char **argv) {
int b = argc, c, d, e, f, g, x[20];
@ -221,6 +242,13 @@ int main (int argc, char **argv) {
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)

View File

@ -0,0 +1,608 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals --global-value-regex ".offload_maptypes.*" ".offload_sizes.*" --global-hex-value-regex ".offload_maptypes.*"
// expected-no-diagnostics
#ifndef HEADER
#define HEADER
// powerpc64le-ibm-linux-gnu
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-extensions \
// RUN: -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ \
// RUN: -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | \
// RUN: FileCheck %s --check-prefixes=CHECK-PPC64LE
// RUN: %clang_cc1 -fopenmp -fopenmp-extensions \
// RUN: -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 \
// RUN: -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp -fopenmp-extensions \
// RUN: -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ \
// RUN: -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t \
// RUN: -verify %s -emit-llvm -o - | \
// RUN: FileCheck %s --check-prefixes=CHECK-PPC64LE
// i386-pc-linux-gnu
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-extensions \
// RUN: -fopenmp-targets=i386-pc-linux-gnu -x c++ \
// RUN: -triple i386-unknown-unknown -emit-llvm %s -o - | \
// RUN: FileCheck %s --check-prefixes=CHECK-I386
// RUN: %clang_cc1 -fopenmp -fopenmp-extensions \
// RUN: -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 \
// RUN: -triple i386-unknown-unknown -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp -fopenmp-extensions \
// RUN: -fopenmp-targets=i386-pc-linux-gnu -x c++ \
// RUN: -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s \
// RUN: -emit-llvm -o - | \
// RUN: FileCheck %s --check-prefixes=CHECK-I386
struct S1 {
int i;
};
struct S2 {
S1 s;
struct S2 *ps;
};
// Map flags used in @.offload_maptypes* below:
//
// TO = 0x1
// FROM = 0x2
// ALWAYS = 0x4
// PTR_AND_OBJ = 0x10
// CLOSE = 0x400
// OMPX_HOLD = 0x2000
// MEMBER_OF_1 = 0x1000000000000
// MEMBER_OF_7 = 0x7000000000000
//.
// CHECK-PPC64LE: @.offload_sizes = private unnamed_addr constant [1 x i64] [i64 20]
// CHECK-PPC64LE: @.offload_maptypes = private unnamed_addr constant [1 x i64] [i64 [[#0x2001]]]
// CHECK-PPC64LE: @.offload_sizes.1 = private unnamed_addr constant [1 x i64] [i64 20]
// CHECK-PPC64LE: @.offload_maptypes.2 = private unnamed_addr constant [1 x i64] [i64 [[#0x2405]]]
// CHECK-PPC64LE: @.offload_sizes.3 = private unnamed_addr constant [1 x i64] [i64 4]
// CHECK-PPC64LE: @.offload_maptypes.4 = private unnamed_addr constant [1 x i64] [i64 [[#0x2003]]]
// CHECK-PPC64LE: @.offload_maptypes.5 = private unnamed_addr constant [11 x i64] [i64 [[#0x2000]], i64 [[#0x1000000002003]], i64 [[#0x1000000002010]], i64 [[#0x2010]], i64 [[#0x2013]], i64 [[#0x3]], i64 [[#0x2000]], i64 [[#0x7000000002003]], i64 [[#0x7000000002010]], i64 [[#0x2010]], i64 [[#0x2013]]]
//.
// CHECK-I386: @.offload_sizes = private unnamed_addr constant [1 x i64] [i64 20]
// CHECK-I386: @.offload_maptypes = private unnamed_addr constant [1 x i64] [i64 [[#0x2001]]]
// CHECK-I386: @.offload_sizes.1 = private unnamed_addr constant [1 x i64] [i64 20]
// CHECK-I386: @.offload_maptypes.2 = private unnamed_addr constant [1 x i64] [i64 [[#0x2405]]]
// CHECK-I386: @.offload_sizes.3 = private unnamed_addr constant [1 x i64] [i64 4]
// CHECK-I386: @.offload_maptypes.4 = private unnamed_addr constant [1 x i64] [i64 [[#0x2003]]]
// CHECK-I386: @.offload_maptypes.5 = private unnamed_addr constant [11 x i64] [i64 [[#0x2000]], i64 [[#0x1000000002003]], i64 [[#0x1000000002010]], i64 [[#0x2010]], i64 [[#0x2013]], i64 [[#0x3]], i64 [[#0x2000]], i64 [[#0x7000000002003]], i64 [[#0x7000000002010]], i64 [[#0x2010]], i64 [[#0x2013]]]
//.
// CHECK-PPC64LE-LABEL: @_Z3fooi(
// CHECK-PPC64LE-NEXT: entry:
// CHECK-PPC64LE-NEXT: [[ARG_ADDR:%.*]] = alloca i32, align 4
// CHECK-PPC64LE-NEXT: [[LB:%.*]] = alloca [5 x float], align 4
// CHECK-PPC64LE-NEXT: [[PS1:%.*]] = alloca %struct.S2*, align 8
// CHECK-PPC64LE-NEXT: [[PS2:%.*]] = alloca %struct.S2*, align 8
// CHECK-PPC64LE-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x i8*], align 8
// CHECK-PPC64LE-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x i8*], align 8
// CHECK-PPC64LE-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x i8*], align 8
// CHECK-PPC64LE-NEXT: [[DOTOFFLOAD_BASEPTRS1:%.*]] = alloca [1 x i8*], align 8
// CHECK-PPC64LE-NEXT: [[DOTOFFLOAD_PTRS2:%.*]] = alloca [1 x i8*], align 8
// CHECK-PPC64LE-NEXT: [[DOTOFFLOAD_MAPPERS3:%.*]] = alloca [1 x i8*], align 8
// CHECK-PPC64LE-NEXT: [[DOTOFFLOAD_BASEPTRS5:%.*]] = alloca [1 x i8*], align 8
// CHECK-PPC64LE-NEXT: [[DOTOFFLOAD_PTRS6:%.*]] = alloca [1 x i8*], align 8
// CHECK-PPC64LE-NEXT: [[DOTOFFLOAD_MAPPERS7:%.*]] = alloca [1 x i8*], align 8
// CHECK-PPC64LE-NEXT: [[DOTOFFLOAD_BASEPTRS29:%.*]] = alloca [11 x i8*], align 8
// CHECK-PPC64LE-NEXT: [[DOTOFFLOAD_PTRS30:%.*]] = alloca [11 x i8*], align 8
// CHECK-PPC64LE-NEXT: [[DOTOFFLOAD_MAPPERS31:%.*]] = alloca [11 x i8*], align 8
// CHECK-PPC64LE-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [11 x i64], align 8
// CHECK-PPC64LE-NEXT: store i32 [[ARG:%.*]], i32* [[ARG_ADDR]], align 4
// CHECK-PPC64LE-NEXT: [[TMP0:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK-PPC64LE-NEXT: [[TMP1:%.*]] = bitcast i8** [[TMP0]] to [5 x float]**
// CHECK-PPC64LE-NEXT: store [5 x float]* [[LB]], [5 x float]** [[TMP1]], align 8
// CHECK-PPC64LE-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK-PPC64LE-NEXT: [[TMP3:%.*]] = bitcast i8** [[TMP2]] to [5 x float]**
// CHECK-PPC64LE-NEXT: store [5 x float]* [[LB]], [5 x float]** [[TMP3]], align 8
// CHECK-PPC64LE-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0
// CHECK-PPC64LE-NEXT: store i8* null, i8** [[TMP4]], align 8
// CHECK-PPC64LE-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK-PPC64LE-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK-PPC64LE-NEXT: call void @__tgt_target_data_begin_mapper(%struct.ident_t* @[[GLOB1:[0-9]+]], i64 -1, i32 1, i8** [[TMP5]], i8** [[TMP6]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null)
// CHECK-PPC64LE-NEXT: [[TMP7:%.*]] = load i32, i32* [[ARG_ADDR]], align 4
// CHECK-PPC64LE-NEXT: [[INC:%.*]] = add nsw i32 [[TMP7]], 1
// CHECK-PPC64LE-NEXT: store i32 [[INC]], i32* [[ARG_ADDR]], align 4
// CHECK-PPC64LE-NEXT: [[TMP8:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK-PPC64LE-NEXT: [[TMP9:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK-PPC64LE-NEXT: call void @__tgt_target_data_end_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i32 1, i8** [[TMP8]], i8** [[TMP9]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null)
// CHECK-PPC64LE-NEXT: [[TMP10:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS1]], i32 0, i32 0
// CHECK-PPC64LE-NEXT: [[TMP11:%.*]] = bitcast i8** [[TMP10]] to [5 x float]**
// CHECK-PPC64LE-NEXT: store [5 x float]* [[LB]], [5 x float]** [[TMP11]], align 8
// CHECK-PPC64LE-NEXT: [[TMP12:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS2]], i32 0, i32 0
// CHECK-PPC64LE-NEXT: [[TMP13:%.*]] = bitcast i8** [[TMP12]] to [5 x float]**
// CHECK-PPC64LE-NEXT: store [5 x float]* [[LB]], [5 x float]** [[TMP13]], align 8
// CHECK-PPC64LE-NEXT: [[TMP14:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_MAPPERS3]], i64 0, i64 0
// CHECK-PPC64LE-NEXT: store i8* null, i8** [[TMP14]], align 8
// CHECK-PPC64LE-NEXT: [[TMP15:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS1]], i32 0, i32 0
// CHECK-PPC64LE-NEXT: [[TMP16:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS2]], i32 0, i32 0
// CHECK-PPC64LE-NEXT: call void @__tgt_target_data_begin_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i32 1, i8** [[TMP15]], i8** [[TMP16]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes.1, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.2, i32 0, i32 0), i8** null, i8** null)
// CHECK-PPC64LE-NEXT: [[TMP17:%.*]] = load i32, i32* [[ARG_ADDR]], align 4
// CHECK-PPC64LE-NEXT: [[INC4:%.*]] = add nsw i32 [[TMP17]], 1
// CHECK-PPC64LE-NEXT: store i32 [[INC4]], i32* [[ARG_ADDR]], align 4
// CHECK-PPC64LE-NEXT: [[TMP18:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS1]], i32 0, i32 0
// CHECK-PPC64LE-NEXT: [[TMP19:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS2]], i32 0, i32 0
// CHECK-PPC64LE-NEXT: call void @__tgt_target_data_end_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i32 1, i8** [[TMP18]], i8** [[TMP19]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes.1, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.2, i32 0, i32 0), i8** null, i8** null)
// CHECK-PPC64LE-NEXT: [[TMP20:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS5]], i32 0, i32 0
// CHECK-PPC64LE-NEXT: [[TMP21:%.*]] = bitcast i8** [[TMP20]] to i32**
// CHECK-PPC64LE-NEXT: store i32* [[ARG_ADDR]], i32** [[TMP21]], align 8
// CHECK-PPC64LE-NEXT: [[TMP22:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS6]], i32 0, i32 0
// CHECK-PPC64LE-NEXT: [[TMP23:%.*]] = bitcast i8** [[TMP22]] to i32**
// CHECK-PPC64LE-NEXT: store i32* [[ARG_ADDR]], i32** [[TMP23]], align 8
// CHECK-PPC64LE-NEXT: [[TMP24:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_MAPPERS7]], i64 0, i64 0
// CHECK-PPC64LE-NEXT: store i8* null, i8** [[TMP24]], align 8
// CHECK-PPC64LE-NEXT: [[TMP25:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS5]], i32 0, i32 0
// CHECK-PPC64LE-NEXT: [[TMP26:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS6]], i32 0, i32 0
// CHECK-PPC64LE-NEXT: call void @__tgt_target_data_begin_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i32 1, i8** [[TMP25]], i8** [[TMP26]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes.3, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.4, i32 0, i32 0), i8** null, i8** null)
// CHECK-PPC64LE-NEXT: [[TMP27:%.*]] = load i32, i32* [[ARG_ADDR]], align 4
// CHECK-PPC64LE-NEXT: [[INC8:%.*]] = add nsw i32 [[TMP27]], 1
// CHECK-PPC64LE-NEXT: store i32 [[INC8]], i32* [[ARG_ADDR]], align 4
// CHECK-PPC64LE-NEXT: [[TMP28:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS5]], i32 0, i32 0
// CHECK-PPC64LE-NEXT: [[TMP29:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS6]], i32 0, i32 0
// CHECK-PPC64LE-NEXT: call void @__tgt_target_data_end_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i32 1, i8** [[TMP28]], i8** [[TMP29]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes.3, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.4, i32 0, i32 0), i8** null, i8** null)
// CHECK-PPC64LE-NEXT: [[TMP30:%.*]] = load %struct.S2*, %struct.S2** [[PS1]], align 8
// CHECK-PPC64LE-NEXT: [[TMP31:%.*]] = load %struct.S2*, %struct.S2** [[PS1]], align 8
// CHECK-PPC64LE-NEXT: [[S:%.*]] = getelementptr inbounds [[STRUCT_S2:%.*]], %struct.S2* [[TMP31]], i32 0, i32 0
// CHECK-PPC64LE-NEXT: [[TMP32:%.*]] = load %struct.S2*, %struct.S2** [[PS1]], align 8
// CHECK-PPC64LE-NEXT: [[TMP33:%.*]] = load %struct.S2*, %struct.S2** [[PS1]], align 8
// CHECK-PPC64LE-NEXT: [[PS:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP33]], i32 0, i32 1
// CHECK-PPC64LE-NEXT: [[TMP34:%.*]] = load %struct.S2*, %struct.S2** [[PS1]], align 8
// CHECK-PPC64LE-NEXT: [[PS9:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP34]], i32 0, i32 1
// CHECK-PPC64LE-NEXT: [[TMP35:%.*]] = load %struct.S2*, %struct.S2** [[PS9]], align 8
// CHECK-PPC64LE-NEXT: [[PS10:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP35]], i32 0, i32 1
// CHECK-PPC64LE-NEXT: [[TMP36:%.*]] = load %struct.S2*, %struct.S2** [[PS1]], align 8
// CHECK-PPC64LE-NEXT: [[PS11:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP36]], i32 0, i32 1
// CHECK-PPC64LE-NEXT: [[TMP37:%.*]] = load %struct.S2*, %struct.S2** [[PS11]], align 8
// CHECK-PPC64LE-NEXT: [[PS12:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP37]], i32 0, i32 1
// CHECK-PPC64LE-NEXT: [[TMP38:%.*]] = load %struct.S2*, %struct.S2** [[PS12]], align 8
// CHECK-PPC64LE-NEXT: [[PS13:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP38]], i32 0, i32 1
// CHECK-PPC64LE-NEXT: [[TMP39:%.*]] = load %struct.S2*, %struct.S2** [[PS1]], align 8
// CHECK-PPC64LE-NEXT: [[PS14:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP39]], i32 0, i32 1
// CHECK-PPC64LE-NEXT: [[TMP40:%.*]] = load %struct.S2*, %struct.S2** [[PS14]], align 8
// CHECK-PPC64LE-NEXT: [[PS15:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP40]], i32 0, i32 1
// CHECK-PPC64LE-NEXT: [[TMP41:%.*]] = load %struct.S2*, %struct.S2** [[PS15]], align 8
// CHECK-PPC64LE-NEXT: [[PS16:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP41]], i32 0, i32 1
// CHECK-PPC64LE-NEXT: [[TMP42:%.*]] = load %struct.S2*, %struct.S2** [[PS16]], align 8
// CHECK-PPC64LE-NEXT: [[S17:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP42]], i32 0, i32 0
// CHECK-PPC64LE-NEXT: [[TMP43:%.*]] = getelementptr %struct.S2*, %struct.S2** [[PS]], i32 1
// CHECK-PPC64LE-NEXT: [[TMP44:%.*]] = bitcast %struct.S1* [[S]] to i8*
// CHECK-PPC64LE-NEXT: [[TMP45:%.*]] = bitcast %struct.S2** [[TMP43]] to i8*
// CHECK-PPC64LE-NEXT: [[TMP46:%.*]] = ptrtoint i8* [[TMP45]] to i64
// CHECK-PPC64LE-NEXT: [[TMP47:%.*]] = ptrtoint i8* [[TMP44]] to i64
// CHECK-PPC64LE-NEXT: [[TMP48:%.*]] = sub i64 [[TMP46]], [[TMP47]]
// CHECK-PPC64LE-NEXT: [[TMP49:%.*]] = sdiv exact i64 [[TMP48]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
// CHECK-PPC64LE-NEXT: [[TMP50:%.*]] = load %struct.S2*, %struct.S2** [[PS2]], align 8
// CHECK-PPC64LE-NEXT: [[TMP51:%.*]] = load %struct.S2*, %struct.S2** [[PS2]], align 8
// CHECK-PPC64LE-NEXT: [[S18:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP51]], i32 0, i32 0
// CHECK-PPC64LE-NEXT: [[TMP52:%.*]] = load %struct.S2*, %struct.S2** [[PS2]], align 8
// CHECK-PPC64LE-NEXT: [[TMP53:%.*]] = load %struct.S2*, %struct.S2** [[PS2]], align 8
// CHECK-PPC64LE-NEXT: [[PS19:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP53]], i32 0, i32 1
// CHECK-PPC64LE-NEXT: [[TMP54:%.*]] = load %struct.S2*, %struct.S2** [[PS2]], align 8
// CHECK-PPC64LE-NEXT: [[PS20:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP54]], i32 0, i32 1
// CHECK-PPC64LE-NEXT: [[TMP55:%.*]] = load %struct.S2*, %struct.S2** [[PS20]], align 8
// CHECK-PPC64LE-NEXT: [[PS21:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP55]], i32 0, i32 1
// CHECK-PPC64LE-NEXT: [[TMP56:%.*]] = load %struct.S2*, %struct.S2** [[PS2]], align 8
// CHECK-PPC64LE-NEXT: [[PS22:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP56]], i32 0, i32 1
// CHECK-PPC64LE-NEXT: [[TMP57:%.*]] = load %struct.S2*, %struct.S2** [[PS22]], align 8
// CHECK-PPC64LE-NEXT: [[PS23:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP57]], i32 0, i32 1
// CHECK-PPC64LE-NEXT: [[TMP58:%.*]] = load %struct.S2*, %struct.S2** [[PS23]], align 8
// CHECK-PPC64LE-NEXT: [[PS24:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP58]], i32 0, i32 1
// CHECK-PPC64LE-NEXT: [[TMP59:%.*]] = load %struct.S2*, %struct.S2** [[PS2]], align 8
// CHECK-PPC64LE-NEXT: [[PS25:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP59]], i32 0, i32 1
// CHECK-PPC64LE-NEXT: [[TMP60:%.*]] = load %struct.S2*, %struct.S2** [[PS25]], align 8
// CHECK-PPC64LE-NEXT: [[PS26:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP60]], i32 0, i32 1
// CHECK-PPC64LE-NEXT: [[TMP61:%.*]] = load %struct.S2*, %struct.S2** [[PS26]], align 8
// CHECK-PPC64LE-NEXT: [[PS27:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP61]], i32 0, i32 1
// CHECK-PPC64LE-NEXT: [[TMP62:%.*]] = load %struct.S2*, %struct.S2** [[PS27]], align 8
// CHECK-PPC64LE-NEXT: [[S28:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP62]], i32 0, i32 0
// CHECK-PPC64LE-NEXT: [[TMP63:%.*]] = getelementptr %struct.S2*, %struct.S2** [[PS19]], i32 1
// CHECK-PPC64LE-NEXT: [[TMP64:%.*]] = bitcast %struct.S1* [[S18]] to i8*
// CHECK-PPC64LE-NEXT: [[TMP65:%.*]] = bitcast %struct.S2** [[TMP63]] to i8*
// CHECK-PPC64LE-NEXT: [[TMP66:%.*]] = ptrtoint i8* [[TMP65]] to i64
// CHECK-PPC64LE-NEXT: [[TMP67:%.*]] = ptrtoint i8* [[TMP64]] to i64
// CHECK-PPC64LE-NEXT: [[TMP68:%.*]] = sub i64 [[TMP66]], [[TMP67]]
// CHECK-PPC64LE-NEXT: [[TMP69:%.*]] = sdiv exact i64 [[TMP68]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
// CHECK-PPC64LE-NEXT: [[TMP70:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 0
// CHECK-PPC64LE-NEXT: [[TMP71:%.*]] = bitcast i8** [[TMP70]] to %struct.S2**
// CHECK-PPC64LE-NEXT: store %struct.S2* [[TMP30]], %struct.S2** [[TMP71]], align 8
// CHECK-PPC64LE-NEXT: [[TMP72:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 0
// CHECK-PPC64LE-NEXT: [[TMP73:%.*]] = bitcast i8** [[TMP72]] to %struct.S1**
// CHECK-PPC64LE-NEXT: store %struct.S1* [[S]], %struct.S1** [[TMP73]], align 8
// CHECK-PPC64LE-NEXT: [[TMP74:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0
// CHECK-PPC64LE-NEXT: store i64 [[TMP49]], i64* [[TMP74]], align 8
// CHECK-PPC64LE-NEXT: [[TMP75:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i64 0, i64 0
// CHECK-PPC64LE-NEXT: store i8* null, i8** [[TMP75]], align 8
// CHECK-PPC64LE-NEXT: [[TMP76:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 1
// CHECK-PPC64LE-NEXT: [[TMP77:%.*]] = bitcast i8** [[TMP76]] to %struct.S2**
// CHECK-PPC64LE-NEXT: store %struct.S2* [[TMP30]], %struct.S2** [[TMP77]], align 8
// CHECK-PPC64LE-NEXT: [[TMP78:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 1
// CHECK-PPC64LE-NEXT: [[TMP79:%.*]] = bitcast i8** [[TMP78]] to %struct.S1**
// CHECK-PPC64LE-NEXT: store %struct.S1* [[S]], %struct.S1** [[TMP79]], align 8
// CHECK-PPC64LE-NEXT: [[TMP80:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 1
// CHECK-PPC64LE-NEXT: store i64 4, i64* [[TMP80]], align 8
// CHECK-PPC64LE-NEXT: [[TMP81:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i64 0, i64 1
// CHECK-PPC64LE-NEXT: store i8* null, i8** [[TMP81]], align 8
// CHECK-PPC64LE-NEXT: [[TMP82:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 2
// CHECK-PPC64LE-NEXT: [[TMP83:%.*]] = bitcast i8** [[TMP82]] to %struct.S2***
// CHECK-PPC64LE-NEXT: store %struct.S2** [[PS]], %struct.S2*** [[TMP83]], align 8
// CHECK-PPC64LE-NEXT: [[TMP84:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 2
// CHECK-PPC64LE-NEXT: [[TMP85:%.*]] = bitcast i8** [[TMP84]] to %struct.S2***
// CHECK-PPC64LE-NEXT: store %struct.S2** [[PS10]], %struct.S2*** [[TMP85]], align 8
// CHECK-PPC64LE-NEXT: [[TMP86:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 2
// CHECK-PPC64LE-NEXT: store i64 8, i64* [[TMP86]], align 8
// CHECK-PPC64LE-NEXT: [[TMP87:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i64 0, i64 2
// CHECK-PPC64LE-NEXT: store i8* null, i8** [[TMP87]], align 8
// CHECK-PPC64LE-NEXT: [[TMP88:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 3
// CHECK-PPC64LE-NEXT: [[TMP89:%.*]] = bitcast i8** [[TMP88]] to %struct.S2***
// CHECK-PPC64LE-NEXT: store %struct.S2** [[PS10]], %struct.S2*** [[TMP89]], align 8
// CHECK-PPC64LE-NEXT: [[TMP90:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 3
// CHECK-PPC64LE-NEXT: [[TMP91:%.*]] = bitcast i8** [[TMP90]] to %struct.S2***
// CHECK-PPC64LE-NEXT: store %struct.S2** [[PS13]], %struct.S2*** [[TMP91]], align 8
// CHECK-PPC64LE-NEXT: [[TMP92:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 3
// CHECK-PPC64LE-NEXT: store i64 8, i64* [[TMP92]], align 8
// CHECK-PPC64LE-NEXT: [[TMP93:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i64 0, i64 3
// CHECK-PPC64LE-NEXT: store i8* null, i8** [[TMP93]], align 8
// CHECK-PPC64LE-NEXT: [[TMP94:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 4
// CHECK-PPC64LE-NEXT: [[TMP95:%.*]] = bitcast i8** [[TMP94]] to %struct.S2***
// CHECK-PPC64LE-NEXT: store %struct.S2** [[PS13]], %struct.S2*** [[TMP95]], align 8
// CHECK-PPC64LE-NEXT: [[TMP96:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 4
// CHECK-PPC64LE-NEXT: [[TMP97:%.*]] = bitcast i8** [[TMP96]] to %struct.S1**
// CHECK-PPC64LE-NEXT: store %struct.S1* [[S17]], %struct.S1** [[TMP97]], align 8
// CHECK-PPC64LE-NEXT: [[TMP98:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 4
// CHECK-PPC64LE-NEXT: store i64 4, i64* [[TMP98]], align 8
// CHECK-PPC64LE-NEXT: [[TMP99:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i64 0, i64 4
// CHECK-PPC64LE-NEXT: store i8* null, i8** [[TMP99]], align 8
// CHECK-PPC64LE-NEXT: [[TMP100:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 5
// CHECK-PPC64LE-NEXT: [[TMP101:%.*]] = bitcast i8** [[TMP100]] to i32**
// CHECK-PPC64LE-NEXT: store i32* [[ARG_ADDR]], i32** [[TMP101]], align 8
// CHECK-PPC64LE-NEXT: [[TMP102:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 5
// CHECK-PPC64LE-NEXT: [[TMP103:%.*]] = bitcast i8** [[TMP102]] to i32**
// CHECK-PPC64LE-NEXT: store i32* [[ARG_ADDR]], i32** [[TMP103]], align 8
// CHECK-PPC64LE-NEXT: [[TMP104:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 5
// CHECK-PPC64LE-NEXT: store i64 4, i64* [[TMP104]], align 8
// CHECK-PPC64LE-NEXT: [[TMP105:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i64 0, i64 5
// CHECK-PPC64LE-NEXT: store i8* null, i8** [[TMP105]], align 8
// CHECK-PPC64LE-NEXT: [[TMP106:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 6
// CHECK-PPC64LE-NEXT: [[TMP107:%.*]] = bitcast i8** [[TMP106]] to %struct.S2**
// CHECK-PPC64LE-NEXT: store %struct.S2* [[TMP50]], %struct.S2** [[TMP107]], align 8
// CHECK-PPC64LE-NEXT: [[TMP108:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 6
// CHECK-PPC64LE-NEXT: [[TMP109:%.*]] = bitcast i8** [[TMP108]] to %struct.S1**
// CHECK-PPC64LE-NEXT: store %struct.S1* [[S18]], %struct.S1** [[TMP109]], align 8
// CHECK-PPC64LE-NEXT: [[TMP110:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 6
// CHECK-PPC64LE-NEXT: store i64 [[TMP69]], i64* [[TMP110]], align 8
// CHECK-PPC64LE-NEXT: [[TMP111:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i64 0, i64 6
// CHECK-PPC64LE-NEXT: store i8* null, i8** [[TMP111]], align 8
// CHECK-PPC64LE-NEXT: [[TMP112:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 7
// CHECK-PPC64LE-NEXT: [[TMP113:%.*]] = bitcast i8** [[TMP112]] to %struct.S2**
// CHECK-PPC64LE-NEXT: store %struct.S2* [[TMP50]], %struct.S2** [[TMP113]], align 8
// CHECK-PPC64LE-NEXT: [[TMP114:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 7
// CHECK-PPC64LE-NEXT: [[TMP115:%.*]] = bitcast i8** [[TMP114]] to %struct.S1**
// CHECK-PPC64LE-NEXT: store %struct.S1* [[S18]], %struct.S1** [[TMP115]], align 8
// CHECK-PPC64LE-NEXT: [[TMP116:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 7
// CHECK-PPC64LE-NEXT: store i64 4, i64* [[TMP116]], align 8
// CHECK-PPC64LE-NEXT: [[TMP117:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i64 0, i64 7
// CHECK-PPC64LE-NEXT: store i8* null, i8** [[TMP117]], align 8
// CHECK-PPC64LE-NEXT: [[TMP118:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 8
// CHECK-PPC64LE-NEXT: [[TMP119:%.*]] = bitcast i8** [[TMP118]] to %struct.S2***
// CHECK-PPC64LE-NEXT: store %struct.S2** [[PS19]], %struct.S2*** [[TMP119]], align 8
// CHECK-PPC64LE-NEXT: [[TMP120:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 8
// CHECK-PPC64LE-NEXT: [[TMP121:%.*]] = bitcast i8** [[TMP120]] to %struct.S2***
// CHECK-PPC64LE-NEXT: store %struct.S2** [[PS21]], %struct.S2*** [[TMP121]], align 8
// CHECK-PPC64LE-NEXT: [[TMP122:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 8
// CHECK-PPC64LE-NEXT: store i64 8, i64* [[TMP122]], align 8
// CHECK-PPC64LE-NEXT: [[TMP123:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i64 0, i64 8
// CHECK-PPC64LE-NEXT: store i8* null, i8** [[TMP123]], align 8
// CHECK-PPC64LE-NEXT: [[TMP124:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 9
// CHECK-PPC64LE-NEXT: [[TMP125:%.*]] = bitcast i8** [[TMP124]] to %struct.S2***
// CHECK-PPC64LE-NEXT: store %struct.S2** [[PS21]], %struct.S2*** [[TMP125]], align 8
// CHECK-PPC64LE-NEXT: [[TMP126:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 9
// CHECK-PPC64LE-NEXT: [[TMP127:%.*]] = bitcast i8** [[TMP126]] to %struct.S2***
// CHECK-PPC64LE-NEXT: store %struct.S2** [[PS24]], %struct.S2*** [[TMP127]], align 8
// CHECK-PPC64LE-NEXT: [[TMP128:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 9
// CHECK-PPC64LE-NEXT: store i64 8, i64* [[TMP128]], align 8
// CHECK-PPC64LE-NEXT: [[TMP129:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i64 0, i64 9
// CHECK-PPC64LE-NEXT: store i8* null, i8** [[TMP129]], align 8
// CHECK-PPC64LE-NEXT: [[TMP130:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 10
// CHECK-PPC64LE-NEXT: [[TMP131:%.*]] = bitcast i8** [[TMP130]] to %struct.S2***
// CHECK-PPC64LE-NEXT: store %struct.S2** [[PS24]], %struct.S2*** [[TMP131]], align 8
// CHECK-PPC64LE-NEXT: [[TMP132:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 10
// CHECK-PPC64LE-NEXT: [[TMP133:%.*]] = bitcast i8** [[TMP132]] to %struct.S1**
// CHECK-PPC64LE-NEXT: store %struct.S1* [[S28]], %struct.S1** [[TMP133]], align 8
// CHECK-PPC64LE-NEXT: [[TMP134:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 10
// CHECK-PPC64LE-NEXT: store i64 4, i64* [[TMP134]], align 8
// CHECK-PPC64LE-NEXT: [[TMP135:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i64 0, i64 10
// CHECK-PPC64LE-NEXT: store i8* null, i8** [[TMP135]], align 8
// CHECK-PPC64LE-NEXT: [[TMP136:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 0
// CHECK-PPC64LE-NEXT: [[TMP137:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 0
// CHECK-PPC64LE-NEXT: [[TMP138:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0
// CHECK-PPC64LE-NEXT: call void @__tgt_target_data_begin_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i32 11, i8** [[TMP136]], i8** [[TMP137]], i64* [[TMP138]], i64* getelementptr inbounds ([11 x i64], [11 x i64]* @.offload_maptypes.5, i32 0, i32 0), i8** null, i8** null)
// CHECK-PPC64LE-NEXT: [[TMP139:%.*]] = load i32, i32* [[ARG_ADDR]], align 4
// CHECK-PPC64LE-NEXT: [[INC32:%.*]] = add nsw i32 [[TMP139]], 1
// CHECK-PPC64LE-NEXT: store i32 [[INC32]], i32* [[ARG_ADDR]], align 4
// CHECK-PPC64LE-NEXT: [[TMP140:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 0
// CHECK-PPC64LE-NEXT: [[TMP141:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 0
// CHECK-PPC64LE-NEXT: [[TMP142:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0
// CHECK-PPC64LE-NEXT: call void @__tgt_target_data_end_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i32 11, i8** [[TMP140]], i8** [[TMP141]], i64* [[TMP142]], i64* getelementptr inbounds ([11 x i64], [11 x i64]* @.offload_maptypes.5, i32 0, i32 0), i8** null, i8** null)
// CHECK-PPC64LE-NEXT: ret void
//
// CHECK-I386-LABEL: @_Z3fooi(
// CHECK-I386-NEXT: entry:
// CHECK-I386-NEXT: [[ARG_ADDR:%.*]] = alloca i32, align 4
// CHECK-I386-NEXT: [[LB:%.*]] = alloca [5 x float], align 4
// CHECK-I386-NEXT: [[PS1:%.*]] = alloca %struct.S2*, align 4
// CHECK-I386-NEXT: [[PS2:%.*]] = alloca %struct.S2*, align 4
// CHECK-I386-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x i8*], align 4
// CHECK-I386-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x i8*], align 4
// CHECK-I386-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x i8*], align 4
// CHECK-I386-NEXT: [[DOTOFFLOAD_BASEPTRS1:%.*]] = alloca [1 x i8*], align 4
// CHECK-I386-NEXT: [[DOTOFFLOAD_PTRS2:%.*]] = alloca [1 x i8*], align 4
// CHECK-I386-NEXT: [[DOTOFFLOAD_MAPPERS3:%.*]] = alloca [1 x i8*], align 4
// CHECK-I386-NEXT: [[DOTOFFLOAD_BASEPTRS5:%.*]] = alloca [1 x i8*], align 4
// CHECK-I386-NEXT: [[DOTOFFLOAD_PTRS6:%.*]] = alloca [1 x i8*], align 4
// CHECK-I386-NEXT: [[DOTOFFLOAD_MAPPERS7:%.*]] = alloca [1 x i8*], align 4
// CHECK-I386-NEXT: [[DOTOFFLOAD_BASEPTRS29:%.*]] = alloca [11 x i8*], align 4
// CHECK-I386-NEXT: [[DOTOFFLOAD_PTRS30:%.*]] = alloca [11 x i8*], align 4
// CHECK-I386-NEXT: [[DOTOFFLOAD_MAPPERS31:%.*]] = alloca [11 x i8*], align 4
// CHECK-I386-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [11 x i64], align 4
// CHECK-I386-NEXT: store i32 [[ARG:%.*]], i32* [[ARG_ADDR]], align 4
// CHECK-I386-NEXT: [[TMP0:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK-I386-NEXT: [[TMP1:%.*]] = bitcast i8** [[TMP0]] to [5 x float]**
// CHECK-I386-NEXT: store [5 x float]* [[LB]], [5 x float]** [[TMP1]], align 4
// CHECK-I386-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK-I386-NEXT: [[TMP3:%.*]] = bitcast i8** [[TMP2]] to [5 x float]**
// CHECK-I386-NEXT: store [5 x float]* [[LB]], [5 x float]** [[TMP3]], align 4
// CHECK-I386-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0
// CHECK-I386-NEXT: store i8* null, i8** [[TMP4]], align 4
// CHECK-I386-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK-I386-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK-I386-NEXT: call void @__tgt_target_data_begin_mapper(%struct.ident_t* @[[GLOB1:[0-9]+]], i64 -1, i32 1, i8** [[TMP5]], i8** [[TMP6]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null)
// CHECK-I386-NEXT: [[TMP7:%.*]] = load i32, i32* [[ARG_ADDR]], align 4
// CHECK-I386-NEXT: [[INC:%.*]] = add nsw i32 [[TMP7]], 1
// CHECK-I386-NEXT: store i32 [[INC]], i32* [[ARG_ADDR]], align 4
// CHECK-I386-NEXT: [[TMP8:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK-I386-NEXT: [[TMP9:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK-I386-NEXT: call void @__tgt_target_data_end_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i32 1, i8** [[TMP8]], i8** [[TMP9]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null)
// CHECK-I386-NEXT: [[TMP10:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS1]], i32 0, i32 0
// CHECK-I386-NEXT: [[TMP11:%.*]] = bitcast i8** [[TMP10]] to [5 x float]**
// CHECK-I386-NEXT: store [5 x float]* [[LB]], [5 x float]** [[TMP11]], align 4
// CHECK-I386-NEXT: [[TMP12:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS2]], i32 0, i32 0
// CHECK-I386-NEXT: [[TMP13:%.*]] = bitcast i8** [[TMP12]] to [5 x float]**
// CHECK-I386-NEXT: store [5 x float]* [[LB]], [5 x float]** [[TMP13]], align 4
// CHECK-I386-NEXT: [[TMP14:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_MAPPERS3]], i32 0, i32 0
// CHECK-I386-NEXT: store i8* null, i8** [[TMP14]], align 4
// CHECK-I386-NEXT: [[TMP15:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS1]], i32 0, i32 0
// CHECK-I386-NEXT: [[TMP16:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS2]], i32 0, i32 0
// CHECK-I386-NEXT: call void @__tgt_target_data_begin_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i32 1, i8** [[TMP15]], i8** [[TMP16]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes.1, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.2, i32 0, i32 0), i8** null, i8** null)
// CHECK-I386-NEXT: [[TMP17:%.*]] = load i32, i32* [[ARG_ADDR]], align 4
// CHECK-I386-NEXT: [[INC4:%.*]] = add nsw i32 [[TMP17]], 1
// CHECK-I386-NEXT: store i32 [[INC4]], i32* [[ARG_ADDR]], align 4
// CHECK-I386-NEXT: [[TMP18:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS1]], i32 0, i32 0
// CHECK-I386-NEXT: [[TMP19:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS2]], i32 0, i32 0
// CHECK-I386-NEXT: call void @__tgt_target_data_end_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i32 1, i8** [[TMP18]], i8** [[TMP19]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes.1, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.2, i32 0, i32 0), i8** null, i8** null)
// CHECK-I386-NEXT: [[TMP20:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS5]], i32 0, i32 0
// CHECK-I386-NEXT: [[TMP21:%.*]] = bitcast i8** [[TMP20]] to i32**
// CHECK-I386-NEXT: store i32* [[ARG_ADDR]], i32** [[TMP21]], align 4
// CHECK-I386-NEXT: [[TMP22:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS6]], i32 0, i32 0
// CHECK-I386-NEXT: [[TMP23:%.*]] = bitcast i8** [[TMP22]] to i32**
// CHECK-I386-NEXT: store i32* [[ARG_ADDR]], i32** [[TMP23]], align 4
// CHECK-I386-NEXT: [[TMP24:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_MAPPERS7]], i32 0, i32 0
// CHECK-I386-NEXT: store i8* null, i8** [[TMP24]], align 4
// CHECK-I386-NEXT: [[TMP25:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS5]], i32 0, i32 0
// CHECK-I386-NEXT: [[TMP26:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS6]], i32 0, i32 0
// CHECK-I386-NEXT: call void @__tgt_target_data_begin_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i32 1, i8** [[TMP25]], i8** [[TMP26]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes.3, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.4, i32 0, i32 0), i8** null, i8** null)
// CHECK-I386-NEXT: [[TMP27:%.*]] = load i32, i32* [[ARG_ADDR]], align 4
// CHECK-I386-NEXT: [[INC8:%.*]] = add nsw i32 [[TMP27]], 1
// CHECK-I386-NEXT: store i32 [[INC8]], i32* [[ARG_ADDR]], align 4
// CHECK-I386-NEXT: [[TMP28:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS5]], i32 0, i32 0
// CHECK-I386-NEXT: [[TMP29:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS6]], i32 0, i32 0
// CHECK-I386-NEXT: call void @__tgt_target_data_end_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i32 1, i8** [[TMP28]], i8** [[TMP29]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes.3, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.4, i32 0, i32 0), i8** null, i8** null)
// CHECK-I386-NEXT: [[TMP30:%.*]] = load %struct.S2*, %struct.S2** [[PS1]], align 4
// CHECK-I386-NEXT: [[TMP31:%.*]] = load %struct.S2*, %struct.S2** [[PS1]], align 4
// CHECK-I386-NEXT: [[S:%.*]] = getelementptr inbounds [[STRUCT_S2:%.*]], %struct.S2* [[TMP31]], i32 0, i32 0
// CHECK-I386-NEXT: [[TMP32:%.*]] = load %struct.S2*, %struct.S2** [[PS1]], align 4
// CHECK-I386-NEXT: [[TMP33:%.*]] = load %struct.S2*, %struct.S2** [[PS1]], align 4
// CHECK-I386-NEXT: [[PS:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP33]], i32 0, i32 1
// CHECK-I386-NEXT: [[TMP34:%.*]] = load %struct.S2*, %struct.S2** [[PS1]], align 4
// CHECK-I386-NEXT: [[PS9:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP34]], i32 0, i32 1
// CHECK-I386-NEXT: [[TMP35:%.*]] = load %struct.S2*, %struct.S2** [[PS9]], align 4
// CHECK-I386-NEXT: [[PS10:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP35]], i32 0, i32 1
// CHECK-I386-NEXT: [[TMP36:%.*]] = load %struct.S2*, %struct.S2** [[PS1]], align 4
// CHECK-I386-NEXT: [[PS11:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP36]], i32 0, i32 1
// CHECK-I386-NEXT: [[TMP37:%.*]] = load %struct.S2*, %struct.S2** [[PS11]], align 4
// CHECK-I386-NEXT: [[PS12:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP37]], i32 0, i32 1
// CHECK-I386-NEXT: [[TMP38:%.*]] = load %struct.S2*, %struct.S2** [[PS12]], align 4
// CHECK-I386-NEXT: [[PS13:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP38]], i32 0, i32 1
// CHECK-I386-NEXT: [[TMP39:%.*]] = load %struct.S2*, %struct.S2** [[PS1]], align 4
// CHECK-I386-NEXT: [[PS14:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP39]], i32 0, i32 1
// CHECK-I386-NEXT: [[TMP40:%.*]] = load %struct.S2*, %struct.S2** [[PS14]], align 4
// CHECK-I386-NEXT: [[PS15:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP40]], i32 0, i32 1
// CHECK-I386-NEXT: [[TMP41:%.*]] = load %struct.S2*, %struct.S2** [[PS15]], align 4
// CHECK-I386-NEXT: [[PS16:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP41]], i32 0, i32 1
// CHECK-I386-NEXT: [[TMP42:%.*]] = load %struct.S2*, %struct.S2** [[PS16]], align 4
// CHECK-I386-NEXT: [[S17:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP42]], i32 0, i32 0
// CHECK-I386-NEXT: [[TMP43:%.*]] = getelementptr %struct.S2*, %struct.S2** [[PS]], i32 1
// CHECK-I386-NEXT: [[TMP44:%.*]] = bitcast %struct.S1* [[S]] to i8*
// CHECK-I386-NEXT: [[TMP45:%.*]] = bitcast %struct.S2** [[TMP43]] to i8*
// CHECK-I386-NEXT: [[TMP46:%.*]] = ptrtoint i8* [[TMP45]] to i64
// CHECK-I386-NEXT: [[TMP47:%.*]] = ptrtoint i8* [[TMP44]] to i64
// CHECK-I386-NEXT: [[TMP48:%.*]] = sub i64 [[TMP46]], [[TMP47]]
// CHECK-I386-NEXT: [[TMP49:%.*]] = sdiv exact i64 [[TMP48]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
// CHECK-I386-NEXT: [[TMP50:%.*]] = load %struct.S2*, %struct.S2** [[PS2]], align 4
// CHECK-I386-NEXT: [[TMP51:%.*]] = load %struct.S2*, %struct.S2** [[PS2]], align 4
// CHECK-I386-NEXT: [[S18:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP51]], i32 0, i32 0
// CHECK-I386-NEXT: [[TMP52:%.*]] = load %struct.S2*, %struct.S2** [[PS2]], align 4
// CHECK-I386-NEXT: [[TMP53:%.*]] = load %struct.S2*, %struct.S2** [[PS2]], align 4
// CHECK-I386-NEXT: [[PS19:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP53]], i32 0, i32 1
// CHECK-I386-NEXT: [[TMP54:%.*]] = load %struct.S2*, %struct.S2** [[PS2]], align 4
// CHECK-I386-NEXT: [[PS20:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP54]], i32 0, i32 1
// CHECK-I386-NEXT: [[TMP55:%.*]] = load %struct.S2*, %struct.S2** [[PS20]], align 4
// CHECK-I386-NEXT: [[PS21:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP55]], i32 0, i32 1
// CHECK-I386-NEXT: [[TMP56:%.*]] = load %struct.S2*, %struct.S2** [[PS2]], align 4
// CHECK-I386-NEXT: [[PS22:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP56]], i32 0, i32 1
// CHECK-I386-NEXT: [[TMP57:%.*]] = load %struct.S2*, %struct.S2** [[PS22]], align 4
// CHECK-I386-NEXT: [[PS23:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP57]], i32 0, i32 1
// CHECK-I386-NEXT: [[TMP58:%.*]] = load %struct.S2*, %struct.S2** [[PS23]], align 4
// CHECK-I386-NEXT: [[PS24:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP58]], i32 0, i32 1
// CHECK-I386-NEXT: [[TMP59:%.*]] = load %struct.S2*, %struct.S2** [[PS2]], align 4
// CHECK-I386-NEXT: [[PS25:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP59]], i32 0, i32 1
// CHECK-I386-NEXT: [[TMP60:%.*]] = load %struct.S2*, %struct.S2** [[PS25]], align 4
// CHECK-I386-NEXT: [[PS26:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP60]], i32 0, i32 1
// CHECK-I386-NEXT: [[TMP61:%.*]] = load %struct.S2*, %struct.S2** [[PS26]], align 4
// CHECK-I386-NEXT: [[PS27:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP61]], i32 0, i32 1
// CHECK-I386-NEXT: [[TMP62:%.*]] = load %struct.S2*, %struct.S2** [[PS27]], align 4
// CHECK-I386-NEXT: [[S28:%.*]] = getelementptr inbounds [[STRUCT_S2]], %struct.S2* [[TMP62]], i32 0, i32 0
// CHECK-I386-NEXT: [[TMP63:%.*]] = getelementptr %struct.S2*, %struct.S2** [[PS19]], i32 1
// CHECK-I386-NEXT: [[TMP64:%.*]] = bitcast %struct.S1* [[S18]] to i8*
// CHECK-I386-NEXT: [[TMP65:%.*]] = bitcast %struct.S2** [[TMP63]] to i8*
// CHECK-I386-NEXT: [[TMP66:%.*]] = ptrtoint i8* [[TMP65]] to i64
// CHECK-I386-NEXT: [[TMP67:%.*]] = ptrtoint i8* [[TMP64]] to i64
// CHECK-I386-NEXT: [[TMP68:%.*]] = sub i64 [[TMP66]], [[TMP67]]
// CHECK-I386-NEXT: [[TMP69:%.*]] = sdiv exact i64 [[TMP68]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
// CHECK-I386-NEXT: [[TMP70:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 0
// CHECK-I386-NEXT: [[TMP71:%.*]] = bitcast i8** [[TMP70]] to %struct.S2**
// CHECK-I386-NEXT: store %struct.S2* [[TMP30]], %struct.S2** [[TMP71]], align 4
// CHECK-I386-NEXT: [[TMP72:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 0
// CHECK-I386-NEXT: [[TMP73:%.*]] = bitcast i8** [[TMP72]] to %struct.S1**
// CHECK-I386-NEXT: store %struct.S1* [[S]], %struct.S1** [[TMP73]], align 4
// CHECK-I386-NEXT: [[TMP74:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0
// CHECK-I386-NEXT: store i64 [[TMP49]], i64* [[TMP74]], align 4
// CHECK-I386-NEXT: [[TMP75:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i32 0, i32 0
// CHECK-I386-NEXT: store i8* null, i8** [[TMP75]], align 4
// CHECK-I386-NEXT: [[TMP76:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 1
// CHECK-I386-NEXT: [[TMP77:%.*]] = bitcast i8** [[TMP76]] to %struct.S2**
// CHECK-I386-NEXT: store %struct.S2* [[TMP30]], %struct.S2** [[TMP77]], align 4
// CHECK-I386-NEXT: [[TMP78:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 1
// CHECK-I386-NEXT: [[TMP79:%.*]] = bitcast i8** [[TMP78]] to %struct.S1**
// CHECK-I386-NEXT: store %struct.S1* [[S]], %struct.S1** [[TMP79]], align 4
// CHECK-I386-NEXT: [[TMP80:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 1
// CHECK-I386-NEXT: store i64 4, i64* [[TMP80]], align 4
// CHECK-I386-NEXT: [[TMP81:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i32 0, i32 1
// CHECK-I386-NEXT: store i8* null, i8** [[TMP81]], align 4
// CHECK-I386-NEXT: [[TMP82:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 2
// CHECK-I386-NEXT: [[TMP83:%.*]] = bitcast i8** [[TMP82]] to %struct.S2***
// CHECK-I386-NEXT: store %struct.S2** [[PS]], %struct.S2*** [[TMP83]], align 4
// CHECK-I386-NEXT: [[TMP84:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 2
// CHECK-I386-NEXT: [[TMP85:%.*]] = bitcast i8** [[TMP84]] to %struct.S2***
// CHECK-I386-NEXT: store %struct.S2** [[PS10]], %struct.S2*** [[TMP85]], align 4
// CHECK-I386-NEXT: [[TMP86:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 2
// CHECK-I386-NEXT: store i64 4, i64* [[TMP86]], align 4
// CHECK-I386-NEXT: [[TMP87:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i32 0, i32 2
// CHECK-I386-NEXT: store i8* null, i8** [[TMP87]], align 4
// CHECK-I386-NEXT: [[TMP88:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 3
// CHECK-I386-NEXT: [[TMP89:%.*]] = bitcast i8** [[TMP88]] to %struct.S2***
// CHECK-I386-NEXT: store %struct.S2** [[PS10]], %struct.S2*** [[TMP89]], align 4
// CHECK-I386-NEXT: [[TMP90:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 3
// CHECK-I386-NEXT: [[TMP91:%.*]] = bitcast i8** [[TMP90]] to %struct.S2***
// CHECK-I386-NEXT: store %struct.S2** [[PS13]], %struct.S2*** [[TMP91]], align 4
// CHECK-I386-NEXT: [[TMP92:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 3
// CHECK-I386-NEXT: store i64 4, i64* [[TMP92]], align 4
// CHECK-I386-NEXT: [[TMP93:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i32 0, i32 3
// CHECK-I386-NEXT: store i8* null, i8** [[TMP93]], align 4
// CHECK-I386-NEXT: [[TMP94:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 4
// CHECK-I386-NEXT: [[TMP95:%.*]] = bitcast i8** [[TMP94]] to %struct.S2***
// CHECK-I386-NEXT: store %struct.S2** [[PS13]], %struct.S2*** [[TMP95]], align 4
// CHECK-I386-NEXT: [[TMP96:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 4
// CHECK-I386-NEXT: [[TMP97:%.*]] = bitcast i8** [[TMP96]] to %struct.S1**
// CHECK-I386-NEXT: store %struct.S1* [[S17]], %struct.S1** [[TMP97]], align 4
// CHECK-I386-NEXT: [[TMP98:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 4
// CHECK-I386-NEXT: store i64 4, i64* [[TMP98]], align 4
// CHECK-I386-NEXT: [[TMP99:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i32 0, i32 4
// CHECK-I386-NEXT: store i8* null, i8** [[TMP99]], align 4
// CHECK-I386-NEXT: [[TMP100:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 5
// CHECK-I386-NEXT: [[TMP101:%.*]] = bitcast i8** [[TMP100]] to i32**
// CHECK-I386-NEXT: store i32* [[ARG_ADDR]], i32** [[TMP101]], align 4
// CHECK-I386-NEXT: [[TMP102:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 5
// CHECK-I386-NEXT: [[TMP103:%.*]] = bitcast i8** [[TMP102]] to i32**
// CHECK-I386-NEXT: store i32* [[ARG_ADDR]], i32** [[TMP103]], align 4
// CHECK-I386-NEXT: [[TMP104:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 5
// CHECK-I386-NEXT: store i64 4, i64* [[TMP104]], align 4
// CHECK-I386-NEXT: [[TMP105:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i32 0, i32 5
// CHECK-I386-NEXT: store i8* null, i8** [[TMP105]], align 4
// CHECK-I386-NEXT: [[TMP106:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 6
// CHECK-I386-NEXT: [[TMP107:%.*]] = bitcast i8** [[TMP106]] to %struct.S2**
// CHECK-I386-NEXT: store %struct.S2* [[TMP50]], %struct.S2** [[TMP107]], align 4
// CHECK-I386-NEXT: [[TMP108:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 6
// CHECK-I386-NEXT: [[TMP109:%.*]] = bitcast i8** [[TMP108]] to %struct.S1**
// CHECK-I386-NEXT: store %struct.S1* [[S18]], %struct.S1** [[TMP109]], align 4
// CHECK-I386-NEXT: [[TMP110:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 6
// CHECK-I386-NEXT: store i64 [[TMP69]], i64* [[TMP110]], align 4
// CHECK-I386-NEXT: [[TMP111:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i32 0, i32 6
// CHECK-I386-NEXT: store i8* null, i8** [[TMP111]], align 4
// CHECK-I386-NEXT: [[TMP112:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 7
// CHECK-I386-NEXT: [[TMP113:%.*]] = bitcast i8** [[TMP112]] to %struct.S2**
// CHECK-I386-NEXT: store %struct.S2* [[TMP50]], %struct.S2** [[TMP113]], align 4
// CHECK-I386-NEXT: [[TMP114:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 7
// CHECK-I386-NEXT: [[TMP115:%.*]] = bitcast i8** [[TMP114]] to %struct.S1**
// CHECK-I386-NEXT: store %struct.S1* [[S18]], %struct.S1** [[TMP115]], align 4
// CHECK-I386-NEXT: [[TMP116:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 7
// CHECK-I386-NEXT: store i64 4, i64* [[TMP116]], align 4
// CHECK-I386-NEXT: [[TMP117:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i32 0, i32 7
// CHECK-I386-NEXT: store i8* null, i8** [[TMP117]], align 4
// CHECK-I386-NEXT: [[TMP118:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 8
// CHECK-I386-NEXT: [[TMP119:%.*]] = bitcast i8** [[TMP118]] to %struct.S2***
// CHECK-I386-NEXT: store %struct.S2** [[PS19]], %struct.S2*** [[TMP119]], align 4
// CHECK-I386-NEXT: [[TMP120:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 8
// CHECK-I386-NEXT: [[TMP121:%.*]] = bitcast i8** [[TMP120]] to %struct.S2***
// CHECK-I386-NEXT: store %struct.S2** [[PS21]], %struct.S2*** [[TMP121]], align 4
// CHECK-I386-NEXT: [[TMP122:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 8
// CHECK-I386-NEXT: store i64 4, i64* [[TMP122]], align 4
// CHECK-I386-NEXT: [[TMP123:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i32 0, i32 8
// CHECK-I386-NEXT: store i8* null, i8** [[TMP123]], align 4
// CHECK-I386-NEXT: [[TMP124:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 9
// CHECK-I386-NEXT: [[TMP125:%.*]] = bitcast i8** [[TMP124]] to %struct.S2***
// CHECK-I386-NEXT: store %struct.S2** [[PS21]], %struct.S2*** [[TMP125]], align 4
// CHECK-I386-NEXT: [[TMP126:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 9
// CHECK-I386-NEXT: [[TMP127:%.*]] = bitcast i8** [[TMP126]] to %struct.S2***
// CHECK-I386-NEXT: store %struct.S2** [[PS24]], %struct.S2*** [[TMP127]], align 4
// CHECK-I386-NEXT: [[TMP128:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 9
// CHECK-I386-NEXT: store i64 4, i64* [[TMP128]], align 4
// CHECK-I386-NEXT: [[TMP129:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i32 0, i32 9
// CHECK-I386-NEXT: store i8* null, i8** [[TMP129]], align 4
// CHECK-I386-NEXT: [[TMP130:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 10
// CHECK-I386-NEXT: [[TMP131:%.*]] = bitcast i8** [[TMP130]] to %struct.S2***
// CHECK-I386-NEXT: store %struct.S2** [[PS24]], %struct.S2*** [[TMP131]], align 4
// CHECK-I386-NEXT: [[TMP132:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 10
// CHECK-I386-NEXT: [[TMP133:%.*]] = bitcast i8** [[TMP132]] to %struct.S1**
// CHECK-I386-NEXT: store %struct.S1* [[S28]], %struct.S1** [[TMP133]], align 4
// CHECK-I386-NEXT: [[TMP134:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 10
// CHECK-I386-NEXT: store i64 4, i64* [[TMP134]], align 4
// CHECK-I386-NEXT: [[TMP135:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_MAPPERS31]], i32 0, i32 10
// CHECK-I386-NEXT: store i8* null, i8** [[TMP135]], align 4
// CHECK-I386-NEXT: [[TMP136:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 0
// CHECK-I386-NEXT: [[TMP137:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 0
// CHECK-I386-NEXT: [[TMP138:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0
// CHECK-I386-NEXT: call void @__tgt_target_data_begin_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i32 11, i8** [[TMP136]], i8** [[TMP137]], i64* [[TMP138]], i64* getelementptr inbounds ([11 x i64], [11 x i64]* @.offload_maptypes.5, i32 0, i32 0), i8** null, i8** null)
// CHECK-I386-NEXT: [[TMP139:%.*]] = load i32, i32* [[ARG_ADDR]], align 4
// CHECK-I386-NEXT: [[INC32:%.*]] = add nsw i32 [[TMP139]], 1
// CHECK-I386-NEXT: store i32 [[INC32]], i32* [[ARG_ADDR]], align 4
// CHECK-I386-NEXT: [[TMP140:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_BASEPTRS29]], i32 0, i32 0
// CHECK-I386-NEXT: [[TMP141:%.*]] = getelementptr inbounds [11 x i8*], [11 x i8*]* [[DOTOFFLOAD_PTRS30]], i32 0, i32 0
// CHECK-I386-NEXT: [[TMP142:%.*]] = getelementptr inbounds [11 x i64], [11 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0
// CHECK-I386-NEXT: call void @__tgt_target_data_end_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i32 11, i8** [[TMP140]], i8** [[TMP141]], i64* [[TMP142]], i64* getelementptr inbounds ([11 x i64], [11 x i64]* @.offload_maptypes.5, i32 0, i32 0), i8** null, i8** null)
// CHECK-I386-NEXT: ret void
//
void foo(int arg) {
float lb[5];
S2 *ps1;
S2 *ps2;
#pragma omp target data map(ompx_hold, to: lb)
{++arg;}
#pragma omp target data map(always close ompx_hold, to: lb)
{++arg;}
#pragma omp target data map(ompx_hold, tofrom : arg)
{++arg;}
// Make sure the struct picks up ompx_hold even if another element of the
// struct doesn't have ompx_hold.
#pragma omp target data map(tofrom : ps1->s, arg) \
map(ompx_hold, tofrom : ps1->ps->ps->ps->s, ps2->s) \
map(tofrom : ps2->ps->ps->ps->s)
{
++(arg);
}
}
#endif

View File

@ -1,8 +1,14 @@
// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp -ferror-limit 100 -o - %s -Wuninitialized
// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp -ferror-limit 100 -o - -x c++ %s -Wuninitialized
// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp -fopenmp -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}}
@ -25,5 +31,12 @@ int main(int argc, char **argv) {
#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;
}

View File

@ -1,8 +1,14 @@
// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp -ferror-limit 100 -o - %s -Wuninitialized
// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp -ferror-limit 100 -o - -x c++ %s -Wuninitialized
// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp -fopenmp -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) {
@ -18,5 +24,15 @@ 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;
}

View File

@ -0,0 +1,928 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" --global-value-regex ".offload_maptypes.*" ".offload_sizes.*" --global-hex-value-regex ".offload_maptypes.*"
// expected-no-diagnostics
#ifndef HEADER
#define HEADER
//--------------------------------------------------
// With -DUSE.
//--------------------------------------------------
// powerpc64le-ibm-linux-gnu
// RUN: %clang_cc1 -DUSE -verify -fopenmp -fopenmp-extensions \
// RUN: -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ \
// RUN: -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | \
// RUN: FileCheck %s --check-prefixes=CHECK-USE-PPC64LE
// RUN: %clang_cc1 -DUSE -fopenmp -fopenmp-extensions \
// RUN: -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 \
// RUN: -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
// RUN: %clang_cc1 -DUSE -fopenmp -fopenmp-extensions \
// RUN: -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ \
// RUN: -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t \
// RUN: -verify %s -emit-llvm -o - | \
// RUN: FileCheck %s --check-prefixes=CHECK-USE-PPC64LE
// i386-pc-linux-gnu
// RUN: %clang_cc1 -DUSE -verify -fopenmp -fopenmp-extensions \
// RUN: -fopenmp-targets=i386-pc-linux-gnu -x c++ \
// RUN: -triple i386-unknown-unknown -emit-llvm %s -o - | \
// RUN: FileCheck %s --check-prefixes=CHECK-USE-I386
// RUN: %clang_cc1 -DUSE -fopenmp -fopenmp-extensions \
// RUN: -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 \
// RUN: -triple i386-unknown-unknown -emit-pch -o %t %s
// RUN: %clang_cc1 -DUSE -fopenmp -fopenmp-extensions \
// RUN: -fopenmp-targets=i386-pc-linux-gnu -x c++ \
// RUN: -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s \
// RUN: -emit-llvm -o - | \
// RUN: FileCheck %s --check-prefixes=CHECK-USE-I386
//--------------------------------------------------
// Without -DUSE.
//--------------------------------------------------
// powerpc64le-ibm-linux-gnu
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-extensions \
// RUN: -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ \
// RUN: -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | \
// RUN: FileCheck %s --check-prefixes=CHECK-NOUSE-PPC64LE
// RUN: %clang_cc1 -fopenmp -fopenmp-extensions \
// RUN: -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 \
// RUN: -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp -fopenmp-extensions \
// RUN: -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ \
// RUN: -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t \
// RUN: -verify %s -emit-llvm -o - | \
// RUN: FileCheck %s --check-prefixes=CHECK-NOUSE-PPC64LE
// i386-pc-linux-gnu
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-extensions \
// RUN: -fopenmp-targets=i386-pc-linux-gnu -x c++ \
// RUN: -triple i386-unknown-unknown -emit-llvm %s -o - | \
// RUN: FileCheck %s --check-prefixes=CHECK-NOUSE-I386
// RUN: %clang_cc1 -fopenmp -fopenmp-extensions \
// RUN: -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 \
// RUN: -triple i386-unknown-unknown -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp -fopenmp-extensions \
// RUN: -fopenmp-targets=i386-pc-linux-gnu -x c++ \
// RUN: -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s \
// RUN: -emit-llvm -o - | \
// RUN: FileCheck %s --check-prefixes=CHECK-NOUSE-I386
// Map flags used in @.offload_maptypes* below:
//
// TO = 0x1
// FROM = 0x2
// ALWAYS = 0x4
// TARGET_PARAM = 0x20
// CLOSE = 0x400
// OMPX_HOLD = 0x2000
// MEMBER_OF_1 = 0x1000000000000
// MEMBER_OF_5 = 0x5000000000000
//.
// CHECK-USE-PPC64LE: @.offload_maptypes = private unnamed_addr constant [7 x i64] [i64 [[#0x2020]], i64 [[#0x1000000002003]], i64 [[#0x1000000002003]], i64 [[#0x2023]], i64 [[#0x2020]], i64 [[#0x5000000002003]], i64 [[#0x5000000002003]]]
// CHECK-USE-PPC64LE: @.offload_sizes = private unnamed_addr constant [1 x i64] [i64 4]
// CHECK-USE-PPC64LE: @.offload_maptypes.1 = private unnamed_addr constant [1 x i64] [i64 [[#0x2427]]]
// CHECK-USE-PPC64LE: @.offload_maptypes.2 = private unnamed_addr constant [3 x i64] [i64 [[#0x2020]], i64 [[#0x1000000002003]], i64 [[#0x1000000002003]]]
//.
// CHECK-USE-I386: @.offload_maptypes = private unnamed_addr constant [7 x i64] [i64 [[#0x2020]], i64 [[#0x1000000002003]], i64 [[#0x1000000002003]], i64 [[#0x2023]], i64 [[#0x2020]], i64 [[#0x5000000002003]], i64 [[#0x5000000002003]]]
// CHECK-USE-I386: @.offload_sizes = private unnamed_addr constant [1 x i64] [i64 4]
// CHECK-USE-I386: @.offload_maptypes.1 = private unnamed_addr constant [1 x i64] [i64 [[#0x2427]]]
// CHECK-USE-I386: @.offload_maptypes.2 = private unnamed_addr constant [3 x i64] [i64 [[#0x2020]], i64 [[#0x1000000002003]], i64 [[#0x1000000002003]]]
//.
// CHECK-NOUSE-PPC64LE: @.offload_maptypes = private unnamed_addr constant [7 x i64] [i64 [[#0x2000]], i64 [[#0x1000000002003]], i64 [[#0x1000000002003]], i64 [[#0x2003]], i64 [[#0x2000]], i64 [[#0x5000000002003]], i64 [[#0x5000000002003]]]
// CHECK-NOUSE-PPC64LE: @.offload_sizes = private unnamed_addr constant [1 x i64] [i64 4]
// CHECK-NOUSE-PPC64LE: @.offload_maptypes.1 = private unnamed_addr constant [1 x i64] [i64 [[#0x2407]]]
// CHECK-NOUSE-PPC64LE: @.offload_maptypes.2 = private unnamed_addr constant [3 x i64] [i64 [[#0x2000]], i64 [[#0x1000000002003]], i64 [[#0x1000000002003]]]
//.
// CHECK-NOUSE-I386: @.offload_maptypes = private unnamed_addr constant [7 x i64] [i64 [[#0x2000]], i64 [[#0x1000000002003]], i64 [[#0x1000000002003]], i64 [[#0x2003]], i64 [[#0x2000]], i64 [[#0x5000000002003]], i64 [[#0x5000000002003]]]
// CHECK-NOUSE-I386: @.offload_sizes = private unnamed_addr constant [1 x i64] [i64 4]
// CHECK-NOUSE-I386: @.offload_maptypes.1 = private unnamed_addr constant [1 x i64] [i64 [[#0x2407]]]
// CHECK-NOUSE-I386: @.offload_maptypes.2 = private unnamed_addr constant [3 x i64] [i64 [[#0x2000]], i64 [[#0x1000000002003]], i64 [[#0x1000000002003]]]
//.
struct ST {
int i;
int j;
void test_present_members();
};
// CHECK-USE-PPC64LE-LABEL: @_Z20explicit_maps_singlei(
// CHECK-USE-PPC64LE-NEXT: entry:
// CHECK-USE-PPC64LE-NEXT: [[II_ADDR:%.*]] = alloca i32, align 4
// CHECK-USE-PPC64LE-NEXT: [[A:%.*]] = alloca i32, align 4
// CHECK-USE-PPC64LE-NEXT: [[ST1:%.*]] = alloca [[STRUCT_ST:%.*]], align 4
// CHECK-USE-PPC64LE-NEXT: [[ST2:%.*]] = alloca [[STRUCT_ST]], align 4
// CHECK-USE-PPC64LE-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [7 x i8*], align 8
// CHECK-USE-PPC64LE-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [7 x i8*], align 8
// CHECK-USE-PPC64LE-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [7 x i8*], align 8
// CHECK-USE-PPC64LE-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [7 x i64], align 8
// CHECK-USE-PPC64LE-NEXT: [[DOTOFFLOAD_BASEPTRS3:%.*]] = alloca [1 x i8*], align 8
// CHECK-USE-PPC64LE-NEXT: [[DOTOFFLOAD_PTRS4:%.*]] = alloca [1 x i8*], align 8
// CHECK-USE-PPC64LE-NEXT: [[DOTOFFLOAD_MAPPERS5:%.*]] = alloca [1 x i8*], align 8
// CHECK-USE-PPC64LE-NEXT: store i32 [[II:%.*]], i32* [[II_ADDR]], align 4
// CHECK-USE-PPC64LE-NEXT: [[TMP0:%.*]] = load i32, i32* [[II_ADDR]], align 4
// CHECK-USE-PPC64LE-NEXT: store i32 [[TMP0]], i32* [[A]], align 4
// CHECK-USE-PPC64LE-NEXT: [[I:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[ST1]], i32 0, i32 0
// CHECK-USE-PPC64LE-NEXT: [[J:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[ST1]], i32 0, i32 1
// CHECK-USE-PPC64LE-NEXT: [[TMP1:%.*]] = getelementptr i32, i32* [[J]], i32 1
// CHECK-USE-PPC64LE-NEXT: [[TMP2:%.*]] = bitcast i32* [[I]] to i8*
// CHECK-USE-PPC64LE-NEXT: [[TMP3:%.*]] = bitcast i32* [[TMP1]] to i8*
// CHECK-USE-PPC64LE-NEXT: [[TMP4:%.*]] = ptrtoint i8* [[TMP3]] to i64
// CHECK-USE-PPC64LE-NEXT: [[TMP5:%.*]] = ptrtoint i8* [[TMP2]] to i64
// CHECK-USE-PPC64LE-NEXT: [[TMP6:%.*]] = sub i64 [[TMP4]], [[TMP5]]
// CHECK-USE-PPC64LE-NEXT: [[TMP7:%.*]] = sdiv exact i64 [[TMP6]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
// CHECK-USE-PPC64LE-NEXT: [[I1:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[ST2]], i32 0, i32 0
// CHECK-USE-PPC64LE-NEXT: [[J2:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[ST2]], i32 0, i32 1
// CHECK-USE-PPC64LE-NEXT: [[TMP8:%.*]] = getelementptr i32, i32* [[J2]], i32 1
// CHECK-USE-PPC64LE-NEXT: [[TMP9:%.*]] = bitcast i32* [[I1]] to i8*
// CHECK-USE-PPC64LE-NEXT: [[TMP10:%.*]] = bitcast i32* [[TMP8]] to i8*
// CHECK-USE-PPC64LE-NEXT: [[TMP11:%.*]] = ptrtoint i8* [[TMP10]] to i64
// CHECK-USE-PPC64LE-NEXT: [[TMP12:%.*]] = ptrtoint i8* [[TMP9]] to i64
// CHECK-USE-PPC64LE-NEXT: [[TMP13:%.*]] = sub i64 [[TMP11]], [[TMP12]]
// CHECK-USE-PPC64LE-NEXT: [[TMP14:%.*]] = sdiv exact i64 [[TMP13]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
// CHECK-USE-PPC64LE-NEXT: [[TMP15:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK-USE-PPC64LE-NEXT: [[TMP16:%.*]] = bitcast i8** [[TMP15]] to %struct.ST**
// CHECK-USE-PPC64LE-NEXT: store %struct.ST* [[ST1]], %struct.ST** [[TMP16]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP17:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK-USE-PPC64LE-NEXT: [[TMP18:%.*]] = bitcast i8** [[TMP17]] to i32**
// CHECK-USE-PPC64LE-NEXT: store i32* [[I]], i32** [[TMP18]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP19:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0
// CHECK-USE-PPC64LE-NEXT: store i64 [[TMP7]], i64* [[TMP19]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP20:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0
// CHECK-USE-PPC64LE-NEXT: store i8* null, i8** [[TMP20]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP21:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1
// CHECK-USE-PPC64LE-NEXT: [[TMP22:%.*]] = bitcast i8** [[TMP21]] to %struct.ST**
// CHECK-USE-PPC64LE-NEXT: store %struct.ST* [[ST1]], %struct.ST** [[TMP22]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP23:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 1
// CHECK-USE-PPC64LE-NEXT: [[TMP24:%.*]] = bitcast i8** [[TMP23]] to i32**
// CHECK-USE-PPC64LE-NEXT: store i32* [[I]], i32** [[TMP24]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP25:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 1
// CHECK-USE-PPC64LE-NEXT: store i64 4, i64* [[TMP25]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP26:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1
// CHECK-USE-PPC64LE-NEXT: store i8* null, i8** [[TMP26]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP27:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2
// CHECK-USE-PPC64LE-NEXT: [[TMP28:%.*]] = bitcast i8** [[TMP27]] to %struct.ST**
// CHECK-USE-PPC64LE-NEXT: store %struct.ST* [[ST1]], %struct.ST** [[TMP28]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP29:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 2
// CHECK-USE-PPC64LE-NEXT: [[TMP30:%.*]] = bitcast i8** [[TMP29]] to i32**
// CHECK-USE-PPC64LE-NEXT: store i32* [[J]], i32** [[TMP30]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP31:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 2
// CHECK-USE-PPC64LE-NEXT: store i64 4, i64* [[TMP31]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP32:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2
// CHECK-USE-PPC64LE-NEXT: store i8* null, i8** [[TMP32]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP33:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 3
// CHECK-USE-PPC64LE-NEXT: [[TMP34:%.*]] = bitcast i8** [[TMP33]] to i32**
// CHECK-USE-PPC64LE-NEXT: store i32* [[A]], i32** [[TMP34]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP35:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 3
// CHECK-USE-PPC64LE-NEXT: [[TMP36:%.*]] = bitcast i8** [[TMP35]] to i32**
// CHECK-USE-PPC64LE-NEXT: store i32* [[A]], i32** [[TMP36]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP37:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 3
// CHECK-USE-PPC64LE-NEXT: store i64 4, i64* [[TMP37]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP38:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 3
// CHECK-USE-PPC64LE-NEXT: store i8* null, i8** [[TMP38]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP39:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 4
// CHECK-USE-PPC64LE-NEXT: [[TMP40:%.*]] = bitcast i8** [[TMP39]] to %struct.ST**
// CHECK-USE-PPC64LE-NEXT: store %struct.ST* [[ST2]], %struct.ST** [[TMP40]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP41:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 4
// CHECK-USE-PPC64LE-NEXT: [[TMP42:%.*]] = bitcast i8** [[TMP41]] to i32**
// CHECK-USE-PPC64LE-NEXT: store i32* [[I1]], i32** [[TMP42]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP43:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 4
// CHECK-USE-PPC64LE-NEXT: store i64 [[TMP14]], i64* [[TMP43]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP44:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 4
// CHECK-USE-PPC64LE-NEXT: store i8* null, i8** [[TMP44]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP45:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 5
// CHECK-USE-PPC64LE-NEXT: [[TMP46:%.*]] = bitcast i8** [[TMP45]] to %struct.ST**
// CHECK-USE-PPC64LE-NEXT: store %struct.ST* [[ST2]], %struct.ST** [[TMP46]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP47:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 5
// CHECK-USE-PPC64LE-NEXT: [[TMP48:%.*]] = bitcast i8** [[TMP47]] to i32**
// CHECK-USE-PPC64LE-NEXT: store i32* [[I1]], i32** [[TMP48]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP49:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 5
// CHECK-USE-PPC64LE-NEXT: store i64 4, i64* [[TMP49]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP50:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 5
// CHECK-USE-PPC64LE-NEXT: store i8* null, i8** [[TMP50]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP51:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 6
// CHECK-USE-PPC64LE-NEXT: [[TMP52:%.*]] = bitcast i8** [[TMP51]] to %struct.ST**
// CHECK-USE-PPC64LE-NEXT: store %struct.ST* [[ST2]], %struct.ST** [[TMP52]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP53:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 6
// CHECK-USE-PPC64LE-NEXT: [[TMP54:%.*]] = bitcast i8** [[TMP53]] to i32**
// CHECK-USE-PPC64LE-NEXT: store i32* [[J2]], i32** [[TMP54]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP55:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 6
// CHECK-USE-PPC64LE-NEXT: store i64 4, i64* [[TMP55]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP56:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 6
// CHECK-USE-PPC64LE-NEXT: store i8* null, i8** [[TMP56]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP57:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK-USE-PPC64LE-NEXT: [[TMP58:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK-USE-PPC64LE-NEXT: [[TMP59:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0
// CHECK-USE-PPC64LE-NEXT: [[TMP60:%.*]] = call i32 @__tgt_target_mapper(%struct.ident_t* @[[GLOB1:[0-9]+]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z20explicit_maps_singlei_l654.region_id, i32 7, i8** [[TMP57]], i8** [[TMP58]], i64* [[TMP59]], i64* getelementptr inbounds ([7 x i64], [7 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null)
// CHECK-USE-PPC64LE-NEXT: [[TMP61:%.*]] = icmp ne i32 [[TMP60]], 0
// CHECK-USE-PPC64LE-NEXT: br i1 [[TMP61]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]
// CHECK-USE-PPC64LE: omp_offload.failed:
// CHECK-USE-PPC64LE-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z20explicit_maps_singlei_l654(%struct.ST* [[ST1]], i32* [[A]], %struct.ST* [[ST2]]) #[[ATTR2:[0-9]+]]
// CHECK-USE-PPC64LE-NEXT: br label [[OMP_OFFLOAD_CONT]]
// CHECK-USE-PPC64LE: omp_offload.cont:
// CHECK-USE-PPC64LE-NEXT: [[TMP62:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0
// CHECK-USE-PPC64LE-NEXT: [[TMP63:%.*]] = bitcast i8** [[TMP62]] to i32**
// CHECK-USE-PPC64LE-NEXT: store i32* [[A]], i32** [[TMP63]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP64:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS4]], i32 0, i32 0
// CHECK-USE-PPC64LE-NEXT: [[TMP65:%.*]] = bitcast i8** [[TMP64]] to i32**
// CHECK-USE-PPC64LE-NEXT: store i32* [[A]], i32** [[TMP65]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP66:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_MAPPERS5]], i64 0, i64 0
// CHECK-USE-PPC64LE-NEXT: store i8* null, i8** [[TMP66]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP67:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0
// CHECK-USE-PPC64LE-NEXT: [[TMP68:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS4]], i32 0, i32 0
// CHECK-USE-PPC64LE-NEXT: [[TMP69:%.*]] = call i32 @__tgt_target_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z20explicit_maps_singlei_l668.region_id, i32 1, i8** [[TMP67]], i8** [[TMP68]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.1, i32 0, i32 0), i8** null, i8** null)
// CHECK-USE-PPC64LE-NEXT: [[TMP70:%.*]] = icmp ne i32 [[TMP69]], 0
// CHECK-USE-PPC64LE-NEXT: br i1 [[TMP70]], label [[OMP_OFFLOAD_FAILED6:%.*]], label [[OMP_OFFLOAD_CONT7:%.*]]
// CHECK-USE-PPC64LE: omp_offload.failed6:
// CHECK-USE-PPC64LE-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z20explicit_maps_singlei_l668(i32* [[A]]) #[[ATTR2]]
// CHECK-USE-PPC64LE-NEXT: br label [[OMP_OFFLOAD_CONT7]]
// CHECK-USE-PPC64LE: omp_offload.cont7:
// CHECK-USE-PPC64LE-NEXT: ret void
//
// CHECK-USE-I386-LABEL: @_Z20explicit_maps_singlei(
// CHECK-USE-I386-NEXT: entry:
// CHECK-USE-I386-NEXT: [[II_ADDR:%.*]] = alloca i32, align 4
// CHECK-USE-I386-NEXT: [[A:%.*]] = alloca i32, align 4
// CHECK-USE-I386-NEXT: [[ST1:%.*]] = alloca [[STRUCT_ST:%.*]], align 4
// CHECK-USE-I386-NEXT: [[ST2:%.*]] = alloca [[STRUCT_ST]], align 4
// CHECK-USE-I386-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [7 x i8*], align 4
// CHECK-USE-I386-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [7 x i8*], align 4
// CHECK-USE-I386-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [7 x i8*], align 4
// CHECK-USE-I386-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [7 x i64], align 4
// CHECK-USE-I386-NEXT: [[DOTOFFLOAD_BASEPTRS3:%.*]] = alloca [1 x i8*], align 4
// CHECK-USE-I386-NEXT: [[DOTOFFLOAD_PTRS4:%.*]] = alloca [1 x i8*], align 4
// CHECK-USE-I386-NEXT: [[DOTOFFLOAD_MAPPERS5:%.*]] = alloca [1 x i8*], align 4
// CHECK-USE-I386-NEXT: store i32 [[II:%.*]], i32* [[II_ADDR]], align 4
// CHECK-USE-I386-NEXT: [[TMP0:%.*]] = load i32, i32* [[II_ADDR]], align 4
// CHECK-USE-I386-NEXT: store i32 [[TMP0]], i32* [[A]], align 4
// CHECK-USE-I386-NEXT: [[I:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[ST1]], i32 0, i32 0
// CHECK-USE-I386-NEXT: [[J:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[ST1]], i32 0, i32 1
// CHECK-USE-I386-NEXT: [[TMP1:%.*]] = getelementptr i32, i32* [[J]], i32 1
// CHECK-USE-I386-NEXT: [[TMP2:%.*]] = bitcast i32* [[I]] to i8*
// CHECK-USE-I386-NEXT: [[TMP3:%.*]] = bitcast i32* [[TMP1]] to i8*
// CHECK-USE-I386-NEXT: [[TMP4:%.*]] = ptrtoint i8* [[TMP3]] to i64
// CHECK-USE-I386-NEXT: [[TMP5:%.*]] = ptrtoint i8* [[TMP2]] to i64
// CHECK-USE-I386-NEXT: [[TMP6:%.*]] = sub i64 [[TMP4]], [[TMP5]]
// CHECK-USE-I386-NEXT: [[TMP7:%.*]] = sdiv exact i64 [[TMP6]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
// CHECK-USE-I386-NEXT: [[I1:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[ST2]], i32 0, i32 0
// CHECK-USE-I386-NEXT: [[J2:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[ST2]], i32 0, i32 1
// CHECK-USE-I386-NEXT: [[TMP8:%.*]] = getelementptr i32, i32* [[J2]], i32 1
// CHECK-USE-I386-NEXT: [[TMP9:%.*]] = bitcast i32* [[I1]] to i8*
// CHECK-USE-I386-NEXT: [[TMP10:%.*]] = bitcast i32* [[TMP8]] to i8*
// CHECK-USE-I386-NEXT: [[TMP11:%.*]] = ptrtoint i8* [[TMP10]] to i64
// CHECK-USE-I386-NEXT: [[TMP12:%.*]] = ptrtoint i8* [[TMP9]] to i64
// CHECK-USE-I386-NEXT: [[TMP13:%.*]] = sub i64 [[TMP11]], [[TMP12]]
// CHECK-USE-I386-NEXT: [[TMP14:%.*]] = sdiv exact i64 [[TMP13]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
// CHECK-USE-I386-NEXT: [[TMP15:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK-USE-I386-NEXT: [[TMP16:%.*]] = bitcast i8** [[TMP15]] to %struct.ST**
// CHECK-USE-I386-NEXT: store %struct.ST* [[ST1]], %struct.ST** [[TMP16]], align 4
// CHECK-USE-I386-NEXT: [[TMP17:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK-USE-I386-NEXT: [[TMP18:%.*]] = bitcast i8** [[TMP17]] to i32**
// CHECK-USE-I386-NEXT: store i32* [[I]], i32** [[TMP18]], align 4
// CHECK-USE-I386-NEXT: [[TMP19:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0
// CHECK-USE-I386-NEXT: store i64 [[TMP7]], i64* [[TMP19]], align 4
// CHECK-USE-I386-NEXT: [[TMP20:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0
// CHECK-USE-I386-NEXT: store i8* null, i8** [[TMP20]], align 4
// CHECK-USE-I386-NEXT: [[TMP21:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1
// CHECK-USE-I386-NEXT: [[TMP22:%.*]] = bitcast i8** [[TMP21]] to %struct.ST**
// CHECK-USE-I386-NEXT: store %struct.ST* [[ST1]], %struct.ST** [[TMP22]], align 4
// CHECK-USE-I386-NEXT: [[TMP23:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 1
// CHECK-USE-I386-NEXT: [[TMP24:%.*]] = bitcast i8** [[TMP23]] to i32**
// CHECK-USE-I386-NEXT: store i32* [[I]], i32** [[TMP24]], align 4
// CHECK-USE-I386-NEXT: [[TMP25:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 1
// CHECK-USE-I386-NEXT: store i64 4, i64* [[TMP25]], align 4
// CHECK-USE-I386-NEXT: [[TMP26:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 1
// CHECK-USE-I386-NEXT: store i8* null, i8** [[TMP26]], align 4
// CHECK-USE-I386-NEXT: [[TMP27:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2
// CHECK-USE-I386-NEXT: [[TMP28:%.*]] = bitcast i8** [[TMP27]] to %struct.ST**
// CHECK-USE-I386-NEXT: store %struct.ST* [[ST1]], %struct.ST** [[TMP28]], align 4
// CHECK-USE-I386-NEXT: [[TMP29:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 2
// CHECK-USE-I386-NEXT: [[TMP30:%.*]] = bitcast i8** [[TMP29]] to i32**
// CHECK-USE-I386-NEXT: store i32* [[J]], i32** [[TMP30]], align 4
// CHECK-USE-I386-NEXT: [[TMP31:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 2
// CHECK-USE-I386-NEXT: store i64 4, i64* [[TMP31]], align 4
// CHECK-USE-I386-NEXT: [[TMP32:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 2
// CHECK-USE-I386-NEXT: store i8* null, i8** [[TMP32]], align 4
// CHECK-USE-I386-NEXT: [[TMP33:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 3
// CHECK-USE-I386-NEXT: [[TMP34:%.*]] = bitcast i8** [[TMP33]] to i32**
// CHECK-USE-I386-NEXT: store i32* [[A]], i32** [[TMP34]], align 4
// CHECK-USE-I386-NEXT: [[TMP35:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 3
// CHECK-USE-I386-NEXT: [[TMP36:%.*]] = bitcast i8** [[TMP35]] to i32**
// CHECK-USE-I386-NEXT: store i32* [[A]], i32** [[TMP36]], align 4
// CHECK-USE-I386-NEXT: [[TMP37:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 3
// CHECK-USE-I386-NEXT: store i64 4, i64* [[TMP37]], align 4
// CHECK-USE-I386-NEXT: [[TMP38:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 3
// CHECK-USE-I386-NEXT: store i8* null, i8** [[TMP38]], align 4
// CHECK-USE-I386-NEXT: [[TMP39:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 4
// CHECK-USE-I386-NEXT: [[TMP40:%.*]] = bitcast i8** [[TMP39]] to %struct.ST**
// CHECK-USE-I386-NEXT: store %struct.ST* [[ST2]], %struct.ST** [[TMP40]], align 4
// CHECK-USE-I386-NEXT: [[TMP41:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 4
// CHECK-USE-I386-NEXT: [[TMP42:%.*]] = bitcast i8** [[TMP41]] to i32**
// CHECK-USE-I386-NEXT: store i32* [[I1]], i32** [[TMP42]], align 4
// CHECK-USE-I386-NEXT: [[TMP43:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 4
// CHECK-USE-I386-NEXT: store i64 [[TMP14]], i64* [[TMP43]], align 4
// CHECK-USE-I386-NEXT: [[TMP44:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 4
// CHECK-USE-I386-NEXT: store i8* null, i8** [[TMP44]], align 4
// CHECK-USE-I386-NEXT: [[TMP45:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 5
// CHECK-USE-I386-NEXT: [[TMP46:%.*]] = bitcast i8** [[TMP45]] to %struct.ST**
// CHECK-USE-I386-NEXT: store %struct.ST* [[ST2]], %struct.ST** [[TMP46]], align 4
// CHECK-USE-I386-NEXT: [[TMP47:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 5
// CHECK-USE-I386-NEXT: [[TMP48:%.*]] = bitcast i8** [[TMP47]] to i32**
// CHECK-USE-I386-NEXT: store i32* [[I1]], i32** [[TMP48]], align 4
// CHECK-USE-I386-NEXT: [[TMP49:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 5
// CHECK-USE-I386-NEXT: store i64 4, i64* [[TMP49]], align 4
// CHECK-USE-I386-NEXT: [[TMP50:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 5
// CHECK-USE-I386-NEXT: store i8* null, i8** [[TMP50]], align 4
// CHECK-USE-I386-NEXT: [[TMP51:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 6
// CHECK-USE-I386-NEXT: [[TMP52:%.*]] = bitcast i8** [[TMP51]] to %struct.ST**
// CHECK-USE-I386-NEXT: store %struct.ST* [[ST2]], %struct.ST** [[TMP52]], align 4
// CHECK-USE-I386-NEXT: [[TMP53:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 6
// CHECK-USE-I386-NEXT: [[TMP54:%.*]] = bitcast i8** [[TMP53]] to i32**
// CHECK-USE-I386-NEXT: store i32* [[J2]], i32** [[TMP54]], align 4
// CHECK-USE-I386-NEXT: [[TMP55:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 6
// CHECK-USE-I386-NEXT: store i64 4, i64* [[TMP55]], align 4
// CHECK-USE-I386-NEXT: [[TMP56:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 6
// CHECK-USE-I386-NEXT: store i8* null, i8** [[TMP56]], align 4
// CHECK-USE-I386-NEXT: [[TMP57:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK-USE-I386-NEXT: [[TMP58:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK-USE-I386-NEXT: [[TMP59:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0
// CHECK-USE-I386-NEXT: [[TMP60:%.*]] = call i32 @__tgt_target_mapper(%struct.ident_t* @[[GLOB1:[0-9]+]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z20explicit_maps_singlei_l654.region_id, i32 7, i8** [[TMP57]], i8** [[TMP58]], i64* [[TMP59]], i64* getelementptr inbounds ([7 x i64], [7 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null)
// CHECK-USE-I386-NEXT: [[TMP61:%.*]] = icmp ne i32 [[TMP60]], 0
// CHECK-USE-I386-NEXT: br i1 [[TMP61]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]
// CHECK-USE-I386: omp_offload.failed:
// CHECK-USE-I386-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z20explicit_maps_singlei_l654(%struct.ST* [[ST1]], i32* [[A]], %struct.ST* [[ST2]]) #[[ATTR2:[0-9]+]]
// CHECK-USE-I386-NEXT: br label [[OMP_OFFLOAD_CONT]]
// CHECK-USE-I386: omp_offload.cont:
// CHECK-USE-I386-NEXT: [[TMP62:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0
// CHECK-USE-I386-NEXT: [[TMP63:%.*]] = bitcast i8** [[TMP62]] to i32**
// CHECK-USE-I386-NEXT: store i32* [[A]], i32** [[TMP63]], align 4
// CHECK-USE-I386-NEXT: [[TMP64:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS4]], i32 0, i32 0
// CHECK-USE-I386-NEXT: [[TMP65:%.*]] = bitcast i8** [[TMP64]] to i32**
// CHECK-USE-I386-NEXT: store i32* [[A]], i32** [[TMP65]], align 4
// CHECK-USE-I386-NEXT: [[TMP66:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_MAPPERS5]], i32 0, i32 0
// CHECK-USE-I386-NEXT: store i8* null, i8** [[TMP66]], align 4
// CHECK-USE-I386-NEXT: [[TMP67:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0
// CHECK-USE-I386-NEXT: [[TMP68:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS4]], i32 0, i32 0
// CHECK-USE-I386-NEXT: [[TMP69:%.*]] = call i32 @__tgt_target_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z20explicit_maps_singlei_l668.region_id, i32 1, i8** [[TMP67]], i8** [[TMP68]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.1, i32 0, i32 0), i8** null, i8** null)
// CHECK-USE-I386-NEXT: [[TMP70:%.*]] = icmp ne i32 [[TMP69]], 0
// CHECK-USE-I386-NEXT: br i1 [[TMP70]], label [[OMP_OFFLOAD_FAILED6:%.*]], label [[OMP_OFFLOAD_CONT7:%.*]]
// CHECK-USE-I386: omp_offload.failed6:
// CHECK-USE-I386-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z20explicit_maps_singlei_l668(i32* [[A]]) #[[ATTR2]]
// CHECK-USE-I386-NEXT: br label [[OMP_OFFLOAD_CONT7]]
// CHECK-USE-I386: omp_offload.cont7:
// CHECK-USE-I386-NEXT: ret void
//
// CHECK-NOUSE-PPC64LE-LABEL: @_Z20explicit_maps_singlei(
// CHECK-NOUSE-PPC64LE-NEXT: entry:
// CHECK-NOUSE-PPC64LE-NEXT: [[II_ADDR:%.*]] = alloca i32, align 4
// CHECK-NOUSE-PPC64LE-NEXT: [[A:%.*]] = alloca i32, align 4
// CHECK-NOUSE-PPC64LE-NEXT: [[ST1:%.*]] = alloca [[STRUCT_ST:%.*]], align 4
// CHECK-NOUSE-PPC64LE-NEXT: [[ST2:%.*]] = alloca [[STRUCT_ST]], align 4
// CHECK-NOUSE-PPC64LE-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [7 x i8*], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [7 x i8*], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [7 x i8*], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [7 x i64], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[DOTOFFLOAD_BASEPTRS3:%.*]] = alloca [1 x i8*], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[DOTOFFLOAD_PTRS4:%.*]] = alloca [1 x i8*], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[DOTOFFLOAD_MAPPERS5:%.*]] = alloca [1 x i8*], align 8
// CHECK-NOUSE-PPC64LE-NEXT: store i32 [[II:%.*]], i32* [[II_ADDR]], align 4
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP0:%.*]] = load i32, i32* [[II_ADDR]], align 4
// CHECK-NOUSE-PPC64LE-NEXT: store i32 [[TMP0]], i32* [[A]], align 4
// CHECK-NOUSE-PPC64LE-NEXT: [[I:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[ST1]], i32 0, i32 0
// CHECK-NOUSE-PPC64LE-NEXT: [[J:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[ST1]], i32 0, i32 1
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP1:%.*]] = getelementptr i32, i32* [[J]], i32 1
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP2:%.*]] = bitcast i32* [[I]] to i8*
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP3:%.*]] = bitcast i32* [[TMP1]] to i8*
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP4:%.*]] = ptrtoint i8* [[TMP3]] to i64
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP5:%.*]] = ptrtoint i8* [[TMP2]] to i64
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP6:%.*]] = sub i64 [[TMP4]], [[TMP5]]
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP7:%.*]] = sdiv exact i64 [[TMP6]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
// CHECK-NOUSE-PPC64LE-NEXT: [[I1:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[ST2]], i32 0, i32 0
// CHECK-NOUSE-PPC64LE-NEXT: [[J2:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[ST2]], i32 0, i32 1
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP8:%.*]] = getelementptr i32, i32* [[J2]], i32 1
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP9:%.*]] = bitcast i32* [[I1]] to i8*
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP10:%.*]] = bitcast i32* [[TMP8]] to i8*
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP11:%.*]] = ptrtoint i8* [[TMP10]] to i64
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP12:%.*]] = ptrtoint i8* [[TMP9]] to i64
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP13:%.*]] = sub i64 [[TMP11]], [[TMP12]]
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP14:%.*]] = sdiv exact i64 [[TMP13]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP15:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP16:%.*]] = bitcast i8** [[TMP15]] to %struct.ST**
// CHECK-NOUSE-PPC64LE-NEXT: store %struct.ST* [[ST1]], %struct.ST** [[TMP16]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP17:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP18:%.*]] = bitcast i8** [[TMP17]] to i32**
// CHECK-NOUSE-PPC64LE-NEXT: store i32* [[I]], i32** [[TMP18]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP19:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0
// CHECK-NOUSE-PPC64LE-NEXT: store i64 [[TMP7]], i64* [[TMP19]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP20:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0
// CHECK-NOUSE-PPC64LE-NEXT: store i8* null, i8** [[TMP20]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP21:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP22:%.*]] = bitcast i8** [[TMP21]] to %struct.ST**
// CHECK-NOUSE-PPC64LE-NEXT: store %struct.ST* [[ST1]], %struct.ST** [[TMP22]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP23:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 1
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP24:%.*]] = bitcast i8** [[TMP23]] to i32**
// CHECK-NOUSE-PPC64LE-NEXT: store i32* [[I]], i32** [[TMP24]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP25:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 1
// CHECK-NOUSE-PPC64LE-NEXT: store i64 4, i64* [[TMP25]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP26:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1
// CHECK-NOUSE-PPC64LE-NEXT: store i8* null, i8** [[TMP26]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP27:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP28:%.*]] = bitcast i8** [[TMP27]] to %struct.ST**
// CHECK-NOUSE-PPC64LE-NEXT: store %struct.ST* [[ST1]], %struct.ST** [[TMP28]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP29:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 2
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP30:%.*]] = bitcast i8** [[TMP29]] to i32**
// CHECK-NOUSE-PPC64LE-NEXT: store i32* [[J]], i32** [[TMP30]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP31:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 2
// CHECK-NOUSE-PPC64LE-NEXT: store i64 4, i64* [[TMP31]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP32:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2
// CHECK-NOUSE-PPC64LE-NEXT: store i8* null, i8** [[TMP32]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP33:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 3
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP34:%.*]] = bitcast i8** [[TMP33]] to i32**
// CHECK-NOUSE-PPC64LE-NEXT: store i32* [[A]], i32** [[TMP34]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP35:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 3
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP36:%.*]] = bitcast i8** [[TMP35]] to i32**
// CHECK-NOUSE-PPC64LE-NEXT: store i32* [[A]], i32** [[TMP36]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP37:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 3
// CHECK-NOUSE-PPC64LE-NEXT: store i64 4, i64* [[TMP37]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP38:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 3
// CHECK-NOUSE-PPC64LE-NEXT: store i8* null, i8** [[TMP38]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP39:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 4
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP40:%.*]] = bitcast i8** [[TMP39]] to %struct.ST**
// CHECK-NOUSE-PPC64LE-NEXT: store %struct.ST* [[ST2]], %struct.ST** [[TMP40]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP41:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 4
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP42:%.*]] = bitcast i8** [[TMP41]] to i32**
// CHECK-NOUSE-PPC64LE-NEXT: store i32* [[I1]], i32** [[TMP42]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP43:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 4
// CHECK-NOUSE-PPC64LE-NEXT: store i64 [[TMP14]], i64* [[TMP43]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP44:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 4
// CHECK-NOUSE-PPC64LE-NEXT: store i8* null, i8** [[TMP44]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP45:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 5
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP46:%.*]] = bitcast i8** [[TMP45]] to %struct.ST**
// CHECK-NOUSE-PPC64LE-NEXT: store %struct.ST* [[ST2]], %struct.ST** [[TMP46]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP47:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 5
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP48:%.*]] = bitcast i8** [[TMP47]] to i32**
// CHECK-NOUSE-PPC64LE-NEXT: store i32* [[I1]], i32** [[TMP48]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP49:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 5
// CHECK-NOUSE-PPC64LE-NEXT: store i64 4, i64* [[TMP49]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP50:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 5
// CHECK-NOUSE-PPC64LE-NEXT: store i8* null, i8** [[TMP50]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP51:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 6
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP52:%.*]] = bitcast i8** [[TMP51]] to %struct.ST**
// CHECK-NOUSE-PPC64LE-NEXT: store %struct.ST* [[ST2]], %struct.ST** [[TMP52]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP53:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 6
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP54:%.*]] = bitcast i8** [[TMP53]] to i32**
// CHECK-NOUSE-PPC64LE-NEXT: store i32* [[J2]], i32** [[TMP54]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP55:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 6
// CHECK-NOUSE-PPC64LE-NEXT: store i64 4, i64* [[TMP55]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP56:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 6
// CHECK-NOUSE-PPC64LE-NEXT: store i8* null, i8** [[TMP56]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP57:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP58:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP59:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP60:%.*]] = call i32 @__tgt_target_mapper(%struct.ident_t* @[[GLOB1:[0-9]+]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z20explicit_maps_singlei_l654.region_id, i32 7, i8** [[TMP57]], i8** [[TMP58]], i64* [[TMP59]], i64* getelementptr inbounds ([7 x i64], [7 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null)
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP61:%.*]] = icmp ne i32 [[TMP60]], 0
// CHECK-NOUSE-PPC64LE-NEXT: br i1 [[TMP61]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]
// CHECK-NOUSE-PPC64LE: omp_offload.failed:
// CHECK-NOUSE-PPC64LE-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z20explicit_maps_singlei_l654() #[[ATTR2:[0-9]+]]
// CHECK-NOUSE-PPC64LE-NEXT: br label [[OMP_OFFLOAD_CONT]]
// CHECK-NOUSE-PPC64LE: omp_offload.cont:
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP62:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP63:%.*]] = bitcast i8** [[TMP62]] to i32**
// CHECK-NOUSE-PPC64LE-NEXT: store i32* [[A]], i32** [[TMP63]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP64:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS4]], i32 0, i32 0
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP65:%.*]] = bitcast i8** [[TMP64]] to i32**
// CHECK-NOUSE-PPC64LE-NEXT: store i32* [[A]], i32** [[TMP65]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP66:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_MAPPERS5]], i64 0, i64 0
// CHECK-NOUSE-PPC64LE-NEXT: store i8* null, i8** [[TMP66]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP67:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP68:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS4]], i32 0, i32 0
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP69:%.*]] = call i32 @__tgt_target_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z20explicit_maps_singlei_l668.region_id, i32 1, i8** [[TMP67]], i8** [[TMP68]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.1, i32 0, i32 0), i8** null, i8** null)
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP70:%.*]] = icmp ne i32 [[TMP69]], 0
// CHECK-NOUSE-PPC64LE-NEXT: br i1 [[TMP70]], label [[OMP_OFFLOAD_FAILED6:%.*]], label [[OMP_OFFLOAD_CONT7:%.*]]
// CHECK-NOUSE-PPC64LE: omp_offload.failed6:
// CHECK-NOUSE-PPC64LE-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z20explicit_maps_singlei_l668() #[[ATTR2]]
// CHECK-NOUSE-PPC64LE-NEXT: br label [[OMP_OFFLOAD_CONT7]]
// CHECK-NOUSE-PPC64LE: omp_offload.cont7:
// CHECK-NOUSE-PPC64LE-NEXT: ret void
//
// CHECK-NOUSE-I386-LABEL: @_Z20explicit_maps_singlei(
// CHECK-NOUSE-I386-NEXT: entry:
// CHECK-NOUSE-I386-NEXT: [[II_ADDR:%.*]] = alloca i32, align 4
// CHECK-NOUSE-I386-NEXT: [[A:%.*]] = alloca i32, align 4
// CHECK-NOUSE-I386-NEXT: [[ST1:%.*]] = alloca [[STRUCT_ST:%.*]], align 4
// CHECK-NOUSE-I386-NEXT: [[ST2:%.*]] = alloca [[STRUCT_ST]], align 4
// CHECK-NOUSE-I386-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [7 x i8*], align 4
// CHECK-NOUSE-I386-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [7 x i8*], align 4
// CHECK-NOUSE-I386-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [7 x i8*], align 4
// CHECK-NOUSE-I386-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [7 x i64], align 4
// CHECK-NOUSE-I386-NEXT: [[DOTOFFLOAD_BASEPTRS3:%.*]] = alloca [1 x i8*], align 4
// CHECK-NOUSE-I386-NEXT: [[DOTOFFLOAD_PTRS4:%.*]] = alloca [1 x i8*], align 4
// CHECK-NOUSE-I386-NEXT: [[DOTOFFLOAD_MAPPERS5:%.*]] = alloca [1 x i8*], align 4
// CHECK-NOUSE-I386-NEXT: store i32 [[II:%.*]], i32* [[II_ADDR]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP0:%.*]] = load i32, i32* [[II_ADDR]], align 4
// CHECK-NOUSE-I386-NEXT: store i32 [[TMP0]], i32* [[A]], align 4
// CHECK-NOUSE-I386-NEXT: [[I:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[ST1]], i32 0, i32 0
// CHECK-NOUSE-I386-NEXT: [[J:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[ST1]], i32 0, i32 1
// CHECK-NOUSE-I386-NEXT: [[TMP1:%.*]] = getelementptr i32, i32* [[J]], i32 1
// CHECK-NOUSE-I386-NEXT: [[TMP2:%.*]] = bitcast i32* [[I]] to i8*
// CHECK-NOUSE-I386-NEXT: [[TMP3:%.*]] = bitcast i32* [[TMP1]] to i8*
// CHECK-NOUSE-I386-NEXT: [[TMP4:%.*]] = ptrtoint i8* [[TMP3]] to i64
// CHECK-NOUSE-I386-NEXT: [[TMP5:%.*]] = ptrtoint i8* [[TMP2]] to i64
// CHECK-NOUSE-I386-NEXT: [[TMP6:%.*]] = sub i64 [[TMP4]], [[TMP5]]
// CHECK-NOUSE-I386-NEXT: [[TMP7:%.*]] = sdiv exact i64 [[TMP6]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
// CHECK-NOUSE-I386-NEXT: [[I1:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[ST2]], i32 0, i32 0
// CHECK-NOUSE-I386-NEXT: [[J2:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[ST2]], i32 0, i32 1
// CHECK-NOUSE-I386-NEXT: [[TMP8:%.*]] = getelementptr i32, i32* [[J2]], i32 1
// CHECK-NOUSE-I386-NEXT: [[TMP9:%.*]] = bitcast i32* [[I1]] to i8*
// CHECK-NOUSE-I386-NEXT: [[TMP10:%.*]] = bitcast i32* [[TMP8]] to i8*
// CHECK-NOUSE-I386-NEXT: [[TMP11:%.*]] = ptrtoint i8* [[TMP10]] to i64
// CHECK-NOUSE-I386-NEXT: [[TMP12:%.*]] = ptrtoint i8* [[TMP9]] to i64
// CHECK-NOUSE-I386-NEXT: [[TMP13:%.*]] = sub i64 [[TMP11]], [[TMP12]]
// CHECK-NOUSE-I386-NEXT: [[TMP14:%.*]] = sdiv exact i64 [[TMP13]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
// CHECK-NOUSE-I386-NEXT: [[TMP15:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK-NOUSE-I386-NEXT: [[TMP16:%.*]] = bitcast i8** [[TMP15]] to %struct.ST**
// CHECK-NOUSE-I386-NEXT: store %struct.ST* [[ST1]], %struct.ST** [[TMP16]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP17:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK-NOUSE-I386-NEXT: [[TMP18:%.*]] = bitcast i8** [[TMP17]] to i32**
// CHECK-NOUSE-I386-NEXT: store i32* [[I]], i32** [[TMP18]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP19:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0
// CHECK-NOUSE-I386-NEXT: store i64 [[TMP7]], i64* [[TMP19]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP20:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0
// CHECK-NOUSE-I386-NEXT: store i8* null, i8** [[TMP20]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP21:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1
// CHECK-NOUSE-I386-NEXT: [[TMP22:%.*]] = bitcast i8** [[TMP21]] to %struct.ST**
// CHECK-NOUSE-I386-NEXT: store %struct.ST* [[ST1]], %struct.ST** [[TMP22]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP23:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 1
// CHECK-NOUSE-I386-NEXT: [[TMP24:%.*]] = bitcast i8** [[TMP23]] to i32**
// CHECK-NOUSE-I386-NEXT: store i32* [[I]], i32** [[TMP24]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP25:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 1
// CHECK-NOUSE-I386-NEXT: store i64 4, i64* [[TMP25]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP26:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 1
// CHECK-NOUSE-I386-NEXT: store i8* null, i8** [[TMP26]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP27:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2
// CHECK-NOUSE-I386-NEXT: [[TMP28:%.*]] = bitcast i8** [[TMP27]] to %struct.ST**
// CHECK-NOUSE-I386-NEXT: store %struct.ST* [[ST1]], %struct.ST** [[TMP28]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP29:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 2
// CHECK-NOUSE-I386-NEXT: [[TMP30:%.*]] = bitcast i8** [[TMP29]] to i32**
// CHECK-NOUSE-I386-NEXT: store i32* [[J]], i32** [[TMP30]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP31:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 2
// CHECK-NOUSE-I386-NEXT: store i64 4, i64* [[TMP31]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP32:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 2
// CHECK-NOUSE-I386-NEXT: store i8* null, i8** [[TMP32]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP33:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 3
// CHECK-NOUSE-I386-NEXT: [[TMP34:%.*]] = bitcast i8** [[TMP33]] to i32**
// CHECK-NOUSE-I386-NEXT: store i32* [[A]], i32** [[TMP34]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP35:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 3
// CHECK-NOUSE-I386-NEXT: [[TMP36:%.*]] = bitcast i8** [[TMP35]] to i32**
// CHECK-NOUSE-I386-NEXT: store i32* [[A]], i32** [[TMP36]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP37:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 3
// CHECK-NOUSE-I386-NEXT: store i64 4, i64* [[TMP37]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP38:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 3
// CHECK-NOUSE-I386-NEXT: store i8* null, i8** [[TMP38]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP39:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 4
// CHECK-NOUSE-I386-NEXT: [[TMP40:%.*]] = bitcast i8** [[TMP39]] to %struct.ST**
// CHECK-NOUSE-I386-NEXT: store %struct.ST* [[ST2]], %struct.ST** [[TMP40]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP41:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 4
// CHECK-NOUSE-I386-NEXT: [[TMP42:%.*]] = bitcast i8** [[TMP41]] to i32**
// CHECK-NOUSE-I386-NEXT: store i32* [[I1]], i32** [[TMP42]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP43:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 4
// CHECK-NOUSE-I386-NEXT: store i64 [[TMP14]], i64* [[TMP43]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP44:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 4
// CHECK-NOUSE-I386-NEXT: store i8* null, i8** [[TMP44]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP45:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 5
// CHECK-NOUSE-I386-NEXT: [[TMP46:%.*]] = bitcast i8** [[TMP45]] to %struct.ST**
// CHECK-NOUSE-I386-NEXT: store %struct.ST* [[ST2]], %struct.ST** [[TMP46]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP47:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 5
// CHECK-NOUSE-I386-NEXT: [[TMP48:%.*]] = bitcast i8** [[TMP47]] to i32**
// CHECK-NOUSE-I386-NEXT: store i32* [[I1]], i32** [[TMP48]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP49:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 5
// CHECK-NOUSE-I386-NEXT: store i64 4, i64* [[TMP49]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP50:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 5
// CHECK-NOUSE-I386-NEXT: store i8* null, i8** [[TMP50]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP51:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 6
// CHECK-NOUSE-I386-NEXT: [[TMP52:%.*]] = bitcast i8** [[TMP51]] to %struct.ST**
// CHECK-NOUSE-I386-NEXT: store %struct.ST* [[ST2]], %struct.ST** [[TMP52]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP53:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 6
// CHECK-NOUSE-I386-NEXT: [[TMP54:%.*]] = bitcast i8** [[TMP53]] to i32**
// CHECK-NOUSE-I386-NEXT: store i32* [[J2]], i32** [[TMP54]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP55:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 6
// CHECK-NOUSE-I386-NEXT: store i64 4, i64* [[TMP55]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP56:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 6
// CHECK-NOUSE-I386-NEXT: store i8* null, i8** [[TMP56]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP57:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK-NOUSE-I386-NEXT: [[TMP58:%.*]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK-NOUSE-I386-NEXT: [[TMP59:%.*]] = getelementptr inbounds [7 x i64], [7 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0
// CHECK-NOUSE-I386-NEXT: [[TMP60:%.*]] = call i32 @__tgt_target_mapper(%struct.ident_t* @[[GLOB1:[0-9]+]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z20explicit_maps_singlei_l654.region_id, i32 7, i8** [[TMP57]], i8** [[TMP58]], i64* [[TMP59]], i64* getelementptr inbounds ([7 x i64], [7 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null)
// CHECK-NOUSE-I386-NEXT: [[TMP61:%.*]] = icmp ne i32 [[TMP60]], 0
// CHECK-NOUSE-I386-NEXT: br i1 [[TMP61]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]
// CHECK-NOUSE-I386: omp_offload.failed:
// CHECK-NOUSE-I386-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z20explicit_maps_singlei_l654() #[[ATTR2:[0-9]+]]
// CHECK-NOUSE-I386-NEXT: br label [[OMP_OFFLOAD_CONT]]
// CHECK-NOUSE-I386: omp_offload.cont:
// CHECK-NOUSE-I386-NEXT: [[TMP62:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0
// CHECK-NOUSE-I386-NEXT: [[TMP63:%.*]] = bitcast i8** [[TMP62]] to i32**
// CHECK-NOUSE-I386-NEXT: store i32* [[A]], i32** [[TMP63]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP64:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS4]], i32 0, i32 0
// CHECK-NOUSE-I386-NEXT: [[TMP65:%.*]] = bitcast i8** [[TMP64]] to i32**
// CHECK-NOUSE-I386-NEXT: store i32* [[A]], i32** [[TMP65]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP66:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_MAPPERS5]], i32 0, i32 0
// CHECK-NOUSE-I386-NEXT: store i8* null, i8** [[TMP66]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP67:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS3]], i32 0, i32 0
// CHECK-NOUSE-I386-NEXT: [[TMP68:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS4]], i32 0, i32 0
// CHECK-NOUSE-I386-NEXT: [[TMP69:%.*]] = call i32 @__tgt_target_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z20explicit_maps_singlei_l668.region_id, i32 1, i8** [[TMP67]], i8** [[TMP68]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.1, i32 0, i32 0), i8** null, i8** null)
// CHECK-NOUSE-I386-NEXT: [[TMP70:%.*]] = icmp ne i32 [[TMP69]], 0
// CHECK-NOUSE-I386-NEXT: br i1 [[TMP70]], label [[OMP_OFFLOAD_FAILED6:%.*]], label [[OMP_OFFLOAD_CONT7:%.*]]
// CHECK-NOUSE-I386: omp_offload.failed6:
// CHECK-NOUSE-I386-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z20explicit_maps_singlei_l668() #[[ATTR2]]
// CHECK-NOUSE-I386-NEXT: br label [[OMP_OFFLOAD_CONT7]]
// CHECK-NOUSE-I386: omp_offload.cont7:
// CHECK-NOUSE-I386-NEXT: ret void
//
void explicit_maps_single(int ii) {
// Map of a scalar.
int a = ii;
struct ST st1;
struct ST st2;
// Make sure the struct picks up ompx_hold even if another element of the
// struct doesn't have ompx_hold.
#pragma omp target map(tofrom : st1.i) \
map(ompx_hold, tofrom : a, st1.j, st2.i) \
map(tofrom : st2.j)
{
#ifdef USE
st1.i++;
a++;
st1.j++;
st2.i++;
st2.j++;
#endif
}
// Always Close Hold.
#pragma omp target map(always close ompx_hold tofrom: a)
{
#ifdef USE
a++;
#endif
}
}
// CHECK-USE-PPC64LE-LABEL: @_ZN2ST20test_present_membersEv(
// CHECK-USE-PPC64LE-NEXT: entry:
// CHECK-USE-PPC64LE-NEXT: [[THIS_ADDR:%.*]] = alloca %struct.ST*, align 8
// CHECK-USE-PPC64LE-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x i8*], align 8
// CHECK-USE-PPC64LE-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x i8*], align 8
// CHECK-USE-PPC64LE-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x i8*], align 8
// CHECK-USE-PPC64LE-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [3 x i64], align 8
// CHECK-USE-PPC64LE-NEXT: store %struct.ST* [[THIS:%.*]], %struct.ST** [[THIS_ADDR]], align 8
// CHECK-USE-PPC64LE-NEXT: [[THIS1:%.*]] = load %struct.ST*, %struct.ST** [[THIS_ADDR]], align 8
// CHECK-USE-PPC64LE-NEXT: [[I:%.*]] = getelementptr inbounds [[STRUCT_ST:%.*]], %struct.ST* [[THIS1]], i32 0, i32 0
// CHECK-USE-PPC64LE-NEXT: [[J:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[THIS1]], i32 0, i32 1
// CHECK-USE-PPC64LE-NEXT: [[TMP0:%.*]] = getelementptr i32, i32* [[J]], i32 1
// CHECK-USE-PPC64LE-NEXT: [[TMP1:%.*]] = bitcast i32* [[I]] to i8*
// CHECK-USE-PPC64LE-NEXT: [[TMP2:%.*]] = bitcast i32* [[TMP0]] to i8*
// CHECK-USE-PPC64LE-NEXT: [[TMP3:%.*]] = ptrtoint i8* [[TMP2]] to i64
// CHECK-USE-PPC64LE-NEXT: [[TMP4:%.*]] = ptrtoint i8* [[TMP1]] to i64
// CHECK-USE-PPC64LE-NEXT: [[TMP5:%.*]] = sub i64 [[TMP3]], [[TMP4]]
// CHECK-USE-PPC64LE-NEXT: [[TMP6:%.*]] = sdiv exact i64 [[TMP5]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
// CHECK-USE-PPC64LE-NEXT: [[TMP7:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK-USE-PPC64LE-NEXT: [[TMP8:%.*]] = bitcast i8** [[TMP7]] to %struct.ST**
// CHECK-USE-PPC64LE-NEXT: store %struct.ST* [[THIS1]], %struct.ST** [[TMP8]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP9:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK-USE-PPC64LE-NEXT: [[TMP10:%.*]] = bitcast i8** [[TMP9]] to i32**
// CHECK-USE-PPC64LE-NEXT: store i32* [[I]], i32** [[TMP10]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP11:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0
// CHECK-USE-PPC64LE-NEXT: store i64 [[TMP6]], i64* [[TMP11]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP12:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0
// CHECK-USE-PPC64LE-NEXT: store i8* null, i8** [[TMP12]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP13:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1
// CHECK-USE-PPC64LE-NEXT: [[TMP14:%.*]] = bitcast i8** [[TMP13]] to %struct.ST**
// CHECK-USE-PPC64LE-NEXT: store %struct.ST* [[THIS1]], %struct.ST** [[TMP14]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP15:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 1
// CHECK-USE-PPC64LE-NEXT: [[TMP16:%.*]] = bitcast i8** [[TMP15]] to i32**
// CHECK-USE-PPC64LE-NEXT: store i32* [[I]], i32** [[TMP16]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP17:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 1
// CHECK-USE-PPC64LE-NEXT: store i64 4, i64* [[TMP17]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP18:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1
// CHECK-USE-PPC64LE-NEXT: store i8* null, i8** [[TMP18]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP19:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2
// CHECK-USE-PPC64LE-NEXT: [[TMP20:%.*]] = bitcast i8** [[TMP19]] to %struct.ST**
// CHECK-USE-PPC64LE-NEXT: store %struct.ST* [[THIS1]], %struct.ST** [[TMP20]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP21:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 2
// CHECK-USE-PPC64LE-NEXT: [[TMP22:%.*]] = bitcast i8** [[TMP21]] to i32**
// CHECK-USE-PPC64LE-NEXT: store i32* [[J]], i32** [[TMP22]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP23:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 2
// CHECK-USE-PPC64LE-NEXT: store i64 4, i64* [[TMP23]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP24:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2
// CHECK-USE-PPC64LE-NEXT: store i8* null, i8** [[TMP24]], align 8
// CHECK-USE-PPC64LE-NEXT: [[TMP25:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK-USE-PPC64LE-NEXT: [[TMP26:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK-USE-PPC64LE-NEXT: [[TMP27:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0
// CHECK-USE-PPC64LE-NEXT: [[TMP28:%.*]] = call i32 @__tgt_target_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2ST20test_present_membersEv_l919.region_id, i32 3, i8** [[TMP25]], i8** [[TMP26]], i64* [[TMP27]], i64* getelementptr inbounds ([3 x i64], [3 x i64]* @.offload_maptypes.2, i32 0, i32 0), i8** null, i8** null)
// CHECK-USE-PPC64LE-NEXT: [[TMP29:%.*]] = icmp ne i32 [[TMP28]], 0
// CHECK-USE-PPC64LE-NEXT: br i1 [[TMP29]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]
// CHECK-USE-PPC64LE: omp_offload.failed:
// CHECK-USE-PPC64LE-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2ST20test_present_membersEv_l919(%struct.ST* [[THIS1]]) #[[ATTR2]]
// CHECK-USE-PPC64LE-NEXT: br label [[OMP_OFFLOAD_CONT]]
// CHECK-USE-PPC64LE: omp_offload.cont:
// CHECK-USE-PPC64LE-NEXT: ret void
//
// CHECK-USE-I386-LABEL: @_ZN2ST20test_present_membersEv(
// CHECK-USE-I386-NEXT: entry:
// CHECK-USE-I386-NEXT: [[THIS_ADDR:%.*]] = alloca %struct.ST*, align 4
// CHECK-USE-I386-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x i8*], align 4
// CHECK-USE-I386-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x i8*], align 4
// CHECK-USE-I386-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x i8*], align 4
// CHECK-USE-I386-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [3 x i64], align 4
// CHECK-USE-I386-NEXT: store %struct.ST* [[THIS:%.*]], %struct.ST** [[THIS_ADDR]], align 4
// CHECK-USE-I386-NEXT: [[THIS1:%.*]] = load %struct.ST*, %struct.ST** [[THIS_ADDR]], align 4
// CHECK-USE-I386-NEXT: [[I:%.*]] = getelementptr inbounds [[STRUCT_ST:%.*]], %struct.ST* [[THIS1]], i32 0, i32 0
// CHECK-USE-I386-NEXT: [[J:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[THIS1]], i32 0, i32 1
// CHECK-USE-I386-NEXT: [[TMP0:%.*]] = getelementptr i32, i32* [[J]], i32 1
// CHECK-USE-I386-NEXT: [[TMP1:%.*]] = bitcast i32* [[I]] to i8*
// CHECK-USE-I386-NEXT: [[TMP2:%.*]] = bitcast i32* [[TMP0]] to i8*
// CHECK-USE-I386-NEXT: [[TMP3:%.*]] = ptrtoint i8* [[TMP2]] to i64
// CHECK-USE-I386-NEXT: [[TMP4:%.*]] = ptrtoint i8* [[TMP1]] to i64
// CHECK-USE-I386-NEXT: [[TMP5:%.*]] = sub i64 [[TMP3]], [[TMP4]]
// CHECK-USE-I386-NEXT: [[TMP6:%.*]] = sdiv exact i64 [[TMP5]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
// CHECK-USE-I386-NEXT: [[TMP7:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK-USE-I386-NEXT: [[TMP8:%.*]] = bitcast i8** [[TMP7]] to %struct.ST**
// CHECK-USE-I386-NEXT: store %struct.ST* [[THIS1]], %struct.ST** [[TMP8]], align 4
// CHECK-USE-I386-NEXT: [[TMP9:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK-USE-I386-NEXT: [[TMP10:%.*]] = bitcast i8** [[TMP9]] to i32**
// CHECK-USE-I386-NEXT: store i32* [[I]], i32** [[TMP10]], align 4
// CHECK-USE-I386-NEXT: [[TMP11:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0
// CHECK-USE-I386-NEXT: store i64 [[TMP6]], i64* [[TMP11]], align 4
// CHECK-USE-I386-NEXT: [[TMP12:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0
// CHECK-USE-I386-NEXT: store i8* null, i8** [[TMP12]], align 4
// CHECK-USE-I386-NEXT: [[TMP13:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1
// CHECK-USE-I386-NEXT: [[TMP14:%.*]] = bitcast i8** [[TMP13]] to %struct.ST**
// CHECK-USE-I386-NEXT: store %struct.ST* [[THIS1]], %struct.ST** [[TMP14]], align 4
// CHECK-USE-I386-NEXT: [[TMP15:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 1
// CHECK-USE-I386-NEXT: [[TMP16:%.*]] = bitcast i8** [[TMP15]] to i32**
// CHECK-USE-I386-NEXT: store i32* [[I]], i32** [[TMP16]], align 4
// CHECK-USE-I386-NEXT: [[TMP17:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 1
// CHECK-USE-I386-NEXT: store i64 4, i64* [[TMP17]], align 4
// CHECK-USE-I386-NEXT: [[TMP18:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 1
// CHECK-USE-I386-NEXT: store i8* null, i8** [[TMP18]], align 4
// CHECK-USE-I386-NEXT: [[TMP19:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2
// CHECK-USE-I386-NEXT: [[TMP20:%.*]] = bitcast i8** [[TMP19]] to %struct.ST**
// CHECK-USE-I386-NEXT: store %struct.ST* [[THIS1]], %struct.ST** [[TMP20]], align 4
// CHECK-USE-I386-NEXT: [[TMP21:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 2
// CHECK-USE-I386-NEXT: [[TMP22:%.*]] = bitcast i8** [[TMP21]] to i32**
// CHECK-USE-I386-NEXT: store i32* [[J]], i32** [[TMP22]], align 4
// CHECK-USE-I386-NEXT: [[TMP23:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 2
// CHECK-USE-I386-NEXT: store i64 4, i64* [[TMP23]], align 4
// CHECK-USE-I386-NEXT: [[TMP24:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 2
// CHECK-USE-I386-NEXT: store i8* null, i8** [[TMP24]], align 4
// CHECK-USE-I386-NEXT: [[TMP25:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK-USE-I386-NEXT: [[TMP26:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK-USE-I386-NEXT: [[TMP27:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0
// CHECK-USE-I386-NEXT: [[TMP28:%.*]] = call i32 @__tgt_target_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2ST20test_present_membersEv_l919.region_id, i32 3, i8** [[TMP25]], i8** [[TMP26]], i64* [[TMP27]], i64* getelementptr inbounds ([3 x i64], [3 x i64]* @.offload_maptypes.2, i32 0, i32 0), i8** null, i8** null)
// CHECK-USE-I386-NEXT: [[TMP29:%.*]] = icmp ne i32 [[TMP28]], 0
// CHECK-USE-I386-NEXT: br i1 [[TMP29]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]
// CHECK-USE-I386: omp_offload.failed:
// CHECK-USE-I386-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2ST20test_present_membersEv_l919(%struct.ST* [[THIS1]]) #[[ATTR2]]
// CHECK-USE-I386-NEXT: br label [[OMP_OFFLOAD_CONT]]
// CHECK-USE-I386: omp_offload.cont:
// CHECK-USE-I386-NEXT: ret void
//
// CHECK-NOUSE-PPC64LE-LABEL: @_ZN2ST20test_present_membersEv(
// CHECK-NOUSE-PPC64LE-NEXT: entry:
// CHECK-NOUSE-PPC64LE-NEXT: [[THIS_ADDR:%.*]] = alloca %struct.ST*, align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x i8*], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x i8*], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x i8*], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [3 x i64], align 8
// CHECK-NOUSE-PPC64LE-NEXT: store %struct.ST* [[THIS:%.*]], %struct.ST** [[THIS_ADDR]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[THIS1:%.*]] = load %struct.ST*, %struct.ST** [[THIS_ADDR]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[I:%.*]] = getelementptr inbounds [[STRUCT_ST:%.*]], %struct.ST* [[THIS1]], i32 0, i32 0
// CHECK-NOUSE-PPC64LE-NEXT: [[J:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[THIS1]], i32 0, i32 1
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP0:%.*]] = getelementptr i32, i32* [[J]], i32 1
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP1:%.*]] = bitcast i32* [[I]] to i8*
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP2:%.*]] = bitcast i32* [[TMP0]] to i8*
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP3:%.*]] = ptrtoint i8* [[TMP2]] to i64
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP4:%.*]] = ptrtoint i8* [[TMP1]] to i64
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP5:%.*]] = sub i64 [[TMP3]], [[TMP4]]
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP6:%.*]] = sdiv exact i64 [[TMP5]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP7:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP8:%.*]] = bitcast i8** [[TMP7]] to %struct.ST**
// CHECK-NOUSE-PPC64LE-NEXT: store %struct.ST* [[THIS1]], %struct.ST** [[TMP8]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP9:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP10:%.*]] = bitcast i8** [[TMP9]] to i32**
// CHECK-NOUSE-PPC64LE-NEXT: store i32* [[I]], i32** [[TMP10]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP11:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0
// CHECK-NOUSE-PPC64LE-NEXT: store i64 [[TMP6]], i64* [[TMP11]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP12:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0
// CHECK-NOUSE-PPC64LE-NEXT: store i8* null, i8** [[TMP12]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP13:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP14:%.*]] = bitcast i8** [[TMP13]] to %struct.ST**
// CHECK-NOUSE-PPC64LE-NEXT: store %struct.ST* [[THIS1]], %struct.ST** [[TMP14]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP15:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 1
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP16:%.*]] = bitcast i8** [[TMP15]] to i32**
// CHECK-NOUSE-PPC64LE-NEXT: store i32* [[I]], i32** [[TMP16]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP17:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 1
// CHECK-NOUSE-PPC64LE-NEXT: store i64 4, i64* [[TMP17]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP18:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1
// CHECK-NOUSE-PPC64LE-NEXT: store i8* null, i8** [[TMP18]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP19:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP20:%.*]] = bitcast i8** [[TMP19]] to %struct.ST**
// CHECK-NOUSE-PPC64LE-NEXT: store %struct.ST* [[THIS1]], %struct.ST** [[TMP20]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP21:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 2
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP22:%.*]] = bitcast i8** [[TMP21]] to i32**
// CHECK-NOUSE-PPC64LE-NEXT: store i32* [[J]], i32** [[TMP22]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP23:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 2
// CHECK-NOUSE-PPC64LE-NEXT: store i64 4, i64* [[TMP23]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP24:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2
// CHECK-NOUSE-PPC64LE-NEXT: store i8* null, i8** [[TMP24]], align 8
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP25:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP26:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP27:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP28:%.*]] = call i32 @__tgt_target_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2ST20test_present_membersEv_l919.region_id, i32 3, i8** [[TMP25]], i8** [[TMP26]], i64* [[TMP27]], i64* getelementptr inbounds ([3 x i64], [3 x i64]* @.offload_maptypes.2, i32 0, i32 0), i8** null, i8** null)
// CHECK-NOUSE-PPC64LE-NEXT: [[TMP29:%.*]] = icmp ne i32 [[TMP28]], 0
// CHECK-NOUSE-PPC64LE-NEXT: br i1 [[TMP29]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]
// CHECK-NOUSE-PPC64LE: omp_offload.failed:
// CHECK-NOUSE-PPC64LE-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2ST20test_present_membersEv_l919() #[[ATTR2]]
// CHECK-NOUSE-PPC64LE-NEXT: br label [[OMP_OFFLOAD_CONT]]
// CHECK-NOUSE-PPC64LE: omp_offload.cont:
// CHECK-NOUSE-PPC64LE-NEXT: ret void
//
// CHECK-NOUSE-I386-LABEL: @_ZN2ST20test_present_membersEv(
// CHECK-NOUSE-I386-NEXT: entry:
// CHECK-NOUSE-I386-NEXT: [[THIS_ADDR:%.*]] = alloca %struct.ST*, align 4
// CHECK-NOUSE-I386-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [3 x i8*], align 4
// CHECK-NOUSE-I386-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [3 x i8*], align 4
// CHECK-NOUSE-I386-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [3 x i8*], align 4
// CHECK-NOUSE-I386-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [3 x i64], align 4
// CHECK-NOUSE-I386-NEXT: store %struct.ST* [[THIS:%.*]], %struct.ST** [[THIS_ADDR]], align 4
// CHECK-NOUSE-I386-NEXT: [[THIS1:%.*]] = load %struct.ST*, %struct.ST** [[THIS_ADDR]], align 4
// CHECK-NOUSE-I386-NEXT: [[I:%.*]] = getelementptr inbounds [[STRUCT_ST:%.*]], %struct.ST* [[THIS1]], i32 0, i32 0
// CHECK-NOUSE-I386-NEXT: [[J:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[THIS1]], i32 0, i32 1
// CHECK-NOUSE-I386-NEXT: [[TMP0:%.*]] = getelementptr i32, i32* [[J]], i32 1
// CHECK-NOUSE-I386-NEXT: [[TMP1:%.*]] = bitcast i32* [[I]] to i8*
// CHECK-NOUSE-I386-NEXT: [[TMP2:%.*]] = bitcast i32* [[TMP0]] to i8*
// CHECK-NOUSE-I386-NEXT: [[TMP3:%.*]] = ptrtoint i8* [[TMP2]] to i64
// CHECK-NOUSE-I386-NEXT: [[TMP4:%.*]] = ptrtoint i8* [[TMP1]] to i64
// CHECK-NOUSE-I386-NEXT: [[TMP5:%.*]] = sub i64 [[TMP3]], [[TMP4]]
// CHECK-NOUSE-I386-NEXT: [[TMP6:%.*]] = sdiv exact i64 [[TMP5]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
// CHECK-NOUSE-I386-NEXT: [[TMP7:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK-NOUSE-I386-NEXT: [[TMP8:%.*]] = bitcast i8** [[TMP7]] to %struct.ST**
// CHECK-NOUSE-I386-NEXT: store %struct.ST* [[THIS1]], %struct.ST** [[TMP8]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP9:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK-NOUSE-I386-NEXT: [[TMP10:%.*]] = bitcast i8** [[TMP9]] to i32**
// CHECK-NOUSE-I386-NEXT: store i32* [[I]], i32** [[TMP10]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP11:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0
// CHECK-NOUSE-I386-NEXT: store i64 [[TMP6]], i64* [[TMP11]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP12:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0
// CHECK-NOUSE-I386-NEXT: store i8* null, i8** [[TMP12]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP13:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1
// CHECK-NOUSE-I386-NEXT: [[TMP14:%.*]] = bitcast i8** [[TMP13]] to %struct.ST**
// CHECK-NOUSE-I386-NEXT: store %struct.ST* [[THIS1]], %struct.ST** [[TMP14]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP15:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 1
// CHECK-NOUSE-I386-NEXT: [[TMP16:%.*]] = bitcast i8** [[TMP15]] to i32**
// CHECK-NOUSE-I386-NEXT: store i32* [[I]], i32** [[TMP16]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP17:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 1
// CHECK-NOUSE-I386-NEXT: store i64 4, i64* [[TMP17]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP18:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 1
// CHECK-NOUSE-I386-NEXT: store i8* null, i8** [[TMP18]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP19:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2
// CHECK-NOUSE-I386-NEXT: [[TMP20:%.*]] = bitcast i8** [[TMP19]] to %struct.ST**
// CHECK-NOUSE-I386-NEXT: store %struct.ST* [[THIS1]], %struct.ST** [[TMP20]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP21:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 2
// CHECK-NOUSE-I386-NEXT: [[TMP22:%.*]] = bitcast i8** [[TMP21]] to i32**
// CHECK-NOUSE-I386-NEXT: store i32* [[J]], i32** [[TMP22]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP23:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 2
// CHECK-NOUSE-I386-NEXT: store i64 4, i64* [[TMP23]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP24:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 2
// CHECK-NOUSE-I386-NEXT: store i8* null, i8** [[TMP24]], align 4
// CHECK-NOUSE-I386-NEXT: [[TMP25:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK-NOUSE-I386-NEXT: [[TMP26:%.*]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK-NOUSE-I386-NEXT: [[TMP27:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[DOTOFFLOAD_SIZES]], i32 0, i32 0
// CHECK-NOUSE-I386-NEXT: [[TMP28:%.*]] = call i32 @__tgt_target_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2ST20test_present_membersEv_l919.region_id, i32 3, i8** [[TMP25]], i8** [[TMP26]], i64* [[TMP27]], i64* getelementptr inbounds ([3 x i64], [3 x i64]* @.offload_maptypes.2, i32 0, i32 0), i8** null, i8** null)
// CHECK-NOUSE-I386-NEXT: [[TMP29:%.*]] = icmp ne i32 [[TMP28]], 0
// CHECK-NOUSE-I386-NEXT: br i1 [[TMP29]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]
// CHECK-NOUSE-I386: omp_offload.failed:
// CHECK-NOUSE-I386-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__ZN2ST20test_present_membersEv_l919() #[[ATTR2]]
// CHECK-NOUSE-I386-NEXT: br label [[OMP_OFFLOAD_CONT]]
// CHECK-NOUSE-I386: omp_offload.cont:
// CHECK-NOUSE-I386-NEXT: ret void
//
void ST::test_present_members() {
// Make sure the struct picks up ompx_hold even if another element of the
// struct doesn't have ompx_hold.
#pragma omp target map(tofrom : i) map(ompx_hold, tofrom : j)
{
#ifdef USE
i++;
j++;
#endif
}
}
#endif

View File

@ -1,16 +1,36 @@
// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp -fopenmp-version=40 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp -fopenmp-version=45 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp -fopenmp-version=50 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp -fopenmp-version=51 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
// RUN: %clang_cc1 -DCCODE -verify -fopenmp -ferror-limit 200 -x c %s -Wno-openmp -Wuninitialized
// -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
// 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-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;
@ -118,38 +138,72 @@ struct SA {
{}
#pragma omp target map(close) // expected-error {{use of undeclared identifier 'close'}}
{}
// lt51-error@+1 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
// lt51-error@+1 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}}
#pragma omp target map(present, tofrom: c,f)
{}
// lt51-error@+1 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
// lt51-error@+1 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}}
#pragma omp target map(present, tofrom: c[1:2],f)
{}
// lt51-error@+1 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
// lt51-error@+1 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}}
#pragma omp target map(present, tofrom: c,f[1:2])
{}
// expected-error@+2 {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
// lt51-error@+1 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
// lt51-error@+1 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}}
#pragma omp target map(present, tofrom: c[:],f)
{}
// expected-error@+2 {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
// lt51-error@+1 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
// lt51-error@+1 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}}
#pragma omp target map(present, tofrom: c,f[:])
{}
// expected-error@+1 {{use of undeclared identifier 'present'}}
#pragma omp target map(present)
{}
// ge51-omp-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}}
// lt51-omp-error@+1 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}}
#pragma omp target map(ompx_hold, tofrom: c,f)
{}
// ge51-omp-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}}
// lt51-omp-error@+1 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}}
#pragma omp target map(ompx_hold, tofrom: c[1:2],f)
{}
// ge51-omp-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}}
// lt51-omp-error@+1 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}}
#pragma omp target map(ompx_hold, tofrom: c,f[1:2])
{}
// expected-error@+3 {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
// ge51-omp-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}}
// lt51-omp-error@+1 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}}
#pragma omp target map(ompx_hold, tofrom: c[:],f)
{}
// expected-error@+3 {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
// ge51-omp-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}}
// lt51-omp-error@+1 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}}
#pragma omp target map(ompx_hold, tofrom: c,f[:])
{}
// expected-error@+1 {{use of undeclared identifier 'ompx_hold'}}
#pragma omp target map(ompx_hold)
{}
#pragma omp target map(close, close, tofrom: a) // expected-error {{same map type modifier has been specified more than once}}
{}
#pragma omp target map(always, close, always, close, tofrom: a) // expected-error 2 {{same map type modifier has been specified more than once}}
{}
// ge51-error@+2 {{same map type modifier has been specified more than once}}
// lt51-error@+1 2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
// lt51-error@+1 2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}}
#pragma omp target map(present, present, tofrom: a)
{}
// expected-error@+3 2 {{same map type modifier has been specified more than once}}
// ge51-error@+2 1 {{same map type modifier has been specified more than once}}
// lt51-error@+1 2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
#pragma omp target map(always, close, present, always, close, present, tofrom: a)
// ompx-error@+3 {{same map type modifier has been specified more than once}}
// ge51-omp-error@+2 2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}}
// lt51-omp-error@+1 2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}}
#pragma omp target map(ompx_hold, ompx_hold, tofrom: a)
{}
// expected-error@+7 2 {{same map type modifier has been specified more than once}}
// ge51-error@+6 {{same map type modifier has been specified more than once}}
// lt51-ompx-error@+5 2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'ompx_hold'}}
// lt51-omp-error@+4 2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}}
// ompx-error@+3 {{same map type modifier has been specified more than once}}
// ge51-omp-error@+2 2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}}
// lt51-omp-error@+1 2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}}
#pragma omp target map(always, close, present, ompx_hold, always, close, present, ompx_hold, tofrom: a)
{}
#pragma omp target map( , tofrom: a) // expected-error {{missing map type modifier}}
{}
@ -157,14 +211,14 @@ struct SA {
{}
#pragma omp target map( , , : a) // expected-error {{missing map type modifier}} expected-error {{missing map type modifier}} expected-error {{missing map type}}
{}
// ge51-error@+3 2 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
// lt51-error@+2 2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
// ge51-error@+3 2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}}
// lt51-error@+2 2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}}
// expected-error@+1 {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
#pragma omp target map( d, f, bf: a)
{}
// expected-error@+4 {{missing map type modifier}}
// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}}
// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}}
// expected-error@+1 {{missing map type}}
#pragma omp target map( , f, : a)
{}
@ -172,13 +226,13 @@ struct SA {
{}
#pragma omp target map(always close bf: a) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
{}
// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}}
// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}}
// expected-error@+1 {{missing map type}}
#pragma omp target map(always tofrom close: a)
{}
// ge51-error@+2 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
// lt51-error@+1 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
// ge51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}}
// lt51-error@+1 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}}
#pragma omp target map(tofrom from: a)
{}
#pragma omp target map(close bf: a) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
@ -600,8 +654,8 @@ T tmain(T argc) {
#pragma omp target data map(always, tofrom: x)
#pragma omp target data map(always: x) // expected-error {{missing map type}}
// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}}
// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}}
// expected-error@+1 {{missing map type}}
#pragma omp target data map(tofrom, always: x)
#pragma omp target data map(always, tofrom: always, tofrom, x)
@ -610,24 +664,24 @@ T tmain(T argc) {
#pragma omp target data map(close, tofrom: x)
#pragma omp target data map(close: x) // expected-error {{missing map type}}
// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}}
// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}}
// expected-error@+1 {{missing map type}}
#pragma omp target data map(tofrom, close: x)
#pragma omp target data map(close, tofrom: close, tofrom, x)
foo();
// lt51-error@+1 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
// lt51-error@+1 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}}
#pragma omp target data map(present, tofrom: x)
// ge51-error@+2 {{missing map type}}
// lt51-error@+1 {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
#pragma omp target data map(present: x)
// ge51-error@+4 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
// lt51-error@+3 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
// ge51-error@+4 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}}
// lt51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}}
// ge51-error@+2 {{missing map type}}
// lt51-error@+1 {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
#pragma omp target data map(tofrom, present: x)
// lt51-error@+1 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
// lt51-error@+1 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}}
#pragma omp target data map(present, tofrom: present, tofrom, x)
foo();
@ -735,8 +789,8 @@ int main(int argc, char **argv) {
#pragma omp target data map(always, tofrom: x)
#pragma omp target data map(always: x) // expected-error {{missing map type}}
// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}}
// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}}
// expected-error@+1 {{missing map type}}
#pragma omp target data map(tofrom, always: x)
#pragma omp target data map(always, tofrom: always, tofrom, x)
@ -744,18 +798,18 @@ int main(int argc, char **argv) {
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)

View File

@ -1,10 +1,10 @@
// 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() {
}
@ -182,8 +182,8 @@ T tmain(T argc) {
for (i = 0; i < argc; ++i) foo();
#pragma omp target parallel for map(always: x) // expected-error {{missing map type}}
for (i = 0; i < argc; ++i) foo();
// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}}
// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}}
// expected-error@+1 {{missing map type}}
#pragma omp target parallel for map(tofrom, always: x)
for (i = 0; i < argc; ++i) foo();
@ -300,8 +300,8 @@ int main(int argc, char **argv) {
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();

View File

@ -1,10 +1,10 @@
// 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() {
}
@ -182,8 +182,8 @@ T tmain(T argc) {
for (i = 0; i < argc; ++i) foo();
#pragma omp target parallel for simd map(always: x) // expected-error {{missing map type}}
for (i = 0; i < argc; ++i) foo();
// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}}
// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}}
// expected-error@+1 {{missing map type}}
#pragma omp target parallel for simd map(tofrom, always: x)
for (i = 0; i < argc; ++i) foo();
@ -300,8 +300,8 @@ int main(int argc, char **argv) {
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();

View File

@ -1,10 +1,10 @@
// 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() {
}
@ -181,8 +181,8 @@ T tmain(T argc) {
foo();
#pragma omp target parallel map(always: x) // expected-error {{missing map type}}
foo();
// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}}
// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}}
// expected-error@+1 {{missing map type}}
#pragma omp target parallel map(tofrom, always: x)
foo();
@ -296,8 +296,8 @@ int main(int argc, char **argv) {
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();

View File

@ -176,8 +176,8 @@ T tmain(T argc) {
for (i = 0; i < argc; ++i) foo();
#pragma omp target simd map(always: x) // expected-error {{missing map type}}
for (i = 0; i < argc; ++i) foo();
// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}}
// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}}
// expected-error@+1 {{missing map type}}
#pragma omp target simd map(tofrom, always: x)
for (i = 0; i < argc; ++i) foo();
@ -293,8 +293,8 @@ int main(int argc, char **argv) {
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();

View File

@ -182,8 +182,8 @@ T tmain(T argc) {
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute map(always: x) // expected-error {{missing map type}}
for (i = 0; i < argc; ++i) foo();
// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}}
// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}}
// expected-error@+1 {{missing map type}}
#pragma omp target teams distribute map(tofrom, always: x)
for (i = 0; i < argc; ++i) foo();
@ -300,8 +300,8 @@ int main(int argc, char **argv) {
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();

View File

@ -1,8 +1,8 @@
// 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() {
}
@ -180,8 +180,8 @@ T tmain(T argc) {
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute parallel for map(always: x) // expected-error {{missing map type}}
for (i = 0; i < argc; ++i) foo();
// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}}
// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}}
// expected-error@+1 {{missing map type}}
#pragma omp target teams distribute parallel for map(tofrom, always: x)
for (i = 0; i < argc; ++i) foo();
@ -302,8 +302,8 @@ int main(int argc, char **argv) {
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();

View File

@ -1,10 +1,10 @@
// 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() {
}
@ -182,8 +182,8 @@ T tmain(T argc) {
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute parallel for simd map(always: x) // expected-error {{missing map type}}
for (i = 0; i < argc; ++i) foo();
// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}}
// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}}
// expected-error@+1 {{missing map type}}
#pragma omp target teams distribute parallel for simd map(tofrom, always: x)
for (i = 0; i < argc; ++i) foo();
@ -299,8 +299,8 @@ int main(int argc, char **argv) {
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();

View File

@ -1,10 +1,10 @@
// 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() {
}
@ -182,8 +182,8 @@ T tmain(T argc) {
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute simd map(always: x) // expected-error {{missing map type}}
for (i = 0; i < argc; ++i) foo();
// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}}
// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}}
// expected-error@+1 {{missing map type}}
#pragma omp target teams distribute simd map(tofrom, always: x)
for (i = 0; i < argc; ++i) foo();
@ -300,8 +300,8 @@ int main(int argc, char **argv) {
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();

View File

@ -1,11 +1,11 @@
// 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 -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 -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-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;
@ -479,8 +479,8 @@ T tmain(T argc) {
#pragma omp target data map(always, tofrom: x)
#pragma omp target data map(always: x) // expected-error {{missing map type}}
// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}}
// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}}
// expected-error@+1 {{missing map type}}
#pragma omp target data map(tofrom, always: x)
#pragma omp target data map(always, tofrom: always, tofrom, x)
@ -562,8 +562,8 @@ int main(int argc, char **argv) {
#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)

View File

@ -28,6 +28,20 @@ A high-level overview of OpenMP in LLVM can be found :doc:`here <design/Overview
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
=========================

View File

@ -0,0 +1,139 @@
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.

View File

@ -0,0 +1,13 @@
OpenACC Support
===============
OpenACC support is under development for both Flang and Clang. For
this purpose, LLVM's OpenMP runtimes are being extended to serve as
OpenACC runtimes.
.. toctree::
:glob:
:hidden:
:maxdepth: 1
OpenMPExtensions