[clang] [clang][OpenMP] 6.0: detect privatization of array section/assumed-size array (PR #152786)
David Pagan via cfe-commits
cfe-commits at lists.llvm.org
Fri Aug 8 13:04:33 PDT 2025
https://github.com/ddpagan created https://github.com/llvm/llvm-project/pull/152786
According to the OpenMP 6.0 specification, array sections with no length and unknown size are considered assumed-size arrays. As of pull request
https://github.com/llvm/llvm-project/pull/148048
these types of array sections are allowed and can be specified in clauses that allow array sections as list items. However, only two clauses explicitly allow array sections that are assumed-size arrays:
- 'map' and 'use_device_addr'.
The other clauses that accept array sections do not explicitly accept assumed-size arrays:
- inclusive, exclusive, has_device_addr, in_reduction, task_reduction, reduction These cases should generate an error. See OpenMP 6.0 specification section 7.4 List Item Privatization, Restrictions, p. 222, L15
Assumed-size arrays must not be privatized
For OpenMP 6.0, function getPrivateItem() now checks for array section list items that are assumed-size arrays and generates an error if they are not allowed for the clause.
Testing
- Updated LIT tests for assumed-size array sections to ensure these clauses generate an error: inclusive, exclusive, has_device_addr, in_reduction, task_reduction, reduction and that this clause is accepted (codegen test): use_device_addr
- check-all
- OpenMP_VV (sollve_vv)
>From 5947d70688474dc8fa86e3593e62f591af26f7dc Mon Sep 17 00:00:00 2001
From: Dave Pagan <dave.pagan at amd.com>
Date: Tue, 29 Jul 2025 16:21:54 -0500
Subject: [PATCH] [clang][OpenMP] 6.0: detect privatization of array
section/assumed-size array
According to the OpenMP 6.0 specification, array sections with no length
and unknown size are considered assumed-size arrays. As of pull request
https://github.com/llvm/llvm-project/pull/148048
these types of array sections are allowed and can be specified in clauses
that allow array sections as list items. However, only two clauses
explicitly allow array sections that are assumed-size arrays:
- 'map' and 'use_device_addr'.
The other clauses that accept array sections do not explicitly accept
assumed-size arrays:
- inclusive, exclusive, has_device_addr, in_reduction, task_reduction,
reduction
These cases should generate an error. See OpenMP 6.0 specification
section 7.4 List Item Privatization, Restrictions, p. 222, L15
Assumed-size arrays must not be privatized
For OpenMP 6.0, function getPrivateItem() now checks for array section list
items that are assumed-size arrays and generates an error if they are not
allowed for the clause.
Testing
- Updated LIT tests for assumed-size array sections to ensure these
clauses generate an error:
inclusive, exclusive, has_device_addr, in_reduction, task_reduction,
reduction
and that this clause is accepted (codegen test):
use_device_addr
- check-all
- OpenMP_VV (sollve_vv)
---
clang/lib/Sema/SemaOpenMP.cpp | 46 +-
clang/test/OpenMP/scan_messages.cpp | 32 +-
.../target_data_use_device_addr_codegen.cpp | 1089 ++++++++++++++---
.../target_has_device_addr_messages.cpp | 26 +-
.../test/OpenMP/task_in_reduction_message.cpp | 55 +-
.../taskgroup_task_reduction_messages.cpp | 22 +-
.../test/OpenMP/teams_reduction_messages.cpp | 26 +-
7 files changed, 1064 insertions(+), 232 deletions(-)
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index 2c5d97c3064ac..2fd3a60fcae0b 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -2780,7 +2780,7 @@ void SemaOpenMP::EndOpenMPClause() {
static std::pair<ValueDecl *, bool>
getPrivateItem(Sema &S, Expr *&RefExpr, SourceLocation &ELoc,
SourceRange &ERange, bool AllowArraySection = false,
- StringRef DiagType = "");
+ bool AllowAssumedSizeArray = false, StringRef DiagType = "");
/// Check consistency of the reduction clauses.
static void checkReductionClauses(Sema &S, DSAStackTy *Stack,
@@ -5148,11 +5148,10 @@ static bool checkIfClauses(Sema &S, OpenMPDirectiveKind Kind,
return ErrorFound;
}
-static std::pair<ValueDecl *, bool> getPrivateItem(Sema &S, Expr *&RefExpr,
- SourceLocation &ELoc,
- SourceRange &ERange,
- bool AllowArraySection,
- StringRef DiagType) {
+static std::pair<ValueDecl *, bool>
+getPrivateItem(Sema &S, Expr *&RefExpr, SourceLocation &ELoc,
+ SourceRange &ERange, bool AllowArraySection,
+ bool AllowAssumedSizeArray, StringRef DiagType) {
if (RefExpr->isTypeDependent() || RefExpr->isValueDependent() ||
RefExpr->containsUnexpandedParameterPack())
return std::make_pair(nullptr, true);
@@ -5162,6 +5161,20 @@ static std::pair<ValueDecl *, bool> getPrivateItem(Sema &S, Expr *&RefExpr,
// OpenMP [2.9.3.3, Restrictions, p.1]
// A variable that is part of another variable (as an array or
// structure element) cannot appear in a private clause.
+ //
+ // OpenMP [6.0]
+ // 5.2.5 Array Sections, p. 166, L28-29
+ // When the length is absent and the size of the dimension is not known,
+ // the array section is an assumed-size array.
+ // 2 Glossary, p. 23, L4-6
+ // assumed-size array
+ // For C/C++, an array section for which the length is absent and the
+ // size of the dimensions is not known.
+ // 5.2.5 Array Sections, p. 168, L11
+ // An assumed-size array can appear only in clauses for which it is
+ // explicitly allowed.
+ // 7.4 List Item Privatization, Restrictions, p. 222, L15
+ // Assumed-size arrays must not be privatized.
RefExpr = RefExpr->IgnoreParens();
enum {
NoArrayExpr = -1,
@@ -5177,6 +5190,17 @@ static std::pair<ValueDecl *, bool> getPrivateItem(Sema &S, Expr *&RefExpr,
IsArrayExpr = ArraySubscript;
} else if (auto *OASE = dyn_cast_or_null<ArraySectionExpr>(RefExpr)) {
Expr *Base = OASE->getBase()->IgnoreParenImpCasts();
+ if (S.getLangOpts().OpenMP >= 60 && !AllowAssumedSizeArray &&
+ OASE->getColonLocFirst().isValid() && !OASE->getLength()) {
+ QualType BaseType = ArraySectionExpr::getBaseOriginalType(Base);
+ if (BaseType.isNull() || (!BaseType->isConstantArrayType() &&
+ !BaseType->isVariableArrayType())) {
+ S.Diag(OASE->getColonLocFirst(),
+ diag::err_omp_section_length_undefined)
+ << (!BaseType.isNull() && BaseType->isArrayType());
+ return std::make_pair(nullptr, false);
+ }
+ }
while (auto *TempOASE = dyn_cast<ArraySectionExpr>(Base))
Base = TempOASE->getBase()->IgnoreParenImpCasts();
while (auto *TempASE = dyn_cast<ArraySubscriptExpr>(Base))
@@ -17294,9 +17318,10 @@ static bool isValidInteropVariable(Sema &SemaRef, Expr *InteropVarExpr,
SourceLocation ELoc;
SourceRange ERange;
Expr *RefExpr = InteropVarExpr;
- auto Res =
- getPrivateItem(SemaRef, RefExpr, ELoc, ERange,
- /*AllowArraySection=*/false, /*DiagType=*/"omp_interop_t");
+ auto Res = getPrivateItem(SemaRef, RefExpr, ELoc, ERange,
+ /*AllowArraySection=*/false,
+ /*AllowAssumedSizeArray=*/false,
+ /*DiagType=*/"omp_interop_t");
if (Res.second) {
// It will be analyzed later.
@@ -23479,7 +23504,8 @@ SemaOpenMP::ActOnOpenMPUseDeviceAddrClause(ArrayRef<Expr *> VarList,
SourceRange ERange;
Expr *SimpleRefExpr = RefExpr;
auto Res = getPrivateItem(SemaRef, SimpleRefExpr, ELoc, ERange,
- /*AllowArraySection=*/true);
+ /*AllowArraySection=*/true,
+ /*AllowAssumedSizeArray=*/true);
if (Res.second) {
// It will be analyzed later.
MVLI.ProcessedVarList.push_back(RefExpr);
diff --git a/clang/test/OpenMP/scan_messages.cpp b/clang/test/OpenMP/scan_messages.cpp
index 0de94898c6571..45516768090a8 100644
--- a/clang/test/OpenMP/scan_messages.cpp
+++ b/clang/test/OpenMP/scan_messages.cpp
@@ -2,9 +2,14 @@
// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 150 %s
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=60 -ferror-limit 150 %s
+
+// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=60 -ferror-limit 150 %s
+
template <class T>
-T tmain() {
+T tfoobar(T ub[]) {
static T argc;
+ T *ptr;
#pragma omp for
for (int i = 0; i < 10; ++i) {
#pragma omp scan // expected-error {{exactly one of 'inclusive' or 'exclusive' clauses is expected}}
@@ -83,12 +88,23 @@ T tmain() {
label1 : {
#pragma omp scan inclusive(argc) // expected-error {{orphaned 'omp scan' directives are prohibited; perhaps you forget to enclose the directive into a for, simd, for simd, parallel for, or parallel for simd region?}}
}}
+#pragma omp simd reduction(inscan, +: argc)
+ for (int i = 0; i < 10; ++i) {
+ #pragma omp scan inclusive(ptr[0:], argc) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+ ;
+ }
+#pragma omp simd reduction(inscan, +: argc)
+ for (int i = 0; i < 10; ++i) {
+ #pragma omp scan exclusive(argc, ub[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is an array of unknown bound}}
+ ;
+ }
return T();
}
-int main() {
+int foobar(int ub[]) {
static int argc;
+ int *ptr;
#pragma omp simd reduction(inscan, +: argc)
for (int i = 0; i < 10; ++i) {
#pragma omp scan inclusive(argc) inclusive(argc) // expected-error {{exactly one of 'inclusive' or 'exclusive' clauses is expected}}
@@ -177,6 +193,16 @@ label1 : {
#pragma omp scan inclusive(argc) // expected-error {{orphaned 'omp scan' directives are prohibited; perhaps you forget to enclose the directive into a for, simd, for simd, parallel for, or parallel for simd region?}}
}
}
+#pragma omp simd reduction(inscan, +: argc)
+ for (int i = 0; i < 10; ++i) {
+ #pragma omp scan inclusive(ptr[0:], argc) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+ ;
+ }
+#pragma omp simd reduction(inscan, +: argc)
+ for (int i = 0; i < 10; ++i) {
+ #pragma omp scan exclusive(argc, ub[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is an array of unknown bound}}
+ ;
+ }
- return tmain<int>(); // expected-note {{in instantiation of function template specialization 'tmain<int>' requested here}}
+ return tfoobar<int>(ub); // expected-note {{in instantiation of function template specialization 'tfoobar<int>' requested here}}
}
diff --git a/clang/test/OpenMP/target_data_use_device_addr_codegen.cpp b/clang/test/OpenMP/target_data_use_device_addr_codegen.cpp
index 19a523ee165c8..a67d3258ea071 100644
--- a/clang/test/OpenMP/target_data_use_device_addr_codegen.cpp
+++ b/clang/test/OpenMP/target_data_use_device_addr_codegen.cpp
@@ -1,23 +1,28 @@
-// RUN: %clang_cc1 -DCK1 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s
-// RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _ --version 5
+// RUN: %clang_cc1 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s
-// RUN: %clang_cc1 -DCK1 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
-// RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -DOMP60 -verify -Wno-vla -fopenmp -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix OMP60 %s
+// RUN: %clang_cc1 -DOMP60 -fopenmp -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -DOMP60 -fopenmp -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck --check-prefix OMP60 %s
+
+// RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
-// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+
+// RUN: %clang_cc1 -DOMP60 -verify -Wno-vla -fopenmp-simd -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0-OMP60 %s
+// RUN: %clang_cc1 -DOMP60 -fopenmp-simd -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -DOMP60 -fopenmp-simd -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0-OMP60 %s
// expected-no-diagnostics
#ifndef HEADER
#define HEADER
-// CHECK-DAG: [[SIZES1:@.+]] = private unnamed_addr constant [6 x i64] [i64 4, i64 16, i64 4, i64 4, i64 0, i64 4]
// 64 = 0x40 = OMP_MAP_RETURN_PARAM
-// CHECK-DAG: [[MAPTYPES1:@.+]] = private unnamed_addr constant [6 x i64] [i64 67, i64 115, i64 51, i64 67, i64 67, i64 67]
-// CHECK-DAG: [[SIZES2:@.+]] = private unnamed_addr constant [6 x i64] [i64 0, i64 4, i64 16, i64 4, i64 4, i64 0]
// 0 = OMP_MAP_NONE
// 281474976710720 = 0x1000000000040 = OMP_MAP_MEMBER_OF | OMP_MAP_RETURN_PARAM
-// CHECK-DAG: [[MAPTYPES2:@.+]] = private unnamed_addr constant [6 x i64] [i64 0, i64 281474976710723, i64 281474976710739, i64 281474976710739, i64 281474976710675, i64 281474976710723]
+
struct S {
int a = 0;
int *ptr = &a;
@@ -25,12 +30,20 @@ struct S {
int arr[4];
S() {}
void foo() {
+#ifdef OMP60
+#pragma omp target data map(tofrom: a, ptr [3:4], ref, ptr[0], arr[:a], ptr[:]) use_device_addr(a, ptr [3:4], ref, ptr[0], arr[:a], ptr[:])
+#else
#pragma omp target data map(tofrom: a, ptr [3:4], ref, ptr[0], arr[:a]) use_device_addr(a, ptr [3:4], ref, ptr[0], arr[:a])
+#endif // OMP60
++a, ++*ptr, ++ref, ++arr[0];
}
};
+#ifdef OMP60
+int main(int argc, char *argv[]) {
+#else
int main() {
+#endif // OMP60
float a = 0;
float *ptr = &a;
float &ref = a;
@@ -38,185 +51,901 @@ int main() {
float vla[(int)a];
S s;
s.foo();
+#ifdef OMP60
+#pragma omp target data map(tofrom: a, ptr [3:4], ref, ptr[0], arr[:(int)a], vla[0], ptr[:], argv[0][:]) use_device_addr(a, ptr [3:4], ref, ptr[0], arr[:(int)a], vla[0], ptr[:], argv[0][:])
+#else
#pragma omp target data map(tofrom: a, ptr [3:4], ref, ptr[0], arr[:(int)a], vla[0]) use_device_addr(a, ptr [3:4], ref, ptr[0], arr[:(int)a], vla[0])
+#endif // OMP60
++a, ++*ptr, ++ref, ++arr[0], ++vla[0];
return a;
}
-// CHECK-LABEL: @main()
-// CHECK: [[A_ADDR:%.+]] = alloca float,
-// CHECK: [[PTR_ADDR:%.+]] = alloca ptr,
-// CHECK: [[REF_ADDR:%.+]] = alloca ptr,
-// CHECK: [[ARR_ADDR:%.+]] = alloca [4 x float],
-// CHECK: [[BPTRS:%.+]] = alloca [6 x ptr],
-// CHECK: [[PTRS:%.+]] = alloca [6 x ptr],
-// CHECK: [[MAP_PTRS:%.+]] = alloca [6 x ptr],
-// CHECK: [[SIZES:%.+]] = alloca [6 x i64],
-// CHECK: [[VLA_ADDR:%.+]] = alloca float, i64 %{{.+}},
-// CHECK: [[PTR:%.+]] = load ptr, ptr [[PTR_ADDR]],
-// CHECK-NEXT: [[ARR_IDX:%.+]] = getelementptr inbounds nuw float, ptr [[PTR]], i64 3
-// CHECK: [[P5:%.+]] = load ptr, ptr [[PTR_ADDR]], align 8
-// CHECK-NEXT: [[ARR_IDX1:%.+]] = getelementptr inbounds float, ptr [[P5]], i64 0
-// CHECK: [[P7:%.+]] = load ptr, ptr [[REF_ADDR]],
-// CHECK-NEXT: [[REF:%.+]] = load ptr, ptr [[REF_ADDR]],
-// CHECK-NEXT: [[ARR_IDX2:%.+]] = getelementptr inbounds nuw [4 x float], ptr [[ARR_ADDR]], i64 0, i64 0
-// CHECK: [[P10:%.+]] = mul nuw i64 {{.+}}, 4
-// CHECK-NEXT: [[ARR_IDX5:%.+]] = getelementptr inbounds float, ptr [[VLA_ADDR]], i64 0
-// CHECK-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[SIZES]], ptr align 8 [[SIZES1]], i64 48, i1 false)
-// CHECK: [[BPTR0:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 0
-// CHECK: store ptr [[A_ADDR]], ptr [[BPTR0]],
-// CHECK: [[PTR0:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 0
-// CHECK: store ptr [[A_ADDR]], ptr [[PTR0]],
-// CHECK: [[BPTR1:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 1
-// CHECK: store ptr [[PTR_ADDR]], ptr [[BPTR1]],
-// CHECK: [[PTR1:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 1
-// CHECK: store ptr [[ARR_IDX]], ptr [[PTR1]],
-// CHECK: [[BPTR2:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 2
-// CHECK: store ptr [[PTR_ADDR]], ptr [[BPTR2]],
-// CHECK: [[PTR2:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 2
-// CHECK: store ptr [[ARR_IDX1]], ptr [[PTR2]],
-// CHECK: [[BPTR3:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 3
-// CHECK: store ptr [[P7]], ptr [[BPTR3]],
-// CHECK: [[PTR3:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 3
-// CHECK: store ptr [[REF]], ptr [[PTR3]],
-// CHECK: [[BPTR4:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 4
-// CHECK: store ptr [[ARR_ADDR]], ptr [[BPTR4]], align
-// CHECK: [[PTR4:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 4
-// CHECK: store ptr [[ARR_IDX2]], ptr [[PTR4]], align 8
-// CHECK: [[SIZE_PTR:%.+]] = getelementptr inbounds [6 x i64], ptr [[SIZES]], i32 0, i32 4
-// CHECK: store i64 [[P10:%.+]], ptr [[SIZE_PTR]], align 8
-// CHECK: [[MAP_PTR:%.+]] = getelementptr inbounds [6 x ptr], ptr [[MAP_PTRS]], i64 0, i64 4
-// CHECK: store ptr null, ptr [[MAP_PTR]], align 8
-// CHECK: [[BPTR5:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 5
-// CHECK: store ptr [[VLA_ADDR]], ptr [[BPTR5]],
-// CHECK: [[PTR5:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 5
-// CHECK: store ptr [[ARR_IDX5]], ptr [[PTR5]],
-// CHECK: [[BPTR:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 0
-// CHECK: [[PTR:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 0
-// CHECK: [[SIZE:%.+]] = getelementptr inbounds [6 x i64], ptr [[SIZES]], i32 0, i32 0
-// CHECK: call void @__tgt_target_data_begin_mapper(ptr @{{.+}}, i64 -1, i32 6, ptr [[BPTR]], ptr [[PTR]], ptr [[SIZE]], ptr [[MAPTYPES1]], ptr null, ptr null)
-// CHECK: [[A_REF:%.+]] = load ptr, ptr [[BPTR0]],
-// CHECK: [[REF_REF:%.+]] = load ptr, ptr [[BPTR3]],
-// CHECK: store ptr [[REF_REF]], ptr [[TMP_REF_ADDR:%.+]],
-// CHECK: [[ARR_REF:%.+]] = load ptr, ptr [[BPTR4]],
-// CHECK: [[VLA_REF:%.+]] = load ptr, ptr [[BPTR5]],
-// CHECK: [[A:%.+]] = load float, ptr [[A_REF]],
-// CHECK: [[INC:%.+]] = fadd float [[A]], 1.000000e+00
-// CHECK: store float [[INC]], ptr [[A_REF]],
-// CHECK: [[PTR:%.+]] = load ptr, ptr [[BPTR1]],
-// CHECK: [[VAL:%.+]] = load float, ptr [[PTR]],
-// CHECK: [[INC:%.+]] = fadd float [[VAL]], 1.000000e+00
-// CHECK: store float [[INC]], ptr [[PTR]],
-// CHECK: [[REF_ADDR:%.+]] = load ptr, ptr [[TMP_REF_ADDR]],
-// CHECK: [[REF:%.+]] = load float, ptr [[REF_ADDR]],
-// CHECK: [[INC:%.+]] = fadd float [[REF]], 1.000000e+00
-// CHECK: store float [[INC]], ptr [[REF_ADDR]],
-// CHECK: [[ARR0_ADDR:%.+]] = getelementptr inbounds [4 x float], ptr [[ARR_REF]], i64 0, i64 0
-// CHECK: [[ARR0:%.+]] = load float, ptr [[ARR0_ADDR]],
-// CHECK: [[INC:%.+]] = fadd float [[ARR0]], 1.000000e+00
-// CHECK: store float [[INC]], ptr [[ARR0_ADDR]],
-// CHECK: [[VLA0_ADDR:%.+]] = getelementptr inbounds float, ptr [[VLA_REF]], i64 0
-// CHECK: [[VLA0:%.+]] = load float, ptr [[VLA0_ADDR]],
-// CHECK: [[INC:%.+]] = fadd float [[VLA0]], 1.000000e+00
-// CHECK: store float [[INC]], ptr [[VLA0_ADDR]],
-// CHECK: [[BPTR:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 0
-// CHECK: [[PTR:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 0
-// CHECK: [[SIZE:%.+]] = getelementptr inbounds [6 x i64], ptr [[SIZES]], i32 0, i32 0
-// CHECK: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 -1, i32 6, ptr [[BPTR]], ptr [[PTR]], ptr [[SIZE]], ptr [[MAPTYPES1]], ptr null, ptr null)
-// CHECK: foo
-// CHECK: [[BPTRS:%.+]] = alloca [6 x ptr],
-// CHECK: [[PTRS:%.+]] = alloca [6 x ptr],
-// CHECK: [[MAP_PTRS:%.+]] = alloca [6 x ptr],
-// CHECK: [[SIZES:%.+]] = alloca [6 x i64],
-// CHECK: [[A_ADDR:%.+]] = getelementptr inbounds nuw %struct.S, ptr [[THIS:%.+]], i32 0, i32 0
-// CHECK: [[PTR_ADDR:%.+]] = getelementptr inbounds nuw %struct.S, ptr [[THIS]], i32 0, i32 1
-// CHECK: [[ARR_IDX:%.+]] = getelementptr inbounds nuw i32, ptr %{{.+}}, i64 3
-// CHECK: [[REF_REF:%.+]] = getelementptr inbounds nuw %struct.S, ptr [[THIS]], i32 0, i32 2
-// CHECK: [[REF_PTR:%.+]] = load ptr, ptr [[REF_REF]],
-// CHECK-NEXT: [[P3:%.+]] = getelementptr inbounds nuw %struct.S, ptr [[THIS]], i32 0, i32 1
-// CHECK: [[ARR_IDX5:%.+]] = getelementptr inbounds i32, ptr {{.+}}, i64 0
-// CHECK: [[ARR_ADDR:%.+]] = getelementptr inbounds nuw %struct.S, ptr [[THIS]], i32 0, i32 3
-// CHECK: [[ARR_IDX6:%.+]] = getelementptr inbounds nuw [4 x i32], ptr [[ARR_ADDR]], i64 0, i64 0
-// CHECK: [[A_ADDR2:%.+]] = getelementptr inbounds nuw %struct.S, ptr [[THIS]], i32 0, i32 0
-// CHECK: [[P4:%.+]] = mul nuw i64 [[CONV:%.+]], 4
-// CHECK: [[A_ADDR3:%.+]] = getelementptr inbounds nuw %struct.S, ptr [[THIS]], i32 0, i32 0
-// CHECK: [[L5:%.+]] = load i32, ptr [[A_ADDR3]]
-// CHECK: [[L6:%.+]] = sext i32 [[L5]] to i64
-// CHECK: [[LB_ADD_LEN:%lb_add_len]] = add nsw i64 -1, [[L6]]
-// CHECK: [[ARR_ADDR9:%.+]] = getelementptr inbounds nuw %struct.S, ptr [[THIS]], i32 0, i32 3
-// CHECK: [[ARR_IDX10:%arrayidx.+]] = getelementptr inbounds nuw [4 x i32], ptr [[ARR_ADDR9]], i64 0, i64 %lb_add_len
-// CHECK: [[ARR_END:%.+]] = getelementptr i32, ptr [[ARR_IDX10]], i32 1
-// CHECK: [[E:%.+]] = ptrtoint ptr [[ARR_END]] to i64
-// CHECK: [[B:%.+]] = ptrtoint ptr [[A_ADDR]] to i64
-// CHECK: [[DIFF:%.+]] = sub i64 [[E]], [[B]]
-// CHECK: [[SZ:%.+]] = sdiv exact i64 [[DIFF]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64)
-// CHECK: [[BPTR0:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 0
-// CHECK: store ptr [[THIS]], ptr [[BPTR0]],
-// CHECK: [[PTR0:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 0
-// CHECK: store ptr [[A_ADDR]], ptr [[PTR0]],
-// CHECK: [[SIZE0:%.+]] = getelementptr inbounds [6 x i64], ptr [[SIZES]], i32 0, i32 0
-// CHECK: store i64 [[SZ]], ptr [[SIZE0]],
-// CHECK: [[BPTR1:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 1
-// CHECK: store ptr [[THIS]], ptr [[BPTR1]]
-// CHECK: [[PTR1:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 1
-// CHECK: store ptr [[A_ADDR]], ptr [[PTR1]],
-// CHECK: [[BPTR2:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 2
-// CHECK: store ptr [[PTR_ADDR]], ptr [[BPTR2]],
-// CHECK: [[PTR2:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 2
-// CHECK: store ptr [[ARR_IDX]], ptr [[PTR2]],
-// CHECK: [[BPTR3:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 3
-// CHECK: store ptr [[THIS]], ptr [[BPTR3]]
-// CHECK: [[PTR3:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 3
-// CHECK: store ptr [[REF_PTR]], ptr [[PTR3]],
-// CHECK: [[BPTR4:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 4
-// CHECK: store ptr [[P3]], ptr [[BPTR4]],
-// CHECK: [[PTR4:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 4
-// CHECK: store ptr [[ARR_IDX5]], ptr [[PTR4]]
-// CHECK: [[BPTR5:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 5
-// CHECK: store ptr [[THIS]], ptr [[BPTR5]], align 8
-// CHECK: [[PTR5:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 5
-// CHECK: store ptr [[ARR_IDX6]], ptr [[PTR5]], align 8
-// CHECK: [[SIZE1:%.+]] = getelementptr inbounds [6 x i64], ptr [[SIZES]], i32 0, i32 5
-// CHECK: store i64 [[P4]], ptr [[SIZE1]], align 8
-// CHECK: [[BPTR:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 0
-// CHECK: [[PTR:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 0
-// CHECK: [[SIZE:%.+]] = getelementptr inbounds [6 x i64], ptr [[SIZES]], i32 0, i32 0
-// CHECK: call void @__tgt_target_data_begin_mapper(ptr @{{.+}}, i64 -1, i32 6, ptr [[BPTR]], ptr [[PTR]], ptr [[SIZE]], ptr [[MAPTYPES2]], ptr null, ptr null)
-// CHECK: [[A_ADDR:%.+]] = load ptr, ptr [[BPTR1]],
-// CHECK: store ptr [[A_ADDR]], ptr [[A_REF:%.+]],
-// CHECK: [[PTR_ADDR:%.+]] = load ptr, ptr [[BPTR2]],
-// CHECK: store ptr [[PTR_ADDR]], ptr [[PTR_REF:%.+]],
-// CHECK: [[REF_PTR:%.+]] = load ptr, ptr [[BPTR3]],
-// CHECK: store ptr [[REF_PTR]], ptr [[REF_REF:%.+]],
-// CHECK: [[PTR_ADDR:%.+]] = load ptr, ptr [[BPTR2]],
-// CHECK: store ptr [[PTR_ADDR]], ptr [[PTR_REF2:%.+]],
-// CHECK: [[ARR_ADDR:%.+]] = load ptr, ptr [[BPTR5]],
-// CHECK: store ptr [[ARR_ADDR]], ptr [[ARR_REF:%.+]],
-// CHECK: [[A_ADDR:%.+]] = load ptr, ptr [[A_REF]],
-// CHECK: [[A:%.+]] = load i32, ptr [[A_ADDR]],
-// CHECK: [[INC:%.+]] = add nsw i32 [[A]], 1
-// CHECK: store i32 [[INC]], ptr [[A_ADDR]],
-// CHECK: [[PTR_PTR:%.+]] = load ptr, ptr [[PTR_REF2]],
-// CHECK: [[PTR:%.+]] = load ptr, ptr [[PTR_PTR]],
-// CHECK: [[VAL:%.+]] = load i32, ptr [[PTR]],
-// CHECK: [[INC:%.+]] = add nsw i32 [[VAL]], 1
-// CHECK: store i32 [[INC]], ptr [[PTR]],
-// CHECK: [[REF_PTR:%.+]] = load ptr, ptr [[REF_REF]],
-// CHECK: [[VAL:%.+]] = load i32, ptr [[REF_PTR]],
-// CHECK: [[INC:%.+]] = add nsw i32 [[VAL]], 1
-// CHECK: store i32 [[INC]], ptr [[REF_PTR]],
-// CHECK: [[ARR_ADDR:%.+]] = load ptr, ptr [[ARR_REF]],
-// CHECK: [[ARR0_ADDR:%.+]] = getelementptr inbounds [4 x i32], ptr [[ARR_ADDR]], i64 0, i64 0
-// CHECK: [[VAL:%.+]] = load i32, ptr [[ARR0_ADDR]],
-// CHECK: [[INC:%.+]] = add nsw i32 [[VAL]], 1
-// CHECK: store i32 [[INC]], ptr [[ARR0_ADDR]],
-// CHECK: [[BPTR:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 0
-// CHECK: [[PTR:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 0
-// CHECK: [[SIZE:%.+]] = getelementptr inbounds [6 x i64], ptr [[SIZES]], i32 0, i32 0
-// CHECK: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 -1, i32 6, ptr [[BPTR]], ptr [[PTR]], ptr [[SIZE]], ptr [[MAPTYPES2]], ptr null, ptr null)
#endif
+// CHECK-LABEL: define dso_local noundef signext i32 @main(
+// CHECK-SAME: ) #[[ATTR0:[0-9]+]] {
+// CHECK-NEXT: [[ENTRY:.*:]]
+// CHECK-NEXT: [[RETVAL:%.*]] = alloca i32, align 4
+// CHECK-NEXT: [[A:%.*]] = alloca float, align 4
+// CHECK-NEXT: [[PTR:%.*]] = alloca ptr, align 8
+// CHECK-NEXT: [[REF:%.*]] = alloca ptr, align 8
+// CHECK-NEXT: [[ARR:%.*]] = alloca [4 x float], align 4
+// CHECK-NEXT: [[SAVED_STACK:%.*]] = alloca ptr, align 8
+// CHECK-NEXT: [[__VLA_EXPR0:%.*]] = alloca i64, align 8
+// CHECK-NEXT: [[S:%.*]] = alloca [[STRUCT_S:%.*]], align 8
+// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [6 x ptr], align 8
+// CHECK-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [6 x ptr], align 8
+// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [6 x ptr], align 8
+// CHECK-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [6 x i64], align 8
+// CHECK-NEXT: [[TMP:%.*]] = alloca ptr, align 8
+// CHECK-NEXT: store i32 0, ptr [[RETVAL]], align 4
+// CHECK-NEXT: store float 0.000000e+00, ptr [[A]], align 4
+// CHECK-NEXT: store ptr [[A]], ptr [[PTR]], align 8
+// CHECK-NEXT: store ptr [[A]], ptr [[REF]], align 8
+// CHECK-NEXT: [[TMP0:%.*]] = load float, ptr [[A]], align 4
+// CHECK-NEXT: [[CONV:%.*]] = fptosi float [[TMP0]] to i32
+// CHECK-NEXT: [[TMP1:%.*]] = zext i32 [[CONV]] to i64
+// CHECK-NEXT: [[TMP2:%.*]] = call ptr @llvm.stacksave.p0()
+// CHECK-NEXT: store ptr [[TMP2]], ptr [[SAVED_STACK]], align 8
+// CHECK-NEXT: [[VLA:%.*]] = alloca float, i64 [[TMP1]], align 4
+// CHECK-NEXT: store i64 [[TMP1]], ptr [[__VLA_EXPR0]], align 8
+// CHECK-NEXT: call void @_ZN1SC1Ev(ptr noundef nonnull align 8 dereferenceable(40) [[S]])
+// CHECK-NEXT: call void @_ZN1S3fooEv(ptr noundef nonnull align 8 dereferenceable(40) [[S]])
+// CHECK-NEXT: [[TMP3:%.*]] = load ptr, ptr [[PTR]], align 8
+// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds nuw float, ptr [[TMP3]], i64 3
+// CHECK-NEXT: [[TMP4:%.*]] = load ptr, ptr [[PTR]], align 8
+// CHECK-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds float, ptr [[TMP4]], i64 0
+// CHECK-NEXT: [[TMP5:%.*]] = load ptr, ptr [[REF]], align 8, !nonnull [[META3:![0-9]+]], !align [[META4:![0-9]+]]
+// CHECK-NEXT: [[TMP6:%.*]] = load ptr, ptr [[REF]], align 8, !nonnull [[META3]], !align [[META4]]
+// CHECK-NEXT: [[ARRAYIDX2:%.*]] = getelementptr inbounds nuw [4 x float], ptr [[ARR]], i64 0, i64 0
+// CHECK-NEXT: [[TMP7:%.*]] = load float, ptr [[A]], align 4
+// CHECK-NEXT: [[CONV3:%.*]] = fptosi float [[TMP7]] to i32
+// CHECK-NEXT: [[CONV4:%.*]] = sext i32 [[CONV3]] to i64
+// CHECK-NEXT: [[TMP8:%.*]] = mul nuw i64 [[CONV4]], 4
+// CHECK-NEXT: [[ARRAYIDX5:%.*]] = getelementptr inbounds float, ptr [[VLA]], i64 0
+// CHECK-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[DOTOFFLOAD_SIZES]], ptr align 8 @.offload_sizes, i64 48, i1 false)
+// CHECK-NEXT: [[TMP9:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
+// CHECK-NEXT: store ptr [[A]], ptr [[TMP9]], align 8
+// CHECK-NEXT: [[TMP10:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// CHECK-NEXT: store ptr [[A]], ptr [[TMP10]], align 8
+// CHECK-NEXT: [[TMP11:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0
+// CHECK-NEXT: store ptr null, ptr [[TMP11]], align 8
+// CHECK-NEXT: [[TMP12:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1
+// CHECK-NEXT: store ptr [[PTR]], ptr [[TMP12]], align 8
+// CHECK-NEXT: [[TMP13:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1
+// CHECK-NEXT: store ptr [[ARRAYIDX]], ptr [[TMP13]], align 8
+// CHECK-NEXT: [[TMP14:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1
+// CHECK-NEXT: store ptr null, ptr [[TMP14]], align 8
+// CHECK-NEXT: [[TMP15:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2
+// CHECK-NEXT: store ptr [[PTR]], ptr [[TMP15]], align 8
+// CHECK-NEXT: [[TMP16:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2
+// CHECK-NEXT: store ptr [[ARRAYIDX1]], ptr [[TMP16]], align 8
+// CHECK-NEXT: [[TMP17:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2
+// CHECK-NEXT: store ptr null, ptr [[TMP17]], align 8
+// CHECK-NEXT: [[TMP18:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 3
+// CHECK-NEXT: store ptr [[TMP5]], ptr [[TMP18]], align 8
+// CHECK-NEXT: [[TMP19:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 3
+// CHECK-NEXT: store ptr [[TMP6]], ptr [[TMP19]], align 8
+// CHECK-NEXT: [[TMP20:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 3
+// CHECK-NEXT: store ptr null, ptr [[TMP20]], align 8
+// CHECK-NEXT: [[TMP21:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 4
+// CHECK-NEXT: store ptr [[ARR]], ptr [[TMP21]], align 8
+// CHECK-NEXT: [[TMP22:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 4
+// CHECK-NEXT: store ptr [[ARRAYIDX2]], ptr [[TMP22]], align 8
+// CHECK-NEXT: [[TMP23:%.*]] = getelementptr inbounds [6 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 4
+// CHECK-NEXT: store i64 [[TMP8]], ptr [[TMP23]], align 8
+// CHECK-NEXT: [[TMP24:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 4
+// CHECK-NEXT: store ptr null, ptr [[TMP24]], align 8
+// CHECK-NEXT: [[TMP25:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 5
+// CHECK-NEXT: store ptr [[VLA]], ptr [[TMP25]], align 8
+// CHECK-NEXT: [[TMP26:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 5
+// CHECK-NEXT: store ptr [[ARRAYIDX5]], ptr [[TMP26]], align 8
+// CHECK-NEXT: [[TMP27:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 5
+// CHECK-NEXT: store ptr null, ptr [[TMP27]], align 8
+// CHECK-NEXT: [[TMP28:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
+// CHECK-NEXT: [[TMP29:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// CHECK-NEXT: [[TMP30:%.*]] = getelementptr inbounds [6 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 0
+// CHECK-NEXT: call void @__tgt_target_data_begin_mapper(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 6, ptr [[TMP28]], ptr [[TMP29]], ptr [[TMP30]], ptr @.offload_maptypes, ptr null, ptr null)
+// CHECK-NEXT: [[TMP31:%.*]] = load ptr, ptr [[TMP9]], align 8
+// CHECK-NEXT: [[TMP32:%.*]] = load ptr, ptr [[TMP18]], align 8
+// CHECK-NEXT: store ptr [[TMP32]], ptr [[TMP]], align 8
+// CHECK-NEXT: [[TMP33:%.*]] = load ptr, ptr [[TMP21]], align 8
+// CHECK-NEXT: [[TMP34:%.*]] = load ptr, ptr [[TMP25]], align 8
+// CHECK-NEXT: [[TMP35:%.*]] = load float, ptr [[TMP31]], align 4
+// CHECK-NEXT: [[INC:%.*]] = fadd float [[TMP35]], 1.000000e+00
+// CHECK-NEXT: store float [[INC]], ptr [[TMP31]], align 4
+// CHECK-NEXT: [[TMP36:%.*]] = load ptr, ptr [[TMP12]], align 8
+// CHECK-NEXT: [[TMP37:%.*]] = load float, ptr [[TMP36]], align 4
+// CHECK-NEXT: [[INC6:%.*]] = fadd float [[TMP37]], 1.000000e+00
+// CHECK-NEXT: store float [[INC6]], ptr [[TMP36]], align 4
+// CHECK-NEXT: [[TMP38:%.*]] = load ptr, ptr [[TMP]], align 8, !nonnull [[META3]], !align [[META4]]
+// CHECK-NEXT: [[TMP39:%.*]] = load float, ptr [[TMP38]], align 4
+// CHECK-NEXT: [[INC7:%.*]] = fadd float [[TMP39]], 1.000000e+00
+// CHECK-NEXT: store float [[INC7]], ptr [[TMP38]], align 4
+// CHECK-NEXT: [[ARRAYIDX8:%.*]] = getelementptr inbounds [4 x float], ptr [[TMP33]], i64 0, i64 0
+// CHECK-NEXT: [[TMP40:%.*]] = load float, ptr [[ARRAYIDX8]], align 4
+// CHECK-NEXT: [[INC9:%.*]] = fadd float [[TMP40]], 1.000000e+00
+// CHECK-NEXT: store float [[INC9]], ptr [[ARRAYIDX8]], align 4
+// CHECK-NEXT: [[ARRAYIDX10:%.*]] = getelementptr inbounds float, ptr [[TMP34]], i64 0
+// CHECK-NEXT: [[TMP41:%.*]] = load float, ptr [[ARRAYIDX10]], align 4
+// CHECK-NEXT: [[INC11:%.*]] = fadd float [[TMP41]], 1.000000e+00
+// CHECK-NEXT: store float [[INC11]], ptr [[ARRAYIDX10]], align 4
+// CHECK-NEXT: [[TMP42:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
+// CHECK-NEXT: [[TMP43:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// CHECK-NEXT: [[TMP44:%.*]] = getelementptr inbounds [6 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 0
+// CHECK-NEXT: call void @__tgt_target_data_end_mapper(ptr @[[GLOB1]], i64 -1, i32 6, ptr [[TMP42]], ptr [[TMP43]], ptr [[TMP44]], ptr @.offload_maptypes, ptr null, ptr null)
+// CHECK-NEXT: [[TMP45:%.*]] = load float, ptr [[A]], align 4
+// CHECK-NEXT: [[CONV12:%.*]] = fptosi float [[TMP45]] to i32
+// CHECK-NEXT: store i32 [[CONV12]], ptr [[RETVAL]], align 4
+// CHECK-NEXT: [[TMP46:%.*]] = load ptr, ptr [[SAVED_STACK]], align 8
+// CHECK-NEXT: call void @llvm.stackrestore.p0(ptr [[TMP46]])
+// CHECK-NEXT: [[TMP47:%.*]] = load i32, ptr [[RETVAL]], align 4
+// CHECK-NEXT: ret i32 [[TMP47]]
+//
+//
+// CHECK-LABEL: define linkonce_odr void @_ZN1SC1Ev(
+// CHECK-SAME: ptr noundef nonnull align 8 dereferenceable(40) [[THIS:%.*]]) unnamed_addr #[[ATTR2:[0-9]+]] comdat {
+// CHECK-NEXT: [[ENTRY:.*:]]
+// CHECK-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8
+// CHECK-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
+// CHECK-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
+// CHECK-NEXT: call void @_ZN1SC2Ev(ptr noundef nonnull align 8 dereferenceable(40) [[THIS1]])
+// CHECK-NEXT: ret void
+//
+//
+// CHECK-LABEL: define linkonce_odr void @_ZN1S3fooEv(
+// CHECK-SAME: ptr noundef nonnull align 8 dereferenceable(40) [[THIS:%.*]]) #[[ATTR2]] comdat {
+// CHECK-NEXT: [[ENTRY:.*:]]
+// CHECK-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8
+// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [6 x ptr], align 8
+// CHECK-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [6 x ptr], align 8
+// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [6 x ptr], align 8
+// CHECK-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [6 x i64], align 8
+// CHECK-NEXT: [[TMP:%.*]] = alloca ptr, align 8
+// CHECK-NEXT: [[_TMP11:%.*]] = alloca ptr, align 8
+// CHECK-NEXT: [[_TMP12:%.*]] = alloca ptr, align 8
+// CHECK-NEXT: [[_TMP13:%.*]] = alloca ptr, align 8
+// CHECK-NEXT: [[_TMP14:%.*]] = alloca ptr, align 8
+// CHECK-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
+// CHECK-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
+// CHECK-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[THIS1]], i32 0, i32 0
+// CHECK-NEXT: [[PTR:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 1
+// CHECK-NEXT: [[PTR2:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 1
+// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[PTR2]], align 8
+// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds nuw i32, ptr [[TMP0]], i64 3
+// CHECK-NEXT: [[REF:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 2
+// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[REF]], align 8, !nonnull [[META3]], !align [[META4]]
+// CHECK-NEXT: [[PTR3:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 1
+// CHECK-NEXT: [[PTR4:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 1
+// CHECK-NEXT: [[TMP2:%.*]] = load ptr, ptr [[PTR4]], align 8
+// CHECK-NEXT: [[ARRAYIDX5:%.*]] = getelementptr inbounds i32, ptr [[TMP2]], i64 0
+// CHECK-NEXT: [[ARR:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 3
+// CHECK-NEXT: [[ARRAYIDX6:%.*]] = getelementptr inbounds nuw [4 x i32], ptr [[ARR]], i64 0, i64 0
+// CHECK-NEXT: [[A7:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 0
+// CHECK-NEXT: [[TMP3:%.*]] = load i32, ptr [[A7]], align 8
+// CHECK-NEXT: [[CONV:%.*]] = sext i32 [[TMP3]] to i64
+// CHECK-NEXT: [[TMP4:%.*]] = mul nuw i64 [[CONV]], 4
+// CHECK-NEXT: [[A8:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 0
+// CHECK-NEXT: [[TMP5:%.*]] = load i32, ptr [[A8]], align 8
+// CHECK-NEXT: [[TMP6:%.*]] = sext i32 [[TMP5]] to i64
+// CHECK-NEXT: [[LB_ADD_LEN:%.*]] = add nsw i64 -1, [[TMP6]]
+// CHECK-NEXT: [[ARR9:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 3
+// CHECK-NEXT: [[ARRAYIDX10:%.*]] = getelementptr inbounds nuw [4 x i32], ptr [[ARR9]], i64 0, i64 [[LB_ADD_LEN]]
+// CHECK-NEXT: [[TMP7:%.*]] = getelementptr i32, ptr [[ARRAYIDX10]], i32 1
+// CHECK-NEXT: [[TMP8:%.*]] = ptrtoint ptr [[TMP7]] to i64
+// CHECK-NEXT: [[TMP9:%.*]] = ptrtoint ptr [[A]] to i64
+// CHECK-NEXT: [[TMP10:%.*]] = sub i64 [[TMP8]], [[TMP9]]
+// CHECK-NEXT: [[TMP11:%.*]] = sdiv exact i64 [[TMP10]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64)
+// CHECK-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[DOTOFFLOAD_SIZES]], ptr align 8 @.offload_sizes.1, i64 48, i1 false)
+// CHECK-NEXT: [[TMP12:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
+// CHECK-NEXT: store ptr [[THIS1]], ptr [[TMP12]], align 8
+// CHECK-NEXT: [[TMP13:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// CHECK-NEXT: store ptr [[A]], ptr [[TMP13]], align 8
+// CHECK-NEXT: [[TMP14:%.*]] = getelementptr inbounds [6 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 0
+// CHECK-NEXT: store i64 [[TMP11]], ptr [[TMP14]], align 8
+// CHECK-NEXT: [[TMP15:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0
+// CHECK-NEXT: store ptr null, ptr [[TMP15]], align 8
+// CHECK-NEXT: [[TMP16:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1
+// CHECK-NEXT: store ptr [[THIS1]], ptr [[TMP16]], align 8
+// CHECK-NEXT: [[TMP17:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1
+// CHECK-NEXT: store ptr [[A]], ptr [[TMP17]], align 8
+// CHECK-NEXT: [[TMP18:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1
+// CHECK-NEXT: store ptr null, ptr [[TMP18]], align 8
+// CHECK-NEXT: [[TMP19:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2
+// CHECK-NEXT: store ptr [[PTR]], ptr [[TMP19]], align 8
+// CHECK-NEXT: [[TMP20:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2
+// CHECK-NEXT: store ptr [[ARRAYIDX]], ptr [[TMP20]], align 8
+// CHECK-NEXT: [[TMP21:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2
+// CHECK-NEXT: store ptr null, ptr [[TMP21]], align 8
+// CHECK-NEXT: [[TMP22:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 3
+// CHECK-NEXT: store ptr [[THIS1]], ptr [[TMP22]], align 8
+// CHECK-NEXT: [[TMP23:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 3
+// CHECK-NEXT: store ptr [[TMP1]], ptr [[TMP23]], align 8
+// CHECK-NEXT: [[TMP24:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 3
+// CHECK-NEXT: store ptr null, ptr [[TMP24]], align 8
+// CHECK-NEXT: [[TMP25:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 4
+// CHECK-NEXT: store ptr [[PTR3]], ptr [[TMP25]], align 8
+// CHECK-NEXT: [[TMP26:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 4
+// CHECK-NEXT: store ptr [[ARRAYIDX5]], ptr [[TMP26]], align 8
+// CHECK-NEXT: [[TMP27:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 4
+// CHECK-NEXT: store ptr null, ptr [[TMP27]], align 8
+// CHECK-NEXT: [[TMP28:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 5
+// CHECK-NEXT: store ptr [[THIS1]], ptr [[TMP28]], align 8
+// CHECK-NEXT: [[TMP29:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 5
+// CHECK-NEXT: store ptr [[ARRAYIDX6]], ptr [[TMP29]], align 8
+// CHECK-NEXT: [[TMP30:%.*]] = getelementptr inbounds [6 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 5
+// CHECK-NEXT: store i64 [[TMP4]], ptr [[TMP30]], align 8
+// CHECK-NEXT: [[TMP31:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 5
+// CHECK-NEXT: store ptr null, ptr [[TMP31]], align 8
+// CHECK-NEXT: [[TMP32:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
+// CHECK-NEXT: [[TMP33:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// CHECK-NEXT: [[TMP34:%.*]] = getelementptr inbounds [6 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 0
+// CHECK-NEXT: call void @__tgt_target_data_begin_mapper(ptr @[[GLOB1]], i64 -1, i32 6, ptr [[TMP32]], ptr [[TMP33]], ptr [[TMP34]], ptr @.offload_maptypes.2, ptr null, ptr null)
+// CHECK-NEXT: [[TMP35:%.*]] = load ptr, ptr [[TMP16]], align 8
+// CHECK-NEXT: store ptr [[TMP35]], ptr [[TMP]], align 8
+// CHECK-NEXT: [[TMP36:%.*]] = load ptr, ptr [[TMP19]], align 8
+// CHECK-NEXT: store ptr [[TMP36]], ptr [[_TMP11]], align 8
+// CHECK-NEXT: [[TMP37:%.*]] = load ptr, ptr [[TMP22]], align 8
+// CHECK-NEXT: store ptr [[TMP37]], ptr [[_TMP12]], align 8
+// CHECK-NEXT: [[TMP38:%.*]] = load ptr, ptr [[TMP19]], align 8
+// CHECK-NEXT: store ptr [[TMP38]], ptr [[_TMP13]], align 8
+// CHECK-NEXT: [[TMP39:%.*]] = load ptr, ptr [[TMP28]], align 8
+// CHECK-NEXT: store ptr [[TMP39]], ptr [[_TMP14]], align 8
+// CHECK-NEXT: [[TMP40:%.*]] = load ptr, ptr [[TMP]], align 8, !nonnull [[META3]], !align [[META4]]
+// CHECK-NEXT: [[TMP41:%.*]] = load i32, ptr [[TMP40]], align 4
+// CHECK-NEXT: [[INC:%.*]] = add nsw i32 [[TMP41]], 1
+// CHECK-NEXT: store i32 [[INC]], ptr [[TMP40]], align 4
+// CHECK-NEXT: [[TMP42:%.*]] = load ptr, ptr [[_TMP13]], align 8, !nonnull [[META3]], !align [[META5:![0-9]+]]
+// CHECK-NEXT: [[TMP43:%.*]] = load ptr, ptr [[TMP42]], align 8
+// CHECK-NEXT: [[TMP44:%.*]] = load i32, ptr [[TMP43]], align 4
+// CHECK-NEXT: [[INC15:%.*]] = add nsw i32 [[TMP44]], 1
+// CHECK-NEXT: store i32 [[INC15]], ptr [[TMP43]], align 4
+// CHECK-NEXT: [[TMP45:%.*]] = load ptr, ptr [[_TMP12]], align 8, !nonnull [[META3]], !align [[META4]]
+// CHECK-NEXT: [[TMP46:%.*]] = load i32, ptr [[TMP45]], align 4
+// CHECK-NEXT: [[INC16:%.*]] = add nsw i32 [[TMP46]], 1
+// CHECK-NEXT: store i32 [[INC16]], ptr [[TMP45]], align 4
+// CHECK-NEXT: [[TMP47:%.*]] = load ptr, ptr [[_TMP14]], align 8, !nonnull [[META3]], !align [[META4]]
+// CHECK-NEXT: [[ARRAYIDX17:%.*]] = getelementptr inbounds [4 x i32], ptr [[TMP47]], i64 0, i64 0
+// CHECK-NEXT: [[TMP48:%.*]] = load i32, ptr [[ARRAYIDX17]], align 4
+// CHECK-NEXT: [[INC18:%.*]] = add nsw i32 [[TMP48]], 1
+// CHECK-NEXT: store i32 [[INC18]], ptr [[ARRAYIDX17]], align 4
+// CHECK-NEXT: [[TMP49:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
+// CHECK-NEXT: [[TMP50:%.*]] = getelementptr inbounds [6 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// CHECK-NEXT: [[TMP51:%.*]] = getelementptr inbounds [6 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 0
+// CHECK-NEXT: call void @__tgt_target_data_end_mapper(ptr @[[GLOB1]], i64 -1, i32 6, ptr [[TMP49]], ptr [[TMP50]], ptr [[TMP51]], ptr @.offload_maptypes.2, ptr null, ptr null)
+// CHECK-NEXT: ret void
+//
+//
+// CHECK-LABEL: define linkonce_odr void @_ZN1SC2Ev(
+// CHECK-SAME: ptr noundef nonnull align 8 dereferenceable(40) [[THIS:%.*]]) unnamed_addr #[[ATTR2]] comdat {
+// CHECK-NEXT: [[ENTRY:.*:]]
+// CHECK-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8
+// CHECK-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
+// CHECK-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
+// CHECK-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[THIS1]], i32 0, i32 0
+// CHECK-NEXT: store i32 0, ptr [[A]], align 8
+// CHECK-NEXT: [[PTR:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 1
+// CHECK-NEXT: [[A2:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 0
+// CHECK-NEXT: store ptr [[A2]], ptr [[PTR]], align 8
+// CHECK-NEXT: [[REF:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 2
+// CHECK-NEXT: [[A3:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 0
+// CHECK-NEXT: store ptr [[A3]], ptr [[REF]], align 8
+// CHECK-NEXT: ret void
+//
+//
+// OMP60-LABEL: define dso_local noundef signext i32 @main(
+// OMP60-SAME: i32 noundef signext [[ARGC:%.*]], ptr noundef [[ARGV:%.*]]) #[[ATTR0:[0-9]+]] {
+// OMP60-NEXT: [[ENTRY:.*:]]
+// OMP60-NEXT: [[RETVAL:%.*]] = alloca i32, align 4
+// OMP60-NEXT: [[ARGC_ADDR:%.*]] = alloca i32, align 4
+// OMP60-NEXT: [[ARGV_ADDR:%.*]] = alloca ptr, align 8
+// OMP60-NEXT: [[A:%.*]] = alloca float, align 4
+// OMP60-NEXT: [[PTR:%.*]] = alloca ptr, align 8
+// OMP60-NEXT: [[REF:%.*]] = alloca ptr, align 8
+// OMP60-NEXT: [[ARR:%.*]] = alloca [4 x float], align 4
+// OMP60-NEXT: [[SAVED_STACK:%.*]] = alloca ptr, align 8
+// OMP60-NEXT: [[__VLA_EXPR0:%.*]] = alloca i64, align 8
+// OMP60-NEXT: [[S:%.*]] = alloca [[STRUCT_S:%.*]], align 8
+// OMP60-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [9 x ptr], align 8
+// OMP60-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [9 x ptr], align 8
+// OMP60-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [9 x ptr], align 8
+// OMP60-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [9 x i64], align 8
+// OMP60-NEXT: [[TMP:%.*]] = alloca ptr, align 8
+// OMP60-NEXT: store i32 0, ptr [[RETVAL]], align 4
+// OMP60-NEXT: store i32 [[ARGC]], ptr [[ARGC_ADDR]], align 4
+// OMP60-NEXT: store ptr [[ARGV]], ptr [[ARGV_ADDR]], align 8
+// OMP60-NEXT: store float 0.000000e+00, ptr [[A]], align 4
+// OMP60-NEXT: store ptr [[A]], ptr [[PTR]], align 8
+// OMP60-NEXT: store ptr [[A]], ptr [[REF]], align 8
+// OMP60-NEXT: [[TMP0:%.*]] = load float, ptr [[A]], align 4
+// OMP60-NEXT: [[CONV:%.*]] = fptosi float [[TMP0]] to i32
+// OMP60-NEXT: [[TMP1:%.*]] = zext i32 [[CONV]] to i64
+// OMP60-NEXT: [[TMP2:%.*]] = call ptr @llvm.stacksave.p0()
+// OMP60-NEXT: store ptr [[TMP2]], ptr [[SAVED_STACK]], align 8
+// OMP60-NEXT: [[VLA:%.*]] = alloca float, i64 [[TMP1]], align 4
+// OMP60-NEXT: store i64 [[TMP1]], ptr [[__VLA_EXPR0]], align 8
+// OMP60-NEXT: call void @_ZN1SC1Ev(ptr noundef nonnull align 8 dereferenceable(40) [[S]])
+// OMP60-NEXT: call void @_ZN1S3fooEv(ptr noundef nonnull align 8 dereferenceable(40) [[S]])
+// OMP60-NEXT: [[TMP3:%.*]] = load ptr, ptr [[PTR]], align 8
+// OMP60-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds nuw float, ptr [[TMP3]], i64 3
+// OMP60-NEXT: [[TMP4:%.*]] = load ptr, ptr [[PTR]], align 8
+// OMP60-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds float, ptr [[TMP4]], i64 0
+// OMP60-NEXT: [[TMP5:%.*]] = load ptr, ptr [[PTR]], align 8
+// OMP60-NEXT: [[ARRAYIDX2:%.*]] = getelementptr inbounds nuw float, ptr [[TMP5]], i64 0
+// OMP60-NEXT: [[TMP6:%.*]] = load ptr, ptr [[REF]], align 8, !nonnull [[META3:![0-9]+]], !align [[META4:![0-9]+]]
+// OMP60-NEXT: [[TMP7:%.*]] = load ptr, ptr [[REF]], align 8, !nonnull [[META3]], !align [[META4]]
+// OMP60-NEXT: [[ARRAYIDX3:%.*]] = getelementptr inbounds nuw [4 x float], ptr [[ARR]], i64 0, i64 0
+// OMP60-NEXT: [[TMP8:%.*]] = load float, ptr [[A]], align 4
+// OMP60-NEXT: [[CONV4:%.*]] = fptosi float [[TMP8]] to i32
+// OMP60-NEXT: [[CONV5:%.*]] = sext i32 [[CONV4]] to i64
+// OMP60-NEXT: [[TMP9:%.*]] = mul nuw i64 [[CONV5]], 4
+// OMP60-NEXT: [[ARRAYIDX6:%.*]] = getelementptr inbounds float, ptr [[VLA]], i64 0
+// OMP60-NEXT: [[TMP10:%.*]] = load ptr, ptr [[ARGV_ADDR]], align 8
+// OMP60-NEXT: [[TMP11:%.*]] = load ptr, ptr [[ARGV_ADDR]], align 8
+// OMP60-NEXT: [[ARRAYIDX7:%.*]] = getelementptr inbounds ptr, ptr [[TMP11]], i64 0
+// OMP60-NEXT: [[TMP12:%.*]] = load ptr, ptr [[ARGV_ADDR]], align 8
+// OMP60-NEXT: [[ARRAYIDX8:%.*]] = getelementptr inbounds ptr, ptr [[TMP12]], i64 0
+// OMP60-NEXT: [[TMP13:%.*]] = load ptr, ptr [[ARRAYIDX8]], align 8
+// OMP60-NEXT: [[ARRAYIDX9:%.*]] = getelementptr inbounds nuw i8, ptr [[TMP13]], i64 0
+// OMP60-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[DOTOFFLOAD_SIZES]], ptr align 8 @.offload_sizes, i64 72, i1 false)
+// OMP60-NEXT: [[TMP14:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
+// OMP60-NEXT: store ptr [[A]], ptr [[TMP14]], align 8
+// OMP60-NEXT: [[TMP15:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// OMP60-NEXT: store ptr [[A]], ptr [[TMP15]], align 8
+// OMP60-NEXT: [[TMP16:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0
+// OMP60-NEXT: store ptr null, ptr [[TMP16]], align 8
+// OMP60-NEXT: [[TMP17:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1
+// OMP60-NEXT: store ptr [[PTR]], ptr [[TMP17]], align 8
+// OMP60-NEXT: [[TMP18:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1
+// OMP60-NEXT: store ptr [[ARRAYIDX]], ptr [[TMP18]], align 8
+// OMP60-NEXT: [[TMP19:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1
+// OMP60-NEXT: store ptr null, ptr [[TMP19]], align 8
+// OMP60-NEXT: [[TMP20:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2
+// OMP60-NEXT: store ptr [[PTR]], ptr [[TMP20]], align 8
+// OMP60-NEXT: [[TMP21:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2
+// OMP60-NEXT: store ptr [[ARRAYIDX1]], ptr [[TMP21]], align 8
+// OMP60-NEXT: [[TMP22:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2
+// OMP60-NEXT: store ptr null, ptr [[TMP22]], align 8
+// OMP60-NEXT: [[TMP23:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 3
+// OMP60-NEXT: store ptr [[PTR]], ptr [[TMP23]], align 8
+// OMP60-NEXT: [[TMP24:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 3
+// OMP60-NEXT: store ptr [[ARRAYIDX2]], ptr [[TMP24]], align 8
+// OMP60-NEXT: [[TMP25:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 3
+// OMP60-NEXT: store ptr null, ptr [[TMP25]], align 8
+// OMP60-NEXT: [[TMP26:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 4
+// OMP60-NEXT: store ptr [[TMP6]], ptr [[TMP26]], align 8
+// OMP60-NEXT: [[TMP27:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 4
+// OMP60-NEXT: store ptr [[TMP7]], ptr [[TMP27]], align 8
+// OMP60-NEXT: [[TMP28:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 4
+// OMP60-NEXT: store ptr null, ptr [[TMP28]], align 8
+// OMP60-NEXT: [[TMP29:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 5
+// OMP60-NEXT: store ptr [[ARR]], ptr [[TMP29]], align 8
+// OMP60-NEXT: [[TMP30:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 5
+// OMP60-NEXT: store ptr [[ARRAYIDX3]], ptr [[TMP30]], align 8
+// OMP60-NEXT: [[TMP31:%.*]] = getelementptr inbounds [9 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 5
+// OMP60-NEXT: store i64 [[TMP9]], ptr [[TMP31]], align 8
+// OMP60-NEXT: [[TMP32:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 5
+// OMP60-NEXT: store ptr null, ptr [[TMP32]], align 8
+// OMP60-NEXT: [[TMP33:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 6
+// OMP60-NEXT: store ptr [[VLA]], ptr [[TMP33]], align 8
+// OMP60-NEXT: [[TMP34:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 6
+// OMP60-NEXT: store ptr [[ARRAYIDX6]], ptr [[TMP34]], align 8
+// OMP60-NEXT: [[TMP35:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 6
+// OMP60-NEXT: store ptr null, ptr [[TMP35]], align 8
+// OMP60-NEXT: [[TMP36:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 7
+// OMP60-NEXT: store ptr [[TMP10]], ptr [[TMP36]], align 8
+// OMP60-NEXT: [[TMP37:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 7
+// OMP60-NEXT: store ptr [[ARRAYIDX7]], ptr [[TMP37]], align 8
+// OMP60-NEXT: [[TMP38:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 7
+// OMP60-NEXT: store ptr null, ptr [[TMP38]], align 8
+// OMP60-NEXT: [[TMP39:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 8
+// OMP60-NEXT: store ptr [[ARRAYIDX7]], ptr [[TMP39]], align 8
+// OMP60-NEXT: [[TMP40:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 8
+// OMP60-NEXT: store ptr [[ARRAYIDX9]], ptr [[TMP40]], align 8
+// OMP60-NEXT: [[TMP41:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 8
+// OMP60-NEXT: store ptr null, ptr [[TMP41]], align 8
+// OMP60-NEXT: [[TMP42:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
+// OMP60-NEXT: [[TMP43:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// OMP60-NEXT: [[TMP44:%.*]] = getelementptr inbounds [9 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 0
+// OMP60-NEXT: call void @__tgt_target_data_begin_mapper(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 9, ptr [[TMP42]], ptr [[TMP43]], ptr [[TMP44]], ptr @.offload_maptypes, ptr null, ptr null)
+// OMP60-NEXT: [[TMP45:%.*]] = load ptr, ptr [[TMP14]], align 8
+// OMP60-NEXT: [[TMP46:%.*]] = load ptr, ptr [[TMP26]], align 8
+// OMP60-NEXT: store ptr [[TMP46]], ptr [[TMP]], align 8
+// OMP60-NEXT: [[TMP47:%.*]] = load ptr, ptr [[TMP29]], align 8
+// OMP60-NEXT: [[TMP48:%.*]] = load ptr, ptr [[TMP33]], align 8
+// OMP60-NEXT: [[TMP49:%.*]] = load float, ptr [[TMP45]], align 4
+// OMP60-NEXT: [[INC:%.*]] = fadd float [[TMP49]], 1.000000e+00
+// OMP60-NEXT: store float [[INC]], ptr [[TMP45]], align 4
+// OMP60-NEXT: [[TMP50:%.*]] = load ptr, ptr [[TMP17]], align 8
+// OMP60-NEXT: [[TMP51:%.*]] = load float, ptr [[TMP50]], align 4
+// OMP60-NEXT: [[INC10:%.*]] = fadd float [[TMP51]], 1.000000e+00
+// OMP60-NEXT: store float [[INC10]], ptr [[TMP50]], align 4
+// OMP60-NEXT: [[TMP52:%.*]] = load ptr, ptr [[TMP]], align 8, !nonnull [[META3]], !align [[META4]]
+// OMP60-NEXT: [[TMP53:%.*]] = load float, ptr [[TMP52]], align 4
+// OMP60-NEXT: [[INC11:%.*]] = fadd float [[TMP53]], 1.000000e+00
+// OMP60-NEXT: store float [[INC11]], ptr [[TMP52]], align 4
+// OMP60-NEXT: [[ARRAYIDX12:%.*]] = getelementptr inbounds [4 x float], ptr [[TMP47]], i64 0, i64 0
+// OMP60-NEXT: [[TMP54:%.*]] = load float, ptr [[ARRAYIDX12]], align 4
+// OMP60-NEXT: [[INC13:%.*]] = fadd float [[TMP54]], 1.000000e+00
+// OMP60-NEXT: store float [[INC13]], ptr [[ARRAYIDX12]], align 4
+// OMP60-NEXT: [[ARRAYIDX14:%.*]] = getelementptr inbounds float, ptr [[TMP48]], i64 0
+// OMP60-NEXT: [[TMP55:%.*]] = load float, ptr [[ARRAYIDX14]], align 4
+// OMP60-NEXT: [[INC15:%.*]] = fadd float [[TMP55]], 1.000000e+00
+// OMP60-NEXT: store float [[INC15]], ptr [[ARRAYIDX14]], align 4
+// OMP60-NEXT: [[TMP56:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
+// OMP60-NEXT: [[TMP57:%.*]] = getelementptr inbounds [9 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// OMP60-NEXT: [[TMP58:%.*]] = getelementptr inbounds [9 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 0
+// OMP60-NEXT: call void @__tgt_target_data_end_mapper(ptr @[[GLOB1]], i64 -1, i32 9, ptr [[TMP56]], ptr [[TMP57]], ptr [[TMP58]], ptr @.offload_maptypes, ptr null, ptr null)
+// OMP60-NEXT: [[TMP59:%.*]] = load float, ptr [[A]], align 4
+// OMP60-NEXT: [[CONV16:%.*]] = fptosi float [[TMP59]] to i32
+// OMP60-NEXT: store i32 [[CONV16]], ptr [[RETVAL]], align 4
+// OMP60-NEXT: [[TMP60:%.*]] = load ptr, ptr [[SAVED_STACK]], align 8
+// OMP60-NEXT: call void @llvm.stackrestore.p0(ptr [[TMP60]])
+// OMP60-NEXT: [[TMP61:%.*]] = load i32, ptr [[RETVAL]], align 4
+// OMP60-NEXT: ret i32 [[TMP61]]
+//
+//
+// OMP60-LABEL: define linkonce_odr void @_ZN1SC1Ev(
+// OMP60-SAME: ptr noundef nonnull align 8 dereferenceable(40) [[THIS:%.*]]) unnamed_addr #[[ATTR2:[0-9]+]] comdat {
+// OMP60-NEXT: [[ENTRY:.*:]]
+// OMP60-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8
+// OMP60-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
+// OMP60-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
+// OMP60-NEXT: call void @_ZN1SC2Ev(ptr noundef nonnull align 8 dereferenceable(40) [[THIS1]])
+// OMP60-NEXT: ret void
+//
+//
+// OMP60-LABEL: define linkonce_odr void @_ZN1S3fooEv(
+// OMP60-SAME: ptr noundef nonnull align 8 dereferenceable(40) [[THIS:%.*]]) #[[ATTR2]] comdat {
+// OMP60-NEXT: [[ENTRY:.*:]]
+// OMP60-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8
+// OMP60-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [7 x ptr], align 8
+// OMP60-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [7 x ptr], align 8
+// OMP60-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [7 x ptr], align 8
+// OMP60-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [7 x i64], align 8
+// OMP60-NEXT: [[TMP:%.*]] = alloca ptr, align 8
+// OMP60-NEXT: [[_TMP14:%.*]] = alloca ptr, align 8
+// OMP60-NEXT: [[_TMP15:%.*]] = alloca ptr, align 8
+// OMP60-NEXT: [[_TMP16:%.*]] = alloca ptr, align 8
+// OMP60-NEXT: [[_TMP17:%.*]] = alloca ptr, align 8
+// OMP60-NEXT: [[_TMP18:%.*]] = alloca ptr, align 8
+// OMP60-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
+// OMP60-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
+// OMP60-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[THIS1]], i32 0, i32 0
+// OMP60-NEXT: [[PTR:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 1
+// OMP60-NEXT: [[PTR2:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 1
+// OMP60-NEXT: [[TMP0:%.*]] = load ptr, ptr [[PTR2]], align 8
+// OMP60-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds nuw i32, ptr [[TMP0]], i64 3
+// OMP60-NEXT: [[REF:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 2
+// OMP60-NEXT: [[TMP1:%.*]] = load ptr, ptr [[REF]], align 8, !nonnull [[META3]], !align [[META4]]
+// OMP60-NEXT: [[PTR3:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 1
+// OMP60-NEXT: [[PTR4:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 1
+// OMP60-NEXT: [[TMP2:%.*]] = load ptr, ptr [[PTR4]], align 8
+// OMP60-NEXT: [[ARRAYIDX5:%.*]] = getelementptr inbounds i32, ptr [[TMP2]], i64 0
+// OMP60-NEXT: [[ARR:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 3
+// OMP60-NEXT: [[ARRAYIDX6:%.*]] = getelementptr inbounds nuw [4 x i32], ptr [[ARR]], i64 0, i64 0
+// OMP60-NEXT: [[A7:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 0
+// OMP60-NEXT: [[TMP3:%.*]] = load i32, ptr [[A7]], align 8
+// OMP60-NEXT: [[CONV:%.*]] = sext i32 [[TMP3]] to i64
+// OMP60-NEXT: [[TMP4:%.*]] = mul nuw i64 [[CONV]], 4
+// OMP60-NEXT: [[A8:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 0
+// OMP60-NEXT: [[TMP5:%.*]] = load i32, ptr [[A8]], align 8
+// OMP60-NEXT: [[TMP6:%.*]] = sext i32 [[TMP5]] to i64
+// OMP60-NEXT: [[LB_ADD_LEN:%.*]] = add nsw i64 -1, [[TMP6]]
+// OMP60-NEXT: [[ARR9:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 3
+// OMP60-NEXT: [[ARRAYIDX10:%.*]] = getelementptr inbounds nuw [4 x i32], ptr [[ARR9]], i64 0, i64 [[LB_ADD_LEN]]
+// OMP60-NEXT: [[PTR11:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 1
+// OMP60-NEXT: [[PTR12:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 1
+// OMP60-NEXT: [[TMP7:%.*]] = load ptr, ptr [[PTR12]], align 8
+// OMP60-NEXT: [[ARRAYIDX13:%.*]] = getelementptr inbounds nuw i32, ptr [[TMP7]], i64 0
+// OMP60-NEXT: [[TMP8:%.*]] = getelementptr i32, ptr [[ARRAYIDX10]], i32 1
+// OMP60-NEXT: [[TMP9:%.*]] = ptrtoint ptr [[TMP8]] to i64
+// OMP60-NEXT: [[TMP10:%.*]] = ptrtoint ptr [[A]] to i64
+// OMP60-NEXT: [[TMP11:%.*]] = sub i64 [[TMP9]], [[TMP10]]
+// OMP60-NEXT: [[TMP12:%.*]] = sdiv exact i64 [[TMP11]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64)
+// OMP60-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[DOTOFFLOAD_SIZES]], ptr align 8 @.offload_sizes.1, i64 56, i1 false)
+// OMP60-NEXT: [[TMP13:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
+// OMP60-NEXT: store ptr [[THIS1]], ptr [[TMP13]], align 8
+// OMP60-NEXT: [[TMP14:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// OMP60-NEXT: store ptr [[A]], ptr [[TMP14]], align 8
+// OMP60-NEXT: [[TMP15:%.*]] = getelementptr inbounds [7 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 0
+// OMP60-NEXT: store i64 [[TMP12]], ptr [[TMP15]], align 8
+// OMP60-NEXT: [[TMP16:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0
+// OMP60-NEXT: store ptr null, ptr [[TMP16]], align 8
+// OMP60-NEXT: [[TMP17:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1
+// OMP60-NEXT: store ptr [[THIS1]], ptr [[TMP17]], align 8
+// OMP60-NEXT: [[TMP18:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1
+// OMP60-NEXT: store ptr [[A]], ptr [[TMP18]], align 8
+// OMP60-NEXT: [[TMP19:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1
+// OMP60-NEXT: store ptr null, ptr [[TMP19]], align 8
+// OMP60-NEXT: [[TMP20:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 2
+// OMP60-NEXT: store ptr [[PTR]], ptr [[TMP20]], align 8
+// OMP60-NEXT: [[TMP21:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 2
+// OMP60-NEXT: store ptr [[ARRAYIDX]], ptr [[TMP21]], align 8
+// OMP60-NEXT: [[TMP22:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 2
+// OMP60-NEXT: store ptr null, ptr [[TMP22]], align 8
+// OMP60-NEXT: [[TMP23:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 3
+// OMP60-NEXT: store ptr [[THIS1]], ptr [[TMP23]], align 8
+// OMP60-NEXT: [[TMP24:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 3
+// OMP60-NEXT: store ptr [[TMP1]], ptr [[TMP24]], align 8
+// OMP60-NEXT: [[TMP25:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 3
+// OMP60-NEXT: store ptr null, ptr [[TMP25]], align 8
+// OMP60-NEXT: [[TMP26:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 4
+// OMP60-NEXT: store ptr [[PTR3]], ptr [[TMP26]], align 8
+// OMP60-NEXT: [[TMP27:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 4
+// OMP60-NEXT: store ptr [[ARRAYIDX5]], ptr [[TMP27]], align 8
+// OMP60-NEXT: [[TMP28:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 4
+// OMP60-NEXT: store ptr null, ptr [[TMP28]], align 8
+// OMP60-NEXT: [[TMP29:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 5
+// OMP60-NEXT: store ptr [[THIS1]], ptr [[TMP29]], align 8
+// OMP60-NEXT: [[TMP30:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 5
+// OMP60-NEXT: store ptr [[ARRAYIDX6]], ptr [[TMP30]], align 8
+// OMP60-NEXT: [[TMP31:%.*]] = getelementptr inbounds [7 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 5
+// OMP60-NEXT: store i64 [[TMP4]], ptr [[TMP31]], align 8
+// OMP60-NEXT: [[TMP32:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 5
+// OMP60-NEXT: store ptr null, ptr [[TMP32]], align 8
+// OMP60-NEXT: [[TMP33:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 6
+// OMP60-NEXT: store ptr [[PTR11]], ptr [[TMP33]], align 8
+// OMP60-NEXT: [[TMP34:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 6
+// OMP60-NEXT: store ptr [[ARRAYIDX13]], ptr [[TMP34]], align 8
+// OMP60-NEXT: [[TMP35:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 6
+// OMP60-NEXT: store ptr null, ptr [[TMP35]], align 8
+// OMP60-NEXT: [[TMP36:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
+// OMP60-NEXT: [[TMP37:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// OMP60-NEXT: [[TMP38:%.*]] = getelementptr inbounds [7 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 0
+// OMP60-NEXT: call void @__tgt_target_data_begin_mapper(ptr @[[GLOB1]], i64 -1, i32 7, ptr [[TMP36]], ptr [[TMP37]], ptr [[TMP38]], ptr @.offload_maptypes.2, ptr null, ptr null)
+// OMP60-NEXT: [[TMP39:%.*]] = load ptr, ptr [[TMP17]], align 8
+// OMP60-NEXT: store ptr [[TMP39]], ptr [[TMP]], align 8
+// OMP60-NEXT: [[TMP40:%.*]] = load ptr, ptr [[TMP20]], align 8
+// OMP60-NEXT: store ptr [[TMP40]], ptr [[_TMP14]], align 8
+// OMP60-NEXT: [[TMP41:%.*]] = load ptr, ptr [[TMP23]], align 8
+// OMP60-NEXT: store ptr [[TMP41]], ptr [[_TMP15]], align 8
+// OMP60-NEXT: [[TMP42:%.*]] = load ptr, ptr [[TMP20]], align 8
+// OMP60-NEXT: store ptr [[TMP42]], ptr [[_TMP16]], align 8
+// OMP60-NEXT: [[TMP43:%.*]] = load ptr, ptr [[TMP29]], align 8
+// OMP60-NEXT: store ptr [[TMP43]], ptr [[_TMP17]], align 8
+// OMP60-NEXT: [[TMP44:%.*]] = load ptr, ptr [[TMP20]], align 8
+// OMP60-NEXT: store ptr [[TMP44]], ptr [[_TMP18]], align 8
+// OMP60-NEXT: [[TMP45:%.*]] = load ptr, ptr [[TMP]], align 8, !nonnull [[META3]], !align [[META4]]
+// OMP60-NEXT: [[TMP46:%.*]] = load i32, ptr [[TMP45]], align 4
+// OMP60-NEXT: [[INC:%.*]] = add nsw i32 [[TMP46]], 1
+// OMP60-NEXT: store i32 [[INC]], ptr [[TMP45]], align 4
+// OMP60-NEXT: [[TMP47:%.*]] = load ptr, ptr [[_TMP18]], align 8, !nonnull [[META3]], !align [[META5:![0-9]+]]
+// OMP60-NEXT: [[TMP48:%.*]] = load ptr, ptr [[TMP47]], align 8
+// OMP60-NEXT: [[TMP49:%.*]] = load i32, ptr [[TMP48]], align 4
+// OMP60-NEXT: [[INC19:%.*]] = add nsw i32 [[TMP49]], 1
+// OMP60-NEXT: store i32 [[INC19]], ptr [[TMP48]], align 4
+// OMP60-NEXT: [[TMP50:%.*]] = load ptr, ptr [[_TMP15]], align 8, !nonnull [[META3]], !align [[META4]]
+// OMP60-NEXT: [[TMP51:%.*]] = load i32, ptr [[TMP50]], align 4
+// OMP60-NEXT: [[INC20:%.*]] = add nsw i32 [[TMP51]], 1
+// OMP60-NEXT: store i32 [[INC20]], ptr [[TMP50]], align 4
+// OMP60-NEXT: [[TMP52:%.*]] = load ptr, ptr [[_TMP17]], align 8, !nonnull [[META3]], !align [[META4]]
+// OMP60-NEXT: [[ARRAYIDX21:%.*]] = getelementptr inbounds [4 x i32], ptr [[TMP52]], i64 0, i64 0
+// OMP60-NEXT: [[TMP53:%.*]] = load i32, ptr [[ARRAYIDX21]], align 4
+// OMP60-NEXT: [[INC22:%.*]] = add nsw i32 [[TMP53]], 1
+// OMP60-NEXT: store i32 [[INC22]], ptr [[ARRAYIDX21]], align 4
+// OMP60-NEXT: [[TMP54:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
+// OMP60-NEXT: [[TMP55:%.*]] = getelementptr inbounds [7 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// OMP60-NEXT: [[TMP56:%.*]] = getelementptr inbounds [7 x i64], ptr [[DOTOFFLOAD_SIZES]], i32 0, i32 0
+// OMP60-NEXT: call void @__tgt_target_data_end_mapper(ptr @[[GLOB1]], i64 -1, i32 7, ptr [[TMP54]], ptr [[TMP55]], ptr [[TMP56]], ptr @.offload_maptypes.2, ptr null, ptr null)
+// OMP60-NEXT: ret void
+//
+//
+// OMP60-LABEL: define linkonce_odr void @_ZN1SC2Ev(
+// OMP60-SAME: ptr noundef nonnull align 8 dereferenceable(40) [[THIS:%.*]]) unnamed_addr #[[ATTR2]] comdat {
+// OMP60-NEXT: [[ENTRY:.*:]]
+// OMP60-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8
+// OMP60-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
+// OMP60-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
+// OMP60-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[THIS1]], i32 0, i32 0
+// OMP60-NEXT: store i32 0, ptr [[A]], align 8
+// OMP60-NEXT: [[PTR:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 1
+// OMP60-NEXT: [[A2:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 0
+// OMP60-NEXT: store ptr [[A2]], ptr [[PTR]], align 8
+// OMP60-NEXT: [[REF:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 2
+// OMP60-NEXT: [[A3:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 0
+// OMP60-NEXT: store ptr [[A3]], ptr [[REF]], align 8
+// OMP60-NEXT: ret void
+//
+//
+// SIMD-ONLY0-LABEL: define dso_local noundef signext i32 @main(
+// SIMD-ONLY0-SAME: ) #[[ATTR0:[0-9]+]] {
+// SIMD-ONLY0-NEXT: [[ENTRY:.*:]]
+// SIMD-ONLY0-NEXT: [[RETVAL:%.*]] = alloca i32, align 4
+// SIMD-ONLY0-NEXT: [[A:%.*]] = alloca float, align 4
+// SIMD-ONLY0-NEXT: [[PTR:%.*]] = alloca ptr, align 8
+// SIMD-ONLY0-NEXT: [[REF:%.*]] = alloca ptr, align 8
+// SIMD-ONLY0-NEXT: [[ARR:%.*]] = alloca [4 x float], align 4
+// SIMD-ONLY0-NEXT: [[SAVED_STACK:%.*]] = alloca ptr, align 8
+// SIMD-ONLY0-NEXT: [[__VLA_EXPR0:%.*]] = alloca i64, align 8
+// SIMD-ONLY0-NEXT: [[S:%.*]] = alloca [[STRUCT_S:%.*]], align 8
+// SIMD-ONLY0-NEXT: [[TMP:%.*]] = alloca ptr, align 8
+// SIMD-ONLY0-NEXT: store i32 0, ptr [[RETVAL]], align 4
+// SIMD-ONLY0-NEXT: store float 0.000000e+00, ptr [[A]], align 4
+// SIMD-ONLY0-NEXT: store ptr [[A]], ptr [[PTR]], align 8
+// SIMD-ONLY0-NEXT: store ptr [[A]], ptr [[REF]], align 8
+// SIMD-ONLY0-NEXT: [[TMP0:%.*]] = load float, ptr [[A]], align 4
+// SIMD-ONLY0-NEXT: [[CONV:%.*]] = fptosi float [[TMP0]] to i32
+// SIMD-ONLY0-NEXT: [[TMP1:%.*]] = zext i32 [[CONV]] to i64
+// SIMD-ONLY0-NEXT: [[TMP2:%.*]] = call ptr @llvm.stacksave.p0()
+// SIMD-ONLY0-NEXT: store ptr [[TMP2]], ptr [[SAVED_STACK]], align 8
+// SIMD-ONLY0-NEXT: [[VLA:%.*]] = alloca float, i64 [[TMP1]], align 4
+// SIMD-ONLY0-NEXT: store i64 [[TMP1]], ptr [[__VLA_EXPR0]], align 8
+// SIMD-ONLY0-NEXT: call void @_ZN1SC1Ev(ptr noundef nonnull align 8 dereferenceable(40) [[S]])
+// SIMD-ONLY0-NEXT: call void @_ZN1S3fooEv(ptr noundef nonnull align 8 dereferenceable(40) [[S]])
+// SIMD-ONLY0-NEXT: [[TMP3:%.*]] = load ptr, ptr [[REF]], align 8, !nonnull [[META2:![0-9]+]], !align [[META3:![0-9]+]]
+// SIMD-ONLY0-NEXT: store ptr [[TMP3]], ptr [[TMP]], align 8
+// SIMD-ONLY0-NEXT: [[TMP4:%.*]] = load float, ptr [[A]], align 4
+// SIMD-ONLY0-NEXT: [[INC:%.*]] = fadd float [[TMP4]], 1.000000e+00
+// SIMD-ONLY0-NEXT: store float [[INC]], ptr [[A]], align 4
+// SIMD-ONLY0-NEXT: [[TMP5:%.*]] = load ptr, ptr [[PTR]], align 8
+// SIMD-ONLY0-NEXT: [[TMP6:%.*]] = load float, ptr [[TMP5]], align 4
+// SIMD-ONLY0-NEXT: [[INC1:%.*]] = fadd float [[TMP6]], 1.000000e+00
+// SIMD-ONLY0-NEXT: store float [[INC1]], ptr [[TMP5]], align 4
+// SIMD-ONLY0-NEXT: [[TMP7:%.*]] = load ptr, ptr [[TMP]], align 8, !nonnull [[META2]], !align [[META3]]
+// SIMD-ONLY0-NEXT: [[TMP8:%.*]] = load float, ptr [[TMP7]], align 4
+// SIMD-ONLY0-NEXT: [[INC2:%.*]] = fadd float [[TMP8]], 1.000000e+00
+// SIMD-ONLY0-NEXT: store float [[INC2]], ptr [[TMP7]], align 4
+// SIMD-ONLY0-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [4 x float], ptr [[ARR]], i64 0, i64 0
+// SIMD-ONLY0-NEXT: [[TMP9:%.*]] = load float, ptr [[ARRAYIDX]], align 4
+// SIMD-ONLY0-NEXT: [[INC3:%.*]] = fadd float [[TMP9]], 1.000000e+00
+// SIMD-ONLY0-NEXT: store float [[INC3]], ptr [[ARRAYIDX]], align 4
+// SIMD-ONLY0-NEXT: [[ARRAYIDX4:%.*]] = getelementptr inbounds float, ptr [[VLA]], i64 0
+// SIMD-ONLY0-NEXT: [[TMP10:%.*]] = load float, ptr [[ARRAYIDX4]], align 4
+// SIMD-ONLY0-NEXT: [[INC5:%.*]] = fadd float [[TMP10]], 1.000000e+00
+// SIMD-ONLY0-NEXT: store float [[INC5]], ptr [[ARRAYIDX4]], align 4
+// SIMD-ONLY0-NEXT: [[TMP11:%.*]] = load float, ptr [[A]], align 4
+// SIMD-ONLY0-NEXT: [[CONV6:%.*]] = fptosi float [[TMP11]] to i32
+// SIMD-ONLY0-NEXT: store i32 [[CONV6]], ptr [[RETVAL]], align 4
+// SIMD-ONLY0-NEXT: [[TMP12:%.*]] = load ptr, ptr [[SAVED_STACK]], align 8
+// SIMD-ONLY0-NEXT: call void @llvm.stackrestore.p0(ptr [[TMP12]])
+// SIMD-ONLY0-NEXT: [[TMP13:%.*]] = load i32, ptr [[RETVAL]], align 4
+// SIMD-ONLY0-NEXT: ret i32 [[TMP13]]
+//
+//
+// SIMD-ONLY0-LABEL: define linkonce_odr void @_ZN1SC1Ev(
+// SIMD-ONLY0-SAME: ptr noundef nonnull align 8 dereferenceable(40) [[THIS:%.*]]) unnamed_addr #[[ATTR2:[0-9]+]] comdat {
+// SIMD-ONLY0-NEXT: [[ENTRY:.*:]]
+// SIMD-ONLY0-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8
+// SIMD-ONLY0-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
+// SIMD-ONLY0-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
+// SIMD-ONLY0-NEXT: call void @_ZN1SC2Ev(ptr noundef nonnull align 8 dereferenceable(40) [[THIS1]])
+// SIMD-ONLY0-NEXT: ret void
+//
+//
+// SIMD-ONLY0-LABEL: define linkonce_odr void @_ZN1S3fooEv(
+// SIMD-ONLY0-SAME: ptr noundef nonnull align 8 dereferenceable(40) [[THIS:%.*]]) #[[ATTR2]] comdat {
+// SIMD-ONLY0-NEXT: [[ENTRY:.*:]]
+// SIMD-ONLY0-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8
+// SIMD-ONLY0-NEXT: [[A:%.*]] = alloca ptr, align 8
+// SIMD-ONLY0-NEXT: [[PTR:%.*]] = alloca ptr, align 8
+// SIMD-ONLY0-NEXT: [[REF:%.*]] = alloca ptr, align 8
+// SIMD-ONLY0-NEXT: [[PTR5:%.*]] = alloca ptr, align 8
+// SIMD-ONLY0-NEXT: [[ARR:%.*]] = alloca ptr, align 8
+// SIMD-ONLY0-NEXT: [[TMP:%.*]] = alloca ptr, align 8
+// SIMD-ONLY0-NEXT: [[_TMP8:%.*]] = alloca ptr, align 8
+// SIMD-ONLY0-NEXT: [[_TMP9:%.*]] = alloca ptr, align 8
+// SIMD-ONLY0-NEXT: [[_TMP10:%.*]] = alloca ptr, align 8
+// SIMD-ONLY0-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
+// SIMD-ONLY0-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
+// SIMD-ONLY0-NEXT: [[A2:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[THIS1]], i32 0, i32 0
+// SIMD-ONLY0-NEXT: store ptr [[A2]], ptr [[A]], align 8
+// SIMD-ONLY0-NEXT: [[PTR3:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 1
+// SIMD-ONLY0-NEXT: store ptr [[PTR3]], ptr [[PTR]], align 8
+// SIMD-ONLY0-NEXT: [[REF4:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 2
+// SIMD-ONLY0-NEXT: [[TMP0:%.*]] = load ptr, ptr [[REF4]], align 8, !nonnull [[META2]], !align [[META3]]
+// SIMD-ONLY0-NEXT: store ptr [[TMP0]], ptr [[REF]], align 8
+// SIMD-ONLY0-NEXT: [[PTR6:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 1
+// SIMD-ONLY0-NEXT: store ptr [[PTR6]], ptr [[PTR5]], align 8
+// SIMD-ONLY0-NEXT: [[ARR7:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 3
+// SIMD-ONLY0-NEXT: store ptr [[ARR7]], ptr [[ARR]], align 8
+// SIMD-ONLY0-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A]], align 8, !nonnull [[META2]], !align [[META3]]
+// SIMD-ONLY0-NEXT: store ptr [[TMP1]], ptr [[TMP]], align 8
+// SIMD-ONLY0-NEXT: [[TMP2:%.*]] = load ptr, ptr [[PTR5]], align 8, !nonnull [[META2]], !align [[META4:![0-9]+]]
+// SIMD-ONLY0-NEXT: store ptr [[TMP2]], ptr [[_TMP8]], align 8
+// SIMD-ONLY0-NEXT: [[TMP3:%.*]] = load ptr, ptr [[REF]], align 8, !nonnull [[META2]], !align [[META3]]
+// SIMD-ONLY0-NEXT: store ptr [[TMP3]], ptr [[_TMP9]], align 8
+// SIMD-ONLY0-NEXT: [[TMP4:%.*]] = load ptr, ptr [[ARR]], align 8, !nonnull [[META2]], !align [[META3]]
+// SIMD-ONLY0-NEXT: store ptr [[TMP4]], ptr [[_TMP10]], align 8
+// SIMD-ONLY0-NEXT: [[TMP5:%.*]] = load ptr, ptr [[TMP]], align 8, !nonnull [[META2]], !align [[META3]]
+// SIMD-ONLY0-NEXT: [[TMP6:%.*]] = load i32, ptr [[TMP5]], align 4
+// SIMD-ONLY0-NEXT: [[INC:%.*]] = add nsw i32 [[TMP6]], 1
+// SIMD-ONLY0-NEXT: store i32 [[INC]], ptr [[TMP5]], align 4
+// SIMD-ONLY0-NEXT: [[TMP7:%.*]] = load ptr, ptr [[_TMP8]], align 8, !nonnull [[META2]], !align [[META4]]
+// SIMD-ONLY0-NEXT: [[TMP8:%.*]] = load ptr, ptr [[TMP7]], align 8
+// SIMD-ONLY0-NEXT: [[TMP9:%.*]] = load i32, ptr [[TMP8]], align 4
+// SIMD-ONLY0-NEXT: [[INC11:%.*]] = add nsw i32 [[TMP9]], 1
+// SIMD-ONLY0-NEXT: store i32 [[INC11]], ptr [[TMP8]], align 4
+// SIMD-ONLY0-NEXT: [[TMP10:%.*]] = load ptr, ptr [[_TMP9]], align 8, !nonnull [[META2]], !align [[META3]]
+// SIMD-ONLY0-NEXT: [[TMP11:%.*]] = load i32, ptr [[TMP10]], align 4
+// SIMD-ONLY0-NEXT: [[INC12:%.*]] = add nsw i32 [[TMP11]], 1
+// SIMD-ONLY0-NEXT: store i32 [[INC12]], ptr [[TMP10]], align 4
+// SIMD-ONLY0-NEXT: [[TMP12:%.*]] = load ptr, ptr [[_TMP10]], align 8, !nonnull [[META2]], !align [[META3]]
+// SIMD-ONLY0-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [4 x i32], ptr [[TMP12]], i64 0, i64 0
+// SIMD-ONLY0-NEXT: [[TMP13:%.*]] = load i32, ptr [[ARRAYIDX]], align 4
+// SIMD-ONLY0-NEXT: [[INC13:%.*]] = add nsw i32 [[TMP13]], 1
+// SIMD-ONLY0-NEXT: store i32 [[INC13]], ptr [[ARRAYIDX]], align 4
+// SIMD-ONLY0-NEXT: ret void
+//
+//
+// SIMD-ONLY0-LABEL: define linkonce_odr void @_ZN1SC2Ev(
+// SIMD-ONLY0-SAME: ptr noundef nonnull align 8 dereferenceable(40) [[THIS:%.*]]) unnamed_addr #[[ATTR2]] comdat {
+// SIMD-ONLY0-NEXT: [[ENTRY:.*:]]
+// SIMD-ONLY0-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8
+// SIMD-ONLY0-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
+// SIMD-ONLY0-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
+// SIMD-ONLY0-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[THIS1]], i32 0, i32 0
+// SIMD-ONLY0-NEXT: store i32 0, ptr [[A]], align 8
+// SIMD-ONLY0-NEXT: [[PTR:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 1
+// SIMD-ONLY0-NEXT: [[A2:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 0
+// SIMD-ONLY0-NEXT: store ptr [[A2]], ptr [[PTR]], align 8
+// SIMD-ONLY0-NEXT: [[REF:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 2
+// SIMD-ONLY0-NEXT: [[A3:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 0
+// SIMD-ONLY0-NEXT: store ptr [[A3]], ptr [[REF]], align 8
+// SIMD-ONLY0-NEXT: ret void
+//
+//
+// SIMD-ONLY0-OMP60-LABEL: define dso_local noundef signext i32 @main(
+// SIMD-ONLY0-OMP60-SAME: i32 noundef signext [[ARGC:%.*]], ptr noundef [[ARGV:%.*]]) #[[ATTR0:[0-9]+]] {
+// SIMD-ONLY0-OMP60-NEXT: [[ENTRY:.*:]]
+// SIMD-ONLY0-OMP60-NEXT: [[RETVAL:%.*]] = alloca i32, align 4
+// SIMD-ONLY0-OMP60-NEXT: [[ARGC_ADDR:%.*]] = alloca i32, align 4
+// SIMD-ONLY0-OMP60-NEXT: [[ARGV_ADDR:%.*]] = alloca ptr, align 8
+// SIMD-ONLY0-OMP60-NEXT: [[A:%.*]] = alloca float, align 4
+// SIMD-ONLY0-OMP60-NEXT: [[PTR:%.*]] = alloca ptr, align 8
+// SIMD-ONLY0-OMP60-NEXT: [[REF:%.*]] = alloca ptr, align 8
+// SIMD-ONLY0-OMP60-NEXT: [[ARR:%.*]] = alloca [4 x float], align 4
+// SIMD-ONLY0-OMP60-NEXT: [[SAVED_STACK:%.*]] = alloca ptr, align 8
+// SIMD-ONLY0-OMP60-NEXT: [[__VLA_EXPR0:%.*]] = alloca i64, align 8
+// SIMD-ONLY0-OMP60-NEXT: [[S:%.*]] = alloca [[STRUCT_S:%.*]], align 8
+// SIMD-ONLY0-OMP60-NEXT: [[TMP:%.*]] = alloca ptr, align 8
+// SIMD-ONLY0-OMP60-NEXT: store i32 0, ptr [[RETVAL]], align 4
+// SIMD-ONLY0-OMP60-NEXT: store i32 [[ARGC]], ptr [[ARGC_ADDR]], align 4
+// SIMD-ONLY0-OMP60-NEXT: store ptr [[ARGV]], ptr [[ARGV_ADDR]], align 8
+// SIMD-ONLY0-OMP60-NEXT: store float 0.000000e+00, ptr [[A]], align 4
+// SIMD-ONLY0-OMP60-NEXT: store ptr [[A]], ptr [[PTR]], align 8
+// SIMD-ONLY0-OMP60-NEXT: store ptr [[A]], ptr [[REF]], align 8
+// SIMD-ONLY0-OMP60-NEXT: [[TMP0:%.*]] = load float, ptr [[A]], align 4
+// SIMD-ONLY0-OMP60-NEXT: [[CONV:%.*]] = fptosi float [[TMP0]] to i32
+// SIMD-ONLY0-OMP60-NEXT: [[TMP1:%.*]] = zext i32 [[CONV]] to i64
+// SIMD-ONLY0-OMP60-NEXT: [[TMP2:%.*]] = call ptr @llvm.stacksave.p0()
+// SIMD-ONLY0-OMP60-NEXT: store ptr [[TMP2]], ptr [[SAVED_STACK]], align 8
+// SIMD-ONLY0-OMP60-NEXT: [[VLA:%.*]] = alloca float, i64 [[TMP1]], align 4
+// SIMD-ONLY0-OMP60-NEXT: store i64 [[TMP1]], ptr [[__VLA_EXPR0]], align 8
+// SIMD-ONLY0-OMP60-NEXT: call void @_ZN1SC1Ev(ptr noundef nonnull align 8 dereferenceable(40) [[S]])
+// SIMD-ONLY0-OMP60-NEXT: call void @_ZN1S3fooEv(ptr noundef nonnull align 8 dereferenceable(40) [[S]])
+// SIMD-ONLY0-OMP60-NEXT: [[TMP3:%.*]] = load ptr, ptr [[REF]], align 8, !nonnull [[META2:![0-9]+]], !align [[META3:![0-9]+]]
+// SIMD-ONLY0-OMP60-NEXT: store ptr [[TMP3]], ptr [[TMP]], align 8
+// SIMD-ONLY0-OMP60-NEXT: [[TMP4:%.*]] = load float, ptr [[A]], align 4
+// SIMD-ONLY0-OMP60-NEXT: [[INC:%.*]] = fadd float [[TMP4]], 1.000000e+00
+// SIMD-ONLY0-OMP60-NEXT: store float [[INC]], ptr [[A]], align 4
+// SIMD-ONLY0-OMP60-NEXT: [[TMP5:%.*]] = load ptr, ptr [[PTR]], align 8
+// SIMD-ONLY0-OMP60-NEXT: [[TMP6:%.*]] = load float, ptr [[TMP5]], align 4
+// SIMD-ONLY0-OMP60-NEXT: [[INC1:%.*]] = fadd float [[TMP6]], 1.000000e+00
+// SIMD-ONLY0-OMP60-NEXT: store float [[INC1]], ptr [[TMP5]], align 4
+// SIMD-ONLY0-OMP60-NEXT: [[TMP7:%.*]] = load ptr, ptr [[TMP]], align 8, !nonnull [[META2]], !align [[META3]]
+// SIMD-ONLY0-OMP60-NEXT: [[TMP8:%.*]] = load float, ptr [[TMP7]], align 4
+// SIMD-ONLY0-OMP60-NEXT: [[INC2:%.*]] = fadd float [[TMP8]], 1.000000e+00
+// SIMD-ONLY0-OMP60-NEXT: store float [[INC2]], ptr [[TMP7]], align 4
+// SIMD-ONLY0-OMP60-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [4 x float], ptr [[ARR]], i64 0, i64 0
+// SIMD-ONLY0-OMP60-NEXT: [[TMP9:%.*]] = load float, ptr [[ARRAYIDX]], align 4
+// SIMD-ONLY0-OMP60-NEXT: [[INC3:%.*]] = fadd float [[TMP9]], 1.000000e+00
+// SIMD-ONLY0-OMP60-NEXT: store float [[INC3]], ptr [[ARRAYIDX]], align 4
+// SIMD-ONLY0-OMP60-NEXT: [[ARRAYIDX4:%.*]] = getelementptr inbounds float, ptr [[VLA]], i64 0
+// SIMD-ONLY0-OMP60-NEXT: [[TMP10:%.*]] = load float, ptr [[ARRAYIDX4]], align 4
+// SIMD-ONLY0-OMP60-NEXT: [[INC5:%.*]] = fadd float [[TMP10]], 1.000000e+00
+// SIMD-ONLY0-OMP60-NEXT: store float [[INC5]], ptr [[ARRAYIDX4]], align 4
+// SIMD-ONLY0-OMP60-NEXT: [[TMP11:%.*]] = load float, ptr [[A]], align 4
+// SIMD-ONLY0-OMP60-NEXT: [[CONV6:%.*]] = fptosi float [[TMP11]] to i32
+// SIMD-ONLY0-OMP60-NEXT: store i32 [[CONV6]], ptr [[RETVAL]], align 4
+// SIMD-ONLY0-OMP60-NEXT: [[TMP12:%.*]] = load ptr, ptr [[SAVED_STACK]], align 8
+// SIMD-ONLY0-OMP60-NEXT: call void @llvm.stackrestore.p0(ptr [[TMP12]])
+// SIMD-ONLY0-OMP60-NEXT: [[TMP13:%.*]] = load i32, ptr [[RETVAL]], align 4
+// SIMD-ONLY0-OMP60-NEXT: ret i32 [[TMP13]]
+//
+//
+// SIMD-ONLY0-OMP60-LABEL: define linkonce_odr void @_ZN1SC1Ev(
+// SIMD-ONLY0-OMP60-SAME: ptr noundef nonnull align 8 dereferenceable(40) [[THIS:%.*]]) unnamed_addr #[[ATTR2:[0-9]+]] comdat {
+// SIMD-ONLY0-OMP60-NEXT: [[ENTRY:.*:]]
+// SIMD-ONLY0-OMP60-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8
+// SIMD-ONLY0-OMP60-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
+// SIMD-ONLY0-OMP60-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
+// SIMD-ONLY0-OMP60-NEXT: call void @_ZN1SC2Ev(ptr noundef nonnull align 8 dereferenceable(40) [[THIS1]])
+// SIMD-ONLY0-OMP60-NEXT: ret void
+//
+//
+// SIMD-ONLY0-OMP60-LABEL: define linkonce_odr void @_ZN1S3fooEv(
+// SIMD-ONLY0-OMP60-SAME: ptr noundef nonnull align 8 dereferenceable(40) [[THIS:%.*]]) #[[ATTR2]] comdat {
+// SIMD-ONLY0-OMP60-NEXT: [[ENTRY:.*:]]
+// SIMD-ONLY0-OMP60-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8
+// SIMD-ONLY0-OMP60-NEXT: [[A:%.*]] = alloca ptr, align 8
+// SIMD-ONLY0-OMP60-NEXT: [[PTR:%.*]] = alloca ptr, align 8
+// SIMD-ONLY0-OMP60-NEXT: [[REF:%.*]] = alloca ptr, align 8
+// SIMD-ONLY0-OMP60-NEXT: [[PTR5:%.*]] = alloca ptr, align 8
+// SIMD-ONLY0-OMP60-NEXT: [[ARR:%.*]] = alloca ptr, align 8
+// SIMD-ONLY0-OMP60-NEXT: [[PTR8:%.*]] = alloca ptr, align 8
+// SIMD-ONLY0-OMP60-NEXT: [[TMP:%.*]] = alloca ptr, align 8
+// SIMD-ONLY0-OMP60-NEXT: [[_TMP10:%.*]] = alloca ptr, align 8
+// SIMD-ONLY0-OMP60-NEXT: [[_TMP11:%.*]] = alloca ptr, align 8
+// SIMD-ONLY0-OMP60-NEXT: [[_TMP12:%.*]] = alloca ptr, align 8
+// SIMD-ONLY0-OMP60-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
+// SIMD-ONLY0-OMP60-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
+// SIMD-ONLY0-OMP60-NEXT: [[A2:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[THIS1]], i32 0, i32 0
+// SIMD-ONLY0-OMP60-NEXT: store ptr [[A2]], ptr [[A]], align 8
+// SIMD-ONLY0-OMP60-NEXT: [[PTR3:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 1
+// SIMD-ONLY0-OMP60-NEXT: store ptr [[PTR3]], ptr [[PTR]], align 8
+// SIMD-ONLY0-OMP60-NEXT: [[REF4:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 2
+// SIMD-ONLY0-OMP60-NEXT: [[TMP0:%.*]] = load ptr, ptr [[REF4]], align 8, !nonnull [[META2]], !align [[META3]]
+// SIMD-ONLY0-OMP60-NEXT: store ptr [[TMP0]], ptr [[REF]], align 8
+// SIMD-ONLY0-OMP60-NEXT: [[PTR6:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 1
+// SIMD-ONLY0-OMP60-NEXT: store ptr [[PTR6]], ptr [[PTR5]], align 8
+// SIMD-ONLY0-OMP60-NEXT: [[ARR7:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 3
+// SIMD-ONLY0-OMP60-NEXT: store ptr [[ARR7]], ptr [[ARR]], align 8
+// SIMD-ONLY0-OMP60-NEXT: [[PTR9:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 1
+// SIMD-ONLY0-OMP60-NEXT: store ptr [[PTR9]], ptr [[PTR8]], align 8
+// SIMD-ONLY0-OMP60-NEXT: [[TMP1:%.*]] = load ptr, ptr [[A]], align 8, !nonnull [[META2]], !align [[META3]]
+// SIMD-ONLY0-OMP60-NEXT: store ptr [[TMP1]], ptr [[TMP]], align 8
+// SIMD-ONLY0-OMP60-NEXT: [[TMP2:%.*]] = load ptr, ptr [[PTR8]], align 8, !nonnull [[META2]], !align [[META4:![0-9]+]]
+// SIMD-ONLY0-OMP60-NEXT: store ptr [[TMP2]], ptr [[_TMP10]], align 8
+// SIMD-ONLY0-OMP60-NEXT: [[TMP3:%.*]] = load ptr, ptr [[REF]], align 8, !nonnull [[META2]], !align [[META3]]
+// SIMD-ONLY0-OMP60-NEXT: store ptr [[TMP3]], ptr [[_TMP11]], align 8
+// SIMD-ONLY0-OMP60-NEXT: [[TMP4:%.*]] = load ptr, ptr [[ARR]], align 8, !nonnull [[META2]], !align [[META3]]
+// SIMD-ONLY0-OMP60-NEXT: store ptr [[TMP4]], ptr [[_TMP12]], align 8
+// SIMD-ONLY0-OMP60-NEXT: [[TMP5:%.*]] = load ptr, ptr [[TMP]], align 8, !nonnull [[META2]], !align [[META3]]
+// SIMD-ONLY0-OMP60-NEXT: [[TMP6:%.*]] = load i32, ptr [[TMP5]], align 4
+// SIMD-ONLY0-OMP60-NEXT: [[INC:%.*]] = add nsw i32 [[TMP6]], 1
+// SIMD-ONLY0-OMP60-NEXT: store i32 [[INC]], ptr [[TMP5]], align 4
+// SIMD-ONLY0-OMP60-NEXT: [[TMP7:%.*]] = load ptr, ptr [[_TMP10]], align 8, !nonnull [[META2]], !align [[META4]]
+// SIMD-ONLY0-OMP60-NEXT: [[TMP8:%.*]] = load ptr, ptr [[TMP7]], align 8
+// SIMD-ONLY0-OMP60-NEXT: [[TMP9:%.*]] = load i32, ptr [[TMP8]], align 4
+// SIMD-ONLY0-OMP60-NEXT: [[INC13:%.*]] = add nsw i32 [[TMP9]], 1
+// SIMD-ONLY0-OMP60-NEXT: store i32 [[INC13]], ptr [[TMP8]], align 4
+// SIMD-ONLY0-OMP60-NEXT: [[TMP10:%.*]] = load ptr, ptr [[_TMP11]], align 8, !nonnull [[META2]], !align [[META3]]
+// SIMD-ONLY0-OMP60-NEXT: [[TMP11:%.*]] = load i32, ptr [[TMP10]], align 4
+// SIMD-ONLY0-OMP60-NEXT: [[INC14:%.*]] = add nsw i32 [[TMP11]], 1
+// SIMD-ONLY0-OMP60-NEXT: store i32 [[INC14]], ptr [[TMP10]], align 4
+// SIMD-ONLY0-OMP60-NEXT: [[TMP12:%.*]] = load ptr, ptr [[_TMP12]], align 8, !nonnull [[META2]], !align [[META3]]
+// SIMD-ONLY0-OMP60-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [4 x i32], ptr [[TMP12]], i64 0, i64 0
+// SIMD-ONLY0-OMP60-NEXT: [[TMP13:%.*]] = load i32, ptr [[ARRAYIDX]], align 4
+// SIMD-ONLY0-OMP60-NEXT: [[INC15:%.*]] = add nsw i32 [[TMP13]], 1
+// SIMD-ONLY0-OMP60-NEXT: store i32 [[INC15]], ptr [[ARRAYIDX]], align 4
+// SIMD-ONLY0-OMP60-NEXT: ret void
+//
+//
+// SIMD-ONLY0-OMP60-LABEL: define linkonce_odr void @_ZN1SC2Ev(
+// SIMD-ONLY0-OMP60-SAME: ptr noundef nonnull align 8 dereferenceable(40) [[THIS:%.*]]) unnamed_addr #[[ATTR2]] comdat {
+// SIMD-ONLY0-OMP60-NEXT: [[ENTRY:.*:]]
+// SIMD-ONLY0-OMP60-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8
+// SIMD-ONLY0-OMP60-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
+// SIMD-ONLY0-OMP60-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
+// SIMD-ONLY0-OMP60-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_S:%.*]], ptr [[THIS1]], i32 0, i32 0
+// SIMD-ONLY0-OMP60-NEXT: store i32 0, ptr [[A]], align 8
+// SIMD-ONLY0-OMP60-NEXT: [[PTR:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 1
+// SIMD-ONLY0-OMP60-NEXT: [[A2:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 0
+// SIMD-ONLY0-OMP60-NEXT: store ptr [[A2]], ptr [[PTR]], align 8
+// SIMD-ONLY0-OMP60-NEXT: [[REF:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 2
+// SIMD-ONLY0-OMP60-NEXT: [[A3:%.*]] = getelementptr inbounds nuw [[STRUCT_S]], ptr [[THIS1]], i32 0, i32 0
+// SIMD-ONLY0-OMP60-NEXT: store ptr [[A3]], ptr [[REF]], align 8
+// SIMD-ONLY0-OMP60-NEXT: ret void
+//
+//.
+// CHECK: [[META3]] = !{}
+// CHECK: [[META4]] = !{i64 4}
+// CHECK: [[META5]] = !{i64 8}
+//.
+// OMP60: [[META3]] = !{}
+// OMP60: [[META4]] = !{i64 4}
+// OMP60: [[META5]] = !{i64 8}
+//.
+// SIMD-ONLY0: [[META2]] = !{}
+// SIMD-ONLY0: [[META3]] = !{i64 4}
+// SIMD-ONLY0: [[META4]] = !{i64 8}
+//.
+// SIMD-ONLY0-OMP60: [[META2]] = !{}
+// SIMD-ONLY0-OMP60: [[META3]] = !{i64 4}
+// SIMD-ONLY0-OMP60: [[META4]] = !{i64 8}
+//.
diff --git a/clang/test/OpenMP/target_has_device_addr_messages.cpp b/clang/test/OpenMP/target_has_device_addr_messages.cpp
index 532ace5caf0a0..bac6fc06d6a16 100644
--- a/clang/test/OpenMP/target_has_device_addr_messages.cpp
+++ b/clang/test/OpenMP/target_has_device_addr_messages.cpp
@@ -4,9 +4,16 @@
// RUN: %clang_cc1 -std=c++11 -verify \
// RUN: -fopenmp-simd -ferror-limit 200 %s -Wuninitialized
+// RUN: %clang_cc1 -std=c++11 -verify \
+// RUN: -fopenmp -fopenmp-version=60 -ferror-limit 200 %s -Wuninitialized
+
+// RUN: %clang_cc1 -std=c++11 -verify \
+// RUN: -fopenmp-simd -fopenmp-version=60 -ferror-limit 200 %s -Wuninitialized
+
struct ST {
int *a;
};
+
typedef int arr[10];
typedef ST STarr[10];
struct SA {
@@ -133,9 +140,9 @@ typedef struct {
int a;
} S6;
-template <typename T, int I>
-T tmain(T argc) {
- const T d = 5;
+template <typename T, typename S, int I>
+T tfoo(T argc, S ub[]) {
+ const T d = I;
const T da[5] = { 0 };
S4 e(4);
S5 g(5);
@@ -200,10 +207,14 @@ T tmain(T argc) {
{}
#pragma omp target private(ps) has_device_addr(ps) // expected-error{{private variable cannot be in a has_device_addr clause in '#pragma omp target' directive}} expected-note{{defined as private}}
{}
+#pragma omp target has_device_addr(ps[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+ {}
+#pragma omp target has_device_addr(ub[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is an array of unknown bound}}
+ {}
return 0;
}
-int main(int argc, char **argv) {
+int foo(int argc, float ub[]) {
const int d = 5;
const int da[5] = { 0 };
S4 e(4);
@@ -269,5 +280,10 @@ int main(int argc, char **argv) {
{}
#pragma omp target private(ps) has_device_addr(ps) // expected-error{{private variable cannot be in a has_device_addr clause in '#pragma omp target' directive}} expected-note{{defined as private}}
{}
- return tmain<int, 3>(argc);
+#pragma omp target has_device_addr(ps[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+ {}
+#pragma omp target has_device_addr(ub[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is an array of unknown bound}}
+ {}
+ return tfoo<int, float, 5>(argc, ub); // expected-note {{in instantiation of function template specialization 'tfoo<int, float, 5>' requested here}}
+
}
diff --git a/clang/test/OpenMP/task_in_reduction_message.cpp b/clang/test/OpenMP/task_in_reduction_message.cpp
index dceefef90c2f2..7979f70aa9699 100644
--- a/clang/test/OpenMP/task_in_reduction_message.cpp
+++ b/clang/test/OpenMP/task_in_reduction_message.cpp
@@ -1,16 +1,17 @@
-// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 150 -o - %s
-// RUN: %clang_cc1 -verify -fopenmp -std=c++98 -ferror-limit 150 -o - %s
-// RUN: %clang_cc1 -verify -fopenmp -std=c++11 -ferror-limit 150 -o - %s
-// RUN: %clang_cc1 -verify=expected,omp52 -fopenmp -fopenmp-version=52 -ferror-limit 150 -o - %s
-// RUN: %clang_cc1 -verify=expected,omp52 -fopenmp -fopenmp-version=52 -std=c++98 -ferror-limit 150 -o - %s
-// RUN: %clang_cc1 -verify=expected,omp52 -fopenmp -fopenmp-version=52 -std=c++11 -ferror-limit 150 -o - %s
+// RUN: %clang_cc1 -verify=expected,omp-lt60 -fopenmp -std=c++98 -ferror-limit 150 -o - %s
+// RUN: %clang_cc1 -verify=expected,omp-lt60 -fopenmp -std=c++11 -ferror-limit 150 -o - %s
+// RUN: %clang_cc1 -verify=expected,omp52,omp-lt60 -fopenmp -fopenmp-version=52 -ferror-limit 150 -o - %s
+// RUN: %clang_cc1 -verify=expected,omp52,omp-lt60 -fopenmp -fopenmp-version=52 -std=c++98 -ferror-limit 150 -o - %s
+// RUN: %clang_cc1 -verify=expected,omp52,omp-lt60 -fopenmp -fopenmp-version=52 -std=c++11 -ferror-limit 150 -o - %s
+// RUN: %clang_cc1 -verify=expected,omp-ge60 -fopenmp -fopenmp-version=60 -ferror-limit 150 -o - %s
-// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 150 -o - %s
-// RUN: %clang_cc1 -verify -fopenmp-simd -std=c++98 -ferror-limit 150 -o - %s
-// RUN: %clang_cc1 -verify -fopenmp-simd -std=c++11 -ferror-limit 150 -o - %s
-// RUN: %clang_cc1 -verify=expected,omp52 -fopenmp-simd -fopenmp-version=52 -ferror-limit 150 -o - %s
-// RUN: %clang_cc1 -verify=expected,omp52 -fopenmp-simd -fopenmp-version=52 -std=c++98 -ferror-limit 150 -o - %s
-// RUN: %clang_cc1 -verify=expected,omp52 -fopenmp-simd -fopenmp-version=52 -std=c++11 -ferror-limit 150 -o - %s
+// RUN: %clang_cc1 -verify=expected,omp-lt60 -fopenmp-simd -ferror-limit 150 -o - %s
+// RUN: %clang_cc1 -verify=expected,omp-lt60 -fopenmp-simd -std=c++98 -ferror-limit 150 -o - %s
+// RUN: %clang_cc1 -verify=expected,omp-lt60 -fopenmp-simd -std=c++11 -ferror-limit 150 -o - %s
+// RUN: %clang_cc1 -verify=expected,omp52,omp-lt60 -fopenmp-simd -fopenmp-version=52 -ferror-limit 150 -o - %s
+// RUN: %clang_cc1 -verify=expected,omp52,omp-lt60 -fopenmp-simd -fopenmp-version=52 -std=c++98 -ferror-limit 150 -o - %s
+// RUN: %clang_cc1 -verify=expected,omp52,omp-lt60 -fopenmp-simd -fopenmp-version=52 -std=c++11 -ferror-limit 150 -o - %s
+// RUN: %clang_cc1 -verify=expected,omp-ge60 -fopenmp-simd -fopenmp-version=60 -std=c++11 -ferror-limit 150 -o - %s
// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
typedef void **omp_allocator_handle_t;
@@ -39,7 +40,7 @@ void foobar(int &ref) {
void foobar1(int &ref) {
#pragma omp taskgroup task_reduction(+:ref)
-#pragma omp task in_reduction(-:ref) // omp52-warning {{minus(-) operator for reductions is deprecated; use + or user defined reduction instead}}
+#pragma omp task in_reduction(-:ref) // omp52-warning {{minus(-) operator for reductions is deprecated; use + or user defined reduction instead}} omp-ge60-error {{incorrect reduction identifier, expected one of '+', '*', '&', '|', '^', '&&', '||', 'min' or 'max' or declare reduction for type 'int}}
foo();
}
@@ -123,8 +124,8 @@ class S6 { // expected-note 3 {{candidate function (the implicit copy assignment
S3 h, k;
#pragma omp threadprivate(h) // expected-note 2 {{defined as threadprivate or thread local}}
-template <class T> // expected-note {{declared here}}
-T tmain(T argc) {
+template <class T, class S> // expected-note {{declared here}}
+T tfoobar5(T argc, S ub[]) {
const T d = T(); // expected-note 4 {{'d' defined here}}
const T da[5] = {T()}; // expected-note 2 {{'da' defined here}}
T qa[5] = {T()};
@@ -134,6 +135,7 @@ T tmain(T argc) {
const T &r = da[(int)i]; // expected-note 2 {{'r' defined here}}
T &q = qa[(int)i];
T fl;
+ T *ptr;
#pragma omp taskgroup task_reduction(+:argc)
#pragma omp task in_reduction // expected-error {{expected '(' after 'in_reduction'}}
foo();
@@ -163,7 +165,7 @@ T tmain(T argc) {
foo();
#pragma omp task in_reduction(|| : argc ? i : argc) // expected-error 2 {{expected variable name, array element or array section}}
foo();
-#pragma omp task in_reduction(foo : argc) //expected-error {{incorrect reduction identifier, expected one of '+', '-', '*', '&', '|', '^', '&&', '||', 'min' or 'max' or declare reduction for type 'float'}} expected-error {{incorrect reduction identifier, expected one of '+', '-', '*', '&', '|', '^', '&&', '||', 'min' or 'max' or declare reduction for type 'int'}}
+#pragma omp task in_reduction(foo : argc) //omp-lt60-error {{incorrect reduction identifier, expected one of '+', '-', '*', '&', '|', '^', '&&', '||', 'min' or 'max' or declare reduction for type 'float'}} omp-lt60-error {{incorrect reduction identifier, expected one of '+', '-', '*', '&', '|', '^', '&&', '||', 'min' or 'max' or declare reduction for type 'int'}} omp-ge60-error {{incorrect reduction identifier, expected one of '+', '*', '&', '|', '^', '&&', '||', 'min' or 'max' or declare reduction for type 'int'}} omp-ge60-error {{incorrect reduction identifier, expected one of '+', '*', '&', '|', '^', '&&', '||', 'min' or 'max' or declare reduction for type 'float'}}
foo();
#pragma omp taskgroup task_reduction(&&:argc)
#pragma omp task in_reduction(&& : argc)
@@ -210,11 +212,15 @@ T tmain(T argc) {
#pragma omp task in_reduction(+ : fl)
foo();
#pragma omp parallel
-#pragma omp for reduction(- : fl) // omp52-warning 3 {{minus(-) operator for reductions is deprecated; use + or user defined reduction instead}}
+#pragma omp for reduction(- : fl) // omp52-warning 3 {{minus(-) operator for reductions is deprecated; use + or user defined reduction instead}} omp-ge60-error {{incorrect reduction identifier, expected one of '+', '*', '&', '|', '^', '&&', '||', 'min' or 'max' or declare reduction for type 'int'}} omp-ge60-error {{incorrect reduction identifier, expected one of '+', '*', '&', '|', '^', '&&', '||', 'min' or 'max' or declare reduction for type 'float'}}
for (int i = 0; i < 10; ++i)
#pragma omp taskgroup task_reduction(+:fl)
#pragma omp task in_reduction(+ : fl)
foo();
+#pragma omp task in_reduction(min:ub[:]) // expected-error 2 {{section length is unspecified and cannot be inferred because subscripted value is an array of unknown bound}}
+ foo();
+#pragma omp task in_reduction(min:ptr[:]) // expected-error 2 {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+ foo();
return T();
}
@@ -227,7 +233,7 @@ namespace B {
using A::x;
}
-int main(int argc, char **argv) {
+int foobar5(int argc, float ub[]) {
const int d = 5; // expected-note 2 {{'d' defined here}}
const int da[5] = {0}; // expected-note {{'da' defined here}}
int qa[5] = {0};
@@ -239,6 +245,7 @@ int main(int argc, char **argv) {
const int &r = da[i]; // expected-note {{'r' defined here}}
int &q = qa[i];
float fl;
+ float *ptr;
#pragma omp task in_reduction // expected-error {{expected '(' after 'in_reduction'}}
foo();
#pragma omp task in_reduction + // expected-error {{expected '(' after 'in_reduction'}} expected-warning {{extra tokens at the end of '#pragma omp task' are ignored}}
@@ -253,7 +260,7 @@ int main(int argc, char **argv) {
foo();
#pragma omp task in_reduction(\) // expected-error {{expected unqualified-id}} expected-warning {{missing ':' after reduction identifier - ignoring}}
foo();
-#pragma omp task in_reduction(foo : argc // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{incorrect reduction identifier, expected one of '+', '-', '*', '&', '|', '^', '&&', '||', 'min' or 'max'}}
+#pragma omp task in_reduction(foo : argc // expected-error {{expected ')'}} expected-note {{to match this '('}} omp-lt60-error {{incorrect reduction identifier, expected one of '+', '-', '*', '&', '|', '^', '&&', '||', 'min' or 'max'}} omp-ge60-error {{incorrect reduction identifier, expected one of '+', '*', '&', '|', '^', '&&', '||', 'min' or 'max'}}
foo();
#pragma omp taskgroup task_reduction(|:argc)
{
@@ -262,10 +269,10 @@ int main(int argc, char **argv) {
#pragma omp task in_reduction(| : argc, // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{in_reduction variable must have the same reduction operation as in a task_reduction clause}}
foo();
}
-#pragma omp task in_reduction(| : argc) allocate , allocate(, allocate(omp_default , allocate(omp_default_mem_alloc, allocate(omp_default_mem_alloc:, allocate(omp_default_mem_alloc: argc, allocate(omp_default_mem_alloc: argv), allocate(argv) // expected-error {{expected '(' after 'allocate'}} expected-error 2 {{expected expression}} expected-error 2 {{expected ')'}} expected-error {{use of undeclared identifier 'omp_default'}} expected-note 2 {{to match this '('}}
+#pragma omp task in_reduction(| : argc) allocate , allocate(, allocate(omp_default , allocate(omp_default_mem_alloc, allocate(omp_default_mem_alloc:, allocate(omp_default_mem_alloc: argc, allocate(omp_default_mem_alloc: ub), allocate(ub) // expected-error {{expected '(' after 'allocate'}} expected-error 2 {{expected expression}} expected-error 2 {{expected ')'}} expected-error {{use of undeclared identifier 'omp_default'}} expected-note 2 {{to match this '('}}
foo();
}
-#pragma omp task in_reduction(|| : argc > 0 ? argv[1] : argv[2]) // expected-error {{expected variable name, array element or array section}}
+#pragma omp task in_reduction(|| : argc > 0 ? ub[1] : ub[2]) // expected-error {{expected variable name, array element or array section}}
foo();
#pragma omp task in_reduction(~ : argc) // expected-error {{expected unqualified-id}}
foo();
@@ -325,6 +332,10 @@ int main(int argc, char **argv) {
#pragma omp taskgroup task_reduction(+:m)
#pragma omp task in_reduction(+ : m) // OK
m++;
+#pragma omp task in_reduction(min:ub[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is an array of unknown bound}}
+ foo();
+#pragma omp task in_reduction(min:ptr[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+ foo();
- return tmain(argc) + tmain(fl); // expected-note {{in instantiation of function template specialization 'tmain<int>' requested here}} expected-note {{in instantiation of function template specialization 'tmain<float>' requested here}}
+ return tfoobar5(argc, ub) + tfoobar5(fl, ub); // expected-note {{in instantiation of function template specialization 'tfoobar5<int, float>' requested here}} expected-note {{in instantiation of function template specialization 'tfoobar5<float, float>' requested here}}
}
diff --git a/clang/test/OpenMP/taskgroup_task_reduction_messages.cpp b/clang/test/OpenMP/taskgroup_task_reduction_messages.cpp
index e038b5edf2066..08b4bf2bbe721 100644
--- a/clang/test/OpenMP/taskgroup_task_reduction_messages.cpp
+++ b/clang/test/OpenMP/taskgroup_task_reduction_messages.cpp
@@ -92,8 +92,8 @@ class S6 { // expected-note 3 {{candidate function (the implicit copy assignment
S3 h, k;
#pragma omp threadprivate(h) // expected-note 2 {{defined as threadprivate or thread local}}
-template <class T> // expected-note {{declared here}}
-T tmain(T argc) {
+template <class T, class S> // expected-note {{declared here}}
+T tfoobar2(T argc, S ub[]) {
const T d = T(); // expected-note 4 {{'d' defined here}}
const T da[5] = {T()}; // expected-note 2 {{'da' defined here}}
T qa[5] = {T()};
@@ -103,6 +103,7 @@ T tmain(T argc) {
const T &r = da[(int)i]; // expected-note 2 {{'r' defined here}}
T &q = qa[(int)i];
T fl;
+ T *ptr;
#pragma omp taskgroup task_reduction // expected-error {{expected '(' after 'task_reduction'}}
foo();
#pragma omp taskgroup task_reduction + // expected-error {{expected '(' after 'task_reduction'}} expected-warning {{extra tokens at the end of '#pragma omp taskgroup' are ignored}}
@@ -125,7 +126,7 @@ T tmain(T argc) {
foo();
#pragma omp taskgroup task_reduction(foo : argc) //omp51-error {{incorrect reduction identifier, expected one of '+', '-', '*', '&', '|', '^', '&&', '||', 'min' or 'max' or declare reduction for type 'float'}} omp51-error {{incorrect reduction identifier, expected one of '+', '-', '*', '&', '|', '^', '&&', '||', 'min' or 'max' or declare reduction for type 'int'}} omp52-error {{incorrect reduction identifier, expected one of '+', '-', '*', '&', '|', '^', '&&', '||', 'min' or 'max' or declare reduction for type 'float'}} omp52-error {{incorrect reduction identifier, expected one of '+', '-', '*', '&', '|', '^', '&&', '||', 'min' or 'max' or declare reduction for type 'int'}} omp60-error {{incorrect reduction identifier, expected one of '+', '*', '&', '|', '^', '&&', '||', 'min' or 'max' or declare reduction for type 'float'}} omp60-error {{incorrect reduction identifier, expected one of '+', '*', '&', '|', '^', '&&', '||', 'min' or 'max' or declare reduction for type 'int'}}
foo();
-#pragma omp taskgroup task_reduction(&& : argc) allocate , allocate(, allocate(omp_default , allocate(omp_default_mem_alloc, allocate(omp_default_mem_alloc:, allocate(omp_default_mem_alloc: argc, allocate(omp_default_mem_alloc: argv), allocate(argv) // expected-error {{expected '(' after 'allocate'}} expected-error 2 {{expected expression}} expected-error 2 {{expected ')'}} expected-error {{use of undeclared identifier 'omp_default'}} expected-note 2 {{to match this '('}}
+#pragma omp taskgroup task_reduction(&& : argc) allocate , allocate(, allocate(omp_default , allocate(omp_default_mem_alloc, allocate(omp_default_mem_alloc:, allocate(omp_default_mem_alloc: argc, allocate(omp_default_mem_alloc: ub), allocate(ub) // expected-error {{expected '(' after 'allocate'}} expected-error 2 {{expected expression}} expected-error 2 {{expected ')'}} expected-error {{use of undeclared identifier 'omp_default'}} expected-note 2 {{to match this '('}}
foo();
#pragma omp taskgroup task_reduction(^ : T) // expected-error {{'T' does not refer to a value}}
foo();
@@ -172,6 +173,10 @@ T tmain(T argc) {
for (int i = 0; i < 10; ++i)
#pragma omp taskgroup task_reduction(+ : fl)
foo();
+#pragma omp taskgroup task_reduction(+:ub[1:]) // expected-error 2 {{section length is unspecified and cannot be inferred because subscripted value is an array of unknown bound}}
+ foo();
+#pragma omp taskgroup task_reduction(+:ptr[:]) // expected-error 2 {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+ foo();
return T();
}
@@ -184,7 +189,7 @@ namespace B {
using A::x;
}
-int main(int argc, char **argv) {
+int foobar2(int argc, float ub[]) {
const int d = 5; // expected-note 2 {{'d' defined here}}
const int da[5] = {0}; // expected-note {{'da' defined here}}
int qa[5] = {0};
@@ -196,6 +201,7 @@ int main(int argc, char **argv) {
const int &r = da[i]; // expected-note {{'r' defined here}}
int &q = qa[i];
float fl;
+ float *ptr;
#pragma omp taskgroup task_reduction // expected-error {{expected '(' after 'task_reduction'}}
foo();
#pragma omp taskgroup task_reduction + // expected-error {{expected '(' after 'task_reduction'}} expected-warning {{extra tokens at the end of '#pragma omp taskgroup' are ignored}}
@@ -214,7 +220,7 @@ int main(int argc, char **argv) {
foo();
#pragma omp taskgroup task_reduction(| : argc, // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
foo();
-#pragma omp taskgroup task_reduction(|| : argc > 0 ? argv[1] : argv[2]) // expected-error {{expected variable name, array element or array section}}
+#pragma omp taskgroup task_reduction(|| : argc > 0 ? ub[1] : ub[2]) // expected-error {{expected variable name, array element or array section}}
foo();
#pragma omp taskgroup task_reduction(~ : argc) // expected-error {{expected unqualified-id}}
foo();
@@ -270,6 +276,10 @@ int main(int argc, char **argv) {
static int m;
#pragma omp taskgroup task_reduction(+ : m) // OK
m++;
+#pragma omp taskgroup task_reduction(+:ub[1:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is an array of unknown bound}}
+ foo();
+#pragma omp taskgroup task_reduction(+:ptr[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+ foo();
- return tmain(argc) + tmain(fl); // expected-note {{in instantiation of function template specialization 'tmain<int>' requested here}} expected-note {{in instantiation of function template specialization 'tmain<float>' requested here}}
+ return tfoobar2(argc, ub) + tfoobar2(fl, ub); // expected-note {{in instantiation of function template specialization 'tfoobar2<int, float>' requested here}} expected-note {{in instantiation of function template specialization 'tfoobar2<float, float>' requested here}}
}
diff --git a/clang/test/OpenMP/teams_reduction_messages.cpp b/clang/test/OpenMP/teams_reduction_messages.cpp
index c980dfa61ef2a..ecf358a0fe8ab 100644
--- a/clang/test/OpenMP/teams_reduction_messages.cpp
+++ b/clang/test/OpenMP/teams_reduction_messages.cpp
@@ -107,8 +107,8 @@ class S6 { // expected-note 3 {{candidate function (the implicit copy assignment
S3 h, k;
#pragma omp threadprivate(h) // expected-note 2 {{defined as threadprivate or thread local}}
-template <class T> // expected-note {{declared here}}
-T tmain(T argc) {
+template <class T, class S> // expected-note {{declared here}}
+T tfoobar2(T argc, S ub[]) {
const T d = T(); // expected-note 4 {{'d' defined here}}
const T da[5] = {T()}; // expected-note 2 {{'da' defined here}}
T qa[5] = {T()};
@@ -118,6 +118,7 @@ T tmain(T argc) {
const T &r = da[(int)i]; // expected-note 2 {{'r' defined here}}
T &q = qa[(int)i]; // expected-note 2 {{'q' defined here}}
T fl;
+ T *ptr;
#pragma omp target
#pragma omp teams reduction // expected-error {{expected '(' after 'reduction'}}
foo();
@@ -152,7 +153,7 @@ T tmain(T argc) {
#pragma omp teams reduction(foo : argc) //omp45-error {{incorrect reduction identifier, expected one of '+', '-', '*', '&', '|', '^', '&&', '||', 'min' or 'max' or declare reduction for type 'float'}} omp45-error {{incorrect reduction identifier, expected one of '+', '-', '*', '&', '|', '^', '&&', '||', 'min' or 'max' or declare reduction for type 'int'}} omp51-error {{incorrect reduction identifier, expected one of '+', '-', '*', '&', '|', '^', '&&', '||', 'min' or 'max' or declare reduction for type 'float'}} omp51-error {{incorrect reduction identifier, expected one of '+', '-', '*', '&', '|', '^', '&&', '||', 'min' or 'max' or declare reduction for type 'int'}} omp52-error {{incorrect reduction identifier, expected one of '+', '-', '*', '&', '|', '^', '&&', '||', 'min' or 'max' or declare reduction for type 'float'}} omp52-error {{incorrect reduction identifier, expected one of '+', '-', '*', '&', '|', '^', '&&', '||', 'min' or 'max' or declare reduction for type 'int'}} omp60-error {{incorrect reduction identifier, expected one of '+', '*', '&', '|', '^', '&&', '||', 'min' or 'max' or declare reduction for type 'float'}} omp60-error {{incorrect reduction identifier, expected one of '+', '*', '&', '|', '^', '&&', '||', 'min' or 'max' or declare reduction for type 'int'}}
foo();
#pragma omp target
-#pragma omp teams reduction(&& : argc) allocate , allocate(, allocate(omp_default , allocate(omp_default_mem_alloc, allocate(omp_default_mem_alloc:, allocate(omp_default_mem_alloc: argc, allocate(omp_default_mem_alloc: argv), allocate(argv) // expected-error {{expected '(' after 'allocate'}} expected-error 2 {{expected expression}} expected-error 2 {{expected ')'}} expected-error {{use of undeclared identifier 'omp_default'}} expected-note 2 {{to match this '('}}
+#pragma omp teams reduction(&& : argc) allocate , allocate(, allocate(omp_default , allocate(omp_default_mem_alloc, allocate(omp_default_mem_alloc:, allocate(omp_default_mem_alloc: argc, allocate(omp_default_mem_alloc: ub), allocate(ub) // expected-error {{expected '(' after 'allocate'}} expected-error 2 {{expected expression}} expected-error 2 {{expected ')'}} expected-error {{use of undeclared identifier 'omp_default'}} expected-note 2 {{to match this '('}}
foo();
#pragma omp target
#pragma omp teams reduction(^ : T) // expected-error {{'T' does not refer to a value}}
@@ -224,6 +225,12 @@ T tmain(T argc) {
#pragma omp target
#pragma omp teams reduction(+ : fl)
foo();
+#pragma omp target
+#pragma omp teams reduction(+ : ptr[0:]) // expected-error 2 {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+ foo();
+#pragma omp target
+#pragma omp teams reduction(* : ub[:]) // expected-error 2 {{section length is unspecified and cannot be inferred because subscripted value is an array of unknown bound}}
+ foo();
return T();
}
@@ -236,7 +243,7 @@ namespace B {
using A::x;
}
-int main(int argc, char **argv) {
+int foobar2(int argc, float ub[]) {
const int d = 5; // expected-note 2 {{'d' defined here}}
const int da[5] = {0}; // expected-note {{'da' defined here}}
int qa[5] = {0};
@@ -248,6 +255,7 @@ int main(int argc, char **argv) {
const int &r = da[i]; // expected-note {{'r' defined here}}
int &q = qa[i]; // expected-note {{'q' defined here}}
float fl;
+ float *ptr;
#pragma omp target
#pragma omp teams reduction // expected-error {{expected '(' after 'reduction'}}
foo();
@@ -276,7 +284,7 @@ int main(int argc, char **argv) {
#pragma omp teams reduction(| : argc, // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
foo();
#pragma omp target
-#pragma omp teams reduction(|| : argc > 0 ? argv[1] : argv[2]) // expected-error {{expected variable name, array element or array section}}
+#pragma omp teams reduction(|| : argc > 0 ? ub[1] : ub[2]) // expected-error {{expected variable name, array element or array section}}
foo();
#pragma omp target
#pragma omp teams reduction(~ : argc) // expected-error {{expected unqualified-id}}
@@ -364,6 +372,12 @@ int main(int argc, char **argv) {
#pragma omp target
#pragma omp teams reduction(task, + : m) // omp45-error 2 {{expected expression}} omp45-warning {{missing ':' after reduction identifier - ignoring}} omp51-error {{'reduction' clause with 'task' modifier allowed only on non-simd parallel or worksharing constructs}} omp52-error {{'reduction' clause with 'task' modifier allowed only on non-simd parallel or worksharing constructs}} omp60-error {{'reduction' clause with 'task' modifier allowed only on non-simd parallel or worksharing constructs}}
foo();
+#pragma omp target
+#pragma omp teams reduction(+ : ptr[0:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+ foo();
+#pragma omp target
+#pragma omp teams reduction(* : ub[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is an array of unknown bound}}
+ foo();
- return tmain(argc) + tmain(fl); // expected-note {{in instantiation of function template specialization 'tmain<int>' requested here}} expected-note {{in instantiation of function template specialization 'tmain<float>' requested here}}
+ return tfoobar2(argc, ub) + tfoobar2(fl, ub); // expected-note {{in instantiation of function template specialization 'tfoobar2<int, float>' requested here}} expected-note {{in instantiation of function template specialization 'tfoobar2<float, float>' requested here}}
}
More information about the cfe-commits
mailing list