1 // 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 _ 2 // expected-no-diagnostics 3 #ifndef HEADER 4 #define HEADER 5 6 ///==========================================================================/// 7 // RUN: %clang_cc1 -DCK4 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix=CHECK1 8 // RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 9 // 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 -allow-deprecated-dag-overlap %s --check-prefix=CHECK2 10 // RUN: %clang_cc1 -DCK4 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix=CHECK3 11 // RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 12 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix=CHECK4 13 14 // RUN: %clang_cc1 -DCK4 -verify -fopenmp -fopenmp-version=45 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix=CHECK5 15 // RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-version=45 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 16 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix=CHECK6 17 // RUN: %clang_cc1 -DCK4 -verify -fopenmp -fopenmp-version=45 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix=CHECK7 18 // RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-version=45 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 19 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix=CHECK8 20 21 // RUN: %clang_cc1 -DCK4 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix=CHECK9 22 // RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 23 // 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 -allow-deprecated-dag-overlap %s --check-prefix=CHECK10 24 // RUN: %clang_cc1 -DCK4 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix=CHECK11 25 // RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 26 // RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix=CHECK12 27 28 // RUN: %clang_cc1 -DCK4 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}" 29 // RUN: %clang_cc1 -DCK4 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 30 // RUN: %clang_cc1 -fopenmp-simd -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 --implicit-check-not="{{__kmpc|__tgt}}" 31 // RUN: %clang_cc1 -DCK4 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}" 32 // RUN: %clang_cc1 -DCK4 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 33 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --implicit-check-not="{{__kmpc|__tgt}}" 34 #ifdef CK4 35 36 37 // Map types: OMP_MAP_PRIVATE_VAL | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 800 38 39 void implicit_maps_nested_integer (int a){ 40 int i = a; 41 42 // The captures in parallel are by reference. Only the capture in target is by 43 // copy. 44 45 #pragma omp parallel 46 { 47 48 #pragma omp target 49 { 50 #pragma omp parallel 51 { 52 ++i; 53 } 54 } 55 } 56 } 57 58 #endif // CK4 59 #endif 60 // CHECK1-LABEL: define {{[^@]+}}@_Z28implicit_maps_nested_integeri 61 // CHECK1-SAME: (i32 signext [[A:%.*]]) #[[ATTR0:[0-9]+]] { 62 // CHECK1-NEXT: entry: 63 // CHECK1-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4 64 // CHECK1-NEXT: [[I:%.*]] = alloca i32, align 4 65 // CHECK1-NEXT: store i32 [[A]], i32* [[A_ADDR]], align 4 66 // CHECK1-NEXT: [[TMP0:%.*]] = load i32, i32* [[A_ADDR]], align 4 67 // CHECK1-NEXT: store i32 [[TMP0]], i32* [[I]], align 4 68 // CHECK1-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* @[[GLOB1:[0-9]+]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined. to void (i32*, i32*, ...)*), i32* [[I]]) 69 // CHECK1-NEXT: ret void 70 // 71 // 72 // CHECK1-LABEL: define {{[^@]+}}@.omp_outlined. 73 // CHECK1-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[I:%.*]]) #[[ATTR1:[0-9]+]] { 74 // CHECK1-NEXT: entry: 75 // CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 76 // CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 77 // CHECK1-NEXT: [[I_ADDR:%.*]] = alloca i32*, align 8 78 // CHECK1-NEXT: [[I_CASTED:%.*]] = alloca i64, align 8 79 // CHECK1-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x i8*], align 8 80 // CHECK1-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x i8*], align 8 81 // CHECK1-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x i8*], align 8 82 // CHECK1-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 83 // CHECK1-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 84 // CHECK1-NEXT: store i32* [[I]], i32** [[I_ADDR]], align 8 85 // CHECK1-NEXT: [[TMP0:%.*]] = load i32*, i32** [[I_ADDR]], align 8 86 // CHECK1-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 87 // CHECK1-NEXT: [[CONV:%.*]] = bitcast i64* [[I_CASTED]] to i32* 88 // CHECK1-NEXT: store i32 [[TMP1]], i32* [[CONV]], align 4 89 // CHECK1-NEXT: [[TMP2:%.*]] = load i64, i64* [[I_CASTED]], align 8 90 // CHECK1-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 91 // CHECK1-NEXT: [[TMP4:%.*]] = bitcast i8** [[TMP3]] to i64* 92 // CHECK1-NEXT: store i64 [[TMP2]], i64* [[TMP4]], align 8 93 // CHECK1-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 94 // CHECK1-NEXT: [[TMP6:%.*]] = bitcast i8** [[TMP5]] to i64* 95 // CHECK1-NEXT: store i64 [[TMP2]], i64* [[TMP6]], align 8 96 // CHECK1-NEXT: [[TMP7:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0 97 // CHECK1-NEXT: store i8* null, i8** [[TMP7]], align 8 98 // CHECK1-NEXT: [[TMP8:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 99 // CHECK1-NEXT: [[TMP9:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 100 // CHECK1-NEXT: [[TMP10:%.*]] = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z28implicit_maps_nested_integeri_l48.region_id, i32 1, i8** [[TMP8]], i8** [[TMP9]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null, i32 1, i32 0) 101 // CHECK1-NEXT: [[TMP11:%.*]] = icmp ne i32 [[TMP10]], 0 102 // CHECK1-NEXT: br i1 [[TMP11]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] 103 // CHECK1: omp_offload.failed: 104 // CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z28implicit_maps_nested_integeri_l48(i64 [[TMP2]]) #[[ATTR3:[0-9]+]] 105 // CHECK1-NEXT: br label [[OMP_OFFLOAD_CONT]] 106 // CHECK1: omp_offload.cont: 107 // CHECK1-NEXT: ret void 108 // 109 // 110 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z28implicit_maps_nested_integeri_l48 111 // CHECK1-SAME: (i64 [[I:%.*]]) #[[ATTR2:[0-9]+]] { 112 // CHECK1-NEXT: entry: 113 // CHECK1-NEXT: [[I_ADDR:%.*]] = alloca i64, align 8 114 // CHECK1-NEXT: store i64 [[I]], i64* [[I_ADDR]], align 8 115 // CHECK1-NEXT: [[CONV:%.*]] = bitcast i64* [[I_ADDR]] to i32* 116 // CHECK1-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* @[[GLOB1]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined..1 to void (i32*, i32*, ...)*), i32* [[CONV]]) 117 // CHECK1-NEXT: ret void 118 // 119 // 120 // CHECK1-LABEL: define {{[^@]+}}@.omp_outlined..1 121 // CHECK1-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[I:%.*]]) #[[ATTR1]] { 122 // CHECK1-NEXT: entry: 123 // CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 124 // CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 125 // CHECK1-NEXT: [[I_ADDR:%.*]] = alloca i32*, align 8 126 // CHECK1-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 127 // CHECK1-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 128 // CHECK1-NEXT: store i32* [[I]], i32** [[I_ADDR]], align 8 129 // CHECK1-NEXT: [[TMP0:%.*]] = load i32*, i32** [[I_ADDR]], align 8 130 // CHECK1-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 131 // CHECK1-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 132 // CHECK1-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4 133 // CHECK1-NEXT: ret void 134 // 135 // 136 // CHECK1-LABEL: define {{[^@]+}}@.omp_offloading.requires_reg 137 // CHECK1-SAME: () #[[ATTR4:[0-9]+]] { 138 // CHECK1-NEXT: entry: 139 // CHECK1-NEXT: call void @__tgt_register_requires(i64 1) 140 // CHECK1-NEXT: ret void 141 // 142 // 143 // CHECK2-LABEL: define {{[^@]+}}@_Z28implicit_maps_nested_integeri 144 // CHECK2-SAME: (i32 signext [[A:%.*]]) #[[ATTR0:[0-9]+]] { 145 // CHECK2-NEXT: entry: 146 // CHECK2-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4 147 // CHECK2-NEXT: [[I:%.*]] = alloca i32, align 4 148 // CHECK2-NEXT: store i32 [[A]], i32* [[A_ADDR]], align 4 149 // CHECK2-NEXT: [[TMP0:%.*]] = load i32, i32* [[A_ADDR]], align 4 150 // CHECK2-NEXT: store i32 [[TMP0]], i32* [[I]], align 4 151 // CHECK2-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* @[[GLOB1:[0-9]+]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined. to void (i32*, i32*, ...)*), i32* [[I]]) 152 // CHECK2-NEXT: ret void 153 // 154 // 155 // CHECK2-LABEL: define {{[^@]+}}@.omp_outlined. 156 // CHECK2-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[I:%.*]]) #[[ATTR1:[0-9]+]] { 157 // CHECK2-NEXT: entry: 158 // CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 159 // CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 160 // CHECK2-NEXT: [[I_ADDR:%.*]] = alloca i32*, align 8 161 // CHECK2-NEXT: [[I_CASTED:%.*]] = alloca i64, align 8 162 // CHECK2-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x i8*], align 8 163 // CHECK2-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x i8*], align 8 164 // CHECK2-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x i8*], align 8 165 // CHECK2-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 166 // CHECK2-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 167 // CHECK2-NEXT: store i32* [[I]], i32** [[I_ADDR]], align 8 168 // CHECK2-NEXT: [[TMP0:%.*]] = load i32*, i32** [[I_ADDR]], align 8 169 // CHECK2-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 170 // CHECK2-NEXT: [[CONV:%.*]] = bitcast i64* [[I_CASTED]] to i32* 171 // CHECK2-NEXT: store i32 [[TMP1]], i32* [[CONV]], align 4 172 // CHECK2-NEXT: [[TMP2:%.*]] = load i64, i64* [[I_CASTED]], align 8 173 // CHECK2-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 174 // CHECK2-NEXT: [[TMP4:%.*]] = bitcast i8** [[TMP3]] to i64* 175 // CHECK2-NEXT: store i64 [[TMP2]], i64* [[TMP4]], align 8 176 // CHECK2-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 177 // CHECK2-NEXT: [[TMP6:%.*]] = bitcast i8** [[TMP5]] to i64* 178 // CHECK2-NEXT: store i64 [[TMP2]], i64* [[TMP6]], align 8 179 // CHECK2-NEXT: [[TMP7:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0 180 // CHECK2-NEXT: store i8* null, i8** [[TMP7]], align 8 181 // CHECK2-NEXT: [[TMP8:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 182 // CHECK2-NEXT: [[TMP9:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 183 // CHECK2-NEXT: [[TMP10:%.*]] = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z28implicit_maps_nested_integeri_l48.region_id, i32 1, i8** [[TMP8]], i8** [[TMP9]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null, i32 1, i32 0) 184 // CHECK2-NEXT: [[TMP11:%.*]] = icmp ne i32 [[TMP10]], 0 185 // CHECK2-NEXT: br i1 [[TMP11]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] 186 // CHECK2: omp_offload.failed: 187 // CHECK2-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z28implicit_maps_nested_integeri_l48(i64 [[TMP2]]) #[[ATTR3:[0-9]+]] 188 // CHECK2-NEXT: br label [[OMP_OFFLOAD_CONT]] 189 // CHECK2: omp_offload.cont: 190 // CHECK2-NEXT: ret void 191 // 192 // 193 // CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z28implicit_maps_nested_integeri_l48 194 // CHECK2-SAME: (i64 [[I:%.*]]) #[[ATTR2:[0-9]+]] { 195 // CHECK2-NEXT: entry: 196 // CHECK2-NEXT: [[I_ADDR:%.*]] = alloca i64, align 8 197 // CHECK2-NEXT: store i64 [[I]], i64* [[I_ADDR]], align 8 198 // CHECK2-NEXT: [[CONV:%.*]] = bitcast i64* [[I_ADDR]] to i32* 199 // CHECK2-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* @[[GLOB1]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined..1 to void (i32*, i32*, ...)*), i32* [[CONV]]) 200 // CHECK2-NEXT: ret void 201 // 202 // 203 // CHECK2-LABEL: define {{[^@]+}}@.omp_outlined..1 204 // CHECK2-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[I:%.*]]) #[[ATTR1]] { 205 // CHECK2-NEXT: entry: 206 // CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 207 // CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 208 // CHECK2-NEXT: [[I_ADDR:%.*]] = alloca i32*, align 8 209 // CHECK2-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 210 // CHECK2-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 211 // CHECK2-NEXT: store i32* [[I]], i32** [[I_ADDR]], align 8 212 // CHECK2-NEXT: [[TMP0:%.*]] = load i32*, i32** [[I_ADDR]], align 8 213 // CHECK2-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 214 // CHECK2-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 215 // CHECK2-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4 216 // CHECK2-NEXT: ret void 217 // 218 // 219 // CHECK2-LABEL: define {{[^@]+}}@.omp_offloading.requires_reg 220 // CHECK2-SAME: () #[[ATTR4:[0-9]+]] { 221 // CHECK2-NEXT: entry: 222 // CHECK2-NEXT: call void @__tgt_register_requires(i64 1) 223 // CHECK2-NEXT: ret void 224 // 225 // 226 // CHECK3-LABEL: define {{[^@]+}}@_Z28implicit_maps_nested_integeri 227 // CHECK3-SAME: (i32 [[A:%.*]]) #[[ATTR0:[0-9]+]] { 228 // CHECK3-NEXT: entry: 229 // CHECK3-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4 230 // CHECK3-NEXT: [[I:%.*]] = alloca i32, align 4 231 // CHECK3-NEXT: store i32 [[A]], i32* [[A_ADDR]], align 4 232 // CHECK3-NEXT: [[TMP0:%.*]] = load i32, i32* [[A_ADDR]], align 4 233 // CHECK3-NEXT: store i32 [[TMP0]], i32* [[I]], align 4 234 // CHECK3-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* @[[GLOB1:[0-9]+]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined. to void (i32*, i32*, ...)*), i32* [[I]]) 235 // CHECK3-NEXT: ret void 236 // 237 // 238 // CHECK3-LABEL: define {{[^@]+}}@.omp_outlined. 239 // CHECK3-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[I:%.*]]) #[[ATTR1:[0-9]+]] { 240 // CHECK3-NEXT: entry: 241 // CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 242 // CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 243 // CHECK3-NEXT: [[I_ADDR:%.*]] = alloca i32*, align 4 244 // CHECK3-NEXT: [[I_CASTED:%.*]] = alloca i32, align 4 245 // CHECK3-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x i8*], align 4 246 // CHECK3-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x i8*], align 4 247 // CHECK3-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x i8*], align 4 248 // CHECK3-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 249 // CHECK3-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 250 // CHECK3-NEXT: store i32* [[I]], i32** [[I_ADDR]], align 4 251 // CHECK3-NEXT: [[TMP0:%.*]] = load i32*, i32** [[I_ADDR]], align 4 252 // CHECK3-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 253 // CHECK3-NEXT: store i32 [[TMP1]], i32* [[I_CASTED]], align 4 254 // CHECK3-NEXT: [[TMP2:%.*]] = load i32, i32* [[I_CASTED]], align 4 255 // CHECK3-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 256 // CHECK3-NEXT: [[TMP4:%.*]] = bitcast i8** [[TMP3]] to i32* 257 // CHECK3-NEXT: store i32 [[TMP2]], i32* [[TMP4]], align 4 258 // CHECK3-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 259 // CHECK3-NEXT: [[TMP6:%.*]] = bitcast i8** [[TMP5]] to i32* 260 // CHECK3-NEXT: store i32 [[TMP2]], i32* [[TMP6]], align 4 261 // CHECK3-NEXT: [[TMP7:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0 262 // CHECK3-NEXT: store i8* null, i8** [[TMP7]], align 4 263 // CHECK3-NEXT: [[TMP8:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 264 // CHECK3-NEXT: [[TMP9:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 265 // CHECK3-NEXT: [[TMP10:%.*]] = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z28implicit_maps_nested_integeri_l48.region_id, i32 1, i8** [[TMP8]], i8** [[TMP9]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null, i32 1, i32 0) 266 // CHECK3-NEXT: [[TMP11:%.*]] = icmp ne i32 [[TMP10]], 0 267 // CHECK3-NEXT: br i1 [[TMP11]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] 268 // CHECK3: omp_offload.failed: 269 // CHECK3-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z28implicit_maps_nested_integeri_l48(i32 [[TMP2]]) #[[ATTR3:[0-9]+]] 270 // CHECK3-NEXT: br label [[OMP_OFFLOAD_CONT]] 271 // CHECK3: omp_offload.cont: 272 // CHECK3-NEXT: ret void 273 // 274 // 275 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z28implicit_maps_nested_integeri_l48 276 // CHECK3-SAME: (i32 [[I:%.*]]) #[[ATTR2:[0-9]+]] { 277 // CHECK3-NEXT: entry: 278 // CHECK3-NEXT: [[I_ADDR:%.*]] = alloca i32, align 4 279 // CHECK3-NEXT: store i32 [[I]], i32* [[I_ADDR]], align 4 280 // CHECK3-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* @[[GLOB1]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined..1 to void (i32*, i32*, ...)*), i32* [[I_ADDR]]) 281 // CHECK3-NEXT: ret void 282 // 283 // 284 // CHECK3-LABEL: define {{[^@]+}}@.omp_outlined..1 285 // CHECK3-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[I:%.*]]) #[[ATTR1]] { 286 // CHECK3-NEXT: entry: 287 // CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 288 // CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 289 // CHECK3-NEXT: [[I_ADDR:%.*]] = alloca i32*, align 4 290 // CHECK3-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 291 // CHECK3-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 292 // CHECK3-NEXT: store i32* [[I]], i32** [[I_ADDR]], align 4 293 // CHECK3-NEXT: [[TMP0:%.*]] = load i32*, i32** [[I_ADDR]], align 4 294 // CHECK3-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 295 // CHECK3-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 296 // CHECK3-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4 297 // CHECK3-NEXT: ret void 298 // 299 // 300 // CHECK3-LABEL: define {{[^@]+}}@.omp_offloading.requires_reg 301 // CHECK3-SAME: () #[[ATTR4:[0-9]+]] { 302 // CHECK3-NEXT: entry: 303 // CHECK3-NEXT: call void @__tgt_register_requires(i64 1) 304 // CHECK3-NEXT: ret void 305 // 306 // 307 // CHECK4-LABEL: define {{[^@]+}}@_Z28implicit_maps_nested_integeri 308 // CHECK4-SAME: (i32 [[A:%.*]]) #[[ATTR0:[0-9]+]] { 309 // CHECK4-NEXT: entry: 310 // CHECK4-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4 311 // CHECK4-NEXT: [[I:%.*]] = alloca i32, align 4 312 // CHECK4-NEXT: store i32 [[A]], i32* [[A_ADDR]], align 4 313 // CHECK4-NEXT: [[TMP0:%.*]] = load i32, i32* [[A_ADDR]], align 4 314 // CHECK4-NEXT: store i32 [[TMP0]], i32* [[I]], align 4 315 // CHECK4-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* @[[GLOB1:[0-9]+]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined. to void (i32*, i32*, ...)*), i32* [[I]]) 316 // CHECK4-NEXT: ret void 317 // 318 // 319 // CHECK4-LABEL: define {{[^@]+}}@.omp_outlined. 320 // CHECK4-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[I:%.*]]) #[[ATTR1:[0-9]+]] { 321 // CHECK4-NEXT: entry: 322 // CHECK4-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 323 // CHECK4-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 324 // CHECK4-NEXT: [[I_ADDR:%.*]] = alloca i32*, align 4 325 // CHECK4-NEXT: [[I_CASTED:%.*]] = alloca i32, align 4 326 // CHECK4-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x i8*], align 4 327 // CHECK4-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x i8*], align 4 328 // CHECK4-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x i8*], align 4 329 // CHECK4-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 330 // CHECK4-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 331 // CHECK4-NEXT: store i32* [[I]], i32** [[I_ADDR]], align 4 332 // CHECK4-NEXT: [[TMP0:%.*]] = load i32*, i32** [[I_ADDR]], align 4 333 // CHECK4-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 334 // CHECK4-NEXT: store i32 [[TMP1]], i32* [[I_CASTED]], align 4 335 // CHECK4-NEXT: [[TMP2:%.*]] = load i32, i32* [[I_CASTED]], align 4 336 // CHECK4-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 337 // CHECK4-NEXT: [[TMP4:%.*]] = bitcast i8** [[TMP3]] to i32* 338 // CHECK4-NEXT: store i32 [[TMP2]], i32* [[TMP4]], align 4 339 // CHECK4-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 340 // CHECK4-NEXT: [[TMP6:%.*]] = bitcast i8** [[TMP5]] to i32* 341 // CHECK4-NEXT: store i32 [[TMP2]], i32* [[TMP6]], align 4 342 // CHECK4-NEXT: [[TMP7:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0 343 // CHECK4-NEXT: store i8* null, i8** [[TMP7]], align 4 344 // CHECK4-NEXT: [[TMP8:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 345 // CHECK4-NEXT: [[TMP9:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 346 // CHECK4-NEXT: [[TMP10:%.*]] = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z28implicit_maps_nested_integeri_l48.region_id, i32 1, i8** [[TMP8]], i8** [[TMP9]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null, i32 1, i32 0) 347 // CHECK4-NEXT: [[TMP11:%.*]] = icmp ne i32 [[TMP10]], 0 348 // CHECK4-NEXT: br i1 [[TMP11]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] 349 // CHECK4: omp_offload.failed: 350 // CHECK4-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z28implicit_maps_nested_integeri_l48(i32 [[TMP2]]) #[[ATTR3:[0-9]+]] 351 // CHECK4-NEXT: br label [[OMP_OFFLOAD_CONT]] 352 // CHECK4: omp_offload.cont: 353 // CHECK4-NEXT: ret void 354 // 355 // 356 // CHECK4-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z28implicit_maps_nested_integeri_l48 357 // CHECK4-SAME: (i32 [[I:%.*]]) #[[ATTR2:[0-9]+]] { 358 // CHECK4-NEXT: entry: 359 // CHECK4-NEXT: [[I_ADDR:%.*]] = alloca i32, align 4 360 // CHECK4-NEXT: store i32 [[I]], i32* [[I_ADDR]], align 4 361 // CHECK4-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* @[[GLOB1]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined..1 to void (i32*, i32*, ...)*), i32* [[I_ADDR]]) 362 // CHECK4-NEXT: ret void 363 // 364 // 365 // CHECK4-LABEL: define {{[^@]+}}@.omp_outlined..1 366 // CHECK4-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[I:%.*]]) #[[ATTR1]] { 367 // CHECK4-NEXT: entry: 368 // CHECK4-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 369 // CHECK4-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 370 // CHECK4-NEXT: [[I_ADDR:%.*]] = alloca i32*, align 4 371 // CHECK4-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 372 // CHECK4-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 373 // CHECK4-NEXT: store i32* [[I]], i32** [[I_ADDR]], align 4 374 // CHECK4-NEXT: [[TMP0:%.*]] = load i32*, i32** [[I_ADDR]], align 4 375 // CHECK4-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 376 // CHECK4-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 377 // CHECK4-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4 378 // CHECK4-NEXT: ret void 379 // 380 // 381 // CHECK4-LABEL: define {{[^@]+}}@.omp_offloading.requires_reg 382 // CHECK4-SAME: () #[[ATTR4:[0-9]+]] { 383 // CHECK4-NEXT: entry: 384 // CHECK4-NEXT: call void @__tgt_register_requires(i64 1) 385 // CHECK4-NEXT: ret void 386 // 387 // 388 // CHECK5-LABEL: define {{[^@]+}}@_Z28implicit_maps_nested_integeri 389 // CHECK5-SAME: (i32 signext [[A:%.*]]) #[[ATTR0:[0-9]+]] { 390 // CHECK5-NEXT: entry: 391 // CHECK5-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4 392 // CHECK5-NEXT: [[I:%.*]] = alloca i32, align 4 393 // CHECK5-NEXT: store i32 [[A]], i32* [[A_ADDR]], align 4 394 // CHECK5-NEXT: [[TMP0:%.*]] = load i32, i32* [[A_ADDR]], align 4 395 // CHECK5-NEXT: store i32 [[TMP0]], i32* [[I]], align 4 396 // CHECK5-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* @[[GLOB1:[0-9]+]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined. to void (i32*, i32*, ...)*), i32* [[I]]) 397 // CHECK5-NEXT: ret void 398 // 399 // 400 // CHECK5-LABEL: define {{[^@]+}}@.omp_outlined. 401 // CHECK5-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[I:%.*]]) #[[ATTR1:[0-9]+]] { 402 // CHECK5-NEXT: entry: 403 // CHECK5-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 404 // CHECK5-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 405 // CHECK5-NEXT: [[I_ADDR:%.*]] = alloca i32*, align 8 406 // CHECK5-NEXT: [[I_CASTED:%.*]] = alloca i64, align 8 407 // CHECK5-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x i8*], align 8 408 // CHECK5-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x i8*], align 8 409 // CHECK5-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x i8*], align 8 410 // CHECK5-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 411 // CHECK5-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 412 // CHECK5-NEXT: store i32* [[I]], i32** [[I_ADDR]], align 8 413 // CHECK5-NEXT: [[TMP0:%.*]] = load i32*, i32** [[I_ADDR]], align 8 414 // CHECK5-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 415 // CHECK5-NEXT: [[CONV:%.*]] = bitcast i64* [[I_CASTED]] to i32* 416 // CHECK5-NEXT: store i32 [[TMP1]], i32* [[CONV]], align 4 417 // CHECK5-NEXT: [[TMP2:%.*]] = load i64, i64* [[I_CASTED]], align 8 418 // CHECK5-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 419 // CHECK5-NEXT: [[TMP4:%.*]] = bitcast i8** [[TMP3]] to i64* 420 // CHECK5-NEXT: store i64 [[TMP2]], i64* [[TMP4]], align 8 421 // CHECK5-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 422 // CHECK5-NEXT: [[TMP6:%.*]] = bitcast i8** [[TMP5]] to i64* 423 // CHECK5-NEXT: store i64 [[TMP2]], i64* [[TMP6]], align 8 424 // CHECK5-NEXT: [[TMP7:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0 425 // CHECK5-NEXT: store i8* null, i8** [[TMP7]], align 8 426 // CHECK5-NEXT: [[TMP8:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 427 // CHECK5-NEXT: [[TMP9:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 428 // CHECK5-NEXT: [[TMP10:%.*]] = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z28implicit_maps_nested_integeri_l48.region_id, i32 1, i8** [[TMP8]], i8** [[TMP9]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null, i32 1, i32 0) 429 // CHECK5-NEXT: [[TMP11:%.*]] = icmp ne i32 [[TMP10]], 0 430 // CHECK5-NEXT: br i1 [[TMP11]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] 431 // CHECK5: omp_offload.failed: 432 // CHECK5-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z28implicit_maps_nested_integeri_l48(i64 [[TMP2]]) #[[ATTR3:[0-9]+]] 433 // CHECK5-NEXT: br label [[OMP_OFFLOAD_CONT]] 434 // CHECK5: omp_offload.cont: 435 // CHECK5-NEXT: ret void 436 // 437 // 438 // CHECK5-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z28implicit_maps_nested_integeri_l48 439 // CHECK5-SAME: (i64 [[I:%.*]]) #[[ATTR2:[0-9]+]] { 440 // CHECK5-NEXT: entry: 441 // CHECK5-NEXT: [[I_ADDR:%.*]] = alloca i64, align 8 442 // CHECK5-NEXT: store i64 [[I]], i64* [[I_ADDR]], align 8 443 // CHECK5-NEXT: [[CONV:%.*]] = bitcast i64* [[I_ADDR]] to i32* 444 // CHECK5-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* @[[GLOB1]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined..1 to void (i32*, i32*, ...)*), i32* [[CONV]]) 445 // CHECK5-NEXT: ret void 446 // 447 // 448 // CHECK5-LABEL: define {{[^@]+}}@.omp_outlined..1 449 // CHECK5-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[I:%.*]]) #[[ATTR1]] { 450 // CHECK5-NEXT: entry: 451 // CHECK5-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 452 // CHECK5-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 453 // CHECK5-NEXT: [[I_ADDR:%.*]] = alloca i32*, align 8 454 // CHECK5-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 455 // CHECK5-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 456 // CHECK5-NEXT: store i32* [[I]], i32** [[I_ADDR]], align 8 457 // CHECK5-NEXT: [[TMP0:%.*]] = load i32*, i32** [[I_ADDR]], align 8 458 // CHECK5-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 459 // CHECK5-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 460 // CHECK5-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4 461 // CHECK5-NEXT: ret void 462 // 463 // 464 // CHECK5-LABEL: define {{[^@]+}}@.omp_offloading.requires_reg 465 // CHECK5-SAME: () #[[ATTR4:[0-9]+]] { 466 // CHECK5-NEXT: entry: 467 // CHECK5-NEXT: call void @__tgt_register_requires(i64 1) 468 // CHECK5-NEXT: ret void 469 // 470 // 471 // CHECK6-LABEL: define {{[^@]+}}@_Z28implicit_maps_nested_integeri 472 // CHECK6-SAME: (i32 signext [[A:%.*]]) #[[ATTR0:[0-9]+]] { 473 // CHECK6-NEXT: entry: 474 // CHECK6-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4 475 // CHECK6-NEXT: [[I:%.*]] = alloca i32, align 4 476 // CHECK6-NEXT: store i32 [[A]], i32* [[A_ADDR]], align 4 477 // CHECK6-NEXT: [[TMP0:%.*]] = load i32, i32* [[A_ADDR]], align 4 478 // CHECK6-NEXT: store i32 [[TMP0]], i32* [[I]], align 4 479 // CHECK6-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* @[[GLOB1:[0-9]+]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined. to void (i32*, i32*, ...)*), i32* [[I]]) 480 // CHECK6-NEXT: ret void 481 // 482 // 483 // CHECK6-LABEL: define {{[^@]+}}@.omp_outlined. 484 // CHECK6-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[I:%.*]]) #[[ATTR1:[0-9]+]] { 485 // CHECK6-NEXT: entry: 486 // CHECK6-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 487 // CHECK6-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 488 // CHECK6-NEXT: [[I_ADDR:%.*]] = alloca i32*, align 8 489 // CHECK6-NEXT: [[I_CASTED:%.*]] = alloca i64, align 8 490 // CHECK6-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x i8*], align 8 491 // CHECK6-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x i8*], align 8 492 // CHECK6-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x i8*], align 8 493 // CHECK6-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 494 // CHECK6-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 495 // CHECK6-NEXT: store i32* [[I]], i32** [[I_ADDR]], align 8 496 // CHECK6-NEXT: [[TMP0:%.*]] = load i32*, i32** [[I_ADDR]], align 8 497 // CHECK6-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 498 // CHECK6-NEXT: [[CONV:%.*]] = bitcast i64* [[I_CASTED]] to i32* 499 // CHECK6-NEXT: store i32 [[TMP1]], i32* [[CONV]], align 4 500 // CHECK6-NEXT: [[TMP2:%.*]] = load i64, i64* [[I_CASTED]], align 8 501 // CHECK6-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 502 // CHECK6-NEXT: [[TMP4:%.*]] = bitcast i8** [[TMP3]] to i64* 503 // CHECK6-NEXT: store i64 [[TMP2]], i64* [[TMP4]], align 8 504 // CHECK6-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 505 // CHECK6-NEXT: [[TMP6:%.*]] = bitcast i8** [[TMP5]] to i64* 506 // CHECK6-NEXT: store i64 [[TMP2]], i64* [[TMP6]], align 8 507 // CHECK6-NEXT: [[TMP7:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0 508 // CHECK6-NEXT: store i8* null, i8** [[TMP7]], align 8 509 // CHECK6-NEXT: [[TMP8:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 510 // CHECK6-NEXT: [[TMP9:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 511 // CHECK6-NEXT: [[TMP10:%.*]] = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z28implicit_maps_nested_integeri_l48.region_id, i32 1, i8** [[TMP8]], i8** [[TMP9]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null, i32 1, i32 0) 512 // CHECK6-NEXT: [[TMP11:%.*]] = icmp ne i32 [[TMP10]], 0 513 // CHECK6-NEXT: br i1 [[TMP11]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] 514 // CHECK6: omp_offload.failed: 515 // CHECK6-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z28implicit_maps_nested_integeri_l48(i64 [[TMP2]]) #[[ATTR3:[0-9]+]] 516 // CHECK6-NEXT: br label [[OMP_OFFLOAD_CONT]] 517 // CHECK6: omp_offload.cont: 518 // CHECK6-NEXT: ret void 519 // 520 // 521 // CHECK6-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z28implicit_maps_nested_integeri_l48 522 // CHECK6-SAME: (i64 [[I:%.*]]) #[[ATTR2:[0-9]+]] { 523 // CHECK6-NEXT: entry: 524 // CHECK6-NEXT: [[I_ADDR:%.*]] = alloca i64, align 8 525 // CHECK6-NEXT: store i64 [[I]], i64* [[I_ADDR]], align 8 526 // CHECK6-NEXT: [[CONV:%.*]] = bitcast i64* [[I_ADDR]] to i32* 527 // CHECK6-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* @[[GLOB1]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined..1 to void (i32*, i32*, ...)*), i32* [[CONV]]) 528 // CHECK6-NEXT: ret void 529 // 530 // 531 // CHECK6-LABEL: define {{[^@]+}}@.omp_outlined..1 532 // CHECK6-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[I:%.*]]) #[[ATTR1]] { 533 // CHECK6-NEXT: entry: 534 // CHECK6-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 535 // CHECK6-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 536 // CHECK6-NEXT: [[I_ADDR:%.*]] = alloca i32*, align 8 537 // CHECK6-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 538 // CHECK6-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 539 // CHECK6-NEXT: store i32* [[I]], i32** [[I_ADDR]], align 8 540 // CHECK6-NEXT: [[TMP0:%.*]] = load i32*, i32** [[I_ADDR]], align 8 541 // CHECK6-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 542 // CHECK6-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 543 // CHECK6-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4 544 // CHECK6-NEXT: ret void 545 // 546 // 547 // CHECK6-LABEL: define {{[^@]+}}@.omp_offloading.requires_reg 548 // CHECK6-SAME: () #[[ATTR4:[0-9]+]] { 549 // CHECK6-NEXT: entry: 550 // CHECK6-NEXT: call void @__tgt_register_requires(i64 1) 551 // CHECK6-NEXT: ret void 552 // 553 // 554 // CHECK7-LABEL: define {{[^@]+}}@_Z28implicit_maps_nested_integeri 555 // CHECK7-SAME: (i32 [[A:%.*]]) #[[ATTR0:[0-9]+]] { 556 // CHECK7-NEXT: entry: 557 // CHECK7-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4 558 // CHECK7-NEXT: [[I:%.*]] = alloca i32, align 4 559 // CHECK7-NEXT: store i32 [[A]], i32* [[A_ADDR]], align 4 560 // CHECK7-NEXT: [[TMP0:%.*]] = load i32, i32* [[A_ADDR]], align 4 561 // CHECK7-NEXT: store i32 [[TMP0]], i32* [[I]], align 4 562 // CHECK7-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* @[[GLOB1:[0-9]+]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined. to void (i32*, i32*, ...)*), i32* [[I]]) 563 // CHECK7-NEXT: ret void 564 // 565 // 566 // CHECK7-LABEL: define {{[^@]+}}@.omp_outlined. 567 // CHECK7-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[I:%.*]]) #[[ATTR1:[0-9]+]] { 568 // CHECK7-NEXT: entry: 569 // CHECK7-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 570 // CHECK7-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 571 // CHECK7-NEXT: [[I_ADDR:%.*]] = alloca i32*, align 4 572 // CHECK7-NEXT: [[I_CASTED:%.*]] = alloca i32, align 4 573 // CHECK7-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x i8*], align 4 574 // CHECK7-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x i8*], align 4 575 // CHECK7-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x i8*], align 4 576 // CHECK7-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 577 // CHECK7-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 578 // CHECK7-NEXT: store i32* [[I]], i32** [[I_ADDR]], align 4 579 // CHECK7-NEXT: [[TMP0:%.*]] = load i32*, i32** [[I_ADDR]], align 4 580 // CHECK7-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 581 // CHECK7-NEXT: store i32 [[TMP1]], i32* [[I_CASTED]], align 4 582 // CHECK7-NEXT: [[TMP2:%.*]] = load i32, i32* [[I_CASTED]], align 4 583 // CHECK7-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 584 // CHECK7-NEXT: [[TMP4:%.*]] = bitcast i8** [[TMP3]] to i32* 585 // CHECK7-NEXT: store i32 [[TMP2]], i32* [[TMP4]], align 4 586 // CHECK7-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 587 // CHECK7-NEXT: [[TMP6:%.*]] = bitcast i8** [[TMP5]] to i32* 588 // CHECK7-NEXT: store i32 [[TMP2]], i32* [[TMP6]], align 4 589 // CHECK7-NEXT: [[TMP7:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0 590 // CHECK7-NEXT: store i8* null, i8** [[TMP7]], align 4 591 // CHECK7-NEXT: [[TMP8:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 592 // CHECK7-NEXT: [[TMP9:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 593 // CHECK7-NEXT: [[TMP10:%.*]] = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z28implicit_maps_nested_integeri_l48.region_id, i32 1, i8** [[TMP8]], i8** [[TMP9]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null, i32 1, i32 0) 594 // CHECK7-NEXT: [[TMP11:%.*]] = icmp ne i32 [[TMP10]], 0 595 // CHECK7-NEXT: br i1 [[TMP11]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] 596 // CHECK7: omp_offload.failed: 597 // CHECK7-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z28implicit_maps_nested_integeri_l48(i32 [[TMP2]]) #[[ATTR3:[0-9]+]] 598 // CHECK7-NEXT: br label [[OMP_OFFLOAD_CONT]] 599 // CHECK7: omp_offload.cont: 600 // CHECK7-NEXT: ret void 601 // 602 // 603 // CHECK7-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z28implicit_maps_nested_integeri_l48 604 // CHECK7-SAME: (i32 [[I:%.*]]) #[[ATTR2:[0-9]+]] { 605 // CHECK7-NEXT: entry: 606 // CHECK7-NEXT: [[I_ADDR:%.*]] = alloca i32, align 4 607 // CHECK7-NEXT: store i32 [[I]], i32* [[I_ADDR]], align 4 608 // CHECK7-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* @[[GLOB1]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined..1 to void (i32*, i32*, ...)*), i32* [[I_ADDR]]) 609 // CHECK7-NEXT: ret void 610 // 611 // 612 // CHECK7-LABEL: define {{[^@]+}}@.omp_outlined..1 613 // CHECK7-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[I:%.*]]) #[[ATTR1]] { 614 // CHECK7-NEXT: entry: 615 // CHECK7-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 616 // CHECK7-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 617 // CHECK7-NEXT: [[I_ADDR:%.*]] = alloca i32*, align 4 618 // CHECK7-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 619 // CHECK7-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 620 // CHECK7-NEXT: store i32* [[I]], i32** [[I_ADDR]], align 4 621 // CHECK7-NEXT: [[TMP0:%.*]] = load i32*, i32** [[I_ADDR]], align 4 622 // CHECK7-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 623 // CHECK7-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 624 // CHECK7-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4 625 // CHECK7-NEXT: ret void 626 // 627 // 628 // CHECK7-LABEL: define {{[^@]+}}@.omp_offloading.requires_reg 629 // CHECK7-SAME: () #[[ATTR4:[0-9]+]] { 630 // CHECK7-NEXT: entry: 631 // CHECK7-NEXT: call void @__tgt_register_requires(i64 1) 632 // CHECK7-NEXT: ret void 633 // 634 // 635 // CHECK8-LABEL: define {{[^@]+}}@_Z28implicit_maps_nested_integeri 636 // CHECK8-SAME: (i32 [[A:%.*]]) #[[ATTR0:[0-9]+]] { 637 // CHECK8-NEXT: entry: 638 // CHECK8-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4 639 // CHECK8-NEXT: [[I:%.*]] = alloca i32, align 4 640 // CHECK8-NEXT: store i32 [[A]], i32* [[A_ADDR]], align 4 641 // CHECK8-NEXT: [[TMP0:%.*]] = load i32, i32* [[A_ADDR]], align 4 642 // CHECK8-NEXT: store i32 [[TMP0]], i32* [[I]], align 4 643 // CHECK8-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* @[[GLOB1:[0-9]+]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined. to void (i32*, i32*, ...)*), i32* [[I]]) 644 // CHECK8-NEXT: ret void 645 // 646 // 647 // CHECK8-LABEL: define {{[^@]+}}@.omp_outlined. 648 // CHECK8-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[I:%.*]]) #[[ATTR1:[0-9]+]] { 649 // CHECK8-NEXT: entry: 650 // CHECK8-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 651 // CHECK8-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 652 // CHECK8-NEXT: [[I_ADDR:%.*]] = alloca i32*, align 4 653 // CHECK8-NEXT: [[I_CASTED:%.*]] = alloca i32, align 4 654 // CHECK8-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x i8*], align 4 655 // CHECK8-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x i8*], align 4 656 // CHECK8-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x i8*], align 4 657 // CHECK8-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 658 // CHECK8-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 659 // CHECK8-NEXT: store i32* [[I]], i32** [[I_ADDR]], align 4 660 // CHECK8-NEXT: [[TMP0:%.*]] = load i32*, i32** [[I_ADDR]], align 4 661 // CHECK8-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 662 // CHECK8-NEXT: store i32 [[TMP1]], i32* [[I_CASTED]], align 4 663 // CHECK8-NEXT: [[TMP2:%.*]] = load i32, i32* [[I_CASTED]], align 4 664 // CHECK8-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 665 // CHECK8-NEXT: [[TMP4:%.*]] = bitcast i8** [[TMP3]] to i32* 666 // CHECK8-NEXT: store i32 [[TMP2]], i32* [[TMP4]], align 4 667 // CHECK8-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 668 // CHECK8-NEXT: [[TMP6:%.*]] = bitcast i8** [[TMP5]] to i32* 669 // CHECK8-NEXT: store i32 [[TMP2]], i32* [[TMP6]], align 4 670 // CHECK8-NEXT: [[TMP7:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0 671 // CHECK8-NEXT: store i8* null, i8** [[TMP7]], align 4 672 // CHECK8-NEXT: [[TMP8:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 673 // CHECK8-NEXT: [[TMP9:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 674 // CHECK8-NEXT: [[TMP10:%.*]] = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z28implicit_maps_nested_integeri_l48.region_id, i32 1, i8** [[TMP8]], i8** [[TMP9]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null, i32 1, i32 0) 675 // CHECK8-NEXT: [[TMP11:%.*]] = icmp ne i32 [[TMP10]], 0 676 // CHECK8-NEXT: br i1 [[TMP11]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] 677 // CHECK8: omp_offload.failed: 678 // CHECK8-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z28implicit_maps_nested_integeri_l48(i32 [[TMP2]]) #[[ATTR3:[0-9]+]] 679 // CHECK8-NEXT: br label [[OMP_OFFLOAD_CONT]] 680 // CHECK8: omp_offload.cont: 681 // CHECK8-NEXT: ret void 682 // 683 // 684 // CHECK8-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z28implicit_maps_nested_integeri_l48 685 // CHECK8-SAME: (i32 [[I:%.*]]) #[[ATTR2:[0-9]+]] { 686 // CHECK8-NEXT: entry: 687 // CHECK8-NEXT: [[I_ADDR:%.*]] = alloca i32, align 4 688 // CHECK8-NEXT: store i32 [[I]], i32* [[I_ADDR]], align 4 689 // CHECK8-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* @[[GLOB1]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined..1 to void (i32*, i32*, ...)*), i32* [[I_ADDR]]) 690 // CHECK8-NEXT: ret void 691 // 692 // 693 // CHECK8-LABEL: define {{[^@]+}}@.omp_outlined..1 694 // CHECK8-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[I:%.*]]) #[[ATTR1]] { 695 // CHECK8-NEXT: entry: 696 // CHECK8-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 697 // CHECK8-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 698 // CHECK8-NEXT: [[I_ADDR:%.*]] = alloca i32*, align 4 699 // CHECK8-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 700 // CHECK8-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 701 // CHECK8-NEXT: store i32* [[I]], i32** [[I_ADDR]], align 4 702 // CHECK8-NEXT: [[TMP0:%.*]] = load i32*, i32** [[I_ADDR]], align 4 703 // CHECK8-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 704 // CHECK8-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 705 // CHECK8-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4 706 // CHECK8-NEXT: ret void 707 // 708 // 709 // CHECK8-LABEL: define {{[^@]+}}@.omp_offloading.requires_reg 710 // CHECK8-SAME: () #[[ATTR4:[0-9]+]] { 711 // CHECK8-NEXT: entry: 712 // CHECK8-NEXT: call void @__tgt_register_requires(i64 1) 713 // CHECK8-NEXT: ret void 714 // 715 // 716 // CHECK9-LABEL: define {{[^@]+}}@_Z28implicit_maps_nested_integeri 717 // CHECK9-SAME: (i32 signext [[A:%.*]]) #[[ATTR0:[0-9]+]] { 718 // CHECK9-NEXT: entry: 719 // CHECK9-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4 720 // CHECK9-NEXT: [[I:%.*]] = alloca i32, align 4 721 // CHECK9-NEXT: store i32 [[A]], i32* [[A_ADDR]], align 4 722 // CHECK9-NEXT: [[TMP0:%.*]] = load i32, i32* [[A_ADDR]], align 4 723 // CHECK9-NEXT: store i32 [[TMP0]], i32* [[I]], align 4 724 // CHECK9-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* @[[GLOB1:[0-9]+]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined. to void (i32*, i32*, ...)*), i32* [[I]]) 725 // CHECK9-NEXT: ret void 726 // 727 // 728 // CHECK9-LABEL: define {{[^@]+}}@.omp_outlined. 729 // CHECK9-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[I:%.*]]) #[[ATTR1:[0-9]+]] { 730 // CHECK9-NEXT: entry: 731 // CHECK9-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 732 // CHECK9-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 733 // CHECK9-NEXT: [[I_ADDR:%.*]] = alloca i32*, align 8 734 // CHECK9-NEXT: [[I_CASTED:%.*]] = alloca i64, align 8 735 // CHECK9-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x i8*], align 8 736 // CHECK9-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x i8*], align 8 737 // CHECK9-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x i8*], align 8 738 // CHECK9-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 739 // CHECK9-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 740 // CHECK9-NEXT: store i32* [[I]], i32** [[I_ADDR]], align 8 741 // CHECK9-NEXT: [[TMP0:%.*]] = load i32*, i32** [[I_ADDR]], align 8 742 // CHECK9-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 743 // CHECK9-NEXT: [[CONV:%.*]] = bitcast i64* [[I_CASTED]] to i32* 744 // CHECK9-NEXT: store i32 [[TMP1]], i32* [[CONV]], align 4 745 // CHECK9-NEXT: [[TMP2:%.*]] = load i64, i64* [[I_CASTED]], align 8 746 // CHECK9-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 747 // CHECK9-NEXT: [[TMP4:%.*]] = bitcast i8** [[TMP3]] to i64* 748 // CHECK9-NEXT: store i64 [[TMP2]], i64* [[TMP4]], align 8 749 // CHECK9-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 750 // CHECK9-NEXT: [[TMP6:%.*]] = bitcast i8** [[TMP5]] to i64* 751 // CHECK9-NEXT: store i64 [[TMP2]], i64* [[TMP6]], align 8 752 // CHECK9-NEXT: [[TMP7:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0 753 // CHECK9-NEXT: store i8* null, i8** [[TMP7]], align 8 754 // CHECK9-NEXT: [[TMP8:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 755 // CHECK9-NEXT: [[TMP9:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 756 // CHECK9-NEXT: [[TMP10:%.*]] = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z28implicit_maps_nested_integeri_l48.region_id, i32 1, i8** [[TMP8]], i8** [[TMP9]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null, i32 1, i32 0) 757 // CHECK9-NEXT: [[TMP11:%.*]] = icmp ne i32 [[TMP10]], 0 758 // CHECK9-NEXT: br i1 [[TMP11]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] 759 // CHECK9: omp_offload.failed: 760 // CHECK9-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z28implicit_maps_nested_integeri_l48(i64 [[TMP2]]) #[[ATTR3:[0-9]+]] 761 // CHECK9-NEXT: br label [[OMP_OFFLOAD_CONT]] 762 // CHECK9: omp_offload.cont: 763 // CHECK9-NEXT: ret void 764 // 765 // 766 // CHECK9-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z28implicit_maps_nested_integeri_l48 767 // CHECK9-SAME: (i64 [[I:%.*]]) #[[ATTR2:[0-9]+]] { 768 // CHECK9-NEXT: entry: 769 // CHECK9-NEXT: [[I_ADDR:%.*]] = alloca i64, align 8 770 // CHECK9-NEXT: store i64 [[I]], i64* [[I_ADDR]], align 8 771 // CHECK9-NEXT: [[CONV:%.*]] = bitcast i64* [[I_ADDR]] to i32* 772 // CHECK9-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* @[[GLOB1]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined..1 to void (i32*, i32*, ...)*), i32* [[CONV]]) 773 // CHECK9-NEXT: ret void 774 // 775 // 776 // CHECK9-LABEL: define {{[^@]+}}@.omp_outlined..1 777 // CHECK9-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[I:%.*]]) #[[ATTR1]] { 778 // CHECK9-NEXT: entry: 779 // CHECK9-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 780 // CHECK9-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 781 // CHECK9-NEXT: [[I_ADDR:%.*]] = alloca i32*, align 8 782 // CHECK9-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 783 // CHECK9-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 784 // CHECK9-NEXT: store i32* [[I]], i32** [[I_ADDR]], align 8 785 // CHECK9-NEXT: [[TMP0:%.*]] = load i32*, i32** [[I_ADDR]], align 8 786 // CHECK9-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 787 // CHECK9-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 788 // CHECK9-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4 789 // CHECK9-NEXT: ret void 790 // 791 // 792 // CHECK9-LABEL: define {{[^@]+}}@.omp_offloading.requires_reg 793 // CHECK9-SAME: () #[[ATTR4:[0-9]+]] { 794 // CHECK9-NEXT: entry: 795 // CHECK9-NEXT: call void @__tgt_register_requires(i64 1) 796 // CHECK9-NEXT: ret void 797 // 798 // 799 // CHECK10-LABEL: define {{[^@]+}}@_Z28implicit_maps_nested_integeri 800 // CHECK10-SAME: (i32 signext [[A:%.*]]) #[[ATTR0:[0-9]+]] { 801 // CHECK10-NEXT: entry: 802 // CHECK10-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4 803 // CHECK10-NEXT: [[I:%.*]] = alloca i32, align 4 804 // CHECK10-NEXT: store i32 [[A]], i32* [[A_ADDR]], align 4 805 // CHECK10-NEXT: [[TMP0:%.*]] = load i32, i32* [[A_ADDR]], align 4 806 // CHECK10-NEXT: store i32 [[TMP0]], i32* [[I]], align 4 807 // CHECK10-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* @[[GLOB1:[0-9]+]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined. to void (i32*, i32*, ...)*), i32* [[I]]) 808 // CHECK10-NEXT: ret void 809 // 810 // 811 // CHECK10-LABEL: define {{[^@]+}}@.omp_outlined. 812 // CHECK10-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[I:%.*]]) #[[ATTR1:[0-9]+]] { 813 // CHECK10-NEXT: entry: 814 // CHECK10-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 815 // CHECK10-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 816 // CHECK10-NEXT: [[I_ADDR:%.*]] = alloca i32*, align 8 817 // CHECK10-NEXT: [[I_CASTED:%.*]] = alloca i64, align 8 818 // CHECK10-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x i8*], align 8 819 // CHECK10-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x i8*], align 8 820 // CHECK10-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x i8*], align 8 821 // CHECK10-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 822 // CHECK10-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 823 // CHECK10-NEXT: store i32* [[I]], i32** [[I_ADDR]], align 8 824 // CHECK10-NEXT: [[TMP0:%.*]] = load i32*, i32** [[I_ADDR]], align 8 825 // CHECK10-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 826 // CHECK10-NEXT: [[CONV:%.*]] = bitcast i64* [[I_CASTED]] to i32* 827 // CHECK10-NEXT: store i32 [[TMP1]], i32* [[CONV]], align 4 828 // CHECK10-NEXT: [[TMP2:%.*]] = load i64, i64* [[I_CASTED]], align 8 829 // CHECK10-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 830 // CHECK10-NEXT: [[TMP4:%.*]] = bitcast i8** [[TMP3]] to i64* 831 // CHECK10-NEXT: store i64 [[TMP2]], i64* [[TMP4]], align 8 832 // CHECK10-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 833 // CHECK10-NEXT: [[TMP6:%.*]] = bitcast i8** [[TMP5]] to i64* 834 // CHECK10-NEXT: store i64 [[TMP2]], i64* [[TMP6]], align 8 835 // CHECK10-NEXT: [[TMP7:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0 836 // CHECK10-NEXT: store i8* null, i8** [[TMP7]], align 8 837 // CHECK10-NEXT: [[TMP8:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 838 // CHECK10-NEXT: [[TMP9:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 839 // CHECK10-NEXT: [[TMP10:%.*]] = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z28implicit_maps_nested_integeri_l48.region_id, i32 1, i8** [[TMP8]], i8** [[TMP9]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null, i32 1, i32 0) 840 // CHECK10-NEXT: [[TMP11:%.*]] = icmp ne i32 [[TMP10]], 0 841 // CHECK10-NEXT: br i1 [[TMP11]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] 842 // CHECK10: omp_offload.failed: 843 // CHECK10-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z28implicit_maps_nested_integeri_l48(i64 [[TMP2]]) #[[ATTR3:[0-9]+]] 844 // CHECK10-NEXT: br label [[OMP_OFFLOAD_CONT]] 845 // CHECK10: omp_offload.cont: 846 // CHECK10-NEXT: ret void 847 // 848 // 849 // CHECK10-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z28implicit_maps_nested_integeri_l48 850 // CHECK10-SAME: (i64 [[I:%.*]]) #[[ATTR2:[0-9]+]] { 851 // CHECK10-NEXT: entry: 852 // CHECK10-NEXT: [[I_ADDR:%.*]] = alloca i64, align 8 853 // CHECK10-NEXT: store i64 [[I]], i64* [[I_ADDR]], align 8 854 // CHECK10-NEXT: [[CONV:%.*]] = bitcast i64* [[I_ADDR]] to i32* 855 // CHECK10-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* @[[GLOB1]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined..1 to void (i32*, i32*, ...)*), i32* [[CONV]]) 856 // CHECK10-NEXT: ret void 857 // 858 // 859 // CHECK10-LABEL: define {{[^@]+}}@.omp_outlined..1 860 // CHECK10-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[I:%.*]]) #[[ATTR1]] { 861 // CHECK10-NEXT: entry: 862 // CHECK10-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 863 // CHECK10-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 864 // CHECK10-NEXT: [[I_ADDR:%.*]] = alloca i32*, align 8 865 // CHECK10-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 866 // CHECK10-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 867 // CHECK10-NEXT: store i32* [[I]], i32** [[I_ADDR]], align 8 868 // CHECK10-NEXT: [[TMP0:%.*]] = load i32*, i32** [[I_ADDR]], align 8 869 // CHECK10-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 870 // CHECK10-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 871 // CHECK10-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4 872 // CHECK10-NEXT: ret void 873 // 874 // 875 // CHECK10-LABEL: define {{[^@]+}}@.omp_offloading.requires_reg 876 // CHECK10-SAME: () #[[ATTR4:[0-9]+]] { 877 // CHECK10-NEXT: entry: 878 // CHECK10-NEXT: call void @__tgt_register_requires(i64 1) 879 // CHECK10-NEXT: ret void 880 // 881 // 882 // CHECK11-LABEL: define {{[^@]+}}@_Z28implicit_maps_nested_integeri 883 // CHECK11-SAME: (i32 [[A:%.*]]) #[[ATTR0:[0-9]+]] { 884 // CHECK11-NEXT: entry: 885 // CHECK11-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4 886 // CHECK11-NEXT: [[I:%.*]] = alloca i32, align 4 887 // CHECK11-NEXT: store i32 [[A]], i32* [[A_ADDR]], align 4 888 // CHECK11-NEXT: [[TMP0:%.*]] = load i32, i32* [[A_ADDR]], align 4 889 // CHECK11-NEXT: store i32 [[TMP0]], i32* [[I]], align 4 890 // CHECK11-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* @[[GLOB1:[0-9]+]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined. to void (i32*, i32*, ...)*), i32* [[I]]) 891 // CHECK11-NEXT: ret void 892 // 893 // 894 // CHECK11-LABEL: define {{[^@]+}}@.omp_outlined. 895 // CHECK11-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[I:%.*]]) #[[ATTR1:[0-9]+]] { 896 // CHECK11-NEXT: entry: 897 // CHECK11-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 898 // CHECK11-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 899 // CHECK11-NEXT: [[I_ADDR:%.*]] = alloca i32*, align 4 900 // CHECK11-NEXT: [[I_CASTED:%.*]] = alloca i32, align 4 901 // CHECK11-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x i8*], align 4 902 // CHECK11-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x i8*], align 4 903 // CHECK11-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x i8*], align 4 904 // CHECK11-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 905 // CHECK11-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 906 // CHECK11-NEXT: store i32* [[I]], i32** [[I_ADDR]], align 4 907 // CHECK11-NEXT: [[TMP0:%.*]] = load i32*, i32** [[I_ADDR]], align 4 908 // CHECK11-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 909 // CHECK11-NEXT: store i32 [[TMP1]], i32* [[I_CASTED]], align 4 910 // CHECK11-NEXT: [[TMP2:%.*]] = load i32, i32* [[I_CASTED]], align 4 911 // CHECK11-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 912 // CHECK11-NEXT: [[TMP4:%.*]] = bitcast i8** [[TMP3]] to i32* 913 // CHECK11-NEXT: store i32 [[TMP2]], i32* [[TMP4]], align 4 914 // CHECK11-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 915 // CHECK11-NEXT: [[TMP6:%.*]] = bitcast i8** [[TMP5]] to i32* 916 // CHECK11-NEXT: store i32 [[TMP2]], i32* [[TMP6]], align 4 917 // CHECK11-NEXT: [[TMP7:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0 918 // CHECK11-NEXT: store i8* null, i8** [[TMP7]], align 4 919 // CHECK11-NEXT: [[TMP8:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 920 // CHECK11-NEXT: [[TMP9:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 921 // CHECK11-NEXT: [[TMP10:%.*]] = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z28implicit_maps_nested_integeri_l48.region_id, i32 1, i8** [[TMP8]], i8** [[TMP9]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null, i32 1, i32 0) 922 // CHECK11-NEXT: [[TMP11:%.*]] = icmp ne i32 [[TMP10]], 0 923 // CHECK11-NEXT: br i1 [[TMP11]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] 924 // CHECK11: omp_offload.failed: 925 // CHECK11-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z28implicit_maps_nested_integeri_l48(i32 [[TMP2]]) #[[ATTR3:[0-9]+]] 926 // CHECK11-NEXT: br label [[OMP_OFFLOAD_CONT]] 927 // CHECK11: omp_offload.cont: 928 // CHECK11-NEXT: ret void 929 // 930 // 931 // CHECK11-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z28implicit_maps_nested_integeri_l48 932 // CHECK11-SAME: (i32 [[I:%.*]]) #[[ATTR2:[0-9]+]] { 933 // CHECK11-NEXT: entry: 934 // CHECK11-NEXT: [[I_ADDR:%.*]] = alloca i32, align 4 935 // CHECK11-NEXT: store i32 [[I]], i32* [[I_ADDR]], align 4 936 // CHECK11-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* @[[GLOB1]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined..1 to void (i32*, i32*, ...)*), i32* [[I_ADDR]]) 937 // CHECK11-NEXT: ret void 938 // 939 // 940 // CHECK11-LABEL: define {{[^@]+}}@.omp_outlined..1 941 // CHECK11-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[I:%.*]]) #[[ATTR1]] { 942 // CHECK11-NEXT: entry: 943 // CHECK11-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 944 // CHECK11-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 945 // CHECK11-NEXT: [[I_ADDR:%.*]] = alloca i32*, align 4 946 // CHECK11-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 947 // CHECK11-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 948 // CHECK11-NEXT: store i32* [[I]], i32** [[I_ADDR]], align 4 949 // CHECK11-NEXT: [[TMP0:%.*]] = load i32*, i32** [[I_ADDR]], align 4 950 // CHECK11-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 951 // CHECK11-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 952 // CHECK11-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4 953 // CHECK11-NEXT: ret void 954 // 955 // 956 // CHECK11-LABEL: define {{[^@]+}}@.omp_offloading.requires_reg 957 // CHECK11-SAME: () #[[ATTR4:[0-9]+]] { 958 // CHECK11-NEXT: entry: 959 // CHECK11-NEXT: call void @__tgt_register_requires(i64 1) 960 // CHECK11-NEXT: ret void 961 // 962 // 963 // CHECK12-LABEL: define {{[^@]+}}@_Z28implicit_maps_nested_integeri 964 // CHECK12-SAME: (i32 [[A:%.*]]) #[[ATTR0:[0-9]+]] { 965 // CHECK12-NEXT: entry: 966 // CHECK12-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4 967 // CHECK12-NEXT: [[I:%.*]] = alloca i32, align 4 968 // CHECK12-NEXT: store i32 [[A]], i32* [[A_ADDR]], align 4 969 // CHECK12-NEXT: [[TMP0:%.*]] = load i32, i32* [[A_ADDR]], align 4 970 // CHECK12-NEXT: store i32 [[TMP0]], i32* [[I]], align 4 971 // CHECK12-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* @[[GLOB1:[0-9]+]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined. to void (i32*, i32*, ...)*), i32* [[I]]) 972 // CHECK12-NEXT: ret void 973 // 974 // 975 // CHECK12-LABEL: define {{[^@]+}}@.omp_outlined. 976 // CHECK12-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[I:%.*]]) #[[ATTR1:[0-9]+]] { 977 // CHECK12-NEXT: entry: 978 // CHECK12-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 979 // CHECK12-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 980 // CHECK12-NEXT: [[I_ADDR:%.*]] = alloca i32*, align 4 981 // CHECK12-NEXT: [[I_CASTED:%.*]] = alloca i32, align 4 982 // CHECK12-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x i8*], align 4 983 // CHECK12-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x i8*], align 4 984 // CHECK12-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x i8*], align 4 985 // CHECK12-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 986 // CHECK12-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 987 // CHECK12-NEXT: store i32* [[I]], i32** [[I_ADDR]], align 4 988 // CHECK12-NEXT: [[TMP0:%.*]] = load i32*, i32** [[I_ADDR]], align 4 989 // CHECK12-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 990 // CHECK12-NEXT: store i32 [[TMP1]], i32* [[I_CASTED]], align 4 991 // CHECK12-NEXT: [[TMP2:%.*]] = load i32, i32* [[I_CASTED]], align 4 992 // CHECK12-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 993 // CHECK12-NEXT: [[TMP4:%.*]] = bitcast i8** [[TMP3]] to i32* 994 // CHECK12-NEXT: store i32 [[TMP2]], i32* [[TMP4]], align 4 995 // CHECK12-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 996 // CHECK12-NEXT: [[TMP6:%.*]] = bitcast i8** [[TMP5]] to i32* 997 // CHECK12-NEXT: store i32 [[TMP2]], i32* [[TMP6]], align 4 998 // CHECK12-NEXT: [[TMP7:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_MAPPERS]], i32 0, i32 0 999 // CHECK12-NEXT: store i8* null, i8** [[TMP7]], align 4 1000 // CHECK12-NEXT: [[TMP8:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 1001 // CHECK12-NEXT: [[TMP9:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_PTRS]], i32 0, i32 0 1002 // CHECK12-NEXT: [[TMP10:%.*]] = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @[[GLOB1]], i64 -1, i8* @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z28implicit_maps_nested_integeri_l48.region_id, i32 1, i8** [[TMP8]], i8** [[TMP9]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null, i32 1, i32 0) 1003 // CHECK12-NEXT: [[TMP11:%.*]] = icmp ne i32 [[TMP10]], 0 1004 // CHECK12-NEXT: br i1 [[TMP11]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] 1005 // CHECK12: omp_offload.failed: 1006 // CHECK12-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z28implicit_maps_nested_integeri_l48(i32 [[TMP2]]) #[[ATTR3:[0-9]+]] 1007 // CHECK12-NEXT: br label [[OMP_OFFLOAD_CONT]] 1008 // CHECK12: omp_offload.cont: 1009 // CHECK12-NEXT: ret void 1010 // 1011 // 1012 // CHECK12-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z28implicit_maps_nested_integeri_l48 1013 // CHECK12-SAME: (i32 [[I:%.*]]) #[[ATTR2:[0-9]+]] { 1014 // CHECK12-NEXT: entry: 1015 // CHECK12-NEXT: [[I_ADDR:%.*]] = alloca i32, align 4 1016 // CHECK12-NEXT: store i32 [[I]], i32* [[I_ADDR]], align 4 1017 // CHECK12-NEXT: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* @[[GLOB1]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined..1 to void (i32*, i32*, ...)*), i32* [[I_ADDR]]) 1018 // CHECK12-NEXT: ret void 1019 // 1020 // 1021 // CHECK12-LABEL: define {{[^@]+}}@.omp_outlined..1 1022 // CHECK12-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]], i32* nonnull align 4 dereferenceable(4) [[I:%.*]]) #[[ATTR1]] { 1023 // CHECK12-NEXT: entry: 1024 // CHECK12-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 1025 // CHECK12-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 1026 // CHECK12-NEXT: [[I_ADDR:%.*]] = alloca i32*, align 4 1027 // CHECK12-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 1028 // CHECK12-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 1029 // CHECK12-NEXT: store i32* [[I]], i32** [[I_ADDR]], align 4 1030 // CHECK12-NEXT: [[TMP0:%.*]] = load i32*, i32** [[I_ADDR]], align 4 1031 // CHECK12-NEXT: [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4 1032 // CHECK12-NEXT: [[INC:%.*]] = add nsw i32 [[TMP1]], 1 1033 // CHECK12-NEXT: store i32 [[INC]], i32* [[TMP0]], align 4 1034 // CHECK12-NEXT: ret void 1035 // 1036 // 1037 // CHECK12-LABEL: define {{[^@]+}}@.omp_offloading.requires_reg 1038 // CHECK12-SAME: () #[[ATTR4:[0-9]+]] { 1039 // CHECK12-NEXT: entry: 1040 // CHECK12-NEXT: call void @__tgt_register_requires(i64 1) 1041 // CHECK12-NEXT: ret void 1042 // 1043