llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-offload Author: Abhinav Gaba (abhinavgaba) <details> <summary>Changes</summary> * Add a few tests for when a mapper does something like `map(s.p[0:10])` where `p` is a pointer. * Add a few tests that require propagation of bits like `present/always` into a mapper. * Fix a few tests that were expecting `p` to be mapped when the mapper only said `map(s.p[0:10])`. The output of a few tests is different from what we expect. They have been annotated with FIXMEs, and the expected output for when the follow-up changes in this stack to propagate the map-type-modifier bits and using attach-style mapping for mappers get merged. --- Patch is 71.93 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/204269.diff 23 Files Affected: - (added) clang/test/OpenMP/target_map_nested_ptr_member_mapper_codegen.cpp (+392) - (modified) offload/test/mapping/declare_mapper_nested_mappers.cpp (+6-3) - (modified) offload/test/mapping/declare_mapper_target.cpp (+3-1) - (modified) offload/test/mapping/declare_mapper_target_data.cpp (+3-1) - (modified) offload/test/mapping/declare_mapper_target_data_enter_exit.cpp (+3-1) - (modified) offload/test/mapping/declare_mapper_target_update.cpp (+3-1) - (added) offload/test/mapping/mapper_enter_data_always_present_ptee.c (+76) - (added) offload/test/mapping/mapper_map_always_from.c (+46) - (added) offload/test/mapping/mapper_map_mbr_ptee_then_present_mbr_ptee.c (+58) - (added) offload/test/mapping/mapper_map_mbr_then_present_mbr_ptee.c (+62) - (added) offload/test/mapping/mapper_map_present_ptee.c (+65) - (added) offload/test/mapping/mapper_map_ptee_only.c (+49) - (added) offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections.c (+61) - (added) offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections_array.c (+93) - (added) offload/test/mapping/mapper_map_ptee_only_2ndlevel.c (+52) - (added) offload/test/mapping/mapper_map_ptee_only_2ndlevel_array.c (+82) - (added) offload/test/mapping/mapper_map_ptee_only_always_array.c (+78) - (added) offload/test/mapping/mapper_map_ptee_only_array.c (+59) - (added) offload/test/mapping/mapper_map_ptee_then_present_absent_mbr.c (+55) - (added) offload/test/mapping/mapper_map_ptr_ptee_nomapper_del_ptee.c (+43) - (added) offload/test/mapping/mapper_map_ptr_ptee_nomapper_del_ptr.c (+43) - (added) offload/test/mapping/mapper_target_update_present_ptee.c (+84) - (added) offload/test/mapping/multiple_deletes_within_one_struct.c (+44) ``````````diff diff --git a/clang/test/OpenMP/target_map_nested_ptr_member_mapper_codegen.cpp b/clang/test/OpenMP/target_map_nested_ptr_member_mapper_codegen.cpp new file mode 100644 index 0000000000000..cfdc9b861090c --- /dev/null +++ b/clang/test/OpenMP/target_map_nested_ptr_member_mapper_codegen.cpp @@ -0,0 +1,392 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --functions ".*mapper.*" --function-signature --check-globals --filter-out-after "getelem.*kernel" --filter-out "= alloca.*" --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _ --global-value-regex "\.offload_.*" --global-hex-value-regex ".offload_maptypes.*" +// RUN: %clang_cc1 -verify -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 %s -emit-llvm -o - | FileCheck %s + +// PRESENT is propagated to pointee (attach-ptr) entries only at OpenMP >= 6.0. +// FIXME: that propagation is done in a follow-on; until then the CHECK-60 +// output below is identical to CHECK (the pointee entries carry map-type mask +// 1036 = ALWAYS|DELETE|CLOSE, without PRESENT). Once PRESENT is propagated, the +// attach-ptr pointee entries should use mask 5132 = ALWAYS|DELETE|CLOSE|PRESENT +// at 6.0. +// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK-60 + +// expected-no-diagnostics +#ifndef HEADER +#define HEADER + +// S2 mapper for: map(to: arr[0:2]) +// Per-element entries (i = array element index; N = __tgt_mapper_num_components()): +// &arr[i], &arr[i].s1p, sizeof(s1p..z), MEMBER_OF(N) | ALLOC +// &arr[i], &arr[i].z, sizeof(int), MEMBER_OF(N+1) | TO | FROM +// &arr[i].s1p, &arr[i].s1p->x, sizeof(int), MEMBER_OF(N+1) | TO | FROM | PTR_AND_OBJ +// &arr[i].s1p, &arr[i].s1p->y, sizeof(int), MEMBER_OF(N+1) | TO | FROM | PTR_AND_OBJ +// FIXME: should use attach-style codegen for s1p->x/y instead of PTR_AND_OBJ, +// which unnecessarily also maps s1p. + +typedef struct { + int x; + int y; +} S1; + +typedef struct { + S1 *s1p; + int z; +} S2; + +#pragma omp declare mapper(default : S2 s2) map(s2.z, s2.s1p->x, s2.s1p->y) + +void foo(S2 *arr) { + // &arr, &arr[0], 2*sizeof(S2), TARGET_PARAM | TO + // (mapper handles individual members) +#pragma omp target enter data map(to: arr[0:2]) + {} +} + +#endif +//. +// CHECK: @.offload_sizes = private unnamed_addr constant [2 x i64] [i64 32, i64 8] +// CHECK: @.offload_maptypes = private unnamed_addr constant [2 x i64] [i64 [[#0x1]], i64 [[#0x4000]]] +//. +// CHECK-60: @.offload_sizes = private unnamed_addr constant [2 x i64] [i64 32, i64 8] +// CHECK-60: @.offload_maptypes = private unnamed_addr constant [2 x i64] [i64 [[#0x1]], i64 [[#0x4000]]] +//. +// CHECK-LABEL: define {{[^@]+}}@_Z3fooP2S2 +// CHECK-SAME: (ptr noundef [[ARR:%.*]]) #[[ATTR0:[0-9]+]] { +// CHECK: entry: +// CHECK: store ptr [[ARR]], ptr [[ARR_ADDR:%.*]], align 8 +// CHECK: [[TMP0:%.*]] = load ptr, ptr [[ARR_ADDR]], align 8 +// CHECK: [[TMP1:%.*]] = load ptr, ptr [[ARR_ADDR]], align 8 +// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds nuw [[STRUCT_S2:%.*]], ptr [[TMP1]], i64 0 +// CHECK: [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0 +// CHECK: store ptr [[TMP0]], ptr [[TMP2]], align 8 +// CHECK: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0 +// CHECK: store ptr [[ARRAYIDX]], ptr [[TMP3]], align 8 +// CHECK: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0 +// CHECK: store ptr @.omp_mapper._ZTS2S2.default, ptr [[TMP4]], align 8 +// CHECK: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 +// CHECK: store ptr [[ARR_ADDR]], ptr [[TMP5]], align 8 +// CHECK: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1 +// CHECK: store ptr [[ARRAYIDX]], ptr [[TMP6]], align 8 +// CHECK: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 +// CHECK: store ptr null, ptr [[TMP7]], align 8 +// CHECK: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK: [[TMP9:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK: call void @__tgt_target_data_begin_mapper(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 2, ptr [[TMP8]], ptr [[TMP9]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr null, ptr [[DOTOFFLOAD_MAPPERS]]) +// CHECK: ret void +// +// +// CHECK-LABEL: define {{[^@]+}}@.omp_mapper._ZTS2S2.default +// CHECK-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]], ptr noundef [[TMP2:%.*]], i64 noundef [[TMP3:%.*]], i64 noundef [[TMP4:%.*]], ptr noundef [[TMP5:%.*]]) #[[ATTR1:[0-9]+]] { +// CHECK: entry: +// CHECK: [[TMP6:%.*]] = udiv exact i64 [[TMP3]], 16 +// CHECK: [[TMP7:%.*]] = getelementptr [[STRUCT_S2:%.*]], ptr [[TMP2]], i64 [[TMP6]] +// CHECK: [[OMP_ARRAYINIT_ISARRAY:%.*]] = icmp sgt i64 [[TMP6]], 1 +// CHECK: [[TMP8:%.*]] = and i64 [[TMP4]], 8 +// CHECK: [[TMP9:%.*]] = icmp ne ptr [[TMP1]], [[TMP2]] +// CHECK: [[TMP10:%.*]] = or i1 [[OMP_ARRAYINIT_ISARRAY]], [[TMP9]] +// CHECK: [[DOTOMP_ARRAY__INIT__DELETE:%.*]] = icmp eq i64 [[TMP8]], 0 +// CHECK: [[TMP11:%.*]] = and i1 [[TMP10]], [[DOTOMP_ARRAY__INIT__DELETE]] +// CHECK: br i1 [[TMP11]], label [[DOTOMP_ARRAY__INIT:%.*]], label [[OMP_ARRAYMAP_HEAD:%.*]] +// CHECK: .omp.array..init: +// CHECK: [[TMP12:%.*]] = mul nuw i64 [[TMP6]], 16 +// CHECK: [[TMP13:%.*]] = and i64 [[TMP4]], -4 +// CHECK: [[TMP14:%.*]] = or i64 [[TMP13]], 512 +// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP12]], i64 [[TMP14]], ptr [[TMP5]]) +// CHECK: br label [[OMP_ARRAYMAP_HEAD]] +// CHECK: omp.arraymap.head: +// CHECK: [[OMP_ARRAYMAP_ISEMPTY:%.*]] = icmp eq ptr [[TMP2]], [[TMP7]] +// CHECK: br i1 [[OMP_ARRAYMAP_ISEMPTY]], label [[OMP_DONE:%.*]], label [[OMP_ARRAYMAP_BODY:%.*]] +// CHECK: omp.arraymap.body: +// CHECK: [[OMP_ARRAYMAP_PTRCURRENT:%.*]] = phi ptr [ [[TMP2]], [[OMP_ARRAYMAP_HEAD]] ], [ [[OMP_ARRAYMAP_NEXT:%.*]], [[OMP_TYPE_END23:%.*]] ] +// CHECK: [[Z:%.*]] = getelementptr inbounds nuw [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 1 +// CHECK: [[S1P:%.*]] = getelementptr inbounds nuw [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0 +// CHECK: [[S1P1:%.*]] = getelementptr inbounds nuw [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0 +// CHECK: [[TMP15:%.*]] = load ptr, ptr [[S1P1]], align 8 +// CHECK: [[X:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP15]], i32 0, i32 0 +// CHECK: [[S1P2:%.*]] = getelementptr inbounds nuw [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0 +// CHECK: [[S1P3:%.*]] = getelementptr inbounds nuw [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0 +// CHECK: [[TMP16:%.*]] = load ptr, ptr [[S1P3]], align 8 +// CHECK: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_S1]], ptr [[TMP16]], i32 0, i32 1 +// CHECK: [[TMP17:%.*]] = getelementptr i32, ptr [[Z]], i32 1 +// CHECK: [[TMP18:%.*]] = ptrtoaddr ptr [[TMP17]] to i64 +// CHECK: [[TMP19:%.*]] = ptrtoaddr ptr [[S1P]] to i64 +// CHECK: [[TMP20:%.*]] = sub i64 [[TMP18]], [[TMP19]] +// CHECK: [[TMP21:%.*]] = call i64 @__tgt_mapper_num_components(ptr [[TMP0]]) +// CHECK: [[TMP22:%.*]] = shl i64 [[TMP21]], 48 +// CHECK: [[TMP23:%.*]] = add nuw i64 0, [[TMP22]] +// CHECK: [[TMP24:%.*]] = and i64 [[TMP4]], 3 +// CHECK: [[TMP25:%.*]] = icmp eq i64 [[TMP24]], 0 +// CHECK: br i1 [[TMP25]], label [[OMP_TYPE_ALLOC:%.*]], label [[OMP_TYPE_ALLOC_ELSE:%.*]] +// CHECK: omp.type.alloc: +// CHECK: [[TMP26:%.*]] = and i64 [[TMP23]], -4 +// CHECK: br label [[OMP_TYPE_END:%.*]] +// CHECK: omp.type.alloc.else: +// CHECK: [[TMP27:%.*]] = icmp eq i64 [[TMP24]], 1 +// CHECK: br i1 [[TMP27]], label [[OMP_TYPE_TO:%.*]], label [[OMP_TYPE_TO_ELSE:%.*]] +// CHECK: omp.type.to: +// CHECK: [[TMP28:%.*]] = and i64 [[TMP23]], -3 +// CHECK: br label [[OMP_TYPE_END]] +// CHECK: omp.type.to.else: +// CHECK: [[TMP29:%.*]] = icmp eq i64 [[TMP24]], 2 +// CHECK: br i1 [[TMP29]], label [[OMP_TYPE_FROM:%.*]], label [[OMP_TYPE_END]] +// CHECK: omp.type.from: +// CHECK: [[TMP30:%.*]] = and i64 [[TMP23]], -2 +// CHECK: br label [[OMP_TYPE_END]] +// CHECK: omp.type.end: +// CHECK: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP26]], [[OMP_TYPE_ALLOC]] ], [ [[TMP28]], [[OMP_TYPE_TO]] ], [ [[TMP30]], [[OMP_TYPE_FROM]] ], [ [[TMP23]], [[OMP_TYPE_TO_ELSE]] ] +// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[S1P]], i64 [[TMP20]], i64 [[OMP_MAPTYPE]], ptr null) +// CHECK: [[TMP31:%.*]] = add nuw i64 281474976710659, [[TMP22]] +// CHECK: [[TMP32:%.*]] = and i64 [[TMP4]], 3 +// CHECK: [[TMP33:%.*]] = icmp eq i64 [[TMP32]], 0 +// CHECK: br i1 [[TMP33]], label [[OMP_TYPE_ALLOC4:%.*]], label [[OMP_TYPE_ALLOC_ELSE5:%.*]] +// CHECK: omp.type.alloc4: +// CHECK: [[TMP34:%.*]] = and i64 [[TMP31]], -4 +// CHECK: br label [[OMP_TYPE_END9:%.*]] +// CHECK: omp.type.alloc.else5: +// CHECK: [[TMP35:%.*]] = icmp eq i64 [[TMP32]], 1 +// CHECK: br i1 [[TMP35]], label [[OMP_TYPE_TO6:%.*]], label [[OMP_TYPE_TO_ELSE7:%.*]] +// CHECK: omp.type.to6: +// CHECK: [[TMP36:%.*]] = and i64 [[TMP31]], -3 +// CHECK: br label [[OMP_TYPE_END9]] +// CHECK: omp.type.to.else7: +// CHECK: [[TMP37:%.*]] = icmp eq i64 [[TMP32]], 2 +// CHECK: br i1 [[TMP37]], label [[OMP_TYPE_FROM8:%.*]], label [[OMP_TYPE_END9]] +// CHECK: omp.type.from8: +// CHECK: [[TMP38:%.*]] = and i64 [[TMP31]], -2 +// CHECK: br label [[OMP_TYPE_END9]] +// CHECK: omp.type.end9: +// CHECK: [[OMP_MAPTYPE10:%.*]] = phi i64 [ [[TMP34]], [[OMP_TYPE_ALLOC4]] ], [ [[TMP36]], [[OMP_TYPE_TO6]] ], [ [[TMP38]], [[OMP_TYPE_FROM8]] ], [ [[TMP31]], [[OMP_TYPE_TO_ELSE7]] ] +// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[Z]], i64 4, i64 [[OMP_MAPTYPE10]], ptr null) +// CHECK: [[TMP39:%.*]] = add nuw i64 281474976710675, [[TMP22]] +// CHECK: [[TMP40:%.*]] = and i64 [[TMP4]], 3 +// CHECK: [[TMP41:%.*]] = icmp eq i64 [[TMP40]], 0 +// CHECK: br i1 [[TMP41]], label [[OMP_TYPE_ALLOC11:%.*]], label [[OMP_TYPE_ALLOC_ELSE12:%.*]] +// CHECK: omp.type.alloc11: +// CHECK: [[TMP42:%.*]] = and i64 [[TMP39]], -4 +// CHECK: br label [[OMP_TYPE_END16:%.*]] +// CHECK: omp.type.alloc.else12: +// CHECK: [[TMP43:%.*]] = icmp eq i64 [[TMP40]], 1 +// CHECK: br i1 [[TMP43]], label [[OMP_TYPE_TO13:%.*]], label [[OMP_TYPE_TO_ELSE14:%.*]] +// CHECK: omp.type.to13: +// CHECK: [[TMP44:%.*]] = and i64 [[TMP39]], -3 +// CHECK: br label [[OMP_TYPE_END16]] +// CHECK: omp.type.to.else14: +// CHECK: [[TMP45:%.*]] = icmp eq i64 [[TMP40]], 2 +// CHECK: br i1 [[TMP45]], label [[OMP_TYPE_FROM15:%.*]], label [[OMP_TYPE_END16]] +// CHECK: omp.type.from15: +// CHECK: [[TMP46:%.*]] = and i64 [[TMP39]], -2 +// CHECK: br label [[OMP_TYPE_END16]] +// CHECK: omp.type.end16: +// CHECK: [[OMP_MAPTYPE17:%.*]] = phi i64 [ [[TMP42]], [[OMP_TYPE_ALLOC11]] ], [ [[TMP44]], [[OMP_TYPE_TO13]] ], [ [[TMP46]], [[OMP_TYPE_FROM15]] ], [ [[TMP39]], [[OMP_TYPE_TO_ELSE14]] ] +// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[S1P]], ptr [[X]], i64 4, i64 [[OMP_MAPTYPE17]], ptr null) +// CHECK: [[TMP47:%.*]] = add nuw i64 281474976710675, [[TMP22]] +// CHECK: [[TMP48:%.*]] = and i64 [[TMP4]], 3 +// CHECK: [[TMP49:%.*]] = icmp eq i64 [[TMP48]], 0 +// CHECK: br i1 [[TMP49]], label [[OMP_TYPE_ALLOC18:%.*]], label [[OMP_TYPE_ALLOC_ELSE19:%.*]] +// CHECK: omp.type.alloc18: +// CHECK: [[TMP50:%.*]] = and i64 [[TMP47]], -4 +// CHECK: br label [[OMP_TYPE_END23]] +// CHECK: omp.type.alloc.else19: +// CHECK: [[TMP51:%.*]] = icmp eq i64 [[TMP48]], 1 +// CHECK: br i1 [[TMP51]], label [[OMP_TYPE_TO20:%.*]], label [[OMP_TYPE_TO_ELSE21:%.*]] +// CHECK: omp.type.to20: +// CHECK: [[TMP52:%.*]] = and i64 [[TMP47]], -3 +// CHECK: br label [[OMP_TYPE_END23]] +// CHECK: omp.type.to.else21: +// CHECK: [[TMP53:%.*]] = icmp eq i64 [[TMP48]], 2 +// CHECK: br i1 [[TMP53]], label [[OMP_TYPE_FROM22:%.*]], label [[OMP_TYPE_END23]] +// CHECK: omp.type.from22: +// CHECK: [[TMP54:%.*]] = and i64 [[TMP47]], -2 +// CHECK: br label [[OMP_TYPE_END23]] +// CHECK: omp.type.end23: +// CHECK: [[OMP_MAPTYPE24:%.*]] = phi i64 [ [[TMP50]], [[OMP_TYPE_ALLOC18]] ], [ [[TMP52]], [[OMP_TYPE_TO20]] ], [ [[TMP54]], [[OMP_TYPE_FROM22]] ], [ [[TMP47]], [[OMP_TYPE_TO_ELSE21]] ] +// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[S1P2]], ptr [[Y]], i64 4, i64 [[OMP_MAPTYPE24]], ptr null) +// CHECK: [[OMP_ARRAYMAP_NEXT]] = getelementptr [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 1 +// CHECK: [[OMP_ARRAYMAP_ISDONE:%.*]] = icmp eq ptr [[OMP_ARRAYMAP_NEXT]], [[TMP7]] +// CHECK: br i1 [[OMP_ARRAYMAP_ISDONE]], label [[OMP_ARRAYMAP_EXIT:%.*]], label [[OMP_ARRAYMAP_BODY]] +// CHECK: omp.arraymap.exit: +// CHECK: [[OMP_ARRAYINIT_ISARRAY25:%.*]] = icmp sgt i64 [[TMP6]], 1 +// CHECK: [[TMP55:%.*]] = and i64 [[TMP4]], 8 +// CHECK: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP55]], 0 +// CHECK: [[TMP56:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY25]], [[DOTOMP_ARRAY__DEL__DELETE]] +// CHECK: br i1 [[TMP56]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]] +// CHECK: .omp.array..del: +// CHECK: [[TMP57:%.*]] = mul nuw i64 [[TMP6]], 16 +// CHECK: [[TMP58:%.*]] = and i64 [[TMP4]], -4 +// CHECK: [[TMP59:%.*]] = or i64 [[TMP58]], 512 +// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP57]], i64 [[TMP59]], ptr [[TMP5]]) +// CHECK: br label [[OMP_DONE]] +// CHECK: omp.done: +// CHECK: ret void +// +// +// CHECK-60-LABEL: define {{[^@]+}}@_Z3fooP2S2 +// CHECK-60-SAME: (ptr noundef [[ARR:%.*]]) #[[ATTR0:[0-9]+]] { +// CHECK-60: entry: +// CHECK-60: store ptr [[ARR]], ptr [[ARR_ADDR:%.*]], align 8 +// CHECK-60: [[TMP0:%.*]] = load ptr, ptr [[ARR_ADDR]], align 8 +// CHECK-60: [[TMP1:%.*]] = load ptr, ptr [[ARR_ADDR]], align 8 +// CHECK-60: [[ARRAYIDX:%.*]] = getelementptr inbounds nuw [[STRUCT_S2:%.*]], ptr [[TMP1]], i64 0 +// CHECK-60: [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0 +// CHECK-60: store ptr [[TMP0]], ptr [[TMP2]], align 8 +// CHECK-60: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0 +// CHECK-60: store ptr [[ARRAYIDX]], ptr [[TMP3]], align 8 +// CHECK-60: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0 +// CHECK-60: store ptr @.omp_mapper._ZTS2S2.default, ptr [[TMP4]], align 8 +// CHECK-60: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1 +// CHECK-60: store ptr [[ARR_ADDR]], ptr [[TMP5]], align 8 +// CHECK-60: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1 +// CHECK-60: store ptr [[ARRAYIDX]], ptr [[TMP6]], align 8 +// CHECK-60: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1 +// CHECK-60: store ptr null, ptr [[TMP7]], align 8 +// CHECK-60: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 +// CHECK-60: [[TMP9:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 +// CHECK-60: call void @__tgt_target_data_begin_mapper(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 2, ptr [[TMP8]], ptr [[TMP9]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr null, ptr [[DOTOFFLOAD_MAPPERS]]) +// CHECK-60: ret void +// +// +// CHECK-60-LABEL: define {{[^@]+}}@.omp_mapper._ZTS2S2.default +// CHECK-60-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]], ptr noundef [[TMP2:%.*]], i64 noundef [[TMP3:%.*]], i64 noundef [[TMP4:%.*]], ptr noundef [[TMP5:%.*]]) #[[ATTR1:[0-9]+]] { +// CHECK-60: entry: +// CHECK-60: [[TMP6:%.*]] = udiv exact i64 [[TMP3]], 16 +// CHECK-60: [[TMP7:%.*]] = getelementptr [[STRUCT_S2:%.*]], ptr [[TMP2]], i64 [[TMP6]] +// CHECK-60: [[OMP_ARRAYINIT_ISARRAY:%.*]] = icmp sgt i64 [[TMP6]], 1 +// CHECK-60: [[TMP8:%.*]] = and i64 [[TMP4]], 8 +// CHECK-60: [[TMP9:%.*]] = icmp ne ptr [[TMP1]], [[TMP2]] +// CHECK-60: [[TMP10:%.*]] = or i1 [[OMP_ARRAYINIT_ISARRAY]], [[TMP9]] +// CHECK-60: [[DOTOMP_ARRAY__INIT__DELETE:%.*]] = icmp eq i64 [[TMP8]], 0 +// CHECK-60: [[TMP11:%.*]] = and i1 [[TMP10]], [[DOTOMP_ARRAY__INIT__DELETE]] +// CHECK-60: br i1 [[TMP11]], label [[DOTOMP_ARRAY__INIT:%.*]], label [[OMP_ARRAYMAP_HEAD:%.*]] +// CHECK-60: .omp.array..init: +// CHECK-60: [[TMP12:%.*]] = mul nuw i64 [[TMP6]], 16 +// CHECK-60: [[TMP13:%.*]] = and i64 [[TMP4]], -4 +// CHECK-60: [[TMP14:%.*]] = or i64 [[TMP13]], 512 +// CHECK-60: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP12]], i64 [[TMP14]], ptr [[TMP5]]) +// CHECK-60: br label [[OMP_ARRAYMAP_HEAD]] +// CHECK-60: omp.arraymap.head: +// CHECK-60: [[OMP_ARRAYMAP_ISEMPTY:%.*]] = icmp eq ptr [[TMP2]], [[TMP7]] +// CHECK-60: br i1 [[OMP_ARRAYMAP_ISEMPTY]], label [[OMP_DONE:%.*]], label [[OMP_ARRAYMAP_BODY:%.*]] +// CHECK-60: omp.arraymap.body: +// CHECK-60: [[OMP_ARRAYMAP_PTRCURRENT:%.*]] = phi ptr [ [[TMP2]], [[OMP_ARRAYMAP_HEAD]] ], [ [[OMP_ARRAYMAP_NEXT:%.*]], [[OMP_TYPE_END23:%.*]] ] +// CHECK-60: [[Z:%.*]] = getelementptr inbounds nuw [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 1 +// CHECK-60: [[S1P:%.*]] = getelementptr inbounds nuw [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0 +// CHECK-60: [[S1P1:%.*]] = getelementptr inbounds nuw [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0 +// CHECK-60: [[TMP15:%.*]] = load ptr, ptr [[S1P1]], align 8 +// CHECK-60: [[X:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP15]], i32 0, i32 0 +// CHECK-60: [[S1P2:%.*]] = getelementptr inbounds nuw [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0 +// CHECK-60: [[S1P3:%.*]] = getelementptr inbounds nuw [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0 +// CHECK-60: [[TMP16:%.*]] = load ptr, ptr [[S1P3]], align 8 +// CHECK-60: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_S1]], ptr [[TMP16]], i32 0, i32 1 +// CHECK-60: [[TMP17:%.*]] = getelementptr i32, ptr [[Z]], i32 1 +// CHECK-60: [[TMP18:%.*]] = ptrtoaddr ptr [[TMP17]] to i64 +// CHECK-60: [[TMP19:%.*]] = ptrtoaddr ptr [[S1P]] to i64 +// CHECK-60: [[TMP20:%.*]] = sub i64 [[TMP18]], [[TMP19]] +// CHECK-60: [[TMP21:%.*]] = call i64 @__tgt_mapper_num_components(ptr [[TMP0]]) +// CHECK-60: [[TMP22:%.*]] = shl i64 [[TMP21]], 48 +// CHECK-60: [[TMP23:%.*]] = add nuw i64 0, [[TMP22]] +// CHECK-60: [[TMP24:%.*]] = and i64 [[TMP4]], 3 +// CHECK-60: [[TMP25:%.*]] = icmp eq i64 [[TMP24]], 0 +// CHECK-60: br i1 [[TMP25]], label [[OMP_TYPE_ALLOC:%.*]], label [[OMP_TYPE_ALLOC_ELSE:%.*]] +// CHECK-60: omp.type.alloc: +// CHECK-60: [[TMP26:%.*]] = and i64 [[TMP23]], -4 +// CHECK-60: br label [[OMP_TYPE_END:%.*]] +// CHECK-60: omp.type.alloc.else: +// CHECK-60: [[TMP27:%.*]] = icmp eq i64 [[TMP24]], 1 +// CHECK-60: br i1 [[TMP27]], label [[OMP_TYPE_TO:%.*]], label [[OMP_TYPE_TO_ELSE:%.*]] +// CHECK-60: omp.type.to: +// CHECK-60: [[TMP28:%.*]] = and i64 [[TMP23]], -3 +// CHECK-60: br label [[OMP_TYP... [truncated] `````````` </details> https://github.com/llvm/llvm-project/pull/204269 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
