Index: clang/docs/OpenMPSupport.rst =================================================================== --- clang/docs/OpenMPSupport.rst +++ clang/docs/OpenMPSupport.rst @@ -360,3 +360,20 @@ +------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+ | task extension | nowait clause on taskwait | :none:`unclaimed` | | +------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+ + +OpenMP Extensions +================= + +The following table provides a quick overview over various OpenMP +extensions and their implementation status. These extensions are not +currently defined by any standard, so links to associated LLVM +documentation are provided. As these extensions mature, they will be +considered for standardization. Please contact *openmp-dev* at +*lists.llvm.org* to provide feedback. + ++------------------------------+----------------------------------------------------------------------+--------------------------+--------------------------------------------------------+ +|Category | Feature | Status | Reviews | ++==============================+======================================================================+==========================+========================================================+ +| device extension | `'hold' map type modifier | :good:`prototyped` | | +| | `_ | | | ++------------------------------+----------------------------------------------------------------------+--------------------------+--------------------------------------------------------+ Index: clang/include/clang/AST/OpenMPClause.h =================================================================== --- clang/include/clang/AST/OpenMPClause.h +++ clang/include/clang/AST/OpenMPClause.h @@ -5606,7 +5606,8 @@ /// Map-type-modifiers for the 'map' clause. OpenMPMapModifierKind MapTypeModifiers[NumberOfOMPMapClauseModifiers] = { OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown, - OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown}; + OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown, + OMPC_MAP_MODIFIER_unknown}; /// Location of map-type-modifiers for the 'map' clause. SourceLocation MapTypeModifiersLoc[NumberOfOMPMapClauseModifiers]; Index: clang/include/clang/Basic/DiagnosticSemaKinds.td =================================================================== --- clang/include/clang/Basic/DiagnosticSemaKinds.td +++ clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -10526,6 +10526,8 @@ "variable already marked as mapped in current construct">; def err_omp_invalid_map_type_for_directive : Error< "%select{map type '%1' is not allowed|map type must be specified}0 for '#pragma omp %2'">; +def err_omp_invalid_map_type_modifier_for_directive : Error< + "map type modifier '%0' is not allowed for '#pragma omp %1'">; def err_omp_no_clause_for_directive : Error< "expected at least one %0 clause for '#pragma omp %1'">; def err_omp_threadprivate_in_clause : Error< Index: clang/include/clang/Basic/OpenMPKinds.def =================================================================== --- clang/include/clang/Basic/OpenMPKinds.def +++ clang/include/clang/Basic/OpenMPKinds.def @@ -123,6 +123,8 @@ OPENMP_MAP_MODIFIER_KIND(close) OPENMP_MAP_MODIFIER_KIND(mapper) OPENMP_MAP_MODIFIER_KIND(present) +// This is an OpenMP extension for the sake of OpenACC support. +OPENMP_MAP_MODIFIER_KIND(hold) // Modifiers for 'to' or 'from' clause. OPENMP_MOTION_MODIFIER_KIND(mapper) Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp =================================================================== --- clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -7106,6 +7106,14 @@ /// 0x800 is reserved for compatibility with XLC. /// Produce a runtime error if the data is not already allocated. OMP_MAP_PRESENT = 0x1000, + // Increment and decrement a separate reference counter so that the data + // cannot be unmapped within the associated region. Thus, this flag is + // intended to be used on 'target' and 'target data' directives because they + // are inherently structured. It is not intended to be used on 'target + // enter data' and 'target exit data' directives because they are inherently + // dynamic. + // This is an OpenMP extension for the sake of OpenACC support. + OMP_MAP_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 @@ -7407,6 +7415,8 @@ llvm::find(MotionModifiers, OMPC_MOTION_MODIFIER_present) != MotionModifiers.end()) Bits |= OMP_MAP_PRESENT; + if (llvm::find(MapModifiers, OMPC_MAP_MODIFIER_hold) != MapModifiers.end()) + Bits |= OMP_MAP_HOLD; if (IsNonContiguous) Bits |= OMP_MAP_NON_CONTIG; return Bits; @@ -8760,6 +8770,20 @@ CombinedInfo.Types.back() |= OMP_MAP_PRESENT; // Remove TARGET_PARAM flag from the first element (*CurTypes.begin()) &= ~OMP_MAP_TARGET_PARAM; + // If any element has the 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_HOLD; + })) { + CombinedInfo.Types.back() |= OMP_MAP_HOLD; + for (auto &M : CurTypes) + M |= OMP_MAP_HOLD; + } // All other current entries will be MEMBER_OF the combined entry // (except for PTR_AND_OBJ entries which do not have a placeholder value Index: clang/lib/Parse/ParseOpenMP.cpp =================================================================== --- clang/lib/Parse/ParseOpenMP.cpp +++ clang/lib/Parse/ParseOpenMP.cpp @@ -3554,7 +3554,8 @@ OpenMPMapModifierKind TypeModifier = isMapModifier(*this); if (TypeModifier == OMPC_MAP_MODIFIER_always || TypeModifier == OMPC_MAP_MODIFIER_close || - TypeModifier == OMPC_MAP_MODIFIER_present) { + TypeModifier == OMPC_MAP_MODIFIER_present || + TypeModifier == OMPC_MAP_MODIFIER_hold) { Data.MapTypeModifiers.push_back(TypeModifier); Data.MapTypeModifiersLoc.push_back(Tok.getLocation()); ConsumeToken(); Index: clang/lib/Sema/SemaOpenMP.cpp =================================================================== --- clang/lib/Sema/SemaOpenMP.cpp +++ clang/lib/Sema/SemaOpenMP.cpp @@ -19321,6 +19321,7 @@ CXXScopeSpec &MapperIdScopeSpec, DeclarationNameInfo MapperId, ArrayRef UnresolvedMappers, OpenMPMapClauseKind MapType = OMPC_MAP_unknown, + ArrayRef Modifiers = None, bool IsMapTypeImplicit = false) { // We only expect mappable expressions in 'to', 'from', and 'map' clauses. assert((CKind == OMPC_map || CKind == OMPC_to || CKind == OMPC_from) && @@ -19342,6 +19343,10 @@ bool UpdateUMIt = false; Expr *UnresolvedMapper = nullptr; + bool HasHoldModifier = + Modifiers.end() != + std::find(Modifiers.begin(), Modifiers.end(), OMPC_MAP_MODIFIER_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. @@ -19540,6 +19545,19 @@ continue; } + // The '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_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] @@ -19614,7 +19632,8 @@ const OMPVarListLocTy &Locs, ArrayRef UnresolvedMappers) { OpenMPMapModifierKind Modifiers[] = { OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown, - OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown}; + OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown, + OMPC_MAP_MODIFIER_unknown}; SourceLocation ModifiersLoc[NumberOfOMPMapClauseModifiers]; // Process map-type-modifiers, flag errors for duplicate modifiers. @@ -19635,7 +19654,7 @@ MappableVarListInfo MVLI(VarList); checkMappableExpressionList(*this, DSAStack, OMPC_map, MVLI, Locs.StartLoc, MapperIdScopeSpec, MapperId, UnresolvedMappers, - MapType, IsMapTypeImplicit); + MapType, Modifiers, IsMapTypeImplicit); // We need to produce a map clause even if we don't have variables so that // other diagnostics related with non-existing map clauses are accurate. Index: clang/test/OpenMP/target_ast_print.cpp =================================================================== --- clang/test/OpenMP/target_ast_print.cpp +++ clang/test/OpenMP/target_ast_print.cpp @@ -1125,4 +1125,76 @@ return tmain(argc, &argc) + tmain(argv[0][0], argv[0]); } #endif // OMP51 + +#ifdef OMPEXT + +// RUN: %clang_cc1 -DOMPEXT -verify -fopenmp -ast-print %s | FileCheck %s --check-prefix=OMPEXT +// RUN: %clang_cc1 -DOMPEXT -fopenmp -x c++ -std=c++11 -emit-pch -o %t %s +// RUN: %clang_cc1 -DOMPEXT -fopenmp -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s --check-prefix=OMPEXT + +// RUN: %clang_cc1 -DOMPEXT -verify -fopenmp-simd -ast-print %s | FileCheck %s --check-prefix=OMPEXT +// RUN: %clang_cc1 -DOMPEXT -fopenmp-simd -x c++ -std=c++11 -emit-pch -o %t %s +// RUN: %clang_cc1 -DOMPEXT -fopenmp-simd -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s --check-prefix=OMPEXT + +void foo() {} + +template +T tmain(T argc, T *argv) { + T i, hold; +#pragma omp target map(hold,alloc: i) + foo(); +#pragma omp target map(hold from: i) + foo(); +#pragma omp target map(hold) + {hold++;} +#pragma omp target map(hold,i) + {hold++;i++;} + return 0; +} + +// OMPEXT: template T tmain(T argc, T *argv) { +// OMPEXT-NEXT: T i, hold; +// OMPEXT-NEXT: #pragma omp target map(hold,alloc: i) +// OMPEXT-NEXT: foo() +// OMPEXT-NEXT: #pragma omp target map(hold,from: i) +// OMPEXT-NEXT: foo() +// OMPEXT-NEXT: #pragma omp target map(tofrom: hold) +// OMPEXT-NEXT: { +// OMPEXT-NEXT: hold++; +// OMPEXT-NEXT: } +// OMPEXT-NEXT: #pragma omp target map(tofrom: hold,i) +// OMPEXT-NEXT: { +// OMPEXT-NEXT: hold++; +// OMPEXT-NEXT: i++; +// OMPEXT-NEXT: } + +// OMPEXT-LABEL: int main(int argc, char **argv) { +// OMPEXT-NEXT: int i, hold; +// OMPEXT-NEXT: #pragma omp target map(hold,alloc: i) +// OMPEXT-NEXT: foo(); +// OMPEXT-NEXT: #pragma omp target map(hold,from: i) +// OMPEXT-NEXT: foo(); +// OMPEXT-NEXT: #pragma omp target map(tofrom: hold) +// OMPEXT-NEXT: { +// OMPEXT-NEXT: hold++; +// OMPEXT-NEXT: } +// OMPEXT-NEXT: #pragma omp target map(tofrom: hold,i) +// OMPEXT-NEXT: { +// OMPEXT-NEXT: hold++; +// OMPEXT-NEXT: i++; +// OMPEXT-NEXT: } +int main (int argc, char **argv) { + int i, hold; + #pragma omp target map(hold,alloc: i) + foo(); + #pragma omp target map(hold from: i) + foo(); + #pragma omp target map(hold) + {hold++;} + #pragma omp target map(hold,i) + {hold++;i++;} + return tmain(argc, &argc) + tmain(argv[0][0], argv[0]); +} + +#endif #endif Index: clang/test/OpenMP/target_data_ast_print.cpp =================================================================== --- clang/test/OpenMP/target_data_ast_print.cpp +++ clang/test/OpenMP/target_data_ast_print.cpp @@ -56,6 +56,9 @@ foo(); #endif +#pragma omp target data map(hold,alloc: e) + foo(); + // nesting a target region #pragma omp target data map(e) { @@ -67,6 +70,8 @@ #pragma omp target map(present, alloc: e) foo(); #endif + #pragma omp target map(hold, alloc: e) + foo(); } return 0; @@ -94,6 +99,8 @@ // CHECK-NEXT: foo(); // OMP51-NEXT: #pragma omp target data map(present,alloc: e) // OMP51-NEXT: foo(); +// CHECK-NEXT: #pragma omp target data map(hold,alloc: e) +// CHECK-NEXT: foo(); // CHECK-NEXT: #pragma omp target data map(tofrom: e) // CHECK-NEXT: { // CHECK-NEXT: #pragma omp target map(always,alloc: e) @@ -102,6 +109,8 @@ // CHECK-NEXT: foo(); // OMP51-NEXT: #pragma omp target map(present,alloc: e) // OMP51-NEXT: foo(); +// CHECK-NEXT: #pragma omp target map(hold,alloc: e) +// CHECK-NEXT: foo(); // CHECK: template<> int tmain(int argc, int *argv) { // CHECK-NEXT: int i, j, b, c, d, e, x[20]; // CHECK-NEXT: #pragma omp target data map(to: c) @@ -124,6 +133,8 @@ // CHECK-NEXT: foo(); // OMP51-NEXT: #pragma omp target data map(present,alloc: e) // OMP51-NEXT: foo(); +// CHECK-NEXT: #pragma omp target data map(hold,alloc: e) +// CHECK-NEXT: foo(); // CHECK-NEXT: #pragma omp target data map(tofrom: e) // CHECK-NEXT: { // CHECK-NEXT: #pragma omp target map(always,alloc: e) @@ -132,6 +143,8 @@ // CHECK-NEXT: foo(); // OMP51-NEXT: #pragma omp target map(present,alloc: e) // OMP51-NEXT: foo(); +// CHECK-NEXT: #pragma omp target map(hold,alloc: e) +// CHECK-NEXT: foo(); // CHECK: template<> char tmain(char argc, char *argv) { // CHECK-NEXT: char i, j, b, c, d, e, x[20]; // CHECK-NEXT: #pragma omp target data map(to: c) @@ -154,6 +167,8 @@ // CHECK-NEXT: foo(); // OMP51-NEXT: #pragma omp target data map(present,alloc: e) // OMP51-NEXT: foo(); +// CHECK-NEXT: #pragma omp target data map(hold,alloc: e) +// CHECK-NEXT: foo(); // CHECK-NEXT: #pragma omp target data map(tofrom: e) // CHECK-NEXT: { // CHECK-NEXT: #pragma omp target map(always,alloc: e) @@ -162,6 +177,8 @@ // CHECK-NEXT: foo(); // OMP51-NEXT: #pragma omp target map(present,alloc: e) // OMP51-NEXT: foo(); +// CHECK-NEXT: #pragma omp target map(hold,alloc: e) +// CHECK-NEXT: foo(); int main (int argc, char **argv) { int b = argc, c, d, e, f, g, x[20]; @@ -221,6 +238,11 @@ foo(); #endif +// CHECK-NEXT: #pragma omp target data map(hold,alloc: e) +// CHECK-NEXT: foo(); +#pragma omp target data map(hold,alloc: e) + foo(); + // nesting a target region #pragma omp target data map(e) // CHECK-NEXT: #pragma omp target data map(tofrom: e) Index: clang/test/OpenMP/target_data_map_codegen_hold.cpp =================================================================== --- /dev/null +++ clang/test/OpenMP/target_data_map_codegen_hold.cpp @@ -0,0 +1,602 @@ +// 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-targets=powerpc64le-ibm-linux-gnu \ +// RUN: -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | \ +// RUN: FileCheck %s --check-prefixes=CHECK-PPC64LE +// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ \ +// RUN: -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -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-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-targets=i386-pc-linux-gnu -x c++ \ +// RUN: -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -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 +// 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(hold, to: lb) + {++arg;} + + #pragma omp target data map(always close hold, to: lb) + {++arg;} + + #pragma omp target data map(hold, tofrom : arg) + {++arg;} + + // Make sure the struct picks up hold even if another element of the struct + // doesn't have hold. + #pragma omp target data map(tofrom : ps1->s, arg) \ + map(hold, tofrom : ps1->ps->ps->ps->s, ps2->s) \ + map(tofrom : ps2->ps->ps->ps->s) + { + ++(arg); + } +} + +#endif Index: clang/test/OpenMP/target_enter_data_map_messages.c =================================================================== --- clang/test/OpenMP/target_enter_data_map_messages.c +++ clang/test/OpenMP/target_enter_data_map_messages.c @@ -25,5 +25,10 @@ #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'}} + // expected-error@+1 {{map type modifier 'hold' is not allowed for '#pragma omp target enter data'}} + #pragma omp target enter data map(hold, alloc: r) + // expected-error@+1 {{map type modifier 'hold' is not allowed for '#pragma omp target enter data'}} + #pragma omp target enter data map(hold, to: r) + return 0; } Index: clang/test/OpenMP/target_exit_data_map_messages.c =================================================================== --- clang/test/OpenMP/target_exit_data_map_messages.c +++ clang/test/OpenMP/target_exit_data_map_messages.c @@ -18,5 +18,12 @@ #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'}} + // expected-error@+1 {{map type modifier 'hold' is not allowed for '#pragma omp target exit data'}} + #pragma omp target exit data map(hold, from: r) + // expected-error@+1 {{map type modifier 'hold' is not allowed for '#pragma omp target exit data'}} + #pragma omp target exit data map(hold, release: r) + // expected-error@+1 {{map type modifier 'hold' is not allowed for '#pragma omp target exit data'}} + #pragma omp target exit data map(hold, delete: r) + return 0; } Index: clang/test/OpenMP/target_map_codegen_hold.cpp =================================================================== --- /dev/null +++ clang/test/OpenMP/target_map_codegen_hold.cpp @@ -0,0 +1,917 @@ +// 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 \ +// 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-targets=powerpc64le-ibm-linux-gnu \ +// RUN: -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch \ +// RUN: -o %t %s +// RUN: %clang_cc1 -DUSE -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu \ +// RUN: -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch \ +// RUN: %t -verify %s -emit-llvm -o - | \ +// RUN: FileCheck %s --check-prefixes=CHECK-USE-PPC64LE + +// i386-pc-linux-gnu + +// RUN: %clang_cc1 -DUSE -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu \ +// RUN: -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | \ +// RUN: FileCheck %s --check-prefixes=CHECK-USE-I386 +// RUN: %clang_cc1 -DUSE -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ \ +// RUN: -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -DUSE -fopenmp -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-targets=powerpc64le-ibm-linux-gnu \ +// RUN: -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | \ +// RUN: FileCheck %s --check-prefixes=CHECK-NOUSE-PPC64LE +// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ \ +// RUN: -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -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-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-targets=i386-pc-linux-gnu -x c++ \ +// RUN: -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -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 +// 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_l644.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_l644(%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_l657.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_l657(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_l644.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_l644(%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_l657.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_l657(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_l644.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_l644() #[[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_l657.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_l657() #[[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_l644.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_l644() #[[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_l657.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_l657() #[[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 hold even if another element of the struct + // doesn't have hold. + #pragma omp target map(tofrom : st1.i) map(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 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_l908.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_l908(%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_l908.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_l908(%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_l908.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_l908() #[[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_l908.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_l908() #[[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 hold even if another element of the + // struct doesn't have hold. + #pragma omp target map(tofrom : i) map(hold, tofrom : j) + { +#ifdef USE + i++; + j++; +#endif + } +} + +#endif Index: clang/test/OpenMP/target_map_messages.cpp =================================================================== --- clang/test/OpenMP/target_map_messages.cpp +++ clang/test/OpenMP/target_map_messages.cpp @@ -138,6 +138,21 @@ // expected-error@+1 {{use of undeclared identifier 'present'}} #pragma omp target map(present) {} + #pragma omp target map(hold, tofrom: c,f) + {} + #pragma omp target map(hold, tofrom: c[1:2],f) + {} + #pragma omp target map(hold, tofrom: c,f[1:2]) + {} + // expected-error@+1 {{section length is unspecified and cannot be inferred because subscripted value is not an array}} + #pragma omp target map(hold, tofrom: c[:],f) + {} + // expected-error@+1 {{section length is unspecified and cannot be inferred because subscripted value is not an array}} + #pragma omp target map(hold, tofrom: c,f[:]) + {} + // expected-error@+1 {{use of undeclared identifier 'hold'}} + #pragma omp target map(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}} @@ -146,10 +161,13 @@ // lt51-error@+1 2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} #pragma omp target map(present, present, tofrom: a) {} - // expected-error@+3 2 {{same map type modifier has been specified more than once}} + // expected-error@+1 {{same map type modifier has been specified more than once}} + #pragma omp target map(hold, hold, tofrom: a) + {} + // expected-error@+3 3 {{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) + #pragma omp target map(always, close, present, hold, always, close, present, hold, tofrom: a) {} #pragma omp target map( , tofrom: a) // expected-error {{missing map type modifier}} {} Index: openmp/docs/index.rst =================================================================== --- openmp/docs/index.rst +++ openmp/docs/index.rst @@ -28,6 +28,21 @@ design/Overview +OpenACC Support +=============== + +:doc:`OpenACC support ` is under development for +both Flang and Clang. For this purpose, LLVM's OpenMP runtimes also +serve as OpenACC runtimes. However, OpenACC requires some +:doc:`runtime functionality ` that OpenMP +does not, and in some cases Clang also supports OpenMP extensions +providing similar functionality. + +.. toctree:: + :hidden: + :maxdepth: 1 + + openacc/Overview LLVM/OpenMP Optimizations ========================= Index: openmp/docs/openacc/OpenMPExtensions.rst =================================================================== --- /dev/null +++ openmp/docs/openacc/OpenMPExtensions.rst @@ -0,0 +1,110 @@ +OpenMP Extensions +================= + +OpenACC requires some runtime functionality that OpenMP does not. In +some cases, Clang also supports OpenMP extensions providing similar +functionality. This section documents those extensions. + +.. _hold: + +``hold`` Map Type Modifier +-------------------------- + +.. _holdExample: + +Example +^^^^^^^ + +.. code-block:: c++ + + #pragma omp target data map(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 ``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. + +.. _holdBehavior: + +Behavior +^^^^^^^^ + +* Stated more generally, the ``hold`` map type modifier specifies + that, regardless of changes in the value of the standard OpenMP + reference count, the specified data is not unmapped until the end of + the construct. If ``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 ``hold`` map type modifier is permitted to appear only on + ``target`` constructs (and associated combined constructs) and + ``target data`` constructs. It is not permitted to appear on + ``target enter data`` or ``target exit data`` directives because + there is no associated statement, so it is not meangingful to hold + onto a mapping until the end of the directive. +* The runtime reports an error if ``omp_target_disassociate_ptr`` is + called for a mapping for which the ``hold`` map type modifier is in + effect. +* Like the ``present`` map type modifier, the ``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 ``hold`` map type + modifier. +* ``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 ``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 ``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 ``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 ``hold`` map type modifier is *in + effect* (see :ref:`Behavior ` above) when the ``hold`` + reference count is greater than zero. + +Relationship with OpenACC +^^^^^^^^^^^^^^^^^^^^^^^^^ + +OpenACC specifies two reference counts for tracking host/device data +mappings. Which reference count is used to implement an OpenACC +directive is determined by the nature of that directive, either +dynamic or structured: + +* The *dynamic reference count* is always used for ``enter data`` and + ``exit data`` directives and corresponding OpenACC routines. +* The *structured reference count* is always used for ``data`` and + compute constructs, which are similar to OpenMP's ``target data`` + and ``target`` constructs. + +Contrast with OpenMP, where the dynamic reference count is always used +unless the application developer specifies an alternate behavior via +our map type modifier extension. We chose the name *hold* for that +map type modifier because, as demonstrated in the above :ref:`example +`, *hold* concisely identifies the desired behavior from +the application developer's perspective without referencing the +implementation of that behavior. + +The ``hold`` reference count is otherwise modeled after OpenACC's +structured reference count. For example, calling ``acc_unmap_data``, +which is similar to ``omp_target_disassociate_ptr``, is an error when +the structured reference count is not zero. + +While Flang and Clang obviously must implement the syntax and +semantics for selecting OpenACC reference counts differently than for +selecting OpenMP reference counts, the implementation is the same at +the runtime level. That is, OpenACC's dynamic reference count is +OpenMP's dynamic reference count, and OpenACC's structured reference +count is our OpenMP ``hold`` reference count extension. Index: openmp/docs/openacc/Overview.rst =================================================================== --- /dev/null +++ openmp/docs/openacc/Overview.rst @@ -0,0 +1,12 @@ +OpenACC Support +=============== + +OpenACC support is under development for both Flang and Clang. For +this purpose, LLVM's OpenMP runtimes also serve as OpenACC runtimes. + +.. toctree:: + :glob: + :hidden: + :maxdepth: 1 + + OpenMPExtensions