This is an archive of the discontinued LLVM Phabricator instance.

Differential D134974

[OpenMP] Add map clause to the LIT test on use_device_addr clause
ClosedPublic

Authored by animeshk-amd on Sep 30 2022, 10:20 AM.

Details

Reviewers
saiislam
JonChesterfield
jdoerfert
Commits
rG0f8e7b43291c: [OpenMP] Add map clause to the LIT test on use_device_addr clause
Summary

As per the OpenMP Spec, "A list item in a use_device_addr clause
must have a corresponding list item in the device data environment"
. Therefore a map clause is added which will make sure that the
respective list items are mapped to the device data environment
before the use_device_addr clause is specified. The CHECK lines
are also modified based on this change.

Diff Detail

Event Timeline

animeshk-amd created this revision.Sep 30 2022, 10:20 AM
Herald added a project: Restricted Project. · View Herald TranscriptSep 30 2022, 10:20 AM
Herald added subscribers: guansong, yaxunl. · View Herald Transcript
animeshk-amd requested review of this revision.Sep 30 2022, 10:20 AM
Herald added a reviewer: jdoerfert. · View Herald TranscriptSep 30 2022, 10:20 AM
Herald added a project: Restricted Project. · View Herald Transcript
Herald added subscribers: cfe-commits, sstefan1. · View Herald Transcript
Harbormaster completed remote builds in B189723: Diff 464320.Sep 30 2022, 10:55 AM
animeshk-amd added a comment.Oct 21 2022, 12:29 PM

ping

saiislam accepted this revision.Oct 21 2022, 12:33 PM

LGTM!

This revision is now accepted and ready to land.Oct 21 2022, 12:33 PM
animeshk-amd updated this revision to Diff 473193.Nov 4 2022, 4:45 AM

The main branch dropped the use of flag -no-opaque-pointers in
this test which generated a different IR than before. This revision
gets rid of the conflicts with the main branch caused by above.

Harbormaster completed remote builds in B196112: Diff 473193.Nov 4 2022, 5:15 AM
Closed by commit rG0f8e7b43291c: [OpenMP] Add map clause to the LIT test on use_device_addr clause (authored by animeshk-amd). · Explain WhyNov 8 2022, 11:00 PM
This revision was automatically updated to reflect the committed changes.
animeshk-amd added a commit: rG0f8e7b43291c: [OpenMP] Add map clause to the LIT test on use_device_addr clause.

Revision Contents

PathSize
clang/
test/
OpenMP/
190 lines

Diff 474161

clang/test/OpenMP/target_data_use_device_addr_codegen.cpp

// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s // RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s
// RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s // RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-version=50 -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-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s // RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
// RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -fopenmp-version=50 -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 -verify -fopenmp-simd -fopenmp-version=50 -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-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s // RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-version=50 -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-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// SIMD-ONLY0-NOT: {{__kmpc|__tgt}} // SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
// expected-no-diagnostics // expected-no-diagnostics
#ifndef HEADER #ifndef HEADER
#define HEADER #define HEADER
// CHECK-DAG: [[SIZES1:@.+]] = private unnamed_addr constant [5 x i64] zeroinitializer // 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 // 64 = 0x40 = OMP_MAP_RETURN_PARAM
// CHECK-DAG: [[MAPTYPES1:@.+]] = private unnamed_addr constant [5 x i64] [i64 64, i64 64, i64 64, i64 64, i64 64] // CHECK-DAG: [[MAPTYPES1:@.+]] = private unnamed_addr constant [6 x i64] [i64 67, i64 67, i64 3, i64 67, i64 67, i64 67]
// CHECK-DAG: [[SIZES2:@.+]] = private unnamed_addr constant [5 x i64] zeroinitializer // 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 // 0 = OMP_MAP_NONE
// 281474976710720 = 0x1000000000040 = OMP_MAP_MEMBER_OF | OMP_MAP_RETURN_PARAM // 281474976710720 = 0x1000000000040 = OMP_MAP_MEMBER_OF | OMP_MAP_RETURN_PARAM
// CHECK-DAG: [[MAPTYPES2:@.+]] = private unnamed_addr constant [5 x i64] [i64 0, i64 281474976710720, i64 281474976710720, i64 281474976710720, i64 281474976710720] // CHECK-DAG: [[MAPTYPES2:@.+]] = private unnamed_addr constant [6 x i64] [i64 0, i64 281474976710723, i64 281474976710739, i64 281474976710739, i64 281474976710675, i64 281474976710723]
struct S { struct S {
int a = 0; int a = 0;
int *ptr = &a; int *ptr = &a;
int &ref = a; int &ref = a;
int arr[4]; int arr[4];
S() {} S() {}
void foo() { void foo() {
#pragma omp target data use_device_addr(a, ptr [3:4], ref, ptr[0], arr[:a]) #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])
++a, ++*ptr, ++ref, ++arr[0]; ++a, ++*ptr, ++ref, ++arr[0];
} }
}; };
int main() { int main() {
float a = 0; float a = 0;
float *ptr = &a; float *ptr = &a;
float &ref = a; float &ref = a;
float arr[4]; float arr[4];
float vla[(int)a]; float vla[(int)a];
S s; S s;
s.foo(); s.foo();
#pragma omp target data use_device_addr(a, ptr [3:4], ref, ptr[0], arr[:(int)a], vla[0]) #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])
++a, ++*ptr, ++ref, ++arr[0], ++vla[0]; ++a, ++*ptr, ++ref, ++arr[0], ++vla[0];
return a; return a;
} }
// CHECK-LABEL: @main() // CHECK-LABEL: @main()
// CHECK: [[A_ADDR:%.+]] = alloca float, // CHECK: [[A_ADDR:%.+]] = alloca float,
// CHECK: [[PTR_ADDR:%.+]] = alloca ptr, // CHECK: [[PTR_ADDR:%.+]] = alloca ptr,
// CHECK: [[REF_ADDR:%.+]] = alloca ptr, // CHECK: [[REF_ADDR:%.+]] = alloca ptr,
// CHECK: [[ARR_ADDR:%.+]] = alloca [4 x float], // CHECK: [[ARR_ADDR:%.+]] = alloca [4 x float],
// CHECK: [[BPTRS:%.+]] = alloca [5 x ptr], // CHECK: [[BPTRS:%.+]] = alloca [6 x ptr],
// CHECK: [[PTRS:%.+]] = alloca [5 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: [[VLA_ADDR:%.+]] = alloca float, i64 %{{.+}},
// CHECK: [[PTR:%.+]] = load ptr, ptr [[PTR_ADDR]], // CHECK: [[PTR:%.+]] = load ptr, ptr [[PTR_ADDR]],
// CHECK: [[REF:%.+]] = load ptr, ptr [[REF_ADDR]], // CHECK-NEXT: [[P4:%.+]] = load ptr, ptr [[PTR_ADDR]], align 8
// CHECK: [[ARR:%.+]] = getelementptr inbounds [4 x float], ptr [[ARR_ADDR]], i64 0, i64 0 // CHECK-NEXT: [[ARR_IDX:%.+]] = getelementptr inbounds float, ptr [[P4]], i64 3
// CHECK: [[BPTR0:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BPTRS]], i32 0, i32 0 // CHECK: [[P5:%.+]] = load ptr, ptr [[PTR_ADDR]], align 8
// CHECK-NEXT: [[P6:%.+]] = load ptr, ptr [[PTR_ADDR]], align 8
// CHECK-NEXT: [[ARR_IDX1:%.+]] = getelementptr inbounds float, ptr [[P6]], i64 0
// CHECK: [[P7:%.+]] = load ptr, ptr [[REF_ADDR]],
// CHECK-NEXT: [[REF:%.+]] = load ptr, ptr [[REF_ADDR]],
// CHECK-NEXT: [[ARR_IDX2:%.+]] = getelementptr inbounds [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: store ptr [[A_ADDR]], ptr [[BPTR0]],
// CHECK: [[PTR0:%.+]] = getelementptr inbounds [5 x ptr], ptr [[PTRS]], i32 0, i32 0 // CHECK: [[PTR0:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 0
// CHECK: store ptr [[A_ADDR]], ptr [[PTR0]], // CHECK: store ptr [[A_ADDR]], ptr [[PTR0]],
// CHECK: [[BPTR1:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BPTRS]], i32 0, i32 1 // CHECK: [[BPTR1:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 1
// CHECK: store ptr [[PTR]], ptr [[BPTR1]], // CHECK: store ptr [[PTR]], ptr [[BPTR1]],
// CHECK: [[PTR1:%.+]] = getelementptr inbounds [5 x ptr], ptr [[PTRS]], i32 0, i32 1 // CHECK: [[PTR1:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 1
// CHECK: store ptr [[PTR]], ptr [[PTR1]], // CHECK: store ptr [[ARR_IDX]], ptr [[PTR1]],
// CHECK: [[BPTR2:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BPTRS]], i32 0, i32 2 // CHECK: [[BPTR2:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 2
// CHECK: store ptr [[REF]], ptr [[BPTR2]], // CHECK: store ptr [[P5]], ptr [[BPTR2]],
// CHECK: [[PTR2:%.+]] = getelementptr inbounds [5 x ptr], ptr [[PTRS]], i32 0, i32 2 // CHECK: [[PTR2:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 2
// CHECK: store ptr [[REF]], ptr [[PTR2]], // CHECK: store ptr [[ARR_IDX1]], ptr [[PTR2]],
// CHECK: [[BPTR3:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BPTRS]], i32 0, i32 3 // CHECK: [[BPTR3:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 3
// CHECK: store ptr [[ARR]], ptr [[BPTR3]], // CHECK: store ptr [[P7]], ptr [[BPTR3]],
// CHECK: [[PTR3:%.+]] = getelementptr inbounds [5 x ptr], ptr [[PTRS]], i32 0, i32 3 // CHECK: [[PTR3:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 3
// CHECK: store ptr [[ARR]], ptr [[PTR3]], // CHECK: store ptr [[REF]], ptr [[PTR3]],
// CHECK: [[BPTR4:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BPTRS]], i32 0, i32 4 // CHECK: [[BPTR4:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 4
// CHECK: store ptr [[VLA_ADDR]], ptr [[BPTR4]], // CHECK: store ptr [[ARR_ADDR]], ptr [[BPTR4]], align
// CHECK: [[PTR4:%.+]] = getelementptr inbounds [5 x ptr], ptr [[PTRS]], i32 0, i32 4 // CHECK: [[PTR4:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 4
// CHECK: store ptr [[VLA_ADDR]], ptr [[PTR4]], // CHECK: store ptr [[ARR_IDX2]], ptr [[PTR4]], align 8
// CHECK: [[BPTR:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BPTRS]], i32 0, i32 0 // CHECK: [[SIZE_PTR:%.+]] = getelementptr inbounds [6 x i64], ptr [[SIZES]], i32 0, i32 4
// CHECK: [[PTR:%.+]] = getelementptr inbounds [5 x ptr], ptr [[PTRS]], i32 0, i32 0 // CHECK: store i64 [[P10:%.+]], ptr [[SIZE_PTR]], align 8
// CHECK: call void @__tgt_target_data_begin_mapper(ptr @{{.+}}, i64 -1, i32 5, ptr [[BPTR]], ptr [[PTR]], ptr [[SIZES1]], ptr [[MAPTYPES1]], ptr null, ptr null) // 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: [[A_REF:%.+]] = load ptr, ptr [[BPTR0]],
// CHECK: [[REF_REF:%.+]] = load ptr, ptr [[BPTR2]], // CHECK: [[REF_REF:%.+]] = load ptr, ptr [[BPTR3]],
// CHECK: store ptr [[REF_REF]], ptr [[TMP_REF_ADDR:%.+]], // CHECK: store ptr [[REF_REF]], ptr [[TMP_REF_ADDR:%.+]],
// CHECK: [[ARR_REF:%.+]] = load ptr, ptr [[BPTR3]], // CHECK: [[ARR_REF:%.+]] = load ptr, ptr [[BPTR4]],
// CHECK: [[VLA_REF:%.+]] = load ptr, ptr [[BPTR4]], // CHECK: [[VLA_REF:%.+]] = load ptr, ptr [[BPTR5]],
// CHECK: [[A:%.+]] = load float, ptr [[A_REF]], // CHECK: [[A:%.+]] = load float, ptr [[A_REF]],
// CHECK: [[INC:%.+]] = fadd float [[A]], 1.000000e+00 // CHECK: [[INC:%.+]] = fadd float [[A]], 1.000000e+00
// CHECK: store float [[INC]], ptr [[A_REF]], // CHECK: store float [[INC]], ptr [[A_REF]],
// CHECK: [[PTR_ADDR:%.+]] = load ptr, ptr [[BPTR1]], // CHECK: [[PTR:%.+]] = load ptr, ptr [[BPTR1]],
// CHECK: [[VAL:%.+]] = load float, ptr [[PTR_ADDR]], // CHECK: [[VAL:%.+]] = load float, ptr [[PTR]],
// CHECK: [[INC:%.+]] = fadd float [[VAL]], 1.000000e+00 // CHECK: [[INC:%.+]] = fadd float [[VAL]], 1.000000e+00
// CHECK: store float [[INC]], ptr [[PTR_ADDR]], // CHECK: store float [[INC]], ptr [[PTR]],
// CHECK: [[REF_ADDR:%.+]] = load ptr, ptr [[TMP_REF_ADDR]], // CHECK: [[REF_ADDR:%.+]] = load ptr, ptr [[TMP_REF_ADDR]],
// CHECK: [[REF:%.+]] = load float, ptr [[REF_ADDR]], // CHECK: [[REF:%.+]] = load float, ptr [[REF_ADDR]],
// CHECK: [[INC:%.+]] = fadd float [[REF]], 1.000000e+00 // CHECK: [[INC:%.+]] = fadd float [[REF]], 1.000000e+00
// CHECK: store float [[INC]], ptr [[REF_ADDR]], // CHECK: store float [[INC]], ptr [[REF_ADDR]],
// CHECK: [[ARR0_ADDR:%.+]] = getelementptr inbounds [4 x float], ptr [[ARR_REF]], i64 0, i64 0 // CHECK: [[ARR0_ADDR:%.+]] = getelementptr inbounds [4 x float], ptr [[ARR_REF]], i64 0, i64 0
// CHECK: [[ARR0:%.+]] = load float, ptr [[ARR0_ADDR]], // CHECK: [[ARR0:%.+]] = load float, ptr [[ARR0_ADDR]],
// CHECK: [[INC:%.+]] = fadd float [[ARR0]], 1.000000e+00 // CHECK: [[INC:%.+]] = fadd float [[ARR0]], 1.000000e+00
// CHECK: store float [[INC]], ptr [[ARR0_ADDR]], // CHECK: store float [[INC]], ptr [[ARR0_ADDR]],
// CHECK: [[VLA0_ADDR:%.+]] = getelementptr inbounds float, ptr [[VLA_REF]], i64 0 // CHECK: [[VLA0_ADDR:%.+]] = getelementptr inbounds float, ptr [[VLA_REF]], i64 0
// CHECK: [[VLA0:%.+]] = load float, ptr [[VLA0_ADDR]], // CHECK: [[VLA0:%.+]] = load float, ptr [[VLA0_ADDR]],
// CHECK: [[INC:%.+]] = fadd float [[VLA0]], 1.000000e+00 // CHECK: [[INC:%.+]] = fadd float [[VLA0]], 1.000000e+00
// CHECK: store float [[INC]], ptr [[VLA0_ADDR]], // CHECK: store float [[INC]], ptr [[VLA0_ADDR]],
// CHECK: [[BPTR:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BPTRS]], i32 0, i32 0 // CHECK: [[BPTR:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 0
// CHECK: [[PTR:%.+]] = getelementptr inbounds [5 x ptr], ptr [[PTRS]], i32 0, i32 0 // CHECK: [[PTR:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 0
// CHECK: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 -1, i32 5, ptr [[BPTR]], ptr [[PTR]], ptr [[SIZES1]], ptr [[MAPTYPES1]], ptr null, ptr null) // 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: foo
// %this.addr = alloca ptr, align 8 // CHECK: [[BPTRS:%.+]] = alloca [6 x ptr],
// CHECK: [[BPTRS:%.+]] = alloca [5 x ptr], // CHECK: [[PTRS:%.+]] = alloca [6 x ptr],
// CHECK: [[PTRS:%.+]] = alloca [5 x ptr], // CHECK: [[MAP_PTRS:%.+]] = alloca [6 x ptr],
// CHECK: [[SIZES:%.+]] = alloca [5 x i64], // CHECK: [[SIZES:%.+]] = alloca [6 x i64],
// %tmp = alloca ptr, align 8
// %tmp6 = alloca ptr, align 8
// %tmp7 = alloca ptr, align 8
// %tmp8 = alloca ptr, align 8
// %tmp9 = alloca ptr, align 8
// store ptr %this, ptr %this.addr, align 8
// %this1 = load ptr, ptr %this.addr, align 8
// CHECK: [[A_ADDR:%.+]] = getelementptr inbounds %struct.S, ptr [[THIS:%.+]], i32 0, i32 0 // CHECK: [[A_ADDR:%.+]] = getelementptr inbounds %struct.S, ptr [[THIS:%.+]], i32 0, i32 0
// %ptr = getelementptr inbounds %struct.S, ptr %this1, i32 0, i32 1
// %ref = getelementptr inbounds %struct.S, ptr %this1, i32 0, i32 2
// %0 = load ptr, ptr %ref, align 8
// CHECK: [[ARR_ADDR:%.+]] = getelementptr inbounds %struct.S, ptr [[THIS]], i32 0, i32 3
// CHECK: [[A_ADDR2:%.+]] = getelementptr inbounds %struct.S, ptr [[THIS]], i32 0, i32 0
// CHECK: [[PTR_ADDR:%.+]] = getelementptr inbounds %struct.S, ptr [[THIS]], i32 0, i32 1 // CHECK: [[PTR_ADDR:%.+]] = getelementptr inbounds %struct.S, ptr [[THIS]], i32 0, i32 1
// CHECK: [[ARR_IDX:%.+]] = getelementptr inbounds i32, ptr %{{.+}}, i64 3
// CHECK: [[REF_REF:%.+]] = getelementptr inbounds %struct.S, ptr [[THIS]], i32 0, i32 2 // CHECK: [[REF_REF:%.+]] = getelementptr inbounds %struct.S, ptr [[THIS]], i32 0, i32 2
// CHECK: [[REF_PTR:%.+]] = load ptr, ptr [[REF_REF]], // CHECK: [[REF_PTR:%.+]] = load ptr, ptr [[REF_REF]],
// CHECK: [[ARR_ADDR2:%.+]] = getelementptr inbounds %struct.S, ptr [[THIS]], i32 0, i32 3 // CHECK-NEXT: [[P3:%.+]] = getelementptr inbounds %struct.S, ptr [[THIS]], i32 0, i32 1
// CHECK: [[ARR_END:%.+]] = getelementptr [4 x i32], ptr [[ARR_ADDR]], i32 1 // CHECK: [[ARR_IDX5:%.+]] = getelementptr inbounds i32, ptr {{.+}}, i64 0
// CHECK: [[ARR_ADDR:%.+]] = getelementptr inbounds %struct.S, ptr [[THIS]], i32 0, i32 3
// CHECK: [[ARR_IDX6:%.+]] = getelementptr inbounds [4 x i32], ptr [[ARR_ADDR]], i64 0, i64 0
// CHECK: [[A_ADDR2:%.+]] = getelementptr inbounds %struct.S, ptr [[THIS]], i32 0, i32 0
// CHECK: [[P4:%.+]] = mul nuw i64 [[CONV:%.+]], 4
// CHECK: [[ARR_END:%.+]] = getelementptr i32, ptr [[ARR_IDX6]], i32 1
// CHECK: [[E:%.+]] = ptrtoint ptr [[ARR_END]] to i64 // CHECK: [[E:%.+]] = ptrtoint ptr [[ARR_END]] to i64
// CHECK: [[B:%.+]] = ptrtoint ptr [[A_ADDR]] to i64 // CHECK: [[B:%.+]] = ptrtoint ptr [[A_ADDR]] to i64
// CHECK: [[DIFF:%.+]] = sub i64 [[E]], [[B]] // CHECK: [[DIFF:%.+]] = sub i64 [[E]], [[B]]
// CHECK: [[SZ:%.+]] = sdiv exact i64 [[DIFF]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64) // CHECK: [[SZ:%.+]] = sdiv exact i64 [[DIFF]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64)
// CHECK: [[BPTR0:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BPTRS]], i32 0, i32 0 // CHECK: [[BPTR0:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 0
// CHECK: store ptr [[THIS]], ptr [[BPTR0]], // CHECK: store ptr [[THIS]], ptr [[BPTR0]],
// CHECK: [[PTR0:%.+]] = getelementptr inbounds [5 x ptr], ptr [[PTRS]], i32 0, i32 0 // CHECK: [[PTR0:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 0
// CHECK: store ptr [[A_ADDR]], ptr [[PTR0]], // CHECK: store ptr [[A_ADDR]], ptr [[PTR0]],
// CHECK: [[SIZE0:%.+]] = getelementptr inbounds [5 x i64], ptr [[SIZES]], i32 0, i32 0 // CHECK: [[SIZE0:%.+]] = getelementptr inbounds [6 x i64], ptr [[SIZES]], i32 0, i32 0
// CHECK: store i64 [[SZ]], ptr [[SIZE0]], // CHECK: store i64 [[SZ]], ptr [[SIZE0]],
// CHECK: [[BPTR1:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BPTRS]], i32 0, i32 1 // CHECK: [[BPTR1:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 1
// CHECK: store ptr [[A_ADDR2]], ptr [[BPTR1]], // CHECK: store ptr [[THIS]], ptr [[BPTR1]]
// CHECK: [[PTR1:%.+]] = getelementptr inbounds [5 x ptr], ptr [[PTRS]], i32 0, i32 1 // CHECK: [[PTR1:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 1
// CHECK: store ptr [[A_ADDR2]], ptr [[PTR1]], // CHECK: store ptr [[A_ADDR]], ptr [[PTR1]],
// CHECK: [[BPTR2:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BPTRS]], i32 0, i32 2 // CHECK: [[BPTR2:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 2
// CHECK: store ptr [[PTR_ADDR]], ptr [[BPTR2]], // CHECK: store ptr [[PTR_ADDR]], ptr [[BPTR2]],
// CHECK: [[PTR2:%.+]] = getelementptr inbounds [5 x ptr], ptr [[PTRS]], i32 0, i32 2 // CHECK: [[PTR2:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 2
// CHECK: store ptr [[PTR_ADDR]], ptr [[PTR2]], // CHECK: store ptr [[ARR_IDX]], ptr [[PTR2]],
// CHECK: [[BPTR3:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BPTRS]], i32 0, i32 3 // CHECK: [[BPTR3:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 3
// CHECK: store ptr [[REF_PTR]], ptr [[BPTR3]], // CHECK: store ptr [[THIS]], ptr [[BPTR3]]
// CHECK: [[PTR3:%.+]] = getelementptr inbounds [5 x ptr], ptr [[PTRS]], i32 0, i32 3 // CHECK: [[PTR3:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 3
// CHECK: store ptr [[REF_PTR]], ptr [[PTR3]], // CHECK: store ptr [[REF_PTR]], ptr [[PTR3]],
// CHECK: [[BPTR4:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BPTRS]], i32 0, i32 4 // CHECK: [[BPTR4:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 4
// CHECK: store ptr [[ARR_ADDR2]], ptr [[BPTR4]], // CHECK: store ptr [[P3]], ptr [[BPTR4]],
// CHECK: [[PTR4:%.+]] = getelementptr inbounds [5 x ptr], ptr [[PTRS]], i32 0, i32 4 // CHECK: [[PTR4:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 4
// CHECK: store ptr [[ARR_ADDR2]], ptr [[PTR4]], // CHECK: store ptr [[ARR_IDX5]], ptr [[PTR4]]
// CHECK: [[BPTR:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BPTRS]], i32 0, i32 0
// CHECK: [[PTR:%.+]] = getelementptr inbounds [5 x ptr], ptr [[PTRS]], i32 0, i32 0 // CHECK: [[BPTR5:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 5
// CHECK: [[SIZE:%.+]] = getelementptr inbounds [5 x i64], ptr [[SIZES]], i32 0, i32 0 // CHECK: store ptr [[THIS]], ptr [[BPTR5]], align 8
// CHECK: call void @__tgt_target_data_begin_mapper(ptr @{{.+}}, i64 -1, i32 5, ptr [[BPTR]], ptr [[PTR]], ptr [[SIZE]], ptr [[MAPTYPES2]], ptr null, ptr null) // 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: [[A_ADDR:%.+]] = load ptr, ptr [[BPTR1]],
// CHECK: store ptr [[A_ADDR]], ptr [[A_REF:%.+]], // CHECK: store ptr [[A_ADDR]], ptr [[A_REF:%.+]],
// CHECK: [[PTR_ADDR:%.+]] = load ptr, ptr [[BPTR2]], // CHECK: [[PTR_ADDR:%.+]] = load ptr, ptr [[BPTR2]],
// CHECK: store ptr [[PTR_ADDR]], ptr [[PTR_REF:%.+]], // CHECK: store ptr [[PTR_ADDR]], ptr [[PTR_REF:%.+]],
// CHECK: [[REF_PTR:%.+]] = load ptr, ptr [[BPTR3]], // CHECK: [[REF_PTR:%.+]] = load ptr, ptr [[BPTR3]],
// CHECK: store ptr [[REF_PTR]], ptr [[REF_REF:%.+]], // CHECK: store ptr [[REF_PTR]], ptr [[REF_REF:%.+]],
// CHECK: [[PTR_ADDR:%.+]] = load ptr, ptr [[BPTR2]], // CHECK: [[PTR_ADDR:%.+]] = load ptr, ptr [[BPTR2]],
// CHECK: store ptr [[PTR_ADDR]], ptr [[PTR_REF2:%.+]], // CHECK: store ptr [[PTR_ADDR]], ptr [[PTR_REF2:%.+]],
// CHECK: [[ARR_ADDR:%.+]] = load ptr, ptr [[BPTR4]], // CHECK: [[ARR_ADDR:%.+]] = load ptr, ptr [[BPTR5]],
// CHECK: store ptr [[ARR_ADDR]], ptr [[ARR_REF:%.+]], // CHECK: store ptr [[ARR_ADDR]], ptr [[ARR_REF:%.+]],
// CHECK: [[A_ADDR:%.+]] = load ptr, ptr [[A_REF]], // CHECK: [[A_ADDR:%.+]] = load ptr, ptr [[A_REF]],
// CHECK: [[A:%.+]] = load i32, ptr [[A_ADDR]], // CHECK: [[A:%.+]] = load i32, ptr [[A_ADDR]],
// CHECK: [[INC:%.+]] = add nsw i32 [[A]], 1 // CHECK: [[INC:%.+]] = add nsw i32 [[A]], 1
// CHECK: store i32 [[INC]], ptr [[A_ADDR]], // CHECK: store i32 [[INC]], ptr [[A_ADDR]],
// CHECK: [[PTR_PTR:%.+]] = load ptr, ptr [[PTR_REF2]], // CHECK: [[PTR_PTR:%.+]] = load ptr, ptr [[PTR_REF2]],
// CHECK: [[PTR:%.+]] = load ptr, ptr [[PTR_PTR]], // CHECK: [[PTR:%.+]] = load ptr, ptr [[PTR_PTR]],
// CHECK: [[VAL:%.+]] = load i32, ptr [[PTR]], // CHECK: [[VAL:%.+]] = load i32, ptr [[PTR]],
// CHECK: [[INC:%.+]] = add nsw i32 [[VAL]], 1 // CHECK: [[INC:%.+]] = add nsw i32 [[VAL]], 1
// CHECK: store i32 [[INC]], ptr [[PTR]], // CHECK: store i32 [[INC]], ptr [[PTR]],
// CHECK: [[REF_PTR:%.+]] = load ptr, ptr [[REF_REF]], // CHECK: [[REF_PTR:%.+]] = load ptr, ptr [[REF_REF]],
// CHECK: [[VAL:%.+]] = load i32, ptr [[REF_PTR]], // CHECK: [[VAL:%.+]] = load i32, ptr [[REF_PTR]],
// CHECK: [[INC:%.+]] = add nsw i32 [[VAL]], 1 // CHECK: [[INC:%.+]] = add nsw i32 [[VAL]], 1
// CHECK: store i32 [[INC]], ptr [[REF_PTR]], // CHECK: store i32 [[INC]], ptr [[REF_PTR]],
// CHECK: [[ARR_ADDR:%.+]] = load ptr, ptr [[ARR_REF]], // CHECK: [[ARR_ADDR:%.+]] = load ptr, ptr [[ARR_REF]],
// CHECK: [[ARR0_ADDR:%.+]] = getelementptr inbounds [4 x i32], ptr [[ARR_ADDR]], i64 0, i64 0 // CHECK: [[ARR0_ADDR:%.+]] = getelementptr inbounds [4 x i32], ptr [[ARR_ADDR]], i64 0, i64 0
// CHECK: [[VAL:%.+]] = load i32, ptr [[ARR0_ADDR]], // CHECK: [[VAL:%.+]] = load i32, ptr [[ARR0_ADDR]],
// CHECK: [[INC:%.+]] = add nsw i32 [[VAL]], 1 // CHECK: [[INC:%.+]] = add nsw i32 [[VAL]], 1
// CHECK: store i32 [[INC]], ptr [[ARR0_ADDR]], // CHECK: store i32 [[INC]], ptr [[ARR0_ADDR]],
// CHECK: [[BPTR:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BPTRS]], i32 0, i32 0 // CHECK: [[BPTR:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 0
// CHECK: [[PTR:%.+]] = getelementptr inbounds [5 x ptr], ptr [[PTRS]], i32 0, i32 0 // CHECK: [[PTR:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 0
// CHECK: [[SIZE:%.+]] = getelementptr inbounds [5 x i64], ptr [[SIZES]], 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 5, ptr [[BPTR]], ptr [[PTR]], ptr [[SIZE]], ptr [[MAPTYPES2]], ptr null, ptr null) // 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 #endif