Skip to content

Use ATTACH maps for array-sections/subscripts on pointers. #1

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Draft
wants to merge 28 commits into
base: tgt-capture-mapped-ptrs-by-ref
Choose a base branch
from
+2,790 −1,472
Draft

Use ATTACH maps for array-sections/subscripts on pointers. #1

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
28 commits
Select commit Hold shift + click to select a range
cc90643
[WIP] Use ATTACH maps for array-sections/subscripts on pointers.
abhinavgaba Jun 11, 2025
6ba75e1
[WIP] Commit some more non-debug print changes.
abhinavgaba Jun 25, 2025
9408873
Merge branch 'tgt-capture-mapped-ptrs-by-ref' into map-ptr-array-sect…
abhinavgaba Jul 8, 2025
3ce181f
Fix findAttachComponent
abhinavgaba Jul 8, 2025
50c8d0e
Fix build issues, keep debug prints
abhinavgaba Jul 8, 2025
c23e01f
Fix handling of array-sections in findAttachComponent
abhinavgaba Jul 8, 2025
9adf9a8
Fix one test.
abhinavgaba Jul 8, 2025
ca0c381
Merge branch 'tgt-capture-mapped-ptrs-by-ref' into map-ptr-array-sect…
abhinavgaba Jul 8, 2025
36f6d9b
Fix star case. Update some tests.
abhinavgaba Jul 9, 2025
8349a49
Refactor AddAttachEntry
abhinavgaba Jul 9, 2025
eaf332f
Delay addition of attach when we populate PartialStruct.
abhinavgaba Jul 9, 2025
da55bf8
Update some more tests.
abhinavgaba Jul 16, 2025
fcdc13a
More test updates
abhinavgaba Jul 16, 2025
c35d7a0
[WIP][Offload] Introduce ATTACH map-type support for pointer attachment.
abhinavgaba Jul 16, 2025
13faca1
Minor formatting changes.
abhinavgaba Jul 16, 2025
54b2ae4
Remove debug prints.
abhinavgaba Jul 16, 2025
00b0767
Limit attach map-type generation to only map-entering constructs.
abhinavgaba Jul 16, 2025
baa9dbb
Minor clean-up
abhinavgaba Jul 16, 2025
c0d20e9
Minor NFC changes.
abhinavgaba Jul 20, 2025
a9b94d4
Use lb instead of base when attaching to a partialstruct.
abhinavgaba Jul 20, 2025
f1acc37
[NFC] Move computation of ptr/ptee addrs outside the loop.
abhinavgaba Jul 21, 2025
976470c
[NFC] Return Expr* from findAttachBasePtr.
abhinavgaba Jul 21, 2025
fd2d077
[NFC] Minor renaming/comment changes.
abhinavgaba Jul 21, 2025
fdced7d
Fix call to getAttachptrPteeAddrs.
abhinavgaba Jul 21, 2025
d791cb3
Remove unnecessary include.
abhinavgaba Jul 21, 2025
df62d2e
No need to compute the pointee address from scratch.
abhinavgaba Jul 22, 2025
63d2ca4
Update tests after avoiding creating pointee addr, and a minor commen…
abhinavgaba Jul 22, 2025
c9f8c0a
Move initialization of FinalLowestElem earlier.
abhinavgaba Jul 22, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
8 changes: 8 additions & 0 deletions clang/include/clang/Basic/OpenMPKinds.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -301,6 +301,14 @@ bool isOpenMPTargetExecutionDirective(OpenMPDirectiveKind DKind);
/// otherwise - false.
bool isOpenMPTargetDataManagementDirective(OpenMPDirectiveKind DKind);

/// Checks if the specified directive is a map-entering target directive.
/// \param DKind Specified directive.
/// \return true - the directive is a map-entering target directive like
/// 'omp target', 'omp target data', 'omp target enter data',
/// 'omp target parallel', etc. (excludes 'omp target exit data', 'omp target
/// update') otherwise - false.
bool isOpenMPTargetMapEnteringDirective(OpenMPDirectiveKind DKind);

