[NFC][OpenMP] Add tests for mapping pointers and their dereferences. #146934
+1,799
−156
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
The output of the compile-and-run tests is incorrect. These will be used for reference in future commits that resolve the issues. Also updated the existing clang LIT test, target_map_both_pointer_pointee_codegen.cpp, with more regions and more narrowed-down update_cc_test_checks filters.
|
✅ With the latest revision this PR passed the C/C++ code formatter. |
abhinavgaba
commented
Jul 3, 2025
| printf("%d %d %d %d\n", p[0], p_mappedptr == &p, x0_mappedptr == &p[-1], | ||
| x0_hostaddr == &p[-1]); | ||
| // EXPECTED: 222 1 1 0 | ||
| // CHECK: 111 0 0 0 |
There was a problem hiding this comment.
I kept the expected output here as a comment, instead of using that for CHECKs, so that we can track the current status of the compiler/runtime, instead of marking this as UNSUPPORTED/XFAIL. I can change that and mark these as XFAIL if that's preferable.
llvmbot
added
clang
|
@llvm/pr-subscribers-offload @llvm/pr-subscribers-clang Author: Abhinav Gaba (abhinavgaba) ChangesThe output of the compile-and-run tests is incorrect. These will be used for reference in future commits that resolve the issues. Also updated the existing clang LIT test, Patch is 108.34 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/146934.diff 12 Files Affected:
diff --git a/clang/test/OpenMP/target_map_both_pointer_pointee_codegen.cpp b/clang/test/OpenMP/target_map_both_pointer_pointee_codegen.cpp
index 87fa7fe462daa..7ad142e51fc09 100644
--- a/clang/test/OpenMP/target_map_both_pointer_pointee_codegen.cpp
+++ b/clang/test/OpenMP/target_map_both_pointer_pointee_codegen.cpp
@@ -1,4 +1,4 @@
-// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --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
@@ -7,168 +7,211 @@
#ifndef HEADER
#define HEADER
-extern void *malloc (int __size) throw () __attribute__ ((__malloc__));
+void f1() {
+ int *ptr;
+ // &ptr, &ptr, sizeof(ptr), TO | PARAM
+ #pragma omp target map(ptr)
+ ptr[1] = 5;
+}
-void foo() {
- int *ptr = (int *) malloc(3 * sizeof(int));
+void f2() {
+ int *ptr;
+ // &ptr[0], &ptr[2], sizeof(ptr[2]), TO | FROM | PARAM
+ #pragma omp target map(ptr[2])
+ ptr[1] = 6;
+}
+void f3() {
+ int *ptr;
+ // &ptr, &ptr[0], sizeof(ptr[0:2]), TO | FROM | PARAM | PTR_AND_OBJ
#pragma omp target map(ptr, ptr[0:2])
- {
- ptr[1] = 6;
- }
+ ptr[1] = 7;
+}
+
+void f4() {
+ int *ptr;
+ // &ptr, &ptr[2], sizeof(ptr[2]), TO | FROM | PARAM | PTR_AND_OBJ
#pragma omp target map(ptr, ptr[2])
- {
ptr[2] = 8;
- }
- #pragma omp target data map(ptr, ptr[2])
- {
+}
+
+void f5() {
+ int *ptr;
+ // &ptr, &ptr[2], sizeof(ptr[2]), TO | FROM | PARAM | PTR_AND_OBJ
+ #pragma omp target data map(ptr[2], ptr)
ptr[2] = 9;
- }
+}
+
+void f6() {
+ int *ptr;
+ // &ptr, &ptr[0], sizeof(ptr[2]), TO | FROM | PARAM | PTR_AND_OBJ
+ #pragma omp target data map(ptr, ptr[2])
+ ptr[2] = 10;
}
#endif
-// CHECK-LABEL: define {{[^@]+}}@_Z3foov
+//.
+// CHECK: @.offload_sizes = private unnamed_addr constant [1 x i64] [i64 8]
+// CHECK: @.offload_maptypes = private unnamed_addr constant [1 x i64] [i64 [[#0x23]]]
+// CHECK: @.offload_sizes.1 = private unnamed_addr constant [1 x i64] [i64 4]
+// CHECK: @.offload_maptypes.2 = private unnamed_addr constant [1 x i64] [i64 [[#0x23]]]
+// CHECK: @.offload_sizes.3 = private unnamed_addr constant [1 x i64] [i64 8]
+// CHECK: @.offload_maptypes.4 = private unnamed_addr constant [1 x i64] [i64 [[#0x33]]]
+// CHECK: @.offload_sizes.5 = private unnamed_addr constant [1 x i64] [i64 4]
+// CHECK: @.offload_maptypes.6 = private unnamed_addr constant [1 x i64] [i64 [[#0x33]]]
+// CHECK: @.offload_sizes.7 = private unnamed_addr constant [1 x i64] [i64 4]
+// CHECK: @.offload_maptypes.8 = private unnamed_addr constant [1 x i64] [i64 [[#0x33]]]
+// CHECK: @.offload_sizes.9 = private unnamed_addr constant [1 x i64] [i64 4]
+// CHECK: @.offload_maptypes.10 = private unnamed_addr constant [1 x i64] [i64 [[#0x33]]]
+//.
+// CHECK-LABEL: define {{[^@]+}}@_Z2f1v
// CHECK-SAME: () #[[ATTR0:[0-9]+]] {
-// CHECK-NEXT: entry:
-// CHECK-NEXT: [[PTR:%.*]] = alloca ptr, align 8
-// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8
-// CHECK-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8
-// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8
-// CHECK-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8
-// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS2:%.*]] = alloca [1 x ptr], align 8
-// CHECK-NEXT: [[DOTOFFLOAD_PTRS3:%.*]] = alloca [1 x ptr], align 8
-// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS4:%.*]] = alloca [1 x ptr], align 8
-// CHECK-NEXT: [[KERNEL_ARGS5:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8
-// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS9:%.*]] = alloca [1 x ptr], align 8
-// CHECK-NEXT: [[DOTOFFLOAD_PTRS10:%.*]] = alloca [1 x ptr], align 8
-// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS11:%.*]] = alloca [1 x ptr], align 8
-// CHECK-NEXT: [[CALL:%.*]] = call noalias noundef ptr @_Z6malloci(i32 noundef signext 12) #[[ATTR3:[0-9]+]]
-// CHECK-NEXT: store ptr [[CALL]], ptr [[PTR]], align 8
-// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[PTR]], align 8
-// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[PTR]], align 8
-// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds nuw i32, ptr [[TMP1]], i64 0
-// CHECK-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
-// CHECK-NEXT: store ptr [[PTR]], ptr [[TMP2]], align 8
-// CHECK-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
-// CHECK-NEXT: store ptr [[ARRAYIDX]], ptr [[TMP3]], align 8
-// CHECK-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0
-// CHECK-NEXT: store ptr null, ptr [[TMP4]], align 8
-// CHECK-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
-// CHECK-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
-// CHECK-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0
-// CHECK-NEXT: store i32 3, ptr [[TMP7]], align 4
-// CHECK-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1
-// CHECK-NEXT: store i32 1, ptr [[TMP8]], align 4
-// CHECK-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2
-// CHECK-NEXT: store ptr [[TMP5]], ptr [[TMP9]], align 8
-// CHECK-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3
-// CHECK-NEXT: store ptr [[TMP6]], ptr [[TMP10]], align 8
-// CHECK-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4
-// CHECK-NEXT: store ptr @.offload_sizes, ptr [[TMP11]], align 8
-// CHECK-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5
-// CHECK-NEXT: store ptr @.offload_maptypes, ptr [[TMP12]], align 8
-// CHECK-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6
-// CHECK-NEXT: store ptr null, ptr [[TMP13]], align 8
-// CHECK-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7
-// CHECK-NEXT: store ptr null, ptr [[TMP14]], align 8
-// CHECK-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8
-// CHECK-NEXT: store i64 0, ptr [[TMP15]], align 8
-// CHECK-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9
-// CHECK-NEXT: store i64 0, ptr [[TMP16]], align 8
-// CHECK-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10
-// CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP17]], align 4
-// CHECK-NEXT: [[TMP18:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11
-// CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP18]], align 4
-// CHECK-NEXT: [[TMP19:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12
-// CHECK-NEXT: store i32 0, ptr [[TMP19]], align 4
-// CHECK-NEXT: [[TMP20:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foov_l15.region_id, ptr [[KERNEL_ARGS]])
-// CHECK-NEXT: [[TMP21:%.*]] = icmp ne i32 [[TMP20]], 0
-// CHECK-NEXT: br i1 [[TMP21]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]
-// CHECK: omp_offload.failed:
-// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foov_l15(ptr [[TMP0]]) #[[ATTR3]]
-// CHECK-NEXT: br label [[OMP_OFFLOAD_CONT]]
-// CHECK: omp_offload.cont:
-// CHECK-NEXT: [[TMP22:%.*]] = load ptr, ptr [[PTR]], align 8
-// CHECK-NEXT: [[TMP23:%.*]] = load ptr, ptr [[PTR]], align 8
-// CHECK-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds i32, ptr [[TMP23]], i64 2
-// CHECK-NEXT: [[TMP24:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 0
-// CHECK-NEXT: store ptr [[PTR]], ptr [[TMP24]], align 8
-// CHECK-NEXT: [[TMP25:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i32 0, i32 0
-// CHECK-NEXT: store ptr [[ARRAYIDX1]], ptr [[TMP25]], align 8
-// CHECK-NEXT: [[TMP26:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS4]], i64 0, i64 0
-// CHECK-NEXT: store ptr null, ptr [[TMP26]], align 8
-// CHECK-NEXT: [[TMP27:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 0
-// CHECK-NEXT: [[TMP28:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i32 0, i32 0
-// CHECK-NEXT: [[TMP29:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 0
-// CHECK-NEXT: store i32 3, ptr [[TMP29]], align 4
-// CHECK-NEXT: [[TMP30:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 1
-// CHECK-NEXT: store i32 1, ptr [[TMP30]], align 4
-// CHECK-NEXT: [[TMP31:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 2
-// CHECK-NEXT: store ptr [[TMP27]], ptr [[TMP31]], align 8
-// CHECK-NEXT: [[TMP32:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 3
-// CHECK-NEXT: store ptr [[TMP28]], ptr [[TMP32]], align 8
-// CHECK-NEXT: [[TMP33:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 4
-// CHECK-NEXT: store ptr @.offload_sizes.1, ptr [[TMP33]], align 8
-// CHECK-NEXT: [[TMP34:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 5
-// CHECK-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP34]], align 8
-// CHECK-NEXT: [[TMP35:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 6
-// CHECK-NEXT: store ptr null, ptr [[TMP35]], align 8
-// CHECK-NEXT: [[TMP36:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 7
-// CHECK-NEXT: store ptr null, ptr [[TMP36]], align 8
-// CHECK-NEXT: [[TMP37:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 8
-// CHECK-NEXT: store i64 0, ptr [[TMP37]], align 8
-// CHECK-NEXT: [[TMP38:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 9
-// CHECK-NEXT: store i64 0, ptr [[TMP38]], align 8
-// CHECK-NEXT: [[TMP39:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 10
-// CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP39]], align 4
-// CHECK-NEXT: [[TMP40:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 11
-// CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP40]], align 4
-// CHECK-NEXT: [[TMP41:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 12
-// CHECK-NEXT: store i32 0, ptr [[TMP41]], align 4
-// CHECK-NEXT: [[TMP42:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foov_l19.region_id, ptr [[KERNEL_ARGS5]])
-// CHECK-NEXT: [[TMP43:%.*]] = icmp ne i32 [[TMP42]], 0
-// CHECK-NEXT: br i1 [[TMP43]], label [[OMP_OFFLOAD_FAILED6:%.*]], label [[OMP_OFFLOAD_CONT7:%.*]]
-// CHECK: omp_offload.failed6:
-// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foov_l19(ptr [[TMP22]]) #[[ATTR3]]
-// CHECK-NEXT: br label [[OMP_OFFLOAD_CONT7]]
-// CHECK: omp_offload.cont7:
-// CHECK-NEXT: [[TMP44:%.*]] = load ptr, ptr [[PTR]], align 8
-// CHECK-NEXT: [[ARRAYIDX8:%.*]] = getelementptr inbounds i32, ptr [[TMP44]], i64 2
-// CHECK-NEXT: [[TMP45:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS9]], i32 0, i32 0
-// CHECK-NEXT: store ptr [[PTR]], ptr [[TMP45]], align 8
-// CHECK-NEXT: [[TMP46:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS10]], i32 0, i32 0
-// CHECK-NEXT: store ptr [[ARRAYIDX8]], ptr [[TMP46]], align 8
-// CHECK-NEXT: [[TMP47:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS11]], i64 0, i64 0
-// CHECK-NEXT: store ptr null, ptr [[TMP47]], align 8
-// CHECK-NEXT: [[TMP48:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS9]], i32 0, i32 0
-// CHECK-NEXT: [[TMP49:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS10]], i32 0, i32 0
-// CHECK-NEXT: call void @__tgt_target_data_begin_mapper(ptr @[[GLOB1]], i64 -1, i32 1, ptr [[TMP48]], ptr [[TMP49]], ptr @.offload_sizes.3, ptr @.offload_maptypes.4, ptr null, ptr null)
-// CHECK-NEXT: [[TMP50:%.*]] = load ptr, ptr [[PTR]], align 8
-// CHECK-NEXT: [[ARRAYIDX12:%.*]] = getelementptr inbounds i32, ptr [[TMP50]], i64 2
-// CHECK-NEXT: store i32 9, ptr [[ARRAYIDX12]], align 4
-// CHECK-NEXT: [[TMP51:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS9]], i32 0, i32 0
-// CHECK-NEXT: [[TMP52:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS10]], i32 0, i32 0
-// CHECK-NEXT: call void @__tgt_target_data_end_mapper(ptr @[[GLOB1]], i64 -1, i32 1, ptr [[TMP51]], ptr [[TMP52]], ptr @.offload_sizes.3, ptr @.offload_maptypes.4, ptr null, ptr null)
-// CHECK-NEXT: ret void
-//
-//
-// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foov_l15
-// CHECK-SAME: (ptr noundef [[PTR:%.*]]) #[[ATTR2:[0-9]+]] {
-// CHECK-NEXT: entry:
-// CHECK-NEXT: [[PTR_ADDR:%.*]] = alloca ptr, align 8
-// CHECK-NEXT: store ptr [[PTR]], ptr [[PTR_ADDR]], align 8
-// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8
-// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 1
-// CHECK-NEXT: store i32 6, ptr [[ARRAYIDX]], align 4
-// CHECK-NEXT: ret void
-//
-//
-// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foov_l19
-// CHECK-SAME: (ptr noundef [[PTR:%.*]]) #[[ATTR2]] {
-// CHECK-NEXT: entry:
-// CHECK-NEXT: [[PTR_ADDR:%.*]] = alloca ptr, align 8
-// CHECK-NEXT: store ptr [[PTR]], ptr [[PTR_ADDR]], align 8
-// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8
-// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 2
-// CHECK-NEXT: store i32 8, ptr [[ARRAYIDX]], align 4
-// CHECK-NEXT: ret void
+// CHECK: entry:
+// CHECK: [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0
+// CHECK: store ptr [[PTR:%.*]], ptr [[TMP0]], align 8
+// CHECK: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0
+// CHECK: store ptr [[PTR]], ptr [[TMP1]], align 8
+// CHECK: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0
+// CHECK: store ptr null, ptr [[TMP2]], align 8
+// CHECK: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
+// CHECK: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// CHECK: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0
+//
+//
+// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f1v_l13
+// CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[PTR:%.*]]) #[[ATTR1:[0-9]+]] {
+// CHECK: entry:
+// CHECK: store ptr [[PTR]], ptr [[PTR_ADDR:%.*]], align 8
+// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8, !nonnull [[META11:![0-9]+]], !align [[META12:![0-9]+]]
+// CHECK: [[TMP1:%.*]] = load ptr, ptr [[TMP0]], align 8
+// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 1
+// CHECK: store i32 5, ptr [[ARRAYIDX]], align 4
+// CHECK: ret void
+//
+//
+// CHECK-LABEL: define {{[^@]+}}@_Z2f2v
+// CHECK-SAME: () #[[ATTR0]] {
+// CHECK: entry:
+// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR:%.*]], align 8
+// CHECK: [[TMP1:%.*]] = load ptr, ptr [[PTR]], align 8
+// CHECK: [[TMP2:%.*]] = load ptr, ptr [[PTR]], align 8
+// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP2]], i64 2
+// CHECK: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0
+// CHECK: store ptr [[TMP1]], ptr [[TMP3]], align 8
+// CHECK: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0
+// CHECK: store ptr [[ARRAYIDX]], ptr [[TMP4]], align 8
+// CHECK: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0
+// CHECK: store ptr null, ptr [[TMP5]], align 8
+// CHECK: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
+// CHECK: [[TMP7:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// CHECK: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0
+//
+//
+// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f2v_l20
+// CHECK-SAME: (ptr noundef [[PTR:%.*]]) #[[ATTR1]] {
+// CHECK: entry:
+// CHECK: store ptr [[PTR]], ptr [[PTR_ADDR:%.*]], align 8
+// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8
+// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 1
+// CHECK: store i32 6, ptr [[ARRAYIDX]], align 4
+// CHECK: ret void
+//
+//
+// CHECK-LABEL: define {{[^@]+}}@_Z2f3v
+// CHECK-SAME: () #[[ATTR0]] {
+// CHECK: entry:
+// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR:%.*]], align 8
+// CHECK: [[TMP1:%.*]] = load ptr, ptr [[PTR]], align 8
+// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds nuw i32, ptr [[TMP1]], i64 0
+// CHECK: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0
+// CHECK: store ptr [[PTR]], ptr [[TMP2]], align 8
+// CHECK: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0
+// CHECK: store ptr [[ARRAYIDX]], ptr [[TMP3]], align 8
+// CHECK: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0
+// CHECK: store ptr null, ptr [[TMP4]], align 8
+// CHECK: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
+// CHECK: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// CHECK: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0
+//
+//
+// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f3v_l27
+// CHECK-SAME: (ptr noundef [[PTR:%.*]]) #[[ATTR1]] {
+// CHECK: entry:
+// CHECK: store ptr [[PTR]], ptr [[PTR_ADDR:%.*]], align 8
+// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8
+// CHECK: [[ARRAYIDX:%.*]] ...
[truncated]
|
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
The output of the compile-and-run tests is incorrect. These will be used for reference in future commits that resolve the issues.
Also updated the existing clang LIT test,
target_map_both_pointer_pointee_codegen.cpp, with more constructs and fewer CHECKs (through more update_cc_test_checks filters).