613 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 -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 -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-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 -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 -fexceptions -fcxx-exceptions -x c++ -triple nvptx-unknown-unknown -aux-triple i386-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 -disable-O0-optnone | FileCheck %s --check-prefix=CHECK27// expected-no-diagnostics8#ifndef HEADER9#define HEADER10 11template<typename tx>12tx ftemplate(int n) {13 tx a = 0;14 short aa = 0;15 tx b[10];16 17 #pragma omp target if(0)18 {19 #pragma omp parallel20 {21 int a = 41;22 }23 a += 1;24 }25 26 #pragma omp target27 {28 #pragma omp parallel29 {30 int a = 42;31 }32 #pragma omp parallel if(0)33 {34 int a = 43;35 }36 #pragma omp parallel if(1)37 {38 int a = 44;39 }40 a += 1;41 }42 43 #pragma omp target if(n>40)44 {45 #pragma omp parallel if(n>1000)46 {47 int a = 45;48#pragma omp barrier49 }50 a += 1;51 aa += 1;52 b[2] += 1;53 }54 55 #pragma omp target56 {57 #pragma omp parallel58 {59 #pragma omp critical60 ++a;61 }62 ++a;63 }64 return a;65}66 67int bar(int n){68 int a = 0;69 70 a += ftemplate<int>(n);71 72 return a;73}74 75#endif76// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l2677// CHECK1-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[A:%.*]]) #[[ATTR0:[0-9]+]] {78// CHECK1-NEXT: entry:79// CHECK1-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 880// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca i64, align 881// CHECK1-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x ptr], align 882// CHECK1-NEXT: [[CAPTURED_VARS_ADDRS1:%.*]] = alloca [0 x ptr], align 883// CHECK1-NEXT: [[CAPTURED_VARS_ADDRS2:%.*]] = alloca [0 x ptr], align 884// CHECK1-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 885// CHECK1-NEXT: store i64 [[A]], ptr [[A_ADDR]], align 886// CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_kernel_environment, ptr [[DYN_PTR]])87// CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -188// CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]89// CHECK1: user_code.entry:90// CHECK1-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1:[0-9]+]])91// CHECK1-NEXT: call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined_wrapper, ptr [[CAPTURED_VARS_ADDRS]], i64 0)92// CHECK1-NEXT: call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP1]], i32 0, i32 -1, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined1_wrapper, ptr [[CAPTURED_VARS_ADDRS1]], i64 0)93// CHECK1-NEXT: call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined2_wrapper, ptr [[CAPTURED_VARS_ADDRS2]], i64 0)94// CHECK1-NEXT: [[TMP2:%.*]] = load i32, ptr [[A_ADDR]], align 495// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP2]], 196// CHECK1-NEXT: store i32 [[ADD]], ptr [[A_ADDR]], align 497// CHECK1-NEXT: call void @__kmpc_target_deinit()98// CHECK1-NEXT: ret void99// CHECK1: worker.exit:100// CHECK1-NEXT: ret void101//102//103// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined104// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1:[0-9]+]] {105// CHECK1-NEXT: entry:106// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8107// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8108// CHECK1-NEXT: [[A:%.*]] = alloca i32, align 4109// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8110// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8111// CHECK1-NEXT: store i32 42, ptr [[A]], align 4112// CHECK1-NEXT: ret void113//114//115// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined_wrapper116// CHECK1-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR2:[0-9]+]] {117// CHECK1-NEXT: entry:118// CHECK1-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2119// CHECK1-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4120// CHECK1-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4121// CHECK1-NEXT: [[GLOBAL_ARGS:%.*]] = alloca ptr, align 8122// CHECK1-NEXT: store i16 [[TMP0]], ptr [[DOTADDR]], align 2123// CHECK1-NEXT: store i32 [[TMP1]], ptr [[DOTADDR1]], align 4124// CHECK1-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 4125// CHECK1-NEXT: call void @__kmpc_get_shared_variables(ptr [[GLOBAL_ARGS]])126// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined(ptr [[DOTADDR1]], ptr [[DOTZERO_ADDR]]) #[[ATTR3:[0-9]+]]127// CHECK1-NEXT: ret void128//129//130// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined1131// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {132// CHECK1-NEXT: entry:133// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8134// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8135// CHECK1-NEXT: [[A:%.*]] = alloca i32, align 4136// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8137// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8138// CHECK1-NEXT: store i32 43, ptr [[A]], align 4139// CHECK1-NEXT: ret void140//141//142// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined1_wrapper143// CHECK1-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR2]] {144// CHECK1-NEXT: entry:145// CHECK1-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2146// CHECK1-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4147// CHECK1-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4148// CHECK1-NEXT: [[GLOBAL_ARGS:%.*]] = alloca ptr, align 8149// CHECK1-NEXT: store i16 [[TMP0]], ptr [[DOTADDR]], align 2150// CHECK1-NEXT: store i32 [[TMP1]], ptr [[DOTADDR1]], align 4151// CHECK1-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 4152// CHECK1-NEXT: call void @__kmpc_get_shared_variables(ptr [[GLOBAL_ARGS]])153// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined1(ptr [[DOTADDR1]], ptr [[DOTZERO_ADDR]]) #[[ATTR3]]154// CHECK1-NEXT: ret void155//156//157// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined2158// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {159// CHECK1-NEXT: entry:160// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8161// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8162// CHECK1-NEXT: [[A:%.*]] = alloca i32, align 4163// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8164// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8165// CHECK1-NEXT: store i32 44, ptr [[A]], align 4166// CHECK1-NEXT: ret void167//168//169// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined2_wrapper170// CHECK1-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR2]] {171// CHECK1-NEXT: entry:172// CHECK1-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2173// CHECK1-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4174// CHECK1-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4175// CHECK1-NEXT: [[GLOBAL_ARGS:%.*]] = alloca ptr, align 8176// CHECK1-NEXT: store i16 [[TMP0]], ptr [[DOTADDR]], align 2177// CHECK1-NEXT: store i32 [[TMP1]], ptr [[DOTADDR1]], align 4178// CHECK1-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 4179// CHECK1-NEXT: call void @__kmpc_get_shared_variables(ptr [[GLOBAL_ARGS]])180// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined2(ptr [[DOTADDR1]], ptr [[DOTZERO_ADDR]]) #[[ATTR3]]181// CHECK1-NEXT: ret void182//183//184// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l43185// CHECK1-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[N:%.*]], i64 noundef [[A:%.*]], i64 noundef [[AA:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]]) #[[ATTR0]] {186// CHECK1-NEXT: entry:187// CHECK1-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8188// CHECK1-NEXT: [[N_ADDR:%.*]] = alloca i64, align 8189// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca i64, align 8190// CHECK1-NEXT: [[AA_ADDR:%.*]] = alloca i64, align 8191// CHECK1-NEXT: [[B_ADDR:%.*]] = alloca ptr, align 8192// CHECK1-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x ptr], align 8193// CHECK1-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8194// CHECK1-NEXT: store i64 [[N]], ptr [[N_ADDR]], align 8195// CHECK1-NEXT: store i64 [[A]], ptr [[A_ADDR]], align 8196// CHECK1-NEXT: store i64 [[AA]], ptr [[AA_ADDR]], align 8197// CHECK1-NEXT: store ptr [[B]], ptr [[B_ADDR]], align 8198// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 8199// CHECK1-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l43_kernel_environment, ptr [[DYN_PTR]])200// CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1201// CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]202// CHECK1: user_code.entry:203// CHECK1-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])204// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[N_ADDR]], align 4205// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 1000206// CHECK1-NEXT: [[TMP4:%.*]] = zext i1 [[CMP]] to i32207// CHECK1-NEXT: call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP2]], i32 [[TMP4]], i32 -1, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l43_omp_outlined, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l43_omp_outlined_wrapper, ptr [[CAPTURED_VARS_ADDRS]], i64 0)208// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[A_ADDR]], align 4209// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP5]], 1210// CHECK1-NEXT: store i32 [[ADD]], ptr [[A_ADDR]], align 4211// CHECK1-NEXT: [[TMP6:%.*]] = load i16, ptr [[AA_ADDR]], align 2212// CHECK1-NEXT: [[CONV:%.*]] = sext i16 [[TMP6]] to i32213// CHECK1-NEXT: [[ADD1:%.*]] = add nsw i32 [[CONV]], 1214// CHECK1-NEXT: [[CONV2:%.*]] = trunc i32 [[ADD1]] to i16215// CHECK1-NEXT: store i16 [[CONV2]], ptr [[AA_ADDR]], align 2216// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i64 0, i64 2217// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[ARRAYIDX]], align 4218// CHECK1-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP7]], 1219// CHECK1-NEXT: store i32 [[ADD3]], ptr [[ARRAYIDX]], align 4220// CHECK1-NEXT: call void @__kmpc_target_deinit()221// CHECK1-NEXT: ret void222// CHECK1: worker.exit:223// CHECK1-NEXT: ret void224//225//226// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l43_omp_outlined227// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {228// CHECK1-NEXT: entry:229// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8230// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8231// CHECK1-NEXT: [[A:%.*]] = alloca i32, align 4232// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8233// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8234// CHECK1-NEXT: store i32 45, ptr [[A]], align 4235// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8236// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4237// CHECK1-NEXT: call void @__kmpc_barrier(ptr @[[GLOB2:[0-9]+]], i32 [[TMP1]])238// CHECK1-NEXT: ret void239//240//241// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l43_omp_outlined_wrapper242// CHECK1-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR2]] {243// CHECK1-NEXT: entry:244// CHECK1-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2245// CHECK1-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4246// CHECK1-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4247// CHECK1-NEXT: [[GLOBAL_ARGS:%.*]] = alloca ptr, align 8248// CHECK1-NEXT: store i16 [[TMP0]], ptr [[DOTADDR]], align 2249// CHECK1-NEXT: store i32 [[TMP1]], ptr [[DOTADDR1]], align 4250// CHECK1-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 4251// CHECK1-NEXT: call void @__kmpc_get_shared_variables(ptr [[GLOBAL_ARGS]])252// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l43_omp_outlined(ptr [[DOTADDR1]], ptr [[DOTZERO_ADDR]]) #[[ATTR3]]253// CHECK1-NEXT: ret void254//255//256// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l55257// CHECK1-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[A:%.*]]) #[[ATTR0]] {258// CHECK1-NEXT: entry:259// CHECK1-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8260// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca i64, align 8261// CHECK1-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x ptr], align 8262// CHECK1-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8263// CHECK1-NEXT: store i64 [[A]], ptr [[A_ADDR]], align 8264// CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l55_kernel_environment, ptr [[DYN_PTR]])265// CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1266// CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]267// CHECK1: user_code.entry:268// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 4269// CHECK1-NEXT: [[A1:%.*]] = call align 16 ptr @__kmpc_alloc_shared(i64 4)270// CHECK1-NEXT: store i32 [[TMP1]], ptr [[A1]], align 4271// CHECK1-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])272// CHECK1-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 0273// CHECK1-NEXT: store ptr [[A1]], ptr [[TMP3]], align 8274// CHECK1-NEXT: call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP2]], i32 1, i32 -1, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l55_omp_outlined, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l55_omp_outlined_wrapper, ptr [[CAPTURED_VARS_ADDRS]], i64 1)275// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[A1]], align 4276// CHECK1-NEXT: [[INC:%.*]] = add nsw i32 [[TMP4]], 1277// CHECK1-NEXT: store i32 [[INC]], ptr [[A1]], align 4278// CHECK1-NEXT: call void @__kmpc_free_shared(ptr [[A1]], i64 4)279// CHECK1-NEXT: call void @__kmpc_target_deinit()280// CHECK1-NEXT: ret void281// CHECK1: worker.exit:282// CHECK1-NEXT: ret void283//284//285// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l55_omp_outlined286// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR1]] {287// CHECK1-NEXT: entry:288// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8289// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8290// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8291// CHECK1-NEXT: [[CRITICAL_COUNTER:%.*]] = alloca i32, align 4292// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8293// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8294// CHECK1-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8295// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8296// CHECK1-NEXT: [[TMP1:%.*]] = call i64 @__kmpc_warp_active_thread_mask()297// CHECK1-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()298// CHECK1-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @__kmpc_get_hardware_num_threads_in_block()299// CHECK1-NEXT: store i32 0, ptr [[CRITICAL_COUNTER]], align 4300// CHECK1-NEXT: br label [[OMP_CRITICAL_LOOP:%.*]]301// CHECK1: omp.critical.loop:302// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[CRITICAL_COUNTER]], align 4303// CHECK1-NEXT: [[TMP4:%.*]] = icmp slt i32 [[TMP3]], [[NVPTX_NUM_THREADS]]304// CHECK1-NEXT: br i1 [[TMP4]], label [[OMP_CRITICAL_TEST:%.*]], label [[OMP_CRITICAL_EXIT:%.*]]305// CHECK1: omp.critical.test:306// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[CRITICAL_COUNTER]], align 4307// CHECK1-NEXT: [[TMP6:%.*]] = icmp eq i32 [[TMP2]], [[TMP5]]308// CHECK1-NEXT: br i1 [[TMP6]], label [[OMP_CRITICAL_BODY:%.*]], label [[OMP_CRITICAL_SYNC:%.*]]309// CHECK1: omp.critical.body:310// CHECK1-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8311// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4312// CHECK1-NEXT: call void @__kmpc_critical(ptr @[[GLOB1]], i32 [[TMP8]], ptr @"_gomp_critical_user_$var")313// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[TMP0]], align 4314// CHECK1-NEXT: [[INC:%.*]] = add nsw i32 [[TMP9]], 1315// CHECK1-NEXT: store i32 [[INC]], ptr [[TMP0]], align 4316// CHECK1-NEXT: call void @__kmpc_end_critical(ptr @[[GLOB1]], i32 [[TMP8]], ptr @"_gomp_critical_user_$var")317// CHECK1-NEXT: br label [[OMP_CRITICAL_SYNC]]318// CHECK1: omp.critical.sync:319// CHECK1-NEXT: call void @__kmpc_syncwarp(i64 [[TMP1]])320// CHECK1-NEXT: [[TMP10:%.*]] = add nsw i32 [[TMP5]], 1321// CHECK1-NEXT: store i32 [[TMP10]], ptr [[CRITICAL_COUNTER]], align 4322// CHECK1-NEXT: br label [[OMP_CRITICAL_LOOP]]323// CHECK1: omp.critical.exit:324// CHECK1-NEXT: ret void325//326//327// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l55_omp_outlined_wrapper328// CHECK1-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR2]] {329// CHECK1-NEXT: entry:330// CHECK1-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2331// CHECK1-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4332// CHECK1-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4333// CHECK1-NEXT: [[GLOBAL_ARGS:%.*]] = alloca ptr, align 8334// CHECK1-NEXT: store i16 [[TMP0]], ptr [[DOTADDR]], align 2335// CHECK1-NEXT: store i32 [[TMP1]], ptr [[DOTADDR1]], align 4336// CHECK1-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 4337// CHECK1-NEXT: call void @__kmpc_get_shared_variables(ptr [[GLOBAL_ARGS]])338// CHECK1-NEXT: [[TMP2:%.*]] = load ptr, ptr [[GLOBAL_ARGS]], align 8339// CHECK1-NEXT: [[TMP3:%.*]] = getelementptr inbounds ptr, ptr [[TMP2]], i64 0340// CHECK1-NEXT: [[TMP4:%.*]] = load ptr, ptr [[TMP3]], align 8341// CHECK1-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l55_omp_outlined(ptr [[DOTADDR1]], ptr [[DOTZERO_ADDR]], ptr [[TMP4]]) #[[ATTR3]]342// CHECK1-NEXT: ret void343//344//345// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26346// CHECK2-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[A:%.*]]) #[[ATTR0:[0-9]+]] {347// CHECK2-NEXT: entry:348// CHECK2-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4349// CHECK2-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4350// CHECK2-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x ptr], align 4351// CHECK2-NEXT: [[CAPTURED_VARS_ADDRS1:%.*]] = alloca [0 x ptr], align 4352// CHECK2-NEXT: [[CAPTURED_VARS_ADDRS2:%.*]] = alloca [0 x ptr], align 4353// CHECK2-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4354// CHECK2-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 4355// CHECK2-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_kernel_environment, ptr [[DYN_PTR]])356// CHECK2-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1357// CHECK2-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]358// CHECK2: user_code.entry:359// CHECK2-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1:[0-9]+]])360// CHECK2-NEXT: call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined_wrapper, ptr [[CAPTURED_VARS_ADDRS]], i32 0)361// CHECK2-NEXT: call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP1]], i32 0, i32 -1, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined1_wrapper, ptr [[CAPTURED_VARS_ADDRS1]], i32 0)362// CHECK2-NEXT: call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined2_wrapper, ptr [[CAPTURED_VARS_ADDRS2]], i32 0)363// CHECK2-NEXT: [[TMP2:%.*]] = load i32, ptr [[A_ADDR]], align 4364// CHECK2-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP2]], 1365// CHECK2-NEXT: store i32 [[ADD]], ptr [[A_ADDR]], align 4366// CHECK2-NEXT: call void @__kmpc_target_deinit()367// CHECK2-NEXT: ret void368// CHECK2: worker.exit:369// CHECK2-NEXT: ret void370//371//372// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined373// CHECK2-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1:[0-9]+]] {374// CHECK2-NEXT: entry:375// CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4376// CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4377// CHECK2-NEXT: [[A:%.*]] = alloca i32, align 4378// CHECK2-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4379// CHECK2-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4380// CHECK2-NEXT: store i32 42, ptr [[A]], align 4381// CHECK2-NEXT: ret void382//383//384// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined_wrapper385// CHECK2-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR1]] {386// CHECK2-NEXT: entry:387// CHECK2-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2388// CHECK2-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4389// CHECK2-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4390// CHECK2-NEXT: [[GLOBAL_ARGS:%.*]] = alloca ptr, align 4391// CHECK2-NEXT: store i16 [[TMP0]], ptr [[DOTADDR]], align 2392// CHECK2-NEXT: store i32 [[TMP1]], ptr [[DOTADDR1]], align 4393// CHECK2-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 4394// CHECK2-NEXT: call void @__kmpc_get_shared_variables(ptr [[GLOBAL_ARGS]])395// CHECK2-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined(ptr [[DOTADDR1]], ptr [[DOTZERO_ADDR]]) #[[ATTR2:[0-9]+]]396// CHECK2-NEXT: ret void397//398//399// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined1400// CHECK2-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {401// CHECK2-NEXT: entry:402// CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4403// CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4404// CHECK2-NEXT: [[A:%.*]] = alloca i32, align 4405// CHECK2-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4406// CHECK2-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4407// CHECK2-NEXT: store i32 43, ptr [[A]], align 4408// CHECK2-NEXT: ret void409//410//411// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined1_wrapper412// CHECK2-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR1]] {413// CHECK2-NEXT: entry:414// CHECK2-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2415// CHECK2-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4416// CHECK2-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4417// CHECK2-NEXT: [[GLOBAL_ARGS:%.*]] = alloca ptr, align 4418// CHECK2-NEXT: store i16 [[TMP0]], ptr [[DOTADDR]], align 2419// CHECK2-NEXT: store i32 [[TMP1]], ptr [[DOTADDR1]], align 4420// CHECK2-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 4421// CHECK2-NEXT: call void @__kmpc_get_shared_variables(ptr [[GLOBAL_ARGS]])422// CHECK2-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined1(ptr [[DOTADDR1]], ptr [[DOTZERO_ADDR]]) #[[ATTR2]]423// CHECK2-NEXT: ret void424//425//426// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined2427// CHECK2-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {428// CHECK2-NEXT: entry:429// CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4430// CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4431// CHECK2-NEXT: [[A:%.*]] = alloca i32, align 4432// CHECK2-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4433// CHECK2-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4434// CHECK2-NEXT: store i32 44, ptr [[A]], align 4435// CHECK2-NEXT: ret void436//437//438// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined2_wrapper439// CHECK2-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR1]] {440// CHECK2-NEXT: entry:441// CHECK2-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2442// CHECK2-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4443// CHECK2-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4444// CHECK2-NEXT: [[GLOBAL_ARGS:%.*]] = alloca ptr, align 4445// CHECK2-NEXT: store i16 [[TMP0]], ptr [[DOTADDR]], align 2446// CHECK2-NEXT: store i32 [[TMP1]], ptr [[DOTADDR1]], align 4447// CHECK2-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 4448// CHECK2-NEXT: call void @__kmpc_get_shared_variables(ptr [[GLOBAL_ARGS]])449// CHECK2-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined2(ptr [[DOTADDR1]], ptr [[DOTZERO_ADDR]]) #[[ATTR2]]450// CHECK2-NEXT: ret void451//452//453// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l43454// CHECK2-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[N:%.*]], i32 noundef [[A:%.*]], i32 noundef [[AA:%.*]], ptr noundef nonnull align 4 dereferenceable(40) [[B:%.*]]) #[[ATTR0]] {455// CHECK2-NEXT: entry:456// CHECK2-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4457// CHECK2-NEXT: [[N_ADDR:%.*]] = alloca i32, align 4458// CHECK2-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4459// CHECK2-NEXT: [[AA_ADDR:%.*]] = alloca i32, align 4460// CHECK2-NEXT: [[B_ADDR:%.*]] = alloca ptr, align 4461// CHECK2-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x ptr], align 4462// CHECK2-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4463// CHECK2-NEXT: store i32 [[N]], ptr [[N_ADDR]], align 4464// CHECK2-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 4465// CHECK2-NEXT: store i32 [[AA]], ptr [[AA_ADDR]], align 4466// CHECK2-NEXT: store ptr [[B]], ptr [[B_ADDR]], align 4467// CHECK2-NEXT: [[TMP0:%.*]] = load ptr, ptr [[B_ADDR]], align 4468// CHECK2-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l43_kernel_environment, ptr [[DYN_PTR]])469// CHECK2-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1470// CHECK2-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]471// CHECK2: user_code.entry:472// CHECK2-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])473// CHECK2-NEXT: [[TMP3:%.*]] = load i32, ptr [[N_ADDR]], align 4474// CHECK2-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 1000475// CHECK2-NEXT: [[TMP4:%.*]] = zext i1 [[CMP]] to i32476// CHECK2-NEXT: call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP2]], i32 [[TMP4]], i32 -1, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l43_omp_outlined, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l43_omp_outlined_wrapper, ptr [[CAPTURED_VARS_ADDRS]], i32 0)477// CHECK2-NEXT: [[TMP5:%.*]] = load i32, ptr [[A_ADDR]], align 4478// CHECK2-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP5]], 1479// CHECK2-NEXT: store i32 [[ADD]], ptr [[A_ADDR]], align 4480// CHECK2-NEXT: [[TMP6:%.*]] = load i16, ptr [[AA_ADDR]], align 2481// CHECK2-NEXT: [[CONV:%.*]] = sext i16 [[TMP6]] to i32482// CHECK2-NEXT: [[ADD1:%.*]] = add nsw i32 [[CONV]], 1483// CHECK2-NEXT: [[CONV2:%.*]] = trunc i32 [[ADD1]] to i16484// CHECK2-NEXT: store i16 [[CONV2]], ptr [[AA_ADDR]], align 2485// CHECK2-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x i32], ptr [[TMP0]], i32 0, i32 2486// CHECK2-NEXT: [[TMP7:%.*]] = load i32, ptr [[ARRAYIDX]], align 4487// CHECK2-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP7]], 1488// CHECK2-NEXT: store i32 [[ADD3]], ptr [[ARRAYIDX]], align 4489// CHECK2-NEXT: call void @__kmpc_target_deinit()490// CHECK2-NEXT: ret void491// CHECK2: worker.exit:492// CHECK2-NEXT: ret void493//494//495// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l43_omp_outlined496// CHECK2-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {497// CHECK2-NEXT: entry:498// CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4499// CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4500// CHECK2-NEXT: [[A:%.*]] = alloca i32, align 4501// CHECK2-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4502// CHECK2-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4503// CHECK2-NEXT: store i32 45, ptr [[A]], align 4504// CHECK2-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4505// CHECK2-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4506// CHECK2-NEXT: call void @__kmpc_barrier(ptr @[[GLOB2:[0-9]+]], i32 [[TMP1]])507// CHECK2-NEXT: ret void508//509//510// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l43_omp_outlined_wrapper511// CHECK2-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR1]] {512// CHECK2-NEXT: entry:513// CHECK2-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2514// CHECK2-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4515// CHECK2-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4516// CHECK2-NEXT: [[GLOBAL_ARGS:%.*]] = alloca ptr, align 4517// CHECK2-NEXT: store i16 [[TMP0]], ptr [[DOTADDR]], align 2518// CHECK2-NEXT: store i32 [[TMP1]], ptr [[DOTADDR1]], align 4519// CHECK2-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 4520// CHECK2-NEXT: call void @__kmpc_get_shared_variables(ptr [[GLOBAL_ARGS]])521// CHECK2-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l43_omp_outlined(ptr [[DOTADDR1]], ptr [[DOTZERO_ADDR]]) #[[ATTR2]]522// CHECK2-NEXT: ret void523//524//525// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l55526// CHECK2-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[A:%.*]]) #[[ATTR0]] {527// CHECK2-NEXT: entry:528// CHECK2-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4529// CHECK2-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4530// CHECK2-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x ptr], align 4531// CHECK2-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4532// CHECK2-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 4533// CHECK2-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l55_kernel_environment, ptr [[DYN_PTR]])534// CHECK2-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1535// CHECK2-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]536// CHECK2: user_code.entry:537// CHECK2-NEXT: [[TMP1:%.*]] = load i32, ptr [[A_ADDR]], align 4538// CHECK2-NEXT: [[A1:%.*]] = call align 4 ptr @__kmpc_alloc_shared(i32 4)539// CHECK2-NEXT: store i32 [[TMP1]], ptr [[A1]], align 4540// CHECK2-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])541// CHECK2-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 0542// CHECK2-NEXT: store ptr [[A1]], ptr [[TMP3]], align 4543// CHECK2-NEXT: call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP2]], i32 1, i32 -1, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l55_omp_outlined, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l55_omp_outlined_wrapper, ptr [[CAPTURED_VARS_ADDRS]], i32 1)544// CHECK2-NEXT: [[TMP4:%.*]] = load i32, ptr [[A1]], align 4545// CHECK2-NEXT: [[INC:%.*]] = add nsw i32 [[TMP4]], 1546// CHECK2-NEXT: store i32 [[INC]], ptr [[A1]], align 4547// CHECK2-NEXT: call void @__kmpc_free_shared(ptr [[A1]], i32 4)548// CHECK2-NEXT: call void @__kmpc_target_deinit()549// CHECK2-NEXT: ret void550// CHECK2: worker.exit:551// CHECK2-NEXT: ret void552//553//554// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l55_omp_outlined555// CHECK2-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]]) #[[ATTR1]] {556// CHECK2-NEXT: entry:557// CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4558// CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4559// CHECK2-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 4560// CHECK2-NEXT: [[CRITICAL_COUNTER:%.*]] = alloca i32, align 4561// CHECK2-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4562// CHECK2-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4563// CHECK2-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 4564// CHECK2-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 4565// CHECK2-NEXT: [[TMP1:%.*]] = call i64 @__kmpc_warp_active_thread_mask()566// CHECK2-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()567// CHECK2-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @__kmpc_get_hardware_num_threads_in_block()568// CHECK2-NEXT: store i32 0, ptr [[CRITICAL_COUNTER]], align 4569// CHECK2-NEXT: br label [[OMP_CRITICAL_LOOP:%.*]]570// CHECK2: omp.critical.loop:571// CHECK2-NEXT: [[TMP3:%.*]] = load i32, ptr [[CRITICAL_COUNTER]], align 4572// CHECK2-NEXT: [[TMP4:%.*]] = icmp slt i32 [[TMP3]], [[NVPTX_NUM_THREADS]]573// CHECK2-NEXT: br i1 [[TMP4]], label [[OMP_CRITICAL_TEST:%.*]], label [[OMP_CRITICAL_EXIT:%.*]]574// CHECK2: omp.critical.test:575// CHECK2-NEXT: [[TMP5:%.*]] = load i32, ptr [[CRITICAL_COUNTER]], align 4576// CHECK2-NEXT: [[TMP6:%.*]] = icmp eq i32 [[TMP2]], [[TMP5]]577// CHECK2-NEXT: br i1 [[TMP6]], label [[OMP_CRITICAL_BODY:%.*]], label [[OMP_CRITICAL_SYNC:%.*]]578// CHECK2: omp.critical.body:579// CHECK2-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4580// CHECK2-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4581// CHECK2-NEXT: call void @__kmpc_critical(ptr @[[GLOB1]], i32 [[TMP8]], ptr @"_gomp_critical_user_$var")582// CHECK2-NEXT: [[TMP9:%.*]] = load i32, ptr [[TMP0]], align 4583// CHECK2-NEXT: [[INC:%.*]] = add nsw i32 [[TMP9]], 1584// CHECK2-NEXT: store i32 [[INC]], ptr [[TMP0]], align 4585// CHECK2-NEXT: call void @__kmpc_end_critical(ptr @[[GLOB1]], i32 [[TMP8]], ptr @"_gomp_critical_user_$var")586// CHECK2-NEXT: br label [[OMP_CRITICAL_SYNC]]587// CHECK2: omp.critical.sync:588// CHECK2-NEXT: call void @__kmpc_syncwarp(i64 [[TMP1]])589// CHECK2-NEXT: [[TMP10:%.*]] = add nsw i32 [[TMP5]], 1590// CHECK2-NEXT: store i32 [[TMP10]], ptr [[CRITICAL_COUNTER]], align 4591// CHECK2-NEXT: br label [[OMP_CRITICAL_LOOP]]592// CHECK2: omp.critical.exit:593// CHECK2-NEXT: ret void594//595//596// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l55_omp_outlined_wrapper597// CHECK2-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR1]] {598// CHECK2-NEXT: entry:599// CHECK2-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2600// CHECK2-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4601// CHECK2-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4602// CHECK2-NEXT: [[GLOBAL_ARGS:%.*]] = alloca ptr, align 4603// CHECK2-NEXT: store i16 [[TMP0]], ptr [[DOTADDR]], align 2604// CHECK2-NEXT: store i32 [[TMP1]], ptr [[DOTADDR1]], align 4605// CHECK2-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 4606// CHECK2-NEXT: call void @__kmpc_get_shared_variables(ptr [[GLOBAL_ARGS]])607// CHECK2-NEXT: [[TMP2:%.*]] = load ptr, ptr [[GLOBAL_ARGS]], align 4608// CHECK2-NEXT: [[TMP3:%.*]] = getelementptr inbounds ptr, ptr [[TMP2]], i32 0609// CHECK2-NEXT: [[TMP4:%.*]] = load ptr, ptr [[TMP3]], align 4610// CHECK2-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l55_omp_outlined(ptr [[DOTADDR1]], ptr [[DOTZERO_ADDR]], ptr [[TMP4]]) #[[ATTR2]]611// CHECK2-NEXT: ret void612//613