/// Checks if the specified composite/combined directive constitutes a teams
/// directive in the outermost nest. For example
/// 'omp teams distribute' or 'omp teams distribute parallel for'.
Expand Down
5 changes: 5 additions & 0 deletions clang/lib/Basic/OpenMPKinds.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -650,6 +650,11 @@ bool clang::isOpenMPTargetDataManagementDirective(OpenMPDirectiveKind DKind) {
DKind == OMPD_target_exit_data || DKind == OMPD_target_update;
}

bool clang::isOpenMPTargetMapEnteringDirective(OpenMPDirectiveKind DKind) {
return DKind == OMPD_target_data || DKind == OMPD_target_enter_data ||
isOpenMPTargetExecutionDirective(DKind);
}

bool clang::isOpenMPNestingTeamsDirective(OpenMPDirectiveKind DKind) {
if (DKind == OMPD_teams)
return true;
Expand Down
319 changes: 237 additions & 82 deletions clang/lib/CodeGen/CGOpenMPRuntime.cpp
View file
Open in desktop

Large diffs are not rendered by default.

318 changes: 174 additions & 144 deletions clang/test/OpenMP/bug60602.cpp
View file
Open in desktop

Large diffs are not rendered by default.

521 changes: 280 additions & 241 deletions clang/test/OpenMP/reduction_implicit_map.cpp
View file
Open in desktop

Large diffs are not rendered by default.

16 changes: 10 additions & 6 deletions clang/test/OpenMP/target_data_codegen.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -625,8 +625,8 @@ void test_present_modifier(int arg) {

// Make sure the struct picks up present even if another element of the struct
// doesn't have present.
// CK8: private unnamed_addr constant [11 x i64] [i64 0, i64 {{4|8}}, i64 {{4|8}}, i64 4, i64 4, i64 4, i64 0, i64 4, i64 {{4|8}}, i64 {{4|8}}, i64 4]
// CK8: private unnamed_addr constant [11 x i64]
// CK8: private unnamed_addr constant [13 x i64] [i64 0, i64 {{4|8}}, i64 {{4|8}}, i64 4, i64 4, i64 {{4|8}}, i64 4, i64 0, i64 4, i64 {{4|8}}, i64 {{4|8}}, i64 4, i64 {{8|4}}]
// CK8: private unnamed_addr constant [13 x i64]

// ps1
//
Expand All @@ -635,9 +635,11 @@ void test_present_modifier(int arg) {
// PRESENT=0x1000 | PTR_AND_OBJ=0x10 = 0x1010
// PRESENT=0x1000 | PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x1013
// MEMBER_OF_1=0x1000000000000 | FROM=0x2 | TO=0x1 = 0x1000000000003
// ATTACH=0x4000
//
// CK8-SAME: {{^}} [i64 [[#0x1000]], i64 [[#0x1000000001010]],
// CK8-SAME: {{^}} i64 [[#0x1010]], i64 [[#0x1013]], i64 [[#0x1000000000003]],
// CK8-SAME: {{^}} i64 [[#0x4000]],

// arg
//
Expand All @@ -648,13 +650,15 @@ void test_present_modifier(int arg) {
// ps2
//
// PRESENT=0x1000 = 0x1000
// MEMBER_OF_7=0x7000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x7000000001003
// MEMBER_OF_7=0x7000000000000 | PTR_AND_OBJ=0x10 = 0x7000000000010
// MEMBER_OF_8=0x8000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x8000000001003
// MEMBER_OF_8=0x8000000000000 | PTR_AND_OBJ=0x10 = 0x8000000000010
// PTR_AND_OBJ=0x10 = 0x10
// PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x13
// ATTACH=0x4000
//
// CK8-SAME: {{^}} i64 [[#0x1000]], i64 [[#0x7000000001003]],
// CK8-SAME: {{^}} i64 [[#0x7000000000010]], i64 [[#0x10]], i64 [[#0x13]]]
// CK8-SAME: {{^}} i64 [[#0x1000]], i64 [[#0x8000000001003]],
// CK8-SAME: {{^}} i64 [[#0x8000000000010]], i64 [[#0x10]], i64 [[#0x13]],
// CK8-SAME: {{^}} i64 [[#0x4000]]]
#pragma omp target data map(tofrom \
: ps1->s) \
map(present, tofrom \
Expand Down
291 changes: 158 additions & 133 deletions clang/test/OpenMP/target_data_map_codegen_hold.cpp
View file
Open in desktop

Large diffs are not rendered by default.

28 changes: 23 additions & 5 deletions clang/test/OpenMP/target_data_map_pointer_array_subscript_codegen.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -33,15 +33,33 @@ typedef struct {
MyObject *objects;
#pragma omp end declare target

// CHECK-DAG: [[SIZES0:@.+]] = private unnamed_addr constant [1 x i64] [i64 {{8|4}}]
// CHECK-DAG: [[MAPS0:@.+]] = private unnamed_addr constant [1 x i64] [i64 17]
// CHECK-DAG: [[SIZES0:@.+]] = private unnamed_addr constant [2 x i64] [i64 {{8|4}}, i64 {{8|4}}]
// CHECK-DAG: [[MAPS0:@.+]] = private unnamed_addr constant [2 x i64] [i64 1, i64 16384]
// CHECK-DAG: [[SIZES1:@.+]] = private unnamed_addr constant [2 x i64] [i64 0, i64 4]
// CHECK-DAG: [[MAPS1:@.+]] = private unnamed_addr constant [2 x i64] [i64 0, i64 281474976710673]
// CHECK: @main
int main(void) {
// CHECK: [[BPTR0:%.+]] = getelementptr inbounds [1 x ptr], ptr %{{.+}}, i32 0, i32 0
// CHECK: store ptr @objects, ptr [[BPTR0]],
// CHECK: call void @__tgt_target_data_begin_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr %{{.+}}, ptr %{{.+}}, ptr [[SIZES0]], ptr [[MAPS0]], ptr null, ptr null)

// &objects[0], &objects[1], 1 * sizeof(objects[0]), TO
// &objects, &objects[1], sizeof(objects), ATTACH

// CHECK-DAG: call void @__tgt_target_data_begin_mapper(ptr @{{.+}}, i64 -1, i32 2, ptr [[BPGEP:%.+]], ptr [[PGEP:%.+]], ptr [[SIZES0]], ptr [[MAPS0]], ptr null, ptr null)
// CHECK-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CHECK-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]

// CHECK-DAG: [[BP0:%.+]] = getelementptr inbounds {{.*}}ptr [[BP]], i32 0, i32 0
// CHECK-DAG: [[P0:%.+]] = getelementptr inbounds {{.*}}ptr [[P]], i32 0, i32 0
// CHECK-DAG: store ptr [[RVAR0:%.+]], ptr [[BP0]]
// CHECK-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
// CHECK-DAG: [[RVAR0]] = load ptr, ptr @objects
// CHECK-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[RVAR00:%.+]], i{{.+}} 1
// CHECK-DAG: [[RVAR00]] = load ptr, ptr @objects

// CHECK-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
// CHECK-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
// CHECK-DAG: store ptr @objects, ptr [[BP1]]
// CHECK-DAG: store ptr [[SEC0]], ptr [[P1]]

#pragma omp target enter data map(to : objects [1:1])
// CHECK: [[OBJ:%.+]] = load ptr, ptr @objects,
// CHECK: [[BPTR0:%.+]] = getelementptr inbounds [2 x ptr], ptr %{{.+}}, i32 0, i32 0
Expand Down
113 changes: 70 additions & 43 deletions clang/test/OpenMP/target_data_use_device_addr_codegen.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -11,9 +11,9 @@
#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]
// CHECK-DAG: [[SIZES1:@.+]] = private unnamed_addr constant [8 x i64] [i64 4, i64 16, i64 8, i64 4, i64 8, 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: [[MAPTYPES1:@.+]] = private unnamed_addr constant [8 x i64] [i64 67, i64 67, i64 16384, i64 3, i64 16384, 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
Expand Down Expand Up @@ -44,63 +44,82 @@ int main() {
}

// CHECK-LABEL: @main()
//
// &a, &a, TO | FROM | RETURN_PARAM
// &ptr[0], &ptr[3], TO | FROM | RETURN_PARAM
// &ptr, &ptr[3], ATTACH
// &ptr[0], &ptr[0], TO | FROM | RETURN_PARAM
// &ptr, &ptr[0], ATTACH
// &ref_ptee(ref), &ref_ptee(ref), TO | FROM | RETURN_PARAM
// &arr, &arr[0], TO | FROM | RETURN_PARAM
//
// 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: [[BPTRS:%.+]] = alloca [8 x ptr],
// CHECK: [[PTRS:%.+]] = alloca [8 x ptr],
// CHECK: [[MAP_PTRS:%.+]] = alloca [8 x ptr],
// CHECK: [[SIZES:%.+]] = alloca [8 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: [[P0:%.+]] = load ptr, ptr [[PTR_ADDR]],
// CHECK: [[P1:%.+]] = load ptr, ptr [[PTR_ADDR]],
// CHECK-NEXT: [[ARR_IDX:%.+]] = getelementptr inbounds nuw float, ptr [[P1]], i64 3
// CHECK: [[P2:%.+]] = load ptr, ptr [[PTR_ADDR]],
// CHECK: [[P5:%.+]] = load ptr, ptr [[PTR_ADDR]],
// 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-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[SIZES]], ptr align 8 [[SIZES1]], i64 64, i1 false)
// CHECK: [[BPTR0:%.+]] = getelementptr inbounds [8 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: [[PTR0:%.+]] = getelementptr inbounds [8 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: [[BPTR1:%.+]] = getelementptr inbounds [8 x ptr], ptr [[BPTRS]], i32 0, i32 1
// CHECK: store ptr [[P0]], ptr [[BPTR1]],
// CHECK: [[PTR1:%.+]] = getelementptr inbounds [8 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: [[BPTR2:%.+]] = getelementptr inbounds [8 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: [[PTR2:%.+]] = getelementptr inbounds [8 x ptr], ptr [[PTRS]], i32 0, i32 2
// CHECK: store ptr [[ARR_IDX]], ptr [[PTR2]],
// CHECK: [[BPTR3:%.+]] = getelementptr inbounds [8 x ptr], ptr [[BPTRS]], i32 0, i32 3
// CHECK: store ptr [[P2]], ptr [[BPTR3]],
// CHECK: [[PTR3:%.+]] = getelementptr inbounds [8 x ptr], ptr [[PTRS]], i32 0, i32 3
// CHECK: store ptr [[ARR_IDX1]], ptr [[PTR3]],
// CHECK: [[BPTR4:%.+]] = getelementptr inbounds [8 x ptr], ptr [[BPTRS]], i32 0, i32 4
// CHECK: store ptr [[PTR_ADDR]], ptr [[BPTR4]], align
// CHECK: [[PTR4:%.+]] = getelementptr inbounds [8 x ptr], ptr [[PTRS]], i32 0, i32 4
// CHECK: store ptr [[ARR_IDX1]], ptr [[PTR4]], align 8
// CHECK: [[BPTR5:%.+]] = getelementptr inbounds [8 x ptr], ptr [[BPTRS]], i32 0, i32 5
// CHECK: store ptr [[P7]], ptr [[BPTR5]], align
// CHECK: [[PTR5:%.+]] = getelementptr inbounds [8 x ptr], ptr [[PTRS]], i32 0, i32 5
// CHECK: store ptr [[REF]], ptr [[PTR5]], align 8
// CHECK: [[BPTR6:%.+]] = getelementptr inbounds [8 x ptr], ptr [[BPTRS]], i32 0, i32 6
// CHECK: store ptr [[ARR_ADDR]], ptr [[BPTR6]],
// CHECK: [[PTR6:%.+]] = getelementptr inbounds [8 x ptr], ptr [[PTRS]], i32 0, i32 6
// CHECK: store ptr [[ARR_IDX2]], ptr [[PTR6]],
// CHECK: [[SIZE_PTR:%.+]] = getelementptr inbounds [8 x i64], ptr [[SIZES]], i32 0, i32 6
// 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: [[MAP_PTR:%.+]] = getelementptr inbounds [8 x ptr], ptr [[MAP_PTRS]], i64 0, i64 6
// 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: [[BPTR7:%.+]] = getelementptr inbounds [8 x ptr], ptr [[BPTRS]], i32 0, i32 7
// CHECK: store ptr [[VLA_ADDR]], ptr [[BPTR7]],
// CHECK: [[PTR7:%.+]] = getelementptr inbounds [8 x ptr], ptr [[PTRS]], i32 0, i32 7
// CHECK: store ptr [[ARR_IDX5]], ptr [[PTR7]],

// 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: [[BPTR:%.+]] = getelementptr inbounds [8 x ptr], ptr [[BPTRS]], i32 0, i32 0
// CHECK: [[PTR:%.+]] = getelementptr inbounds [8 x ptr], ptr [[PTRS]], i32 0, i32 0
// CHECK: [[SIZE:%.+]] = getelementptr inbounds [8 x i64], ptr [[SIZES]], i32 0, i32 0
// CHECK: call void @__tgt_target_data_begin_mapper(ptr @{{.+}}, i64 -1, i32 8, 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: [[REF_REF:%.+]] = load ptr, ptr [[BPTR5]],
// CHECK: store ptr [[REF_REF]], ptr [[TMP_REF_ADDR:%.+]],
// CHECK: [[ARR_REF:%.+]] = load ptr, ptr [[BPTR4]],
// CHECK: [[VLA_REF:%.+]] = load ptr, ptr [[BPTR5]],
// CHECK: [[ARR_REF:%.+]] = load ptr, ptr [[BPTR6]],
// CHECK: [[VLA_REF:%.+]] = load ptr, ptr [[BPTR7]],
// CHECK: [[A:%.+]] = load float, ptr [[A_REF]],
// CHECK: [[INC:%.+]] = fadd float [[A]], 1.000000e+00
// CHECK: store float [[INC]], ptr [[A_REF]],
Expand All @@ -120,12 +139,20 @@ int main() {
// 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: [[BPTR:%.+]] = getelementptr inbounds [8 x ptr], ptr [[BPTRS]], i32 0, i32 0
// CHECK: [[PTR:%.+]] = getelementptr inbounds [8 x ptr], ptr [[PTRS]], i32 0, i32 0
// CHECK: [[SIZE:%.+]] = getelementptr inbounds [8 x i64], ptr [[SIZES]], i32 0, i32 0
// CHECK: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 -1, i32 8, ptr [[BPTR]], ptr [[PTR]], ptr [[SIZE]], ptr [[MAPTYPES1]], ptr null, ptr null)

// CHECK: foo
//
// &this[0], &this->a, sizeof(this[0].(a-to-arr[a]) | ALLOC
// &this[0], &this->a, sizeof(a), TO | FROM | RETURN_PARAM | MEMBER_OF(1)
// &this->ptr, &this->ptr[3], 4 * sizeof(ptr[0], TO | FROM | PTR_AND_OBJ | RETURN_PARAM | MEMBER_OF(1)
// &this[0], &ref_ptee(this->ref), sizeof(this->ref[0]), TO | FROM | PTR_AND_OBJ | RETURN_PARAM | MEMBER_OF(1)
// &this->ptr, &this->ptr[0], sizeof(ptr[0], TO | FROM | PTR_AND_OBJ | MEMBER_OF(1)
// &this, &this->arr[0], 4 * sizeof(arr[0]), TO | FROM | RETURN_PARAM | MEMBER_OF(1)
//
// CHECK: [[BPTRS:%.+]] = alloca [6 x ptr],
// CHECK: [[PTRS:%.+]] = alloca [6 x ptr],
// CHECK: [[MAP_PTRS:%.+]] = alloca [6 x ptr],
Expand Down
Loading