259 lines · cpp
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.bc4// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -disable-llvm-optzns | FileCheck %s --check-prefix=CHECK15// 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.bc6// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - -disable-llvm-optzns | FileCheck %s --check-prefix=CHECK27// 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-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - -disable-llvm-optzns | FileCheck %s --check-prefix=CHECK28 9// expected-no-diagnostics10#ifndef HEADER11#define HEADER12 13void work(int *C) {14 #pragma omp atomic15 ++(*C);16}17 18void use(int *C) {19 #pragma omp parallel num_threads(2)20 work(C);21}22 23int main() {24 int C = 0;25 #pragma omp target map(C)26 {27 use(&C);28 #pragma omp parallel num_threads(2)29 use(&C);30 }31 32 return C;33}34 35#endif36// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l2537// CHECK1-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[C:%.*]]) #[[ATTR0:[0-9]+]] {38// CHECK1-NEXT: entry:39// CHECK1-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 840// CHECK1-NEXT: [[C_ADDR:%.*]] = alloca ptr, align 841// CHECK1-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x ptr], align 842// CHECK1-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 843// CHECK1-NEXT: store ptr [[C]], ptr [[C_ADDR]], align 844// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[C_ADDR]], align 845// CHECK1-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l25_kernel_environment, ptr [[DYN_PTR]])46// CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -147// CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]48// CHECK1: user_code.entry:49// CHECK1-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1:[0-9]+]])50// CHECK1-NEXT: call void @_Z3usePi(ptr noundef [[TMP0]]) #[[ATTR6:[0-9]+]]51// CHECK1-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 052// CHECK1-NEXT: store ptr [[TMP0]], ptr [[TMP3]], align 853// CHECK1-NEXT: call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP2]], i32 1, i32 2, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l25_omp_outlined, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l25_omp_outlined_wrapper, ptr [[CAPTURED_VARS_ADDRS]], i64 1)54// CHECK1-NEXT: call void @__kmpc_target_deinit()55// CHECK1-NEXT: ret void56// CHECK1: worker.exit:57// CHECK1-NEXT: ret void58//59//60// CHECK1-LABEL: define {{[^@]+}}@_Z3usePi61// CHECK1-SAME: (ptr noundef [[C:%.*]]) #[[ATTR1:[0-9]+]] {62// CHECK1-NEXT: entry:63// CHECK1-NEXT: [[C_ADDR:%.*]] = alloca ptr, align 864// CHECK1-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x ptr], align 865// CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])66// CHECK1-NEXT: store ptr [[C]], ptr [[C_ADDR]], align 867// CHECK1-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 068// CHECK1-NEXT: store ptr [[C_ADDR]], ptr [[TMP1]], align 869// CHECK1-NEXT: call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP0]], i32 1, i32 2, i32 -1, ptr @_Z3usePi_omp_outlined, ptr @_Z3usePi_omp_outlined_wrapper, ptr [[CAPTURED_VARS_ADDRS]], i64 1)70// CHECK1-NEXT: ret void71//72//73// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l25_omp_outlined74// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[C:%.*]]) #[[ATTR2:[0-9]+]] {75// CHECK1-NEXT: entry:76// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 877// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 878// CHECK1-NEXT: [[C_ADDR:%.*]] = alloca ptr, align 879// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 880// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 881// CHECK1-NEXT: store ptr [[C]], ptr [[C_ADDR]], align 882// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[C_ADDR]], align 883// CHECK1-NEXT: call void @_Z3usePi(ptr noundef [[TMP0]]) #[[ATTR6]]84// CHECK1-NEXT: ret void85//86//87// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l25_omp_outlined_wrapper88// CHECK1-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR3:[0-9]+]] {89// CHECK1-NEXT: entry:90// CHECK1-NEXT: [[DOTADDR:%.*]] = alloca i16, align 291// CHECK1-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 492// CHECK1-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 493// CHECK1-NEXT: [[GLOBAL_ARGS:%.*]] = alloca ptr, align 894// CHECK1-NEXT: store i16 [[TMP0]], ptr [[DOTADDR]], align 295// CHECK1-NEXT: store i32 [[TMP1]], ptr [[DOTADDR1]], align 496// CHECK1-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 497// CHECK1-NEXT: call void @__kmpc_get_shared_variables(ptr [[GLOBAL_ARGS]])98// CHECK1-NEXT: [[TMP2:%.*]] = load ptr, ptr [[GLOBAL_ARGS]], align 899// CHECK1-NEXT: [[TMP3:%.*]] = getelementptr inbounds ptr, ptr [[TMP2]], i64 0100// CHECK1-NEXT: [[TMP4:%.*]] = load ptr, ptr [[TMP3]], align 8101// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l25_omp_outlined(ptr [[DOTADDR1]], ptr [[DOTZERO_ADDR]], ptr [[TMP4]]) #[[ATTR4:[0-9]+]]102// CHECK1-NEXT: ret void103//104//105// CHECK1-LABEL: define {{[^@]+}}@_Z3usePi_omp_outlined106// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[C:%.*]]) #[[ATTR2]] {107// CHECK1-NEXT: entry:108// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8109// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8110// CHECK1-NEXT: [[C_ADDR:%.*]] = alloca ptr, align 8111// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8112// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8113// CHECK1-NEXT: store ptr [[C]], ptr [[C_ADDR]], align 8114// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[C_ADDR]], align 8115// CHECK1-NEXT: [[TMP1:%.*]] = load ptr, ptr [[TMP0]], align 8116// CHECK1-NEXT: call void @_Z4workPi(ptr noundef [[TMP1]]) #[[ATTR6]]117// CHECK1-NEXT: ret void118//119//120// CHECK1-LABEL: define {{[^@]+}}@_Z4workPi121// CHECK1-SAME: (ptr noundef [[C:%.*]]) #[[ATTR1]] {122// CHECK1-NEXT: entry:123// CHECK1-NEXT: [[C_ADDR:%.*]] = alloca ptr, align 8124// CHECK1-NEXT: store ptr [[C]], ptr [[C_ADDR]], align 8125// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[C_ADDR]], align 8126// CHECK1-NEXT: [[TMP1:%.*]] = atomicrmw add ptr [[TMP0]], i32 1 monotonic, align 4127// CHECK1-NEXT: ret void128//129//130// CHECK1-LABEL: define {{[^@]+}}@_Z3usePi_omp_outlined_wrapper131// CHECK1-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR3]] {132// CHECK1-NEXT: entry:133// CHECK1-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2134// CHECK1-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4135// CHECK1-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4136// CHECK1-NEXT: [[GLOBAL_ARGS:%.*]] = alloca ptr, align 8137// CHECK1-NEXT: store i16 [[TMP0]], ptr [[DOTADDR]], align 2138// CHECK1-NEXT: store i32 [[TMP1]], ptr [[DOTADDR1]], align 4139// CHECK1-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 4140// CHECK1-NEXT: call void @__kmpc_get_shared_variables(ptr [[GLOBAL_ARGS]])141// CHECK1-NEXT: [[TMP2:%.*]] = load ptr, ptr [[GLOBAL_ARGS]], align 8142// CHECK1-NEXT: [[TMP3:%.*]] = getelementptr inbounds ptr, ptr [[TMP2]], i64 0143// CHECK1-NEXT: [[TMP4:%.*]] = load ptr, ptr [[TMP3]], align 8144// CHECK1-NEXT: call void @_Z3usePi_omp_outlined(ptr [[DOTADDR1]], ptr [[DOTZERO_ADDR]], ptr [[TMP4]]) #[[ATTR4]]145// CHECK1-NEXT: ret void146//147//148// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l25149// CHECK2-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[C:%.*]]) #[[ATTR0:[0-9]+]] {150// CHECK2-NEXT: entry:151// CHECK2-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4152// CHECK2-NEXT: [[C_ADDR:%.*]] = alloca ptr, align 4153// CHECK2-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x ptr], align 4154// CHECK2-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4155// CHECK2-NEXT: store ptr [[C]], ptr [[C_ADDR]], align 4156// CHECK2-NEXT: [[TMP0:%.*]] = load ptr, ptr [[C_ADDR]], align 4157// CHECK2-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l25_kernel_environment, ptr [[DYN_PTR]])158// CHECK2-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1159// CHECK2-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]160// CHECK2: user_code.entry:161// CHECK2-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1:[0-9]+]])162// CHECK2-NEXT: call void @_Z3usePi(ptr noundef [[TMP0]]) #[[ATTR6:[0-9]+]]163// CHECK2-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 0164// CHECK2-NEXT: store ptr [[TMP0]], ptr [[TMP3]], align 4165// CHECK2-NEXT: call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP2]], i32 1, i32 2, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l25_omp_outlined, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l25_omp_outlined_wrapper, ptr [[CAPTURED_VARS_ADDRS]], i32 1)166// CHECK2-NEXT: call void @__kmpc_target_deinit()167// CHECK2-NEXT: ret void168// CHECK2: worker.exit:169// CHECK2-NEXT: ret void170//171//172// CHECK2-LABEL: define {{[^@]+}}@_Z3usePi173// CHECK2-SAME: (ptr noundef [[C:%.*]]) #[[ATTR1:[0-9]+]] {174// CHECK2-NEXT: entry:175// CHECK2-NEXT: [[C_ADDR:%.*]] = alloca ptr, align 4176// CHECK2-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x ptr], align 4177// CHECK2-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])178// CHECK2-NEXT: store ptr [[C]], ptr [[C_ADDR]], align 4179// CHECK2-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 0180// CHECK2-NEXT: store ptr [[C_ADDR]], ptr [[TMP1]], align 4181// CHECK2-NEXT: call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP0]], i32 1, i32 2, i32 -1, ptr @_Z3usePi_omp_outlined, ptr @_Z3usePi_omp_outlined_wrapper, ptr [[CAPTURED_VARS_ADDRS]], i32 1)182// CHECK2-NEXT: ret void183//184//185// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l25_omp_outlined186// CHECK2-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[C:%.*]]) #[[ATTR2:[0-9]+]] {187// CHECK2-NEXT: entry:188// CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4189// CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4190// CHECK2-NEXT: [[C_ADDR:%.*]] = alloca ptr, align 4191// CHECK2-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4192// CHECK2-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4193// CHECK2-NEXT: store ptr [[C]], ptr [[C_ADDR]], align 4194// CHECK2-NEXT: [[TMP0:%.*]] = load ptr, ptr [[C_ADDR]], align 4195// CHECK2-NEXT: call void @_Z3usePi(ptr noundef [[TMP0]]) #[[ATTR6]]196// CHECK2-NEXT: ret void197//198//199// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l25_omp_outlined_wrapper200// CHECK2-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR3:[0-9]+]] {201// CHECK2-NEXT: entry:202// CHECK2-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2203// CHECK2-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4204// CHECK2-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4205// CHECK2-NEXT: [[GLOBAL_ARGS:%.*]] = alloca ptr, align 4206// CHECK2-NEXT: store i16 [[TMP0]], ptr [[DOTADDR]], align 2207// CHECK2-NEXT: store i32 [[TMP1]], ptr [[DOTADDR1]], align 4208// CHECK2-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 4209// CHECK2-NEXT: call void @__kmpc_get_shared_variables(ptr [[GLOBAL_ARGS]])210// CHECK2-NEXT: [[TMP2:%.*]] = load ptr, ptr [[GLOBAL_ARGS]], align 4211// CHECK2-NEXT: [[TMP3:%.*]] = getelementptr inbounds ptr, ptr [[TMP2]], i32 0212// CHECK2-NEXT: [[TMP4:%.*]] = load ptr, ptr [[TMP3]], align 4213// CHECK2-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l25_omp_outlined(ptr [[DOTADDR1]], ptr [[DOTZERO_ADDR]], ptr [[TMP4]]) #[[ATTR4:[0-9]+]]214// CHECK2-NEXT: ret void215//216//217// CHECK2-LABEL: define {{[^@]+}}@_Z3usePi_omp_outlined218// CHECK2-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[C:%.*]]) #[[ATTR2]] {219// CHECK2-NEXT: entry:220// CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4221// CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4222// CHECK2-NEXT: [[C_ADDR:%.*]] = alloca ptr, align 4223// CHECK2-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4224// CHECK2-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4225// CHECK2-NEXT: store ptr [[C]], ptr [[C_ADDR]], align 4226// CHECK2-NEXT: [[TMP0:%.*]] = load ptr, ptr [[C_ADDR]], align 4227// CHECK2-NEXT: [[TMP1:%.*]] = load ptr, ptr [[TMP0]], align 4228// CHECK2-NEXT: call void @_Z4workPi(ptr noundef [[TMP1]]) #[[ATTR6]]229// CHECK2-NEXT: ret void230//231//232// CHECK2-LABEL: define {{[^@]+}}@_Z4workPi233// CHECK2-SAME: (ptr noundef [[C:%.*]]) #[[ATTR1]] {234// CHECK2-NEXT: entry:235// CHECK2-NEXT: [[C_ADDR:%.*]] = alloca ptr, align 4236// CHECK2-NEXT: store ptr [[C]], ptr [[C_ADDR]], align 4237// CHECK2-NEXT: [[TMP0:%.*]] = load ptr, ptr [[C_ADDR]], align 4238// CHECK2-NEXT: [[TMP1:%.*]] = atomicrmw add ptr [[TMP0]], i32 1 monotonic, align 4239// CHECK2-NEXT: ret void240//241//242// CHECK2-LABEL: define {{[^@]+}}@_Z3usePi_omp_outlined_wrapper243// CHECK2-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR3]] {244// CHECK2-NEXT: entry:245// CHECK2-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2246// CHECK2-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4247// CHECK2-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4248// CHECK2-NEXT: [[GLOBAL_ARGS:%.*]] = alloca ptr, align 4249// CHECK2-NEXT: store i16 [[TMP0]], ptr [[DOTADDR]], align 2250// CHECK2-NEXT: store i32 [[TMP1]], ptr [[DOTADDR1]], align 4251// CHECK2-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 4252// CHECK2-NEXT: call void @__kmpc_get_shared_variables(ptr [[GLOBAL_ARGS]])253// CHECK2-NEXT: [[TMP2:%.*]] = load ptr, ptr [[GLOBAL_ARGS]], align 4254// CHECK2-NEXT: [[TMP3:%.*]] = getelementptr inbounds ptr, ptr [[TMP2]], i32 0255// CHECK2-NEXT: [[TMP4:%.*]] = load ptr, ptr [[TMP3]], align 4256// CHECK2-NEXT: call void @_Z3usePi_omp_outlined(ptr [[DOTADDR1]], ptr [[DOTZERO_ADDR]], ptr [[TMP4]]) #[[ATTR4]]257// CHECK2-NEXT: ret void258//259