[clang][OpenMP] 6.0: detect privatization of array section/assumed-size array by ddpagan · Pull Request #152786 · llvm/llvm-project (original) (raw)
@llvm/pr-subscribers-clang
Author: David Pagan (ddpagan)
Changes
According to the OpenMP 6.0 specification, array sections with no length and unknown size are considered assumed-size arrays. As of pull request
#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)
Patch is 115.25 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/152786.diff
7 Files Affected:
- (modified) clang/lib/Sema/SemaOpenMP.cpp (+36-10)
- (modified) clang/test/OpenMP/scan_messages.cpp (+29-3)
- (modified) clang/test/OpenMP/target_data_use_device_addr_codegen.cpp (+909-180)
- (modified) clang/test/OpenMP/target_has_device_addr_messages.cpp (+21-5)
- (modified) clang/test/OpenMP/task_in_reduction_message.cpp (+33-22)
- (modified) clang/test/OpenMP/taskgroup_task_reduction_messages.cpp (+16-6)
- (modified) clang/test/OpenMP/teams_reduction_messages.cpp (+20-6)
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,
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.bool AllowAssumedSizeArray, StringRef DiagType) {- //
- // 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(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,
if (Res.second) { // It will be analyzed later./*DiagType=*/"omp_interop_t");
@@ -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,
if (Res.second) { // It will be analyzed later. MVLI.ProcessedVarList.push_back(RefExpr);/*AllowAssumedSizeArray=*/true);
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 -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(); // expected-note {{in instantiation of function template specialization 'tmain' requested here}}
- return tfoobar(ub); // expected-note {{in instantiation of function template specialization 'tfoobar' 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]], p... [truncated]