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 // Test target codegen - host bc file has to be created first. 3 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc 4 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -disable-llvm-optzns | FileCheck %s --check-prefix=CHECK1 5 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc 6 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - -disable-llvm-optzns | FileCheck %s --check-prefix=CHECK2 7 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -fexceptions -fcxx-exceptions -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - -disable-llvm-optzns | FileCheck %s --check-prefix=CHECK3 8 9 // expected-no-diagnostics 10 #ifndef HEADER 11 #define HEADER 12 13 void work(); 14 15 void use() { 16 #pragma omp parallel 17 work(); 18 } 19 20 int main() { 21 #pragma omp target parallel 22 { use(); } 23 #pragma omp target 24 { use(); } 25 } 26 27 #endif 28 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l21 29 // CHECK1-SAME: () #[[ATTR0:[0-9]+]] { 30 // CHECK1-NEXT: entry: 31 // CHECK1-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8 32 // CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1:[0-9]+]], i1 true, i1 false, i1 true) 33 // CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1 34 // CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] 35 // CHECK1: user_code.entry: 36 // CHECK1-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB2:[0-9]+]]) 37 // CHECK1-NEXT: [[TMP2:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** 38 // CHECK1-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB2]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*)* @__omp_outlined__ to i8*), i8* null, i8** [[TMP2]], i64 0) 39 // CHECK1-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i1 true, i1 true) 40 // CHECK1-NEXT: ret void 41 // CHECK1: worker.exit: 42 // CHECK1-NEXT: ret void 43 // 44 // 45 // CHECK1-LABEL: define {{[^@]+}}@__omp_outlined__ 46 // CHECK1-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] { 47 // CHECK1-NEXT: entry: 48 // CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 49 // CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 50 // CHECK1-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 51 // CHECK1-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 52 // CHECK1-NEXT: call void @_Z3usev() #[[ATTR5:[0-9]+]] 53 // CHECK1-NEXT: ret void 54 // 55 // 56 // CHECK1-LABEL: define {{[^@]+}}@_Z3usev 57 // CHECK1-SAME: () #[[ATTR1:[0-9]+]] { 58 // CHECK1-NEXT: entry: 59 // CHECK1-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8 60 // CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB3:[0-9]+]]) 61 // CHECK1-NEXT: [[TMP1:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** 62 // CHECK1-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB3]], i32 [[TMP0]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*)* @__omp_outlined__1 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__1_wrapper to i8*), i8** [[TMP1]], i64 0) 63 // CHECK1-NEXT: ret void 64 // 65 // 66 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l23 67 // CHECK1-SAME: () #[[ATTR0]] { 68 // CHECK1-NEXT: entry: 69 // CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i1 false, i1 true, i1 true) 70 // CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1 71 // CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] 72 // CHECK1: user_code.entry: 73 // CHECK1-NEXT: call void @_Z3usev() #[[ATTR5]] 74 // CHECK1-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i1 false, i1 true) 75 // CHECK1-NEXT: ret void 76 // CHECK1: worker.exit: 77 // CHECK1-NEXT: ret void 78 // 79 // 80 // CHECK1-LABEL: define {{[^@]+}}@__omp_outlined__1 81 // CHECK1-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] { 82 // CHECK1-NEXT: entry: 83 // CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 84 // CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 85 // CHECK1-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 86 // CHECK1-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 87 // CHECK1-NEXT: call void @_Z4workv() #[[ATTR5]] 88 // CHECK1-NEXT: ret void 89 // 90 // 91 // CHECK1-LABEL: define {{[^@]+}}@__omp_outlined__1_wrapper 92 // CHECK1-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR4:[0-9]+]] { 93 // CHECK1-NEXT: entry: 94 // CHECK1-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2 95 // CHECK1-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4 96 // CHECK1-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 97 // CHECK1-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 8 98 // CHECK1-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 99 // CHECK1-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2 100 // CHECK1-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4 101 // CHECK1-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]]) 102 // CHECK1-NEXT: call void @__omp_outlined__1(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]]) #[[ATTR2:[0-9]+]] 103 // CHECK1-NEXT: ret void 104 // 105 // 106 // CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l21 107 // CHECK2-SAME: () #[[ATTR0:[0-9]+]] { 108 // CHECK2-NEXT: entry: 109 // CHECK2-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 4 110 // CHECK2-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1:[0-9]+]], i1 true, i1 false, i1 true) 111 // CHECK2-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1 112 // CHECK2-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] 113 // CHECK2: user_code.entry: 114 // CHECK2-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB2:[0-9]+]]) 115 // CHECK2-NEXT: [[TMP2:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** 116 // CHECK2-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB2]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*)* @__omp_outlined__ to i8*), i8* null, i8** [[TMP2]], i32 0) 117 // CHECK2-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i1 true, i1 true) 118 // CHECK2-NEXT: ret void 119 // CHECK2: worker.exit: 120 // CHECK2-NEXT: ret void 121 // 122 // 123 // CHECK2-LABEL: define {{[^@]+}}@__omp_outlined__ 124 // CHECK2-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] { 125 // CHECK2-NEXT: entry: 126 // CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 127 // CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 128 // CHECK2-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 129 // CHECK2-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 130 // CHECK2-NEXT: call void @_Z3usev() #[[ATTR5:[0-9]+]] 131 // CHECK2-NEXT: ret void 132 // 133 // 134 // CHECK2-LABEL: define {{[^@]+}}@_Z3usev 135 // CHECK2-SAME: () #[[ATTR1:[0-9]+]] { 136 // CHECK2-NEXT: entry: 137 // CHECK2-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 4 138 // CHECK2-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB3:[0-9]+]]) 139 // CHECK2-NEXT: [[TMP1:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** 140 // CHECK2-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB3]], i32 [[TMP0]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*)* @__omp_outlined__1 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__1_wrapper to i8*), i8** [[TMP1]], i32 0) 141 // CHECK2-NEXT: ret void 142 // 143 // 144 // CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l23 145 // CHECK2-SAME: () #[[ATTR0]] { 146 // CHECK2-NEXT: entry: 147 // CHECK2-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i1 false, i1 true, i1 true) 148 // CHECK2-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1 149 // CHECK2-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] 150 // CHECK2: user_code.entry: 151 // CHECK2-NEXT: call void @_Z3usev() #[[ATTR5]] 152 // CHECK2-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i1 false, i1 true) 153 // CHECK2-NEXT: ret void 154 // CHECK2: worker.exit: 155 // CHECK2-NEXT: ret void 156 // 157 // 158 // CHECK2-LABEL: define {{[^@]+}}@__omp_outlined__1 159 // CHECK2-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] { 160 // CHECK2-NEXT: entry: 161 // CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 162 // CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 163 // CHECK2-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 164 // CHECK2-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 165 // CHECK2-NEXT: call void @_Z4workv() #[[ATTR5]] 166 // CHECK2-NEXT: ret void 167 // 168 // 169 // CHECK2-LABEL: define {{[^@]+}}@__omp_outlined__1_wrapper 170 // CHECK2-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR4:[0-9]+]] { 171 // CHECK2-NEXT: entry: 172 // CHECK2-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2 173 // CHECK2-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4 174 // CHECK2-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 175 // CHECK2-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 4 176 // CHECK2-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 177 // CHECK2-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2 178 // CHECK2-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4 179 // CHECK2-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]]) 180 // CHECK2-NEXT: call void @__omp_outlined__1(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]]) #[[ATTR2:[0-9]+]] 181 // CHECK2-NEXT: ret void 182 // 183 // 184 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l21 185 // CHECK3-SAME: () #[[ATTR0:[0-9]+]] { 186 // CHECK3-NEXT: entry: 187 // CHECK3-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 4 188 // CHECK3-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1:[0-9]+]], i1 true, i1 false, i1 true) 189 // CHECK3-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1 190 // CHECK3-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] 191 // CHECK3: user_code.entry: 192 // CHECK3-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB2:[0-9]+]]) 193 // CHECK3-NEXT: [[TMP2:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** 194 // CHECK3-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB2]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*)* @__omp_outlined__ to i8*), i8* null, i8** [[TMP2]], i32 0) 195 // CHECK3-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i1 true, i1 true) 196 // CHECK3-NEXT: ret void 197 // CHECK3: worker.exit: 198 // CHECK3-NEXT: ret void 199 // 200 // 201 // CHECK3-LABEL: define {{[^@]+}}@__omp_outlined__ 202 // CHECK3-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] { 203 // CHECK3-NEXT: entry: 204 // CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 205 // CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 206 // CHECK3-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 207 // CHECK3-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 208 // CHECK3-NEXT: call void @_Z3usev() #[[ATTR5:[0-9]+]] 209 // CHECK3-NEXT: ret void 210 // 211 // 212 // CHECK3-LABEL: define {{[^@]+}}@_Z3usev 213 // CHECK3-SAME: () #[[ATTR1:[0-9]+]] { 214 // CHECK3-NEXT: entry: 215 // CHECK3-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 4 216 // CHECK3-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @[[GLOB3:[0-9]+]]) 217 // CHECK3-NEXT: [[TMP1:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** 218 // CHECK3-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* @[[GLOB3]], i32 [[TMP0]], i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*)* @__omp_outlined__1 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__1_wrapper to i8*), i8** [[TMP1]], i32 0) 219 // CHECK3-NEXT: ret void 220 // 221 // 222 // CHECK3-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l23 223 // CHECK3-SAME: () #[[ATTR0]] { 224 // CHECK3-NEXT: entry: 225 // CHECK3-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* @[[GLOB1]], i1 false, i1 true, i1 true) 226 // CHECK3-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1 227 // CHECK3-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] 228 // CHECK3: user_code.entry: 229 // CHECK3-NEXT: call void @_Z3usev() #[[ATTR5]] 230 // CHECK3-NEXT: call void @__kmpc_target_deinit(%struct.ident_t* @[[GLOB1]], i1 false, i1 true) 231 // CHECK3-NEXT: ret void 232 // CHECK3: worker.exit: 233 // CHECK3-NEXT: ret void 234 // 235 // 236 // CHECK3-LABEL: define {{[^@]+}}@__omp_outlined__1 237 // CHECK3-SAME: (i32* noalias [[DOTGLOBAL_TID_:%.*]], i32* noalias [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] { 238 // CHECK3-NEXT: entry: 239 // CHECK3-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 4 240 // CHECK3-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 4 241 // CHECK3-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 4 242 // CHECK3-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 4 243 // CHECK3-NEXT: call void @_Z4workv() #[[ATTR5]] 244 // CHECK3-NEXT: ret void 245 // 246 // 247 // CHECK3-LABEL: define {{[^@]+}}@__omp_outlined__1_wrapper 248 // CHECK3-SAME: (i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR4:[0-9]+]] { 249 // CHECK3-NEXT: entry: 250 // CHECK3-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2 251 // CHECK3-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4 252 // CHECK3-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 253 // CHECK3-NEXT: [[GLOBAL_ARGS:%.*]] = alloca i8**, align 4 254 // CHECK3-NEXT: store i32 0, i32* [[DOTZERO_ADDR]], align 4 255 // CHECK3-NEXT: store i16 [[TMP0]], i16* [[DOTADDR]], align 2 256 // CHECK3-NEXT: store i32 [[TMP1]], i32* [[DOTADDR1]], align 4 257 // CHECK3-NEXT: call void @__kmpc_get_shared_variables(i8*** [[GLOBAL_ARGS]]) 258 // CHECK3-NEXT: call void @__omp_outlined__1(i32* [[DOTADDR1]], i32* [[DOTZERO_ADDR]]) #[[ATTR2:[0-9]+]] 259 // CHECK3-NEXT: ret void 260 // 261