2121 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 - | FileCheck %s --check-prefix=CHECK-645// 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 - | FileCheck %s --check-prefix=CHECK-327// 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 - | FileCheck %s --check-prefix=CHECK-32-EX8// expected-no-diagnostics9#ifndef HEADER10#define HEADER11 12// Check for the data transfer medium in shared memory to transfer the reduction list to the first warp.13 14// Check that the execution mode of all 3 target regions is set to Spmd Mode.15 16template<typename tx>17tx ftemplate(int n) {18 int a;19 short b;20 tx c;21 float d;22 double e;23 24 #pragma omp target parallel reduction(+: e)25 {26 e += 5;27 }28 29 #pragma omp target parallel reduction(^: c) reduction(*: d)30 {31 c ^= 2;32 d *= 33;33 }34 35 #pragma omp target parallel reduction(|: a) reduction(max: b)36 {37 a |= 1;38 b = 99 > b ? 99 : b;39 }40 41 return a+b+c+d+e;42}43 44int bar(int n){45 int a = 0;46 47 a += ftemplate<char>(n);48 49 return a;50}51 52// define internal void [[PFN]](53 54 55// Reduction function56 57// Shuffle and reduce function58// Condition to reduce59// Now check if we should just copy over the remote reduction list60 61// Inter warp copy function62// [[DO_COPY]]63// Barrier after copy to shared memory storage medium.64// Read into warp 0.65 66// define internal void [[PFN1]](67 68// Reduction function69 70// Shuffle and reduce function71// Condition to reduce72// Now check if we should just copy over the remote reduction list73 74// Inter warp copy function75// [[DO_COPY]]76// Barrier after copy to shared memory storage medium.77// Read into warp 0.78// [[DO_COPY]]79// Barrier after copy to shared memory storage medium.80// Read into warp 0.81 82// define internal void [[PFN2]](83 84 85// Reduction function86 87// Shuffle and reduce function88// Condition to reduce89// Now check if we should just copy over the remote reduction list90 91// Inter warp copy function92// [[DO_COPY]]93// Barrier after copy to shared memory storage medium.94// Read into warp 0.95// [[DO_COPY]]96// Barrier after copy to shared memory storage medium.97// Read into warp 0.98 99#endif100// CHECK-64-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l24101// CHECK-64-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[E:%.*]]) #[[ATTR0:[0-9]+]] {102// CHECK-64-NEXT: entry:103// CHECK-64-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8104// CHECK-64-NEXT: [[E_ADDR:%.*]] = alloca ptr, align 8105// CHECK-64-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x ptr], align 8106// CHECK-64-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8107// CHECK-64-NEXT: store ptr [[E]], ptr [[E_ADDR]], align 8108// CHECK-64-NEXT: [[TMP0:%.*]] = load ptr, ptr [[E_ADDR]], align 8109// CHECK-64-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l24_kernel_environment, ptr [[DYN_PTR]])110// CHECK-64-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1111// CHECK-64-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]112// CHECK-64: user_code.entry:113// CHECK-64-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1:[0-9]+]])114// CHECK-64-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 0115// CHECK-64-NEXT: store ptr [[TMP0]], ptr [[TMP3]], align 8116// CHECK-64-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]+}}__Z9ftemplateIcET_i_l24_omp_outlined, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i64 1)117// CHECK-64-NEXT: call void @__kmpc_target_deinit()118// CHECK-64-NEXT: ret void119// CHECK-64: worker.exit:120// CHECK-64-NEXT: ret void121//122//123// CHECK-64-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l24_omp_outlined124// CHECK-64-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[E:%.*]]) #[[ATTR1:[0-9]+]] {125// CHECK-64-NEXT: entry:126// CHECK-64-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8127// CHECK-64-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8128// CHECK-64-NEXT: [[E_ADDR:%.*]] = alloca ptr, align 8129// CHECK-64-NEXT: [[E1:%.*]] = alloca double, align 8130// CHECK-64-NEXT: [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 8131// CHECK-64-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8132// CHECK-64-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8133// CHECK-64-NEXT: store ptr [[E]], ptr [[E_ADDR]], align 8134// CHECK-64-NEXT: [[TMP0:%.*]] = load ptr, ptr [[E_ADDR]], align 8135// CHECK-64-NEXT: store double 0.000000e+00, ptr [[E1]], align 8136// CHECK-64-NEXT: [[TMP1:%.*]] = load double, ptr [[E1]], align 8137// CHECK-64-NEXT: [[ADD:%.*]] = fadd double [[TMP1]], 5.000000e+00138// CHECK-64-NEXT: store double [[ADD]], ptr [[E1]], align 8139// CHECK-64-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i64 0, i64 0140// CHECK-64-NEXT: store ptr [[E1]], ptr [[TMP2]], align 8141// CHECK-64-NEXT: [[TMP3:%.*]] = call i32 @__kmpc_nvptx_parallel_reduce_nowait_v2(ptr @[[GLOB1]], i64 8, ptr [[DOTOMP_REDUCTION_RED_LIST]], ptr @_omp_reduction_shuffle_and_reduce_func, ptr @_omp_reduction_inter_warp_copy_func)142// CHECK-64-NEXT: [[TMP4:%.*]] = icmp eq i32 [[TMP3]], 1143// CHECK-64-NEXT: br i1 [[TMP4]], label [[DOTOMP_REDUCTION_THEN:%.*]], label [[DOTOMP_REDUCTION_DONE:%.*]]144// CHECK-64: .omp.reduction.then:145// CHECK-64-NEXT: [[TMP5:%.*]] = load double, ptr [[TMP0]], align 8146// CHECK-64-NEXT: [[TMP6:%.*]] = load double, ptr [[E1]], align 8147// CHECK-64-NEXT: [[ADD2:%.*]] = fadd double [[TMP5]], [[TMP6]]148// CHECK-64-NEXT: store double [[ADD2]], ptr [[TMP0]], align 8149// CHECK-64-NEXT: br label [[DOTOMP_REDUCTION_DONE]]150// CHECK-64: .omp.reduction.done:151// CHECK-64-NEXT: ret void152//153//154// CHECK-64-LABEL: define {{[^@]+}}@_omp_reduction_shuffle_and_reduce_func155// CHECK-64-SAME: (ptr noundef [[TMP0:%.*]], i16 noundef signext [[TMP1:%.*]], i16 noundef signext [[TMP2:%.*]], i16 noundef signext [[TMP3:%.*]]) #[[ATTR2:[0-9]+]] {156// CHECK-64-NEXT: entry:157// CHECK-64-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 8158// CHECK-64-NEXT: [[DOTADDR1:%.*]] = alloca i16, align 2159// CHECK-64-NEXT: [[DOTADDR2:%.*]] = alloca i16, align 2160// CHECK-64-NEXT: [[DOTADDR3:%.*]] = alloca i16, align 2161// CHECK-64-NEXT: [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST:%.*]] = alloca [1 x ptr], align 8162// CHECK-64-NEXT: [[DOTOMP_REDUCTION_ELEMENT:%.*]] = alloca double, align 8163// CHECK-64-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 8164// CHECK-64-NEXT: store i16 [[TMP1]], ptr [[DOTADDR1]], align 2165// CHECK-64-NEXT: store i16 [[TMP2]], ptr [[DOTADDR2]], align 2166// CHECK-64-NEXT: store i16 [[TMP3]], ptr [[DOTADDR3]], align 2167// CHECK-64-NEXT: [[TMP4:%.*]] = load ptr, ptr [[DOTADDR]], align 8168// CHECK-64-NEXT: [[TMP5:%.*]] = load i16, ptr [[DOTADDR1]], align 2169// CHECK-64-NEXT: [[TMP6:%.*]] = load i16, ptr [[DOTADDR2]], align 2170// CHECK-64-NEXT: [[TMP7:%.*]] = load i16, ptr [[DOTADDR3]], align 2171// CHECK-64-NEXT: [[TMP8:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP4]], i64 0, i64 0172// CHECK-64-NEXT: [[TMP9:%.*]] = load ptr, ptr [[TMP8]], align 8173// CHECK-64-NEXT: [[TMP10:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]], i64 0, i64 0174// CHECK-64-NEXT: [[TMP11:%.*]] = getelementptr double, ptr [[TMP9]], i64 1175// CHECK-64-NEXT: [[TMP12:%.*]] = load i64, ptr [[TMP9]], align 8176// CHECK-64-NEXT: [[TMP13:%.*]] = call i32 @__kmpc_get_warp_size()177// CHECK-64-NEXT: [[TMP14:%.*]] = trunc i32 [[TMP13]] to i16178// CHECK-64-NEXT: [[TMP15:%.*]] = call i64 @__kmpc_shuffle_int64(i64 [[TMP12]], i16 [[TMP6]], i16 [[TMP14]])179// CHECK-64-NEXT: store i64 [[TMP15]], ptr [[DOTOMP_REDUCTION_ELEMENT]], align 8180// CHECK-64-NEXT: [[TMP16:%.*]] = getelementptr i64, ptr [[TMP9]], i64 1181// CHECK-64-NEXT: [[TMP17:%.*]] = getelementptr i64, ptr [[DOTOMP_REDUCTION_ELEMENT]], i64 1182// CHECK-64-NEXT: store ptr [[DOTOMP_REDUCTION_ELEMENT]], ptr [[TMP10]], align 8183// CHECK-64-NEXT: [[TMP18:%.*]] = icmp eq i16 [[TMP7]], 0184// CHECK-64-NEXT: [[TMP19:%.*]] = icmp eq i16 [[TMP7]], 1185// CHECK-64-NEXT: [[TMP20:%.*]] = icmp ult i16 [[TMP5]], [[TMP6]]186// CHECK-64-NEXT: [[TMP21:%.*]] = and i1 [[TMP19]], [[TMP20]]187// CHECK-64-NEXT: [[TMP22:%.*]] = icmp eq i16 [[TMP7]], 2188// CHECK-64-NEXT: [[TMP23:%.*]] = and i16 [[TMP5]], 1189// CHECK-64-NEXT: [[TMP24:%.*]] = icmp eq i16 [[TMP23]], 0190// CHECK-64-NEXT: [[TMP25:%.*]] = and i1 [[TMP22]], [[TMP24]]191// CHECK-64-NEXT: [[TMP26:%.*]] = icmp sgt i16 [[TMP6]], 0192// CHECK-64-NEXT: [[TMP27:%.*]] = and i1 [[TMP25]], [[TMP26]]193// CHECK-64-NEXT: [[TMP28:%.*]] = or i1 [[TMP18]], [[TMP21]]194// CHECK-64-NEXT: [[TMP29:%.*]] = or i1 [[TMP28]], [[TMP27]]195// CHECK-64-NEXT: br i1 [[TMP29]], label [[THEN:%.*]], label [[ELSE:%.*]]196// CHECK-64: then:197// CHECK-64-NEXT: call void @"{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l24_omp_outlined_omp$reduction$reduction_func"(ptr [[TMP4]], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]]) #[[ATTR3:[0-9]+]]198// CHECK-64-NEXT: br label [[IFCONT:%.*]]199// CHECK-64: else:200// CHECK-64-NEXT: br label [[IFCONT]]201// CHECK-64: ifcont:202// CHECK-64-NEXT: [[TMP30:%.*]] = icmp eq i16 [[TMP7]], 1203// CHECK-64-NEXT: [[TMP31:%.*]] = icmp uge i16 [[TMP5]], [[TMP6]]204// CHECK-64-NEXT: [[TMP32:%.*]] = and i1 [[TMP30]], [[TMP31]]205// CHECK-64-NEXT: br i1 [[TMP32]], label [[THEN4:%.*]], label [[ELSE5:%.*]]206// CHECK-64: then4:207// CHECK-64-NEXT: [[TMP33:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]], i64 0, i64 0208// CHECK-64-NEXT: [[TMP34:%.*]] = load ptr, ptr [[TMP33]], align 8209// CHECK-64-NEXT: [[TMP35:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP4]], i64 0, i64 0210// CHECK-64-NEXT: [[TMP36:%.*]] = load ptr, ptr [[TMP35]], align 8211// CHECK-64-NEXT: [[TMP37:%.*]] = load double, ptr [[TMP34]], align 8212// CHECK-64-NEXT: store double [[TMP37]], ptr [[TMP36]], align 8213// CHECK-64-NEXT: br label [[IFCONT6:%.*]]214// CHECK-64: else5:215// CHECK-64-NEXT: br label [[IFCONT6]]216// CHECK-64: ifcont6:217// CHECK-64-NEXT: ret void218//219//220// CHECK-64-LABEL: define {{[^@]+}}@_omp_reduction_inter_warp_copy_func221// CHECK-64-SAME: (ptr noundef [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR2]] {222// CHECK-64-NEXT: entry:223// CHECK-64-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 8224// CHECK-64-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4225// CHECK-64-NEXT: [[DOTCNT_ADDR:%.*]] = alloca i32, align 4226// CHECK-64-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 8227// CHECK-64-NEXT: store i32 [[TMP1]], ptr [[DOTADDR1]], align 4228// CHECK-64-NEXT: [[TMP3:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()229// CHECK-64-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()230// CHECK-64-NEXT: [[NVPTX_LANE_ID:%.*]] = and i32 [[TMP4]], 31231// CHECK-64-NEXT: [[TMP5:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()232// CHECK-64-NEXT: [[NVPTX_WARP_ID:%.*]] = ashr i32 [[TMP5]], 5233// CHECK-64-NEXT: [[TMP6:%.*]] = load ptr, ptr [[DOTADDR]], align 8234// CHECK-64-NEXT: store i32 0, ptr [[DOTCNT_ADDR]], align 4235// CHECK-64-NEXT: br label [[PRECOND:%.*]]236// CHECK-64: precond:237// CHECK-64-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTCNT_ADDR]], align 4238// CHECK-64-NEXT: [[TMP8:%.*]] = icmp ult i32 [[TMP7]], 2239// CHECK-64-NEXT: br i1 [[TMP8]], label [[BODY:%.*]], label [[EXIT:%.*]]240// CHECK-64: body:241// CHECK-64-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])242// CHECK-64-NEXT: call void @__kmpc_barrier(ptr @[[GLOB2:[0-9]+]], i32 [[TMP2]])243// CHECK-64-NEXT: [[WARP_MASTER:%.*]] = icmp eq i32 [[NVPTX_LANE_ID]], 0244// CHECK-64-NEXT: br i1 [[WARP_MASTER]], label [[THEN:%.*]], label [[ELSE:%.*]]245// CHECK-64: then:246// CHECK-64-NEXT: [[TMP9:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP6]], i64 0, i64 0247// CHECK-64-NEXT: [[TMP10:%.*]] = load ptr, ptr [[TMP9]], align 8248// CHECK-64-NEXT: [[TMP11:%.*]] = getelementptr i32, ptr [[TMP10]], i32 [[TMP7]]249// CHECK-64-NEXT: [[TMP12:%.*]] = getelementptr inbounds [32 x i32], ptr addrspace(3) @__openmp_nvptx_data_transfer_temporary_storage, i64 0, i32 [[NVPTX_WARP_ID]]250// CHECK-64-NEXT: [[TMP13:%.*]] = load i32, ptr [[TMP11]], align 4251// CHECK-64-NEXT: store volatile i32 [[TMP13]], ptr addrspace(3) [[TMP12]], align 4252// CHECK-64-NEXT: br label [[IFCONT:%.*]]253// CHECK-64: else:254// CHECK-64-NEXT: br label [[IFCONT]]255// CHECK-64: ifcont:256// CHECK-64-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])257// CHECK-64-NEXT: call void @__kmpc_barrier(ptr @[[GLOB2]], i32 [[TMP2]])258// CHECK-64-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTADDR1]], align 4259// CHECK-64-NEXT: [[IS_ACTIVE_THREAD:%.*]] = icmp ult i32 [[TMP3]], [[TMP14]]260// CHECK-64-NEXT: br i1 [[IS_ACTIVE_THREAD]], label [[THEN2:%.*]], label [[ELSE3:%.*]]261// CHECK-64: then3:262// CHECK-64-NEXT: [[TMP15:%.*]] = getelementptr inbounds [32 x i32], ptr addrspace(3) @__openmp_nvptx_data_transfer_temporary_storage, i64 0, i32 [[TMP3]]263// CHECK-64-NEXT: [[TMP16:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP6]], i64 0, i64 0264// CHECK-64-NEXT: [[TMP17:%.*]] = load ptr, ptr [[TMP16]], align 8265// CHECK-64-NEXT: [[TMP18:%.*]] = getelementptr i32, ptr [[TMP17]], i32 [[TMP7]]266// CHECK-64-NEXT: [[TMP19:%.*]] = load volatile i32, ptr addrspace(3) [[TMP15]], align 4267// CHECK-64-NEXT: store i32 [[TMP19]], ptr [[TMP18]], align 4268// CHECK-64-NEXT: br label [[IFCONT4:%.*]]269// CHECK-64: else4:270// CHECK-64-NEXT: br label [[IFCONT4]]271// CHECK-64: ifcont5:272// CHECK-64-NEXT: [[TMP20:%.*]] = add nsw i32 [[TMP7]], 1273// CHECK-64-NEXT: store i32 [[TMP20]], ptr [[DOTCNT_ADDR]], align 4274// CHECK-64-NEXT: br label [[PRECOND]]275// CHECK-64: exit:276// CHECK-64-NEXT: ret void277//278//279// CHECK-64-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l29280// CHECK-64-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef nonnull align 1 dereferenceable(1) [[C:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[D:%.*]]) #[[ATTR0]] {281// CHECK-64-NEXT: entry:282// CHECK-64-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8283// CHECK-64-NEXT: [[C_ADDR:%.*]] = alloca ptr, align 8284// CHECK-64-NEXT: [[D_ADDR:%.*]] = alloca ptr, align 8285// CHECK-64-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [2 x ptr], align 8286// CHECK-64-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8287// CHECK-64-NEXT: store ptr [[C]], ptr [[C_ADDR]], align 8288// CHECK-64-NEXT: store ptr [[D]], ptr [[D_ADDR]], align 8289// CHECK-64-NEXT: [[TMP0:%.*]] = load ptr, ptr [[C_ADDR]], align 8290// CHECK-64-NEXT: [[TMP1:%.*]] = load ptr, ptr [[D_ADDR]], align 8291// CHECK-64-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l29_kernel_environment, ptr [[DYN_PTR]])292// CHECK-64-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP2]], -1293// CHECK-64-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]294// CHECK-64: user_code.entry:295// CHECK-64-NEXT: [[TMP3:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])296// CHECK-64-NEXT: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 0297// CHECK-64-NEXT: store ptr [[TMP0]], ptr [[TMP4]], align 8298// CHECK-64-NEXT: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 1299// CHECK-64-NEXT: store ptr [[TMP1]], ptr [[TMP5]], align 8300// CHECK-64-NEXT: call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP3]], i32 1, i32 -1, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l29_omp_outlined, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i64 2)301// CHECK-64-NEXT: call void @__kmpc_target_deinit()302// CHECK-64-NEXT: ret void303// CHECK-64: worker.exit:304// CHECK-64-NEXT: ret void305//306//307// CHECK-64-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l29_omp_outlined308// CHECK-64-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 1 dereferenceable(1) [[C:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[D:%.*]]) #[[ATTR1]] {309// CHECK-64-NEXT: entry:310// CHECK-64-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8311// CHECK-64-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8312// CHECK-64-NEXT: [[C_ADDR:%.*]] = alloca ptr, align 8313// CHECK-64-NEXT: [[D_ADDR:%.*]] = alloca ptr, align 8314// CHECK-64-NEXT: [[C1:%.*]] = alloca i8, align 1315// CHECK-64-NEXT: [[D2:%.*]] = alloca float, align 4316// CHECK-64-NEXT: [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [2 x ptr], align 8317// CHECK-64-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8318// CHECK-64-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8319// CHECK-64-NEXT: store ptr [[C]], ptr [[C_ADDR]], align 8320// CHECK-64-NEXT: store ptr [[D]], ptr [[D_ADDR]], align 8321// CHECK-64-NEXT: [[TMP0:%.*]] = load ptr, ptr [[C_ADDR]], align 8322// CHECK-64-NEXT: [[TMP1:%.*]] = load ptr, ptr [[D_ADDR]], align 8323// CHECK-64-NEXT: store i8 0, ptr [[C1]], align 1324// CHECK-64-NEXT: store float 1.000000e+00, ptr [[D2]], align 4325// CHECK-64-NEXT: [[TMP2:%.*]] = load i8, ptr [[C1]], align 1326// CHECK-64-NEXT: [[CONV:%.*]] = sext i8 [[TMP2]] to i32327// CHECK-64-NEXT: [[XOR:%.*]] = xor i32 [[CONV]], 2328// CHECK-64-NEXT: [[CONV3:%.*]] = trunc i32 [[XOR]] to i8329// CHECK-64-NEXT: store i8 [[CONV3]], ptr [[C1]], align 1330// CHECK-64-NEXT: [[TMP3:%.*]] = load float, ptr [[D2]], align 4331// CHECK-64-NEXT: [[MUL:%.*]] = fmul float [[TMP3]], 3.300000e+01332// CHECK-64-NEXT: store float [[MUL]], ptr [[D2]], align 4333// CHECK-64-NEXT: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i64 0, i64 0334// CHECK-64-NEXT: store ptr [[C1]], ptr [[TMP4]], align 8335// CHECK-64-NEXT: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i64 0, i64 1336// CHECK-64-NEXT: store ptr [[D2]], ptr [[TMP5]], align 8337// CHECK-64-NEXT: [[TMP6:%.*]] = call i32 @__kmpc_nvptx_parallel_reduce_nowait_v2(ptr @[[GLOB1]], i64 8, ptr [[DOTOMP_REDUCTION_RED_LIST]], ptr @_omp_reduction_shuffle_and_reduce_func1, ptr @_omp_reduction_inter_warp_copy_func2)338// CHECK-64-NEXT: [[TMP7:%.*]] = icmp eq i32 [[TMP6]], 1339// CHECK-64-NEXT: br i1 [[TMP7]], label [[DOTOMP_REDUCTION_THEN:%.*]], label [[DOTOMP_REDUCTION_DONE:%.*]]340// CHECK-64: .omp.reduction.then:341// CHECK-64-NEXT: [[TMP8:%.*]] = load i8, ptr [[TMP0]], align 1342// CHECK-64-NEXT: [[CONV4:%.*]] = sext i8 [[TMP8]] to i32343// CHECK-64-NEXT: [[TMP9:%.*]] = load i8, ptr [[C1]], align 1344// CHECK-64-NEXT: [[CONV5:%.*]] = sext i8 [[TMP9]] to i32345// CHECK-64-NEXT: [[XOR6:%.*]] = xor i32 [[CONV4]], [[CONV5]]346// CHECK-64-NEXT: [[CONV7:%.*]] = trunc i32 [[XOR6]] to i8347// CHECK-64-NEXT: store i8 [[CONV7]], ptr [[TMP0]], align 1348// CHECK-64-NEXT: [[TMP10:%.*]] = load float, ptr [[TMP1]], align 4349// CHECK-64-NEXT: [[TMP11:%.*]] = load float, ptr [[D2]], align 4350// CHECK-64-NEXT: [[MUL8:%.*]] = fmul float [[TMP10]], [[TMP11]]351// CHECK-64-NEXT: store float [[MUL8]], ptr [[TMP1]], align 4352// CHECK-64-NEXT: br label [[DOTOMP_REDUCTION_DONE]]353// CHECK-64: .omp.reduction.done:354// CHECK-64-NEXT: ret void355//356//357// CHECK-64-LABEL: define {{[^@]+}}@_omp_reduction_shuffle_and_reduce_func1358// CHECK-64-SAME: (ptr noundef [[TMP0:%.*]], i16 noundef signext [[TMP1:%.*]], i16 noundef signext [[TMP2:%.*]], i16 noundef signext [[TMP3:%.*]]) #[[ATTR2]] {359// CHECK-64-NEXT: entry:360// CHECK-64-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 8361// CHECK-64-NEXT: [[DOTADDR1:%.*]] = alloca i16, align 2362// CHECK-64-NEXT: [[DOTADDR2:%.*]] = alloca i16, align 2363// CHECK-64-NEXT: [[DOTADDR3:%.*]] = alloca i16, align 2364// CHECK-64-NEXT: [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST:%.*]] = alloca [2 x ptr], align 8365// CHECK-64-NEXT: [[DOTOMP_REDUCTION_ELEMENT:%.*]] = alloca i8, align 1366// CHECK-64-NEXT: [[DOTOMP_REDUCTION_ELEMENT4:%.*]] = alloca float, align 4367// CHECK-64-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 8368// CHECK-64-NEXT: store i16 [[TMP1]], ptr [[DOTADDR1]], align 2369// CHECK-64-NEXT: store i16 [[TMP2]], ptr [[DOTADDR2]], align 2370// CHECK-64-NEXT: store i16 [[TMP3]], ptr [[DOTADDR3]], align 2371// CHECK-64-NEXT: [[TMP4:%.*]] = load ptr, ptr [[DOTADDR]], align 8372// CHECK-64-NEXT: [[TMP5:%.*]] = load i16, ptr [[DOTADDR1]], align 2373// CHECK-64-NEXT: [[TMP6:%.*]] = load i16, ptr [[DOTADDR2]], align 2374// CHECK-64-NEXT: [[TMP7:%.*]] = load i16, ptr [[DOTADDR3]], align 2375// CHECK-64-NEXT: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP4]], i64 0, i64 0376// CHECK-64-NEXT: [[TMP9:%.*]] = load ptr, ptr [[TMP8]], align 8377// CHECK-64-NEXT: [[TMP10:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]], i64 0, i64 0378// CHECK-64-NEXT: [[TMP11:%.*]] = getelementptr i8, ptr [[TMP9]], i64 1379// CHECK-64-NEXT: [[TMP12:%.*]] = load i8, ptr [[TMP9]], align 1380// CHECK-64-NEXT: [[TMP13:%.*]] = sext i8 [[TMP12]] to i32381// CHECK-64-NEXT: [[TMP14:%.*]] = call i32 @__kmpc_get_warp_size()382// CHECK-64-NEXT: [[TMP15:%.*]] = trunc i32 [[TMP14]] to i16383// CHECK-64-NEXT: [[TMP16:%.*]] = call i32 @__kmpc_shuffle_int32(i32 [[TMP13]], i16 [[TMP6]], i16 [[TMP15]])384// CHECK-64-NEXT: [[TMP17:%.*]] = trunc i32 [[TMP16]] to i8385// CHECK-64-NEXT: store i8 [[TMP17]], ptr [[DOTOMP_REDUCTION_ELEMENT]], align 1386// CHECK-64-NEXT: [[TMP18:%.*]] = getelementptr i8, ptr [[TMP9]], i64 1387// CHECK-64-NEXT: [[TMP19:%.*]] = getelementptr i8, ptr [[DOTOMP_REDUCTION_ELEMENT]], i64 1388// CHECK-64-NEXT: store ptr [[DOTOMP_REDUCTION_ELEMENT]], ptr [[TMP10]], align 8389// CHECK-64-NEXT: [[TMP20:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP4]], i64 0, i64 1390// CHECK-64-NEXT: [[TMP21:%.*]] = load ptr, ptr [[TMP20]], align 8391// CHECK-64-NEXT: [[TMP22:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]], i64 0, i64 1392// CHECK-64-NEXT: [[TMP23:%.*]] = getelementptr float, ptr [[TMP21]], i64 1393// CHECK-64-NEXT: [[TMP24:%.*]] = load i32, ptr [[TMP21]], align 4394// CHECK-64-NEXT: [[TMP25:%.*]] = call i32 @__kmpc_get_warp_size()395// CHECK-64-NEXT: [[TMP26:%.*]] = trunc i32 [[TMP25]] to i16396// CHECK-64-NEXT: [[TMP27:%.*]] = call i32 @__kmpc_shuffle_int32(i32 [[TMP24]], i16 [[TMP6]], i16 [[TMP26]])397// CHECK-64-NEXT: store i32 [[TMP27]], ptr [[DOTOMP_REDUCTION_ELEMENT4]], align 4398// CHECK-64-NEXT: [[TMP28:%.*]] = getelementptr i32, ptr [[TMP21]], i64 1399// CHECK-64-NEXT: [[TMP29:%.*]] = getelementptr i32, ptr [[DOTOMP_REDUCTION_ELEMENT4]], i64 1400// CHECK-64-NEXT: store ptr [[DOTOMP_REDUCTION_ELEMENT4]], ptr [[TMP22]], align 8401// CHECK-64-NEXT: [[TMP30:%.*]] = icmp eq i16 [[TMP7]], 0402// CHECK-64-NEXT: [[TMP31:%.*]] = icmp eq i16 [[TMP7]], 1403// CHECK-64-NEXT: [[TMP32:%.*]] = icmp ult i16 [[TMP5]], [[TMP6]]404// CHECK-64-NEXT: [[TMP33:%.*]] = and i1 [[TMP31]], [[TMP32]]405// CHECK-64-NEXT: [[TMP34:%.*]] = icmp eq i16 [[TMP7]], 2406// CHECK-64-NEXT: [[TMP35:%.*]] = and i16 [[TMP5]], 1407// CHECK-64-NEXT: [[TMP36:%.*]] = icmp eq i16 [[TMP35]], 0408// CHECK-64-NEXT: [[TMP37:%.*]] = and i1 [[TMP34]], [[TMP36]]409// CHECK-64-NEXT: [[TMP38:%.*]] = icmp sgt i16 [[TMP6]], 0410// CHECK-64-NEXT: [[TMP39:%.*]] = and i1 [[TMP37]], [[TMP38]]411// CHECK-64-NEXT: [[TMP40:%.*]] = or i1 [[TMP30]], [[TMP33]]412// CHECK-64-NEXT: [[TMP41:%.*]] = or i1 [[TMP40]], [[TMP39]]413// CHECK-64-NEXT: br i1 [[TMP41]], label [[THEN:%.*]], label [[ELSE:%.*]]414// CHECK-64: then:415// CHECK-64-NEXT: call void @"{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l29_omp_outlined_omp$reduction$reduction_func"(ptr [[TMP4]], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]]) #[[ATTR3]]416// CHECK-64-NEXT: br label [[IFCONT:%.*]]417// CHECK-64: else:418// CHECK-64-NEXT: br label [[IFCONT]]419// CHECK-64: ifcont:420// CHECK-64-NEXT: [[TMP42:%.*]] = icmp eq i16 [[TMP7]], 1421// CHECK-64-NEXT: [[TMP43:%.*]] = icmp uge i16 [[TMP5]], [[TMP6]]422// CHECK-64-NEXT: [[TMP44:%.*]] = and i1 [[TMP42]], [[TMP43]]423// CHECK-64-NEXT: br i1 [[TMP44]], label [[THEN5:%.*]], label [[ELSE6:%.*]]424// CHECK-64: then5:425// CHECK-64-NEXT: [[TMP45:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]], i64 0, i64 0426// CHECK-64-NEXT: [[TMP46:%.*]] = load ptr, ptr [[TMP45]], align 8427// CHECK-64-NEXT: [[TMP47:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP4]], i64 0, i64 0428// CHECK-64-NEXT: [[TMP48:%.*]] = load ptr, ptr [[TMP47]], align 8429// CHECK-64-NEXT: [[TMP49:%.*]] = load i8, ptr [[TMP46]], align 1430// CHECK-64-NEXT: store i8 [[TMP49]], ptr [[TMP48]], align 1431// CHECK-64-NEXT: [[TMP50:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]], i64 0, i64 1432// CHECK-64-NEXT: [[TMP51:%.*]] = load ptr, ptr [[TMP50]], align 8433// CHECK-64-NEXT: [[TMP52:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP4]], i64 0, i64 1434// CHECK-64-NEXT: [[TMP53:%.*]] = load ptr, ptr [[TMP52]], align 8435// CHECK-64-NEXT: [[TMP54:%.*]] = load float, ptr [[TMP51]], align 4436// CHECK-64-NEXT: store float [[TMP54]], ptr [[TMP53]], align 4437// CHECK-64-NEXT: br label [[IFCONT7:%.*]]438// CHECK-64: else6:439// CHECK-64-NEXT: br label [[IFCONT7]]440// CHECK-64: ifcont7:441// CHECK-64-NEXT: ret void442//443//444// CHECK-64-LABEL: define {{[^@]+}}@_omp_reduction_inter_warp_copy_func2445// CHECK-64-SAME: (ptr noundef [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR2]] {446// CHECK-64-NEXT: entry:447// CHECK-64-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 8448// CHECK-64-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4449// CHECK-64-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 8450// CHECK-64-NEXT: store i32 [[TMP1]], ptr [[DOTADDR1]], align 4451// CHECK-64-NEXT: [[TMP3:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()452// CHECK-64-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()453// CHECK-64-NEXT: [[NVPTX_LANE_ID:%.*]] = and i32 [[TMP4]], 31454// CHECK-64-NEXT: [[TMP5:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()455// CHECK-64-NEXT: [[NVPTX_WARP_ID:%.*]] = ashr i32 [[TMP5]], 5456// CHECK-64-NEXT: [[TMP6:%.*]] = load ptr, ptr [[DOTADDR]], align 8457// CHECK-64-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])458// CHECK-64-NEXT: call void @__kmpc_barrier(ptr @[[GLOB2]], i32 [[TMP2]])459// CHECK-64-NEXT: [[WARP_MASTER:%.*]] = icmp eq i32 [[NVPTX_LANE_ID]], 0460// CHECK-64-NEXT: br i1 [[WARP_MASTER]], label [[THEN:%.*]], label [[ELSE:%.*]]461// CHECK-64: then:462// CHECK-64-NEXT: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP6]], i64 0, i64 0463// CHECK-64-NEXT: [[TMP8:%.*]] = load ptr, ptr [[TMP7]], align 8464// CHECK-64-NEXT: [[TMP9:%.*]] = getelementptr inbounds [32 x i32], ptr addrspace(3) @__openmp_nvptx_data_transfer_temporary_storage, i64 0, i32 [[NVPTX_WARP_ID]]465// CHECK-64-NEXT: [[TMP10:%.*]] = load i8, ptr [[TMP8]], align 1466// CHECK-64-NEXT: store volatile i8 [[TMP10]], ptr addrspace(3) [[TMP9]], align 1467// CHECK-64-NEXT: br label [[IFCONT:%.*]]468// CHECK-64: else:469// CHECK-64-NEXT: br label [[IFCONT]]470// CHECK-64: ifcont:471// CHECK-64-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])472// CHECK-64-NEXT: call void @__kmpc_barrier(ptr @[[GLOB2]], i32 [[TMP2]])473// CHECK-64-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTADDR1]], align 4474// CHECK-64-NEXT: [[IS_ACTIVE_THREAD:%.*]] = icmp ult i32 [[TMP3]], [[TMP11]]475// CHECK-64-NEXT: br i1 [[IS_ACTIVE_THREAD]], label [[THEN2:%.*]], label [[ELSE3:%.*]]476// CHECK-64: then3:477// CHECK-64-NEXT: [[TMP12:%.*]] = getelementptr inbounds [32 x i32], ptr addrspace(3) @__openmp_nvptx_data_transfer_temporary_storage, i64 0, i32 [[TMP3]]478// CHECK-64-NEXT: [[TMP13:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP6]], i64 0, i64 0479// CHECK-64-NEXT: [[TMP14:%.*]] = load ptr, ptr [[TMP13]], align 8480// CHECK-64-NEXT: [[TMP15:%.*]] = load volatile i8, ptr addrspace(3) [[TMP12]], align 1481// CHECK-64-NEXT: store i8 [[TMP15]], ptr [[TMP14]], align 1482// CHECK-64-NEXT: br label [[IFCONT4:%.*]]483// CHECK-64: else4:484// CHECK-64-NEXT: br label [[IFCONT4]]485// CHECK-64: ifcont5:486// CHECK-64-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])487// CHECK-64-NEXT: call void @__kmpc_barrier(ptr @[[GLOB2]], i32 [[TMP2]])488// CHECK-64-NEXT: [[WARP_MASTER5:%.*]] = icmp eq i32 [[NVPTX_LANE_ID]], 0489// CHECK-64-NEXT: br i1 [[WARP_MASTER5]], label [[THEN6:%.*]], label [[ELSE7:%.*]]490// CHECK-64: then8:491// CHECK-64-NEXT: [[TMP16:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP6]], i64 0, i64 1492// CHECK-64-NEXT: [[TMP17:%.*]] = load ptr, ptr [[TMP16]], align 8493// CHECK-64-NEXT: [[TMP18:%.*]] = getelementptr inbounds [32 x i32], ptr addrspace(3) @__openmp_nvptx_data_transfer_temporary_storage, i64 0, i32 [[NVPTX_WARP_ID]]494// CHECK-64-NEXT: [[TMP19:%.*]] = load i32, ptr [[TMP17]], align 4495// CHECK-64-NEXT: store volatile i32 [[TMP19]], ptr addrspace(3) [[TMP18]], align 4496// CHECK-64-NEXT: br label [[IFCONT8:%.*]]497// CHECK-64: else9:498// CHECK-64-NEXT: br label [[IFCONT8]]499// CHECK-64: ifcont10:500// CHECK-64-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])501// CHECK-64-NEXT: call void @__kmpc_barrier(ptr @[[GLOB2]], i32 [[TMP2]])502// CHECK-64-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTADDR1]], align 4503// CHECK-64-NEXT: [[IS_ACTIVE_THREAD9:%.*]] = icmp ult i32 [[TMP3]], [[TMP20]]504// CHECK-64-NEXT: br i1 [[IS_ACTIVE_THREAD9]], label [[THEN10:%.*]], label [[ELSE11:%.*]]505// CHECK-64: then13:506// CHECK-64-NEXT: [[TMP21:%.*]] = getelementptr inbounds [32 x i32], ptr addrspace(3) @__openmp_nvptx_data_transfer_temporary_storage, i64 0, i32 [[TMP3]]507// CHECK-64-NEXT: [[TMP22:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP6]], i64 0, i64 1508// CHECK-64-NEXT: [[TMP23:%.*]] = load ptr, ptr [[TMP22]], align 8509// CHECK-64-NEXT: [[TMP24:%.*]] = load volatile i32, ptr addrspace(3) [[TMP21]], align 4510// CHECK-64-NEXT: store i32 [[TMP24]], ptr [[TMP23]], align 4511// CHECK-64-NEXT: br label [[IFCONT12:%.*]]512// CHECK-64: else14:513// CHECK-64-NEXT: br label [[IFCONT12]]514// CHECK-64: ifcont15:515// CHECK-64-NEXT: ret void516//517//518// CHECK-64-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l35519// CHECK-64-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[B:%.*]]) #[[ATTR0]] {520// CHECK-64-NEXT: entry:521// CHECK-64-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8522// CHECK-64-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8523// CHECK-64-NEXT: [[B_ADDR:%.*]] = alloca ptr, align 8524// CHECK-64-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [2 x ptr], align 8525// CHECK-64-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8526// CHECK-64-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8527// CHECK-64-NEXT: store ptr [[B]], ptr [[B_ADDR]], align 8528// CHECK-64-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8529// CHECK-64-NEXT: [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 8530// CHECK-64-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l35_kernel_environment, ptr [[DYN_PTR]])531// CHECK-64-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP2]], -1532// CHECK-64-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]533// CHECK-64: user_code.entry:534// CHECK-64-NEXT: [[TMP3:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])535// CHECK-64-NEXT: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 0536// CHECK-64-NEXT: store ptr [[TMP0]], ptr [[TMP4]], align 8537// CHECK-64-NEXT: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 1538// CHECK-64-NEXT: store ptr [[TMP1]], ptr [[TMP5]], align 8539// CHECK-64-NEXT: call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP3]], i32 1, i32 -1, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l35_omp_outlined, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i64 2)540// CHECK-64-NEXT: call void @__kmpc_target_deinit()541// CHECK-64-NEXT: ret void542// CHECK-64: worker.exit:543// CHECK-64-NEXT: ret void544//545//546// CHECK-64-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l35_omp_outlined547// CHECK-64-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[B:%.*]]) #[[ATTR1]] {548// CHECK-64-NEXT: entry:549// CHECK-64-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8550// CHECK-64-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8551// CHECK-64-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8552// CHECK-64-NEXT: [[B_ADDR:%.*]] = alloca ptr, align 8553// CHECK-64-NEXT: [[A1:%.*]] = alloca i32, align 4554// CHECK-64-NEXT: [[B2:%.*]] = alloca i16, align 2555// CHECK-64-NEXT: [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [2 x ptr], align 8556// CHECK-64-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8557// CHECK-64-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8558// CHECK-64-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8559// CHECK-64-NEXT: store ptr [[B]], ptr [[B_ADDR]], align 8560// CHECK-64-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8561// CHECK-64-NEXT: [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 8562// CHECK-64-NEXT: store i32 0, ptr [[A1]], align 4563// CHECK-64-NEXT: store i16 -32768, ptr [[B2]], align 2564// CHECK-64-NEXT: [[TMP2:%.*]] = load i32, ptr [[A1]], align 4565// CHECK-64-NEXT: [[OR:%.*]] = or i32 [[TMP2]], 1566// CHECK-64-NEXT: store i32 [[OR]], ptr [[A1]], align 4567// CHECK-64-NEXT: [[TMP3:%.*]] = load i16, ptr [[B2]], align 2568// CHECK-64-NEXT: [[CONV:%.*]] = sext i16 [[TMP3]] to i32569// CHECK-64-NEXT: [[CMP:%.*]] = icmp sgt i32 99, [[CONV]]570// CHECK-64-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]571// CHECK-64: cond.true:572// CHECK-64-NEXT: br label [[COND_END:%.*]]573// CHECK-64: cond.false:574// CHECK-64-NEXT: [[TMP4:%.*]] = load i16, ptr [[B2]], align 2575// CHECK-64-NEXT: [[CONV3:%.*]] = sext i16 [[TMP4]] to i32576// CHECK-64-NEXT: br label [[COND_END]]577// CHECK-64: cond.end:578// CHECK-64-NEXT: [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[CONV3]], [[COND_FALSE]] ]579// CHECK-64-NEXT: [[CONV4:%.*]] = trunc i32 [[COND]] to i16580// CHECK-64-NEXT: store i16 [[CONV4]], ptr [[B2]], align 2581// CHECK-64-NEXT: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i64 0, i64 0582// CHECK-64-NEXT: store ptr [[A1]], ptr [[TMP5]], align 8583// CHECK-64-NEXT: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i64 0, i64 1584// CHECK-64-NEXT: store ptr [[B2]], ptr [[TMP6]], align 8585// CHECK-64-NEXT: [[TMP7:%.*]] = call i32 @__kmpc_nvptx_parallel_reduce_nowait_v2(ptr @[[GLOB1]], i64 8, ptr [[DOTOMP_REDUCTION_RED_LIST]], ptr @_omp_reduction_shuffle_and_reduce_func3, ptr @_omp_reduction_inter_warp_copy_func4)586// CHECK-64-NEXT: [[TMP8:%.*]] = icmp eq i32 [[TMP7]], 1587// CHECK-64-NEXT: br i1 [[TMP8]], label [[DOTOMP_REDUCTION_THEN:%.*]], label [[DOTOMP_REDUCTION_DONE:%.*]]588// CHECK-64: .omp.reduction.then:589// CHECK-64-NEXT: [[TMP9:%.*]] = load i32, ptr [[TMP0]], align 4590// CHECK-64-NEXT: [[TMP10:%.*]] = load i32, ptr [[A1]], align 4591// CHECK-64-NEXT: [[OR5:%.*]] = or i32 [[TMP9]], [[TMP10]]592// CHECK-64-NEXT: store i32 [[OR5]], ptr [[TMP0]], align 4593// CHECK-64-NEXT: [[TMP11:%.*]] = load i16, ptr [[TMP1]], align 2594// CHECK-64-NEXT: [[CONV6:%.*]] = sext i16 [[TMP11]] to i32595// CHECK-64-NEXT: [[TMP12:%.*]] = load i16, ptr [[B2]], align 2596// CHECK-64-NEXT: [[CONV7:%.*]] = sext i16 [[TMP12]] to i32597// CHECK-64-NEXT: [[CMP8:%.*]] = icmp sgt i32 [[CONV6]], [[CONV7]]598// CHECK-64-NEXT: br i1 [[CMP8]], label [[COND_TRUE9:%.*]], label [[COND_FALSE10:%.*]]599// CHECK-64: cond.true9:600// CHECK-64-NEXT: [[TMP13:%.*]] = load i16, ptr [[TMP1]], align 2601// CHECK-64-NEXT: br label [[COND_END11:%.*]]602// CHECK-64: cond.false10:603// CHECK-64-NEXT: [[TMP14:%.*]] = load i16, ptr [[B2]], align 2604// CHECK-64-NEXT: br label [[COND_END11]]605// CHECK-64: cond.end11:606// CHECK-64-NEXT: [[COND12:%.*]] = phi i16 [ [[TMP13]], [[COND_TRUE9]] ], [ [[TMP14]], [[COND_FALSE10]] ]607// CHECK-64-NEXT: store i16 [[COND12]], ptr [[TMP1]], align 2608// CHECK-64-NEXT: br label [[DOTOMP_REDUCTION_DONE]]609// CHECK-64: .omp.reduction.done:610// CHECK-64-NEXT: ret void611//612//613// CHECK-64-LABEL: define {{[^@]+}}@_omp_reduction_shuffle_and_reduce_func3614// CHECK-64-SAME: (ptr noundef [[TMP0:%.*]], i16 noundef signext [[TMP1:%.*]], i16 noundef signext [[TMP2:%.*]], i16 noundef signext [[TMP3:%.*]]) #[[ATTR2]] {615// CHECK-64-NEXT: entry:616// CHECK-64-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 8617// CHECK-64-NEXT: [[DOTADDR1:%.*]] = alloca i16, align 2618// CHECK-64-NEXT: [[DOTADDR2:%.*]] = alloca i16, align 2619// CHECK-64-NEXT: [[DOTADDR3:%.*]] = alloca i16, align 2620// CHECK-64-NEXT: [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST:%.*]] = alloca [2 x ptr], align 8621// CHECK-64-NEXT: [[DOTOMP_REDUCTION_ELEMENT:%.*]] = alloca i32, align 4622// CHECK-64-NEXT: [[DOTOMP_REDUCTION_ELEMENT4:%.*]] = alloca i16, align 2623// CHECK-64-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 8624// CHECK-64-NEXT: store i16 [[TMP1]], ptr [[DOTADDR1]], align 2625// CHECK-64-NEXT: store i16 [[TMP2]], ptr [[DOTADDR2]], align 2626// CHECK-64-NEXT: store i16 [[TMP3]], ptr [[DOTADDR3]], align 2627// CHECK-64-NEXT: [[TMP4:%.*]] = load ptr, ptr [[DOTADDR]], align 8628// CHECK-64-NEXT: [[TMP5:%.*]] = load i16, ptr [[DOTADDR1]], align 2629// CHECK-64-NEXT: [[TMP6:%.*]] = load i16, ptr [[DOTADDR2]], align 2630// CHECK-64-NEXT: [[TMP7:%.*]] = load i16, ptr [[DOTADDR3]], align 2631// CHECK-64-NEXT: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP4]], i64 0, i64 0632// CHECK-64-NEXT: [[TMP9:%.*]] = load ptr, ptr [[TMP8]], align 8633// CHECK-64-NEXT: [[TMP10:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]], i64 0, i64 0634// CHECK-64-NEXT: [[TMP11:%.*]] = getelementptr i32, ptr [[TMP9]], i64 1635// CHECK-64-NEXT: [[TMP12:%.*]] = load i32, ptr [[TMP9]], align 4636// CHECK-64-NEXT: [[TMP13:%.*]] = call i32 @__kmpc_get_warp_size()637// CHECK-64-NEXT: [[TMP14:%.*]] = trunc i32 [[TMP13]] to i16638// CHECK-64-NEXT: [[TMP15:%.*]] = call i32 @__kmpc_shuffle_int32(i32 [[TMP12]], i16 [[TMP6]], i16 [[TMP14]])639// CHECK-64-NEXT: store i32 [[TMP15]], ptr [[DOTOMP_REDUCTION_ELEMENT]], align 4640// CHECK-64-NEXT: [[TMP16:%.*]] = getelementptr i32, ptr [[TMP9]], i64 1641// CHECK-64-NEXT: [[TMP17:%.*]] = getelementptr i32, ptr [[DOTOMP_REDUCTION_ELEMENT]], i64 1642// CHECK-64-NEXT: store ptr [[DOTOMP_REDUCTION_ELEMENT]], ptr [[TMP10]], align 8643// CHECK-64-NEXT: [[TMP18:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP4]], i64 0, i64 1644// CHECK-64-NEXT: [[TMP19:%.*]] = load ptr, ptr [[TMP18]], align 8645// CHECK-64-NEXT: [[TMP20:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]], i64 0, i64 1646// CHECK-64-NEXT: [[TMP21:%.*]] = getelementptr i16, ptr [[TMP19]], i64 1647// CHECK-64-NEXT: [[TMP22:%.*]] = load i16, ptr [[TMP19]], align 2648// CHECK-64-NEXT: [[TMP23:%.*]] = sext i16 [[TMP22]] to i32649// CHECK-64-NEXT: [[TMP24:%.*]] = call i32 @__kmpc_get_warp_size()650// CHECK-64-NEXT: [[TMP25:%.*]] = trunc i32 [[TMP24]] to i16651// CHECK-64-NEXT: [[TMP26:%.*]] = call i32 @__kmpc_shuffle_int32(i32 [[TMP23]], i16 [[TMP6]], i16 [[TMP25]])652// CHECK-64-NEXT: [[TMP27:%.*]] = trunc i32 [[TMP26]] to i16653// CHECK-64-NEXT: store i16 [[TMP27]], ptr [[DOTOMP_REDUCTION_ELEMENT4]], align 2654// CHECK-64-NEXT: [[TMP28:%.*]] = getelementptr i16, ptr [[TMP19]], i64 1655// CHECK-64-NEXT: [[TMP29:%.*]] = getelementptr i16, ptr [[DOTOMP_REDUCTION_ELEMENT4]], i64 1656// CHECK-64-NEXT: store ptr [[DOTOMP_REDUCTION_ELEMENT4]], ptr [[TMP20]], align 8657// CHECK-64-NEXT: [[TMP30:%.*]] = icmp eq i16 [[TMP7]], 0658// CHECK-64-NEXT: [[TMP31:%.*]] = icmp eq i16 [[TMP7]], 1659// CHECK-64-NEXT: [[TMP32:%.*]] = icmp ult i16 [[TMP5]], [[TMP6]]660// CHECK-64-NEXT: [[TMP33:%.*]] = and i1 [[TMP31]], [[TMP32]]661// CHECK-64-NEXT: [[TMP34:%.*]] = icmp eq i16 [[TMP7]], 2662// CHECK-64-NEXT: [[TMP35:%.*]] = and i16 [[TMP5]], 1663// CHECK-64-NEXT: [[TMP36:%.*]] = icmp eq i16 [[TMP35]], 0664// CHECK-64-NEXT: [[TMP37:%.*]] = and i1 [[TMP34]], [[TMP36]]665// CHECK-64-NEXT: [[TMP38:%.*]] = icmp sgt i16 [[TMP6]], 0666// CHECK-64-NEXT: [[TMP39:%.*]] = and i1 [[TMP37]], [[TMP38]]667// CHECK-64-NEXT: [[TMP40:%.*]] = or i1 [[TMP30]], [[TMP33]]668// CHECK-64-NEXT: [[TMP41:%.*]] = or i1 [[TMP40]], [[TMP39]]669// CHECK-64-NEXT: br i1 [[TMP41]], label [[THEN:%.*]], label [[ELSE:%.*]]670// CHECK-64: then:671// CHECK-64-NEXT: call void @"{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l35_omp_outlined_omp$reduction$reduction_func"(ptr [[TMP4]], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]]) #[[ATTR3]]672// CHECK-64-NEXT: br label [[IFCONT:%.*]]673// CHECK-64: else:674// CHECK-64-NEXT: br label [[IFCONT]]675// CHECK-64: ifcont:676// CHECK-64-NEXT: [[TMP42:%.*]] = icmp eq i16 [[TMP7]], 1677// CHECK-64-NEXT: [[TMP43:%.*]] = icmp uge i16 [[TMP5]], [[TMP6]]678// CHECK-64-NEXT: [[TMP44:%.*]] = and i1 [[TMP42]], [[TMP43]]679// CHECK-64-NEXT: br i1 [[TMP44]], label [[THEN5:%.*]], label [[ELSE6:%.*]]680// CHECK-64: then5:681// CHECK-64-NEXT: [[TMP45:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]], i64 0, i64 0682// CHECK-64-NEXT: [[TMP46:%.*]] = load ptr, ptr [[TMP45]], align 8683// CHECK-64-NEXT: [[TMP47:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP4]], i64 0, i64 0684// CHECK-64-NEXT: [[TMP48:%.*]] = load ptr, ptr [[TMP47]], align 8685// CHECK-64-NEXT: [[TMP49:%.*]] = load i32, ptr [[TMP46]], align 4686// CHECK-64-NEXT: store i32 [[TMP49]], ptr [[TMP48]], align 4687// CHECK-64-NEXT: [[TMP50:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]], i64 0, i64 1688// CHECK-64-NEXT: [[TMP51:%.*]] = load ptr, ptr [[TMP50]], align 8689// CHECK-64-NEXT: [[TMP52:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP4]], i64 0, i64 1690// CHECK-64-NEXT: [[TMP53:%.*]] = load ptr, ptr [[TMP52]], align 8691// CHECK-64-NEXT: [[TMP54:%.*]] = load i16, ptr [[TMP51]], align 2692// CHECK-64-NEXT: store i16 [[TMP54]], ptr [[TMP53]], align 2693// CHECK-64-NEXT: br label [[IFCONT7:%.*]]694// CHECK-64: else6:695// CHECK-64-NEXT: br label [[IFCONT7]]696// CHECK-64: ifcont7:697// CHECK-64-NEXT: ret void698//699//700// CHECK-64-LABEL: define {{[^@]+}}@_omp_reduction_inter_warp_copy_func4701// CHECK-64-SAME: (ptr noundef [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR2]] {702// CHECK-64-NEXT: entry:703// CHECK-64-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 8704// CHECK-64-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4705// CHECK-64-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 8706// CHECK-64-NEXT: store i32 [[TMP1]], ptr [[DOTADDR1]], align 4707// CHECK-64-NEXT: [[TMP3:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()708// CHECK-64-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()709// CHECK-64-NEXT: [[NVPTX_LANE_ID:%.*]] = and i32 [[TMP4]], 31710// CHECK-64-NEXT: [[TMP5:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()711// CHECK-64-NEXT: [[NVPTX_WARP_ID:%.*]] = ashr i32 [[TMP5]], 5712// CHECK-64-NEXT: [[TMP6:%.*]] = load ptr, ptr [[DOTADDR]], align 8713// CHECK-64-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])714// CHECK-64-NEXT: call void @__kmpc_barrier(ptr @[[GLOB2]], i32 [[TMP2]])715// CHECK-64-NEXT: [[WARP_MASTER:%.*]] = icmp eq i32 [[NVPTX_LANE_ID]], 0716// CHECK-64-NEXT: br i1 [[WARP_MASTER]], label [[THEN:%.*]], label [[ELSE:%.*]]717// CHECK-64: then:718// CHECK-64-NEXT: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP6]], i64 0, i64 0719// CHECK-64-NEXT: [[TMP8:%.*]] = load ptr, ptr [[TMP7]], align 8720// CHECK-64-NEXT: [[TMP9:%.*]] = getelementptr inbounds [32 x i32], ptr addrspace(3) @__openmp_nvptx_data_transfer_temporary_storage, i64 0, i32 [[NVPTX_WARP_ID]]721// CHECK-64-NEXT: [[TMP10:%.*]] = load i32, ptr [[TMP8]], align 4722// CHECK-64-NEXT: store volatile i32 [[TMP10]], ptr addrspace(3) [[TMP9]], align 4723// CHECK-64-NEXT: br label [[IFCONT:%.*]]724// CHECK-64: else:725// CHECK-64-NEXT: br label [[IFCONT]]726// CHECK-64: ifcont:727// CHECK-64-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])728// CHECK-64-NEXT: call void @__kmpc_barrier(ptr @[[GLOB2]], i32 [[TMP2]])729// CHECK-64-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTADDR1]], align 4730// CHECK-64-NEXT: [[IS_ACTIVE_THREAD:%.*]] = icmp ult i32 [[TMP3]], [[TMP11]]731// CHECK-64-NEXT: br i1 [[IS_ACTIVE_THREAD]], label [[THEN2:%.*]], label [[ELSE3:%.*]]732// CHECK-64: then3:733// CHECK-64-NEXT: [[TMP12:%.*]] = getelementptr inbounds [32 x i32], ptr addrspace(3) @__openmp_nvptx_data_transfer_temporary_storage, i64 0, i32 [[TMP3]]734// CHECK-64-NEXT: [[TMP13:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP6]], i64 0, i64 0735// CHECK-64-NEXT: [[TMP14:%.*]] = load ptr, ptr [[TMP13]], align 8736// CHECK-64-NEXT: [[TMP15:%.*]] = load volatile i32, ptr addrspace(3) [[TMP12]], align 4737// CHECK-64-NEXT: store i32 [[TMP15]], ptr [[TMP14]], align 4738// CHECK-64-NEXT: br label [[IFCONT4:%.*]]739// CHECK-64: else4:740// CHECK-64-NEXT: br label [[IFCONT4]]741// CHECK-64: ifcont5:742// CHECK-64-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])743// CHECK-64-NEXT: call void @__kmpc_barrier(ptr @[[GLOB2]], i32 [[TMP2]])744// CHECK-64-NEXT: [[WARP_MASTER5:%.*]] = icmp eq i32 [[NVPTX_LANE_ID]], 0745// CHECK-64-NEXT: br i1 [[WARP_MASTER5]], label [[THEN6:%.*]], label [[ELSE7:%.*]]746// CHECK-64: then8:747// CHECK-64-NEXT: [[TMP16:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP6]], i64 0, i64 1748// CHECK-64-NEXT: [[TMP17:%.*]] = load ptr, ptr [[TMP16]], align 8749// CHECK-64-NEXT: [[TMP18:%.*]] = getelementptr inbounds [32 x i32], ptr addrspace(3) @__openmp_nvptx_data_transfer_temporary_storage, i64 0, i32 [[NVPTX_WARP_ID]]750// CHECK-64-NEXT: [[TMP19:%.*]] = load i16, ptr [[TMP17]], align 2751// CHECK-64-NEXT: store volatile i16 [[TMP19]], ptr addrspace(3) [[TMP18]], align 2752// CHECK-64-NEXT: br label [[IFCONT8:%.*]]753// CHECK-64: else9:754// CHECK-64-NEXT: br label [[IFCONT8]]755// CHECK-64: ifcont10:756// CHECK-64-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])757// CHECK-64-NEXT: call void @__kmpc_barrier(ptr @[[GLOB2]], i32 [[TMP2]])758// CHECK-64-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTADDR1]], align 4759// CHECK-64-NEXT: [[IS_ACTIVE_THREAD9:%.*]] = icmp ult i32 [[TMP3]], [[TMP20]]760// CHECK-64-NEXT: br i1 [[IS_ACTIVE_THREAD9]], label [[THEN10:%.*]], label [[ELSE11:%.*]]761// CHECK-64: then13:762// CHECK-64-NEXT: [[TMP21:%.*]] = getelementptr inbounds [32 x i32], ptr addrspace(3) @__openmp_nvptx_data_transfer_temporary_storage, i64 0, i32 [[TMP3]]763// CHECK-64-NEXT: [[TMP22:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP6]], i64 0, i64 1764// CHECK-64-NEXT: [[TMP23:%.*]] = load ptr, ptr [[TMP22]], align 8765// CHECK-64-NEXT: [[TMP24:%.*]] = load volatile i16, ptr addrspace(3) [[TMP21]], align 2766// CHECK-64-NEXT: store i16 [[TMP24]], ptr [[TMP23]], align 2767// CHECK-64-NEXT: br label [[IFCONT12:%.*]]768// CHECK-64: else14:769// CHECK-64-NEXT: br label [[IFCONT12]]770// CHECK-64: ifcont15:771// CHECK-64-NEXT: ret void772//773//774// CHECK-32-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l24775// CHECK-32-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[E:%.*]]) #[[ATTR0:[0-9]+]] {776// CHECK-32-NEXT: entry:777// CHECK-32-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4778// CHECK-32-NEXT: [[E_ADDR:%.*]] = alloca ptr, align 4779// CHECK-32-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x ptr], align 4780// CHECK-32-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4781// CHECK-32-NEXT: store ptr [[E]], ptr [[E_ADDR]], align 4782// CHECK-32-NEXT: [[TMP0:%.*]] = load ptr, ptr [[E_ADDR]], align 4783// CHECK-32-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l24_kernel_environment, ptr [[DYN_PTR]])784// CHECK-32-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1785// CHECK-32-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]786// CHECK-32: user_code.entry:787// CHECK-32-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1:[0-9]+]])788// CHECK-32-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 0789// CHECK-32-NEXT: store ptr [[TMP0]], ptr [[TMP3]], align 4790// CHECK-32-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]+}}__Z9ftemplateIcET_i_l24_omp_outlined, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i32 1)791// CHECK-32-NEXT: call void @__kmpc_target_deinit()792// CHECK-32-NEXT: ret void793// CHECK-32: worker.exit:794// CHECK-32-NEXT: ret void795//796//797// CHECK-32-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l24_omp_outlined798// CHECK-32-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[E:%.*]]) #[[ATTR1:[0-9]+]] {799// CHECK-32-NEXT: entry:800// CHECK-32-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4801// CHECK-32-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4802// CHECK-32-NEXT: [[E_ADDR:%.*]] = alloca ptr, align 4803// CHECK-32-NEXT: [[E1:%.*]] = alloca double, align 8804// CHECK-32-NEXT: [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 4805// CHECK-32-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4806// CHECK-32-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4807// CHECK-32-NEXT: store ptr [[E]], ptr [[E_ADDR]], align 4808// CHECK-32-NEXT: [[TMP0:%.*]] = load ptr, ptr [[E_ADDR]], align 4809// CHECK-32-NEXT: store double 0.000000e+00, ptr [[E1]], align 8810// CHECK-32-NEXT: [[TMP1:%.*]] = load double, ptr [[E1]], align 8811// CHECK-32-NEXT: [[ADD:%.*]] = fadd double [[TMP1]], 5.000000e+00812// CHECK-32-NEXT: store double [[ADD]], ptr [[E1]], align 8813// CHECK-32-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i32 0, i32 0814// CHECK-32-NEXT: store ptr [[E1]], ptr [[TMP2]], align 4815// CHECK-32-NEXT: [[TMP3:%.*]] = call i32 @__kmpc_nvptx_parallel_reduce_nowait_v2(ptr @[[GLOB1]], i64 8, ptr [[DOTOMP_REDUCTION_RED_LIST]], ptr @_omp_reduction_shuffle_and_reduce_func, ptr @_omp_reduction_inter_warp_copy_func)816// CHECK-32-NEXT: [[TMP4:%.*]] = icmp eq i32 [[TMP3]], 1817// CHECK-32-NEXT: br i1 [[TMP4]], label [[DOTOMP_REDUCTION_THEN:%.*]], label [[DOTOMP_REDUCTION_DONE:%.*]]818// CHECK-32: .omp.reduction.then:819// CHECK-32-NEXT: [[TMP5:%.*]] = load double, ptr [[TMP0]], align 8820// CHECK-32-NEXT: [[TMP6:%.*]] = load double, ptr [[E1]], align 8821// CHECK-32-NEXT: [[ADD2:%.*]] = fadd double [[TMP5]], [[TMP6]]822// CHECK-32-NEXT: store double [[ADD2]], ptr [[TMP0]], align 8823// CHECK-32-NEXT: br label [[DOTOMP_REDUCTION_DONE]]824// CHECK-32: .omp.reduction.done:825// CHECK-32-NEXT: ret void826//827//828// CHECK-32-LABEL: define {{[^@]+}}@_omp_reduction_shuffle_and_reduce_func829// CHECK-32-SAME: (ptr noundef [[TMP0:%.*]], i16 noundef signext [[TMP1:%.*]], i16 noundef signext [[TMP2:%.*]], i16 noundef signext [[TMP3:%.*]]) #[[ATTR2:[0-9]+]] {830// CHECK-32-NEXT: entry:831// CHECK-32-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 4832// CHECK-32-NEXT: [[DOTADDR1:%.*]] = alloca i16, align 2833// CHECK-32-NEXT: [[DOTADDR2:%.*]] = alloca i16, align 2834// CHECK-32-NEXT: [[DOTADDR3:%.*]] = alloca i16, align 2835// CHECK-32-NEXT: [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST:%.*]] = alloca [1 x ptr], align 4836// CHECK-32-NEXT: [[DOTOMP_REDUCTION_ELEMENT:%.*]] = alloca double, align 8837// CHECK-32-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 4838// CHECK-32-NEXT: store i16 [[TMP1]], ptr [[DOTADDR1]], align 2839// CHECK-32-NEXT: store i16 [[TMP2]], ptr [[DOTADDR2]], align 2840// CHECK-32-NEXT: store i16 [[TMP3]], ptr [[DOTADDR3]], align 2841// CHECK-32-NEXT: [[TMP4:%.*]] = load ptr, ptr [[DOTADDR]], align 4842// CHECK-32-NEXT: [[TMP5:%.*]] = load i16, ptr [[DOTADDR1]], align 2843// CHECK-32-NEXT: [[TMP6:%.*]] = load i16, ptr [[DOTADDR2]], align 2844// CHECK-32-NEXT: [[TMP7:%.*]] = load i16, ptr [[DOTADDR3]], align 2845// CHECK-32-NEXT: [[TMP8:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP4]], i32 0, i32 0846// CHECK-32-NEXT: [[TMP9:%.*]] = load ptr, ptr [[TMP8]], align 4847// CHECK-32-NEXT: [[TMP10:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]], i32 0, i32 0848// CHECK-32-NEXT: [[TMP11:%.*]] = getelementptr double, ptr [[TMP9]], i32 1849// CHECK-32-NEXT: [[TMP12:%.*]] = load i64, ptr [[TMP9]], align 8850// CHECK-32-NEXT: [[TMP13:%.*]] = call i32 @__kmpc_get_warp_size()851// CHECK-32-NEXT: [[TMP14:%.*]] = trunc i32 [[TMP13]] to i16852// CHECK-32-NEXT: [[TMP15:%.*]] = call i64 @__kmpc_shuffle_int64(i64 [[TMP12]], i16 [[TMP6]], i16 [[TMP14]])853// CHECK-32-NEXT: store i64 [[TMP15]], ptr [[DOTOMP_REDUCTION_ELEMENT]], align 8854// CHECK-32-NEXT: [[TMP16:%.*]] = getelementptr i64, ptr [[TMP9]], i32 1855// CHECK-32-NEXT: [[TMP17:%.*]] = getelementptr i64, ptr [[DOTOMP_REDUCTION_ELEMENT]], i32 1856// CHECK-32-NEXT: store ptr [[DOTOMP_REDUCTION_ELEMENT]], ptr [[TMP10]], align 4857// CHECK-32-NEXT: [[TMP18:%.*]] = icmp eq i16 [[TMP7]], 0858// CHECK-32-NEXT: [[TMP19:%.*]] = icmp eq i16 [[TMP7]], 1859// CHECK-32-NEXT: [[TMP20:%.*]] = icmp ult i16 [[TMP5]], [[TMP6]]860// CHECK-32-NEXT: [[TMP21:%.*]] = and i1 [[TMP19]], [[TMP20]]861// CHECK-32-NEXT: [[TMP22:%.*]] = icmp eq i16 [[TMP7]], 2862// CHECK-32-NEXT: [[TMP23:%.*]] = and i16 [[TMP5]], 1863// CHECK-32-NEXT: [[TMP24:%.*]] = icmp eq i16 [[TMP23]], 0864// CHECK-32-NEXT: [[TMP25:%.*]] = and i1 [[TMP22]], [[TMP24]]865// CHECK-32-NEXT: [[TMP26:%.*]] = icmp sgt i16 [[TMP6]], 0866// CHECK-32-NEXT: [[TMP27:%.*]] = and i1 [[TMP25]], [[TMP26]]867// CHECK-32-NEXT: [[TMP28:%.*]] = or i1 [[TMP18]], [[TMP21]]868// CHECK-32-NEXT: [[TMP29:%.*]] = or i1 [[TMP28]], [[TMP27]]869// CHECK-32-NEXT: br i1 [[TMP29]], label [[THEN:%.*]], label [[ELSE:%.*]]870// CHECK-32: then:871// CHECK-32-NEXT: call void @"{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l24_omp_outlined_omp$reduction$reduction_func"(ptr [[TMP4]], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]]) #[[ATTR3:[0-9]+]]872// CHECK-32-NEXT: br label [[IFCONT:%.*]]873// CHECK-32: else:874// CHECK-32-NEXT: br label [[IFCONT]]875// CHECK-32: ifcont:876// CHECK-32-NEXT: [[TMP30:%.*]] = icmp eq i16 [[TMP7]], 1877// CHECK-32-NEXT: [[TMP31:%.*]] = icmp uge i16 [[TMP5]], [[TMP6]]878// CHECK-32-NEXT: [[TMP32:%.*]] = and i1 [[TMP30]], [[TMP31]]879// CHECK-32-NEXT: br i1 [[TMP32]], label [[THEN4:%.*]], label [[ELSE5:%.*]]880// CHECK-32: then4:881// CHECK-32-NEXT: [[TMP33:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]], i32 0, i32 0882// CHECK-32-NEXT: [[TMP34:%.*]] = load ptr, ptr [[TMP33]], align 4883// CHECK-32-NEXT: [[TMP35:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP4]], i32 0, i32 0884// CHECK-32-NEXT: [[TMP36:%.*]] = load ptr, ptr [[TMP35]], align 4885// CHECK-32-NEXT: [[TMP37:%.*]] = load double, ptr [[TMP34]], align 8886// CHECK-32-NEXT: store double [[TMP37]], ptr [[TMP36]], align 8887// CHECK-32-NEXT: br label [[IFCONT6:%.*]]888// CHECK-32: else5:889// CHECK-32-NEXT: br label [[IFCONT6]]890// CHECK-32: ifcont6:891// CHECK-32-NEXT: ret void892//893//894// CHECK-32-LABEL: define {{[^@]+}}@_omp_reduction_inter_warp_copy_func895// CHECK-32-SAME: (ptr noundef [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR2]] {896// CHECK-32-NEXT: entry:897// CHECK-32-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 4898// CHECK-32-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4899// CHECK-32-NEXT: [[DOTCNT_ADDR:%.*]] = alloca i32, align 4900// CHECK-32-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 4901// CHECK-32-NEXT: store i32 [[TMP1]], ptr [[DOTADDR1]], align 4902// CHECK-32-NEXT: [[TMP3:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()903// CHECK-32-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()904// CHECK-32-NEXT: [[NVPTX_LANE_ID:%.*]] = and i32 [[TMP4]], 31905// CHECK-32-NEXT: [[TMP5:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()906// CHECK-32-NEXT: [[NVPTX_WARP_ID:%.*]] = ashr i32 [[TMP5]], 5907// CHECK-32-NEXT: [[TMP6:%.*]] = load ptr, ptr [[DOTADDR]], align 4908// CHECK-32-NEXT: store i32 0, ptr [[DOTCNT_ADDR]], align 4909// CHECK-32-NEXT: br label [[PRECOND:%.*]]910// CHECK-32: precond:911// CHECK-32-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTCNT_ADDR]], align 4912// CHECK-32-NEXT: [[TMP8:%.*]] = icmp ult i32 [[TMP7]], 2913// CHECK-32-NEXT: br i1 [[TMP8]], label [[BODY:%.*]], label [[EXIT:%.*]]914// CHECK-32: body:915// CHECK-32-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])916// CHECK-32-NEXT: call void @__kmpc_barrier(ptr @[[GLOB2:[0-9]+]], i32 [[TMP2]])917// CHECK-32-NEXT: [[WARP_MASTER:%.*]] = icmp eq i32 [[NVPTX_LANE_ID]], 0918// CHECK-32-NEXT: br i1 [[WARP_MASTER]], label [[THEN:%.*]], label [[ELSE:%.*]]919// CHECK-32: then:920// CHECK-32-NEXT: [[TMP9:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP6]], i32 0, i32 0921// CHECK-32-NEXT: [[TMP10:%.*]] = load ptr, ptr [[TMP9]], align 4922// CHECK-32-NEXT: [[TMP11:%.*]] = getelementptr i32, ptr [[TMP10]], i32 [[TMP7]]923// CHECK-32-NEXT: [[TMP12:%.*]] = getelementptr inbounds [32 x i32], ptr addrspace(3) @__openmp_nvptx_data_transfer_temporary_storage, i64 0, i32 [[NVPTX_WARP_ID]]924// CHECK-32-NEXT: [[TMP13:%.*]] = load i32, ptr [[TMP11]], align 4925// CHECK-32-NEXT: store volatile i32 [[TMP13]], ptr addrspace(3) [[TMP12]], align 4926// CHECK-32-NEXT: br label [[IFCONT:%.*]]927// CHECK-32: else:928// CHECK-32-NEXT: br label [[IFCONT]]929// CHECK-32: ifcont:930// CHECK-32-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])931// CHECK-32-NEXT: call void @__kmpc_barrier(ptr @[[GLOB2]], i32 [[TMP2]])932// CHECK-32-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTADDR1]], align 4933// CHECK-32-NEXT: [[IS_ACTIVE_THREAD:%.*]] = icmp ult i32 [[TMP3]], [[TMP14]]934// CHECK-32-NEXT: br i1 [[IS_ACTIVE_THREAD]], label [[THEN2:%.*]], label [[ELSE3:%.*]]935// CHECK-32: then3:936// CHECK-32-NEXT: [[TMP15:%.*]] = getelementptr inbounds [32 x i32], ptr addrspace(3) @__openmp_nvptx_data_transfer_temporary_storage, i64 0, i32 [[TMP3]]937// CHECK-32-NEXT: [[TMP16:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP6]], i32 0, i32 0938// CHECK-32-NEXT: [[TMP17:%.*]] = load ptr, ptr [[TMP16]], align 4939// CHECK-32-NEXT: [[TMP18:%.*]] = getelementptr i32, ptr [[TMP17]], i32 [[TMP7]]940// CHECK-32-NEXT: [[TMP19:%.*]] = load volatile i32, ptr addrspace(3) [[TMP15]], align 4941// CHECK-32-NEXT: store i32 [[TMP19]], ptr [[TMP18]], align 4942// CHECK-32-NEXT: br label [[IFCONT4:%.*]]943// CHECK-32: else4:944// CHECK-32-NEXT: br label [[IFCONT4]]945// CHECK-32: ifcont5:946// CHECK-32-NEXT: [[TMP20:%.*]] = add nsw i32 [[TMP7]], 1947// CHECK-32-NEXT: store i32 [[TMP20]], ptr [[DOTCNT_ADDR]], align 4948// CHECK-32-NEXT: br label [[PRECOND]]949// CHECK-32: exit:950// CHECK-32-NEXT: ret void951//952//953// CHECK-32-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l29954// CHECK-32-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef nonnull align 1 dereferenceable(1) [[C:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[D:%.*]]) #[[ATTR0]] {955// CHECK-32-NEXT: entry:956// CHECK-32-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4957// CHECK-32-NEXT: [[C_ADDR:%.*]] = alloca ptr, align 4958// CHECK-32-NEXT: [[D_ADDR:%.*]] = alloca ptr, align 4959// CHECK-32-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [2 x ptr], align 4960// CHECK-32-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4961// CHECK-32-NEXT: store ptr [[C]], ptr [[C_ADDR]], align 4962// CHECK-32-NEXT: store ptr [[D]], ptr [[D_ADDR]], align 4963// CHECK-32-NEXT: [[TMP0:%.*]] = load ptr, ptr [[C_ADDR]], align 4964// CHECK-32-NEXT: [[TMP1:%.*]] = load ptr, ptr [[D_ADDR]], align 4965// CHECK-32-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l29_kernel_environment, ptr [[DYN_PTR]])966// CHECK-32-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP2]], -1967// CHECK-32-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]968// CHECK-32: user_code.entry:969// CHECK-32-NEXT: [[TMP3:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])970// CHECK-32-NEXT: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 0971// CHECK-32-NEXT: store ptr [[TMP0]], ptr [[TMP4]], align 4972// CHECK-32-NEXT: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 1973// CHECK-32-NEXT: store ptr [[TMP1]], ptr [[TMP5]], align 4974// CHECK-32-NEXT: call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP3]], i32 1, i32 -1, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l29_omp_outlined, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i32 2)975// CHECK-32-NEXT: call void @__kmpc_target_deinit()976// CHECK-32-NEXT: ret void977// CHECK-32: worker.exit:978// CHECK-32-NEXT: ret void979//980//981// CHECK-32-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l29_omp_outlined982// CHECK-32-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 1 dereferenceable(1) [[C:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[D:%.*]]) #[[ATTR1]] {983// CHECK-32-NEXT: entry:984// CHECK-32-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4985// CHECK-32-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4986// CHECK-32-NEXT: [[C_ADDR:%.*]] = alloca ptr, align 4987// CHECK-32-NEXT: [[D_ADDR:%.*]] = alloca ptr, align 4988// CHECK-32-NEXT: [[C1:%.*]] = alloca i8, align 1989// CHECK-32-NEXT: [[D2:%.*]] = alloca float, align 4990// CHECK-32-NEXT: [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [2 x ptr], align 4991// CHECK-32-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4992// CHECK-32-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4993// CHECK-32-NEXT: store ptr [[C]], ptr [[C_ADDR]], align 4994// CHECK-32-NEXT: store ptr [[D]], ptr [[D_ADDR]], align 4995// CHECK-32-NEXT: [[TMP0:%.*]] = load ptr, ptr [[C_ADDR]], align 4996// CHECK-32-NEXT: [[TMP1:%.*]] = load ptr, ptr [[D_ADDR]], align 4997// CHECK-32-NEXT: store i8 0, ptr [[C1]], align 1998// CHECK-32-NEXT: store float 1.000000e+00, ptr [[D2]], align 4999// CHECK-32-NEXT: [[TMP2:%.*]] = load i8, ptr [[C1]], align 11000// CHECK-32-NEXT: [[CONV:%.*]] = sext i8 [[TMP2]] to i321001// CHECK-32-NEXT: [[XOR:%.*]] = xor i32 [[CONV]], 21002// CHECK-32-NEXT: [[CONV3:%.*]] = trunc i32 [[XOR]] to i81003// CHECK-32-NEXT: store i8 [[CONV3]], ptr [[C1]], align 11004// CHECK-32-NEXT: [[TMP3:%.*]] = load float, ptr [[D2]], align 41005// CHECK-32-NEXT: [[MUL:%.*]] = fmul float [[TMP3]], 3.300000e+011006// CHECK-32-NEXT: store float [[MUL]], ptr [[D2]], align 41007// CHECK-32-NEXT: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i32 0, i32 01008// CHECK-32-NEXT: store ptr [[C1]], ptr [[TMP4]], align 41009// CHECK-32-NEXT: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i32 0, i32 11010// CHECK-32-NEXT: store ptr [[D2]], ptr [[TMP5]], align 41011// CHECK-32-NEXT: [[TMP6:%.*]] = call i32 @__kmpc_nvptx_parallel_reduce_nowait_v2(ptr @[[GLOB1]], i64 8, ptr [[DOTOMP_REDUCTION_RED_LIST]], ptr @_omp_reduction_shuffle_and_reduce_func1, ptr @_omp_reduction_inter_warp_copy_func2)1012// CHECK-32-NEXT: [[TMP7:%.*]] = icmp eq i32 [[TMP6]], 11013// CHECK-32-NEXT: br i1 [[TMP7]], label [[DOTOMP_REDUCTION_THEN:%.*]], label [[DOTOMP_REDUCTION_DONE:%.*]]1014// CHECK-32: .omp.reduction.then:1015// CHECK-32-NEXT: [[TMP8:%.*]] = load i8, ptr [[TMP0]], align 11016// CHECK-32-NEXT: [[CONV4:%.*]] = sext i8 [[TMP8]] to i321017// CHECK-32-NEXT: [[TMP9:%.*]] = load i8, ptr [[C1]], align 11018// CHECK-32-NEXT: [[CONV5:%.*]] = sext i8 [[TMP9]] to i321019// CHECK-32-NEXT: [[XOR6:%.*]] = xor i32 [[CONV4]], [[CONV5]]1020// CHECK-32-NEXT: [[CONV7:%.*]] = trunc i32 [[XOR6]] to i81021// CHECK-32-NEXT: store i8 [[CONV7]], ptr [[TMP0]], align 11022// CHECK-32-NEXT: [[TMP10:%.*]] = load float, ptr [[TMP1]], align 41023// CHECK-32-NEXT: [[TMP11:%.*]] = load float, ptr [[D2]], align 41024// CHECK-32-NEXT: [[MUL8:%.*]] = fmul float [[TMP10]], [[TMP11]]1025// CHECK-32-NEXT: store float [[MUL8]], ptr [[TMP1]], align 41026// CHECK-32-NEXT: br label [[DOTOMP_REDUCTION_DONE]]1027// CHECK-32: .omp.reduction.done:1028// CHECK-32-NEXT: ret void1029//1030//1031// CHECK-32-LABEL: define {{[^@]+}}@_omp_reduction_shuffle_and_reduce_func11032// CHECK-32-SAME: (ptr noundef [[TMP0:%.*]], i16 noundef signext [[TMP1:%.*]], i16 noundef signext [[TMP2:%.*]], i16 noundef signext [[TMP3:%.*]]) #[[ATTR2]] {1033// CHECK-32-NEXT: entry:1034// CHECK-32-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 41035// CHECK-32-NEXT: [[DOTADDR1:%.*]] = alloca i16, align 21036// CHECK-32-NEXT: [[DOTADDR2:%.*]] = alloca i16, align 21037// CHECK-32-NEXT: [[DOTADDR3:%.*]] = alloca i16, align 21038// CHECK-32-NEXT: [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST:%.*]] = alloca [2 x ptr], align 41039// CHECK-32-NEXT: [[DOTOMP_REDUCTION_ELEMENT:%.*]] = alloca i8, align 11040// CHECK-32-NEXT: [[DOTOMP_REDUCTION_ELEMENT4:%.*]] = alloca float, align 41041// CHECK-32-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 41042// CHECK-32-NEXT: store i16 [[TMP1]], ptr [[DOTADDR1]], align 21043// CHECK-32-NEXT: store i16 [[TMP2]], ptr [[DOTADDR2]], align 21044// CHECK-32-NEXT: store i16 [[TMP3]], ptr [[DOTADDR3]], align 21045// CHECK-32-NEXT: [[TMP4:%.*]] = load ptr, ptr [[DOTADDR]], align 41046// CHECK-32-NEXT: [[TMP5:%.*]] = load i16, ptr [[DOTADDR1]], align 21047// CHECK-32-NEXT: [[TMP6:%.*]] = load i16, ptr [[DOTADDR2]], align 21048// CHECK-32-NEXT: [[TMP7:%.*]] = load i16, ptr [[DOTADDR3]], align 21049// CHECK-32-NEXT: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP4]], i32 0, i32 01050// CHECK-32-NEXT: [[TMP9:%.*]] = load ptr, ptr [[TMP8]], align 41051// CHECK-32-NEXT: [[TMP10:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]], i32 0, i32 01052// CHECK-32-NEXT: [[TMP11:%.*]] = getelementptr i8, ptr [[TMP9]], i32 11053// CHECK-32-NEXT: [[TMP12:%.*]] = load i8, ptr [[TMP9]], align 11054// CHECK-32-NEXT: [[TMP13:%.*]] = sext i8 [[TMP12]] to i321055// CHECK-32-NEXT: [[TMP14:%.*]] = call i32 @__kmpc_get_warp_size()1056// CHECK-32-NEXT: [[TMP15:%.*]] = trunc i32 [[TMP14]] to i161057// CHECK-32-NEXT: [[TMP16:%.*]] = call i32 @__kmpc_shuffle_int32(i32 [[TMP13]], i16 [[TMP6]], i16 [[TMP15]])1058// CHECK-32-NEXT: [[TMP17:%.*]] = trunc i32 [[TMP16]] to i81059// CHECK-32-NEXT: store i8 [[TMP17]], ptr [[DOTOMP_REDUCTION_ELEMENT]], align 11060// CHECK-32-NEXT: [[TMP18:%.*]] = getelementptr i8, ptr [[TMP9]], i32 11061// CHECK-32-NEXT: [[TMP19:%.*]] = getelementptr i8, ptr [[DOTOMP_REDUCTION_ELEMENT]], i32 11062// CHECK-32-NEXT: store ptr [[DOTOMP_REDUCTION_ELEMENT]], ptr [[TMP10]], align 41063// CHECK-32-NEXT: [[TMP20:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP4]], i32 0, i32 11064// CHECK-32-NEXT: [[TMP21:%.*]] = load ptr, ptr [[TMP20]], align 41065// CHECK-32-NEXT: [[TMP22:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]], i32 0, i32 11066// CHECK-32-NEXT: [[TMP23:%.*]] = getelementptr float, ptr [[TMP21]], i32 11067// CHECK-32-NEXT: [[TMP24:%.*]] = load i32, ptr [[TMP21]], align 41068// CHECK-32-NEXT: [[TMP25:%.*]] = call i32 @__kmpc_get_warp_size()1069// CHECK-32-NEXT: [[TMP26:%.*]] = trunc i32 [[TMP25]] to i161070// CHECK-32-NEXT: [[TMP27:%.*]] = call i32 @__kmpc_shuffle_int32(i32 [[TMP24]], i16 [[TMP6]], i16 [[TMP26]])1071// CHECK-32-NEXT: store i32 [[TMP27]], ptr [[DOTOMP_REDUCTION_ELEMENT4]], align 41072// CHECK-32-NEXT: [[TMP28:%.*]] = getelementptr i32, ptr [[TMP21]], i32 11073// CHECK-32-NEXT: [[TMP29:%.*]] = getelementptr i32, ptr [[DOTOMP_REDUCTION_ELEMENT4]], i32 11074// CHECK-32-NEXT: store ptr [[DOTOMP_REDUCTION_ELEMENT4]], ptr [[TMP22]], align 41075// CHECK-32-NEXT: [[TMP30:%.*]] = icmp eq i16 [[TMP7]], 01076// CHECK-32-NEXT: [[TMP31:%.*]] = icmp eq i16 [[TMP7]], 11077// CHECK-32-NEXT: [[TMP32:%.*]] = icmp ult i16 [[TMP5]], [[TMP6]]1078// CHECK-32-NEXT: [[TMP33:%.*]] = and i1 [[TMP31]], [[TMP32]]1079// CHECK-32-NEXT: [[TMP34:%.*]] = icmp eq i16 [[TMP7]], 21080// CHECK-32-NEXT: [[TMP35:%.*]] = and i16 [[TMP5]], 11081// CHECK-32-NEXT: [[TMP36:%.*]] = icmp eq i16 [[TMP35]], 01082// CHECK-32-NEXT: [[TMP37:%.*]] = and i1 [[TMP34]], [[TMP36]]1083// CHECK-32-NEXT: [[TMP38:%.*]] = icmp sgt i16 [[TMP6]], 01084// CHECK-32-NEXT: [[TMP39:%.*]] = and i1 [[TMP37]], [[TMP38]]1085// CHECK-32-NEXT: [[TMP40:%.*]] = or i1 [[TMP30]], [[TMP33]]1086// CHECK-32-NEXT: [[TMP41:%.*]] = or i1 [[TMP40]], [[TMP39]]1087// CHECK-32-NEXT: br i1 [[TMP41]], label [[THEN:%.*]], label [[ELSE:%.*]]1088// CHECK-32: then:1089// CHECK-32-NEXT: call void @"{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l29_omp_outlined_omp$reduction$reduction_func"(ptr [[TMP4]], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]]) #[[ATTR3]]1090// CHECK-32-NEXT: br label [[IFCONT:%.*]]1091// CHECK-32: else:1092// CHECK-32-NEXT: br label [[IFCONT]]1093// CHECK-32: ifcont:1094// CHECK-32-NEXT: [[TMP42:%.*]] = icmp eq i16 [[TMP7]], 11095// CHECK-32-NEXT: [[TMP43:%.*]] = icmp uge i16 [[TMP5]], [[TMP6]]1096// CHECK-32-NEXT: [[TMP44:%.*]] = and i1 [[TMP42]], [[TMP43]]1097// CHECK-32-NEXT: br i1 [[TMP44]], label [[THEN5:%.*]], label [[ELSE6:%.*]]1098// CHECK-32: then5:1099// CHECK-32-NEXT: [[TMP45:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]], i32 0, i32 01100// CHECK-32-NEXT: [[TMP46:%.*]] = load ptr, ptr [[TMP45]], align 41101// CHECK-32-NEXT: [[TMP47:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP4]], i32 0, i32 01102// CHECK-32-NEXT: [[TMP48:%.*]] = load ptr, ptr [[TMP47]], align 41103// CHECK-32-NEXT: [[TMP49:%.*]] = load i8, ptr [[TMP46]], align 11104// CHECK-32-NEXT: store i8 [[TMP49]], ptr [[TMP48]], align 11105// CHECK-32-NEXT: [[TMP50:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]], i32 0, i32 11106// CHECK-32-NEXT: [[TMP51:%.*]] = load ptr, ptr [[TMP50]], align 41107// CHECK-32-NEXT: [[TMP52:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP4]], i32 0, i32 11108// CHECK-32-NEXT: [[TMP53:%.*]] = load ptr, ptr [[TMP52]], align 41109// CHECK-32-NEXT: [[TMP54:%.*]] = load float, ptr [[TMP51]], align 41110// CHECK-32-NEXT: store float [[TMP54]], ptr [[TMP53]], align 41111// CHECK-32-NEXT: br label [[IFCONT7:%.*]]1112// CHECK-32: else6:1113// CHECK-32-NEXT: br label [[IFCONT7]]1114// CHECK-32: ifcont7:1115// CHECK-32-NEXT: ret void1116//1117//1118// CHECK-32-LABEL: define {{[^@]+}}@_omp_reduction_inter_warp_copy_func21119// CHECK-32-SAME: (ptr noundef [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR2]] {1120// CHECK-32-NEXT: entry:1121// CHECK-32-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 41122// CHECK-32-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 41123// CHECK-32-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 41124// CHECK-32-NEXT: store i32 [[TMP1]], ptr [[DOTADDR1]], align 41125// CHECK-32-NEXT: [[TMP3:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()1126// CHECK-32-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()1127// CHECK-32-NEXT: [[NVPTX_LANE_ID:%.*]] = and i32 [[TMP4]], 311128// CHECK-32-NEXT: [[TMP5:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()1129// CHECK-32-NEXT: [[NVPTX_WARP_ID:%.*]] = ashr i32 [[TMP5]], 51130// CHECK-32-NEXT: [[TMP6:%.*]] = load ptr, ptr [[DOTADDR]], align 41131// CHECK-32-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1132// CHECK-32-NEXT: call void @__kmpc_barrier(ptr @[[GLOB2]], i32 [[TMP2]])1133// CHECK-32-NEXT: [[WARP_MASTER:%.*]] = icmp eq i32 [[NVPTX_LANE_ID]], 01134// CHECK-32-NEXT: br i1 [[WARP_MASTER]], label [[THEN:%.*]], label [[ELSE:%.*]]1135// CHECK-32: then:1136// CHECK-32-NEXT: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP6]], i32 0, i32 01137// CHECK-32-NEXT: [[TMP8:%.*]] = load ptr, ptr [[TMP7]], align 41138// CHECK-32-NEXT: [[TMP9:%.*]] = getelementptr inbounds [32 x i32], ptr addrspace(3) @__openmp_nvptx_data_transfer_temporary_storage, i64 0, i32 [[NVPTX_WARP_ID]]1139// CHECK-32-NEXT: [[TMP10:%.*]] = load i8, ptr [[TMP8]], align 11140// CHECK-32-NEXT: store volatile i8 [[TMP10]], ptr addrspace(3) [[TMP9]], align 11141// CHECK-32-NEXT: br label [[IFCONT:%.*]]1142// CHECK-32: else:1143// CHECK-32-NEXT: br label [[IFCONT]]1144// CHECK-32: ifcont:1145// CHECK-32-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1146// CHECK-32-NEXT: call void @__kmpc_barrier(ptr @[[GLOB2]], i32 [[TMP2]])1147// CHECK-32-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTADDR1]], align 41148// CHECK-32-NEXT: [[IS_ACTIVE_THREAD:%.*]] = icmp ult i32 [[TMP3]], [[TMP11]]1149// CHECK-32-NEXT: br i1 [[IS_ACTIVE_THREAD]], label [[THEN2:%.*]], label [[ELSE3:%.*]]1150// CHECK-32: then3:1151// CHECK-32-NEXT: [[TMP12:%.*]] = getelementptr inbounds [32 x i32], ptr addrspace(3) @__openmp_nvptx_data_transfer_temporary_storage, i64 0, i32 [[TMP3]]1152// CHECK-32-NEXT: [[TMP13:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP6]], i32 0, i32 01153// CHECK-32-NEXT: [[TMP14:%.*]] = load ptr, ptr [[TMP13]], align 41154// CHECK-32-NEXT: [[TMP15:%.*]] = load volatile i8, ptr addrspace(3) [[TMP12]], align 11155// CHECK-32-NEXT: store i8 [[TMP15]], ptr [[TMP14]], align 11156// CHECK-32-NEXT: br label [[IFCONT4:%.*]]1157// CHECK-32: else4:1158// CHECK-32-NEXT: br label [[IFCONT4]]1159// CHECK-32: ifcont5:1160// CHECK-32-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1161// CHECK-32-NEXT: call void @__kmpc_barrier(ptr @[[GLOB2]], i32 [[TMP2]])1162// CHECK-32-NEXT: [[WARP_MASTER5:%.*]] = icmp eq i32 [[NVPTX_LANE_ID]], 01163// CHECK-32-NEXT: br i1 [[WARP_MASTER5]], label [[THEN6:%.*]], label [[ELSE7:%.*]]1164// CHECK-32: then8:1165// CHECK-32-NEXT: [[TMP16:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP6]], i32 0, i32 11166// CHECK-32-NEXT: [[TMP17:%.*]] = load ptr, ptr [[TMP16]], align 41167// CHECK-32-NEXT: [[TMP18:%.*]] = getelementptr inbounds [32 x i32], ptr addrspace(3) @__openmp_nvptx_data_transfer_temporary_storage, i64 0, i32 [[NVPTX_WARP_ID]]1168// CHECK-32-NEXT: [[TMP19:%.*]] = load i32, ptr [[TMP17]], align 41169// CHECK-32-NEXT: store volatile i32 [[TMP19]], ptr addrspace(3) [[TMP18]], align 41170// CHECK-32-NEXT: br label [[IFCONT8:%.*]]1171// CHECK-32: else9:1172// CHECK-32-NEXT: br label [[IFCONT8]]1173// CHECK-32: ifcont10:1174// CHECK-32-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1175// CHECK-32-NEXT: call void @__kmpc_barrier(ptr @[[GLOB2]], i32 [[TMP2]])1176// CHECK-32-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTADDR1]], align 41177// CHECK-32-NEXT: [[IS_ACTIVE_THREAD9:%.*]] = icmp ult i32 [[TMP3]], [[TMP20]]1178// CHECK-32-NEXT: br i1 [[IS_ACTIVE_THREAD9]], label [[THEN10:%.*]], label [[ELSE11:%.*]]1179// CHECK-32: then13:1180// CHECK-32-NEXT: [[TMP21:%.*]] = getelementptr inbounds [32 x i32], ptr addrspace(3) @__openmp_nvptx_data_transfer_temporary_storage, i64 0, i32 [[TMP3]]1181// CHECK-32-NEXT: [[TMP22:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP6]], i32 0, i32 11182// CHECK-32-NEXT: [[TMP23:%.*]] = load ptr, ptr [[TMP22]], align 41183// CHECK-32-NEXT: [[TMP24:%.*]] = load volatile i32, ptr addrspace(3) [[TMP21]], align 41184// CHECK-32-NEXT: store i32 [[TMP24]], ptr [[TMP23]], align 41185// CHECK-32-NEXT: br label [[IFCONT12:%.*]]1186// CHECK-32: else14:1187// CHECK-32-NEXT: br label [[IFCONT12]]1188// CHECK-32: ifcont15:1189// CHECK-32-NEXT: ret void1190//1191//1192// CHECK-32-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l351193// CHECK-32-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[B:%.*]]) #[[ATTR0]] {1194// CHECK-32-NEXT: entry:1195// CHECK-32-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 41196// CHECK-32-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 41197// CHECK-32-NEXT: [[B_ADDR:%.*]] = alloca ptr, align 41198// CHECK-32-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [2 x ptr], align 41199// CHECK-32-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 41200// CHECK-32-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 41201// CHECK-32-NEXT: store ptr [[B]], ptr [[B_ADDR]], align 41202// CHECK-32-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 41203// CHECK-32-NEXT: [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 41204// CHECK-32-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l35_kernel_environment, ptr [[DYN_PTR]])1205// CHECK-32-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP2]], -11206// CHECK-32-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]1207// CHECK-32: user_code.entry:1208// CHECK-32-NEXT: [[TMP3:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1209// CHECK-32-NEXT: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 01210// CHECK-32-NEXT: store ptr [[TMP0]], ptr [[TMP4]], align 41211// CHECK-32-NEXT: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 11212// CHECK-32-NEXT: store ptr [[TMP1]], ptr [[TMP5]], align 41213// CHECK-32-NEXT: call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP3]], i32 1, i32 -1, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l35_omp_outlined, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i32 2)1214// CHECK-32-NEXT: call void @__kmpc_target_deinit()1215// CHECK-32-NEXT: ret void1216// CHECK-32: worker.exit:1217// CHECK-32-NEXT: ret void1218//1219//1220// CHECK-32-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l35_omp_outlined1221// CHECK-32-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[B:%.*]]) #[[ATTR1]] {1222// CHECK-32-NEXT: entry:1223// CHECK-32-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41224// CHECK-32-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41225// CHECK-32-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 41226// CHECK-32-NEXT: [[B_ADDR:%.*]] = alloca ptr, align 41227// CHECK-32-NEXT: [[A1:%.*]] = alloca i32, align 41228// CHECK-32-NEXT: [[B2:%.*]] = alloca i16, align 21229// CHECK-32-NEXT: [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [2 x ptr], align 41230// CHECK-32-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41231// CHECK-32-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41232// CHECK-32-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 41233// CHECK-32-NEXT: store ptr [[B]], ptr [[B_ADDR]], align 41234// CHECK-32-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 41235// CHECK-32-NEXT: [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 41236// CHECK-32-NEXT: store i32 0, ptr [[A1]], align 41237// CHECK-32-NEXT: store i16 -32768, ptr [[B2]], align 21238// CHECK-32-NEXT: [[TMP2:%.*]] = load i32, ptr [[A1]], align 41239// CHECK-32-NEXT: [[OR:%.*]] = or i32 [[TMP2]], 11240// CHECK-32-NEXT: store i32 [[OR]], ptr [[A1]], align 41241// CHECK-32-NEXT: [[TMP3:%.*]] = load i16, ptr [[B2]], align 21242// CHECK-32-NEXT: [[CONV:%.*]] = sext i16 [[TMP3]] to i321243// CHECK-32-NEXT: [[CMP:%.*]] = icmp sgt i32 99, [[CONV]]1244// CHECK-32-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1245// CHECK-32: cond.true:1246// CHECK-32-NEXT: br label [[COND_END:%.*]]1247// CHECK-32: cond.false:1248// CHECK-32-NEXT: [[TMP4:%.*]] = load i16, ptr [[B2]], align 21249// CHECK-32-NEXT: [[CONV3:%.*]] = sext i16 [[TMP4]] to i321250// CHECK-32-NEXT: br label [[COND_END]]1251// CHECK-32: cond.end:1252// CHECK-32-NEXT: [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[CONV3]], [[COND_FALSE]] ]1253// CHECK-32-NEXT: [[CONV4:%.*]] = trunc i32 [[COND]] to i161254// CHECK-32-NEXT: store i16 [[CONV4]], ptr [[B2]], align 21255// CHECK-32-NEXT: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i32 0, i32 01256// CHECK-32-NEXT: store ptr [[A1]], ptr [[TMP5]], align 41257// CHECK-32-NEXT: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i32 0, i32 11258// CHECK-32-NEXT: store ptr [[B2]], ptr [[TMP6]], align 41259// CHECK-32-NEXT: [[TMP7:%.*]] = call i32 @__kmpc_nvptx_parallel_reduce_nowait_v2(ptr @[[GLOB1]], i64 8, ptr [[DOTOMP_REDUCTION_RED_LIST]], ptr @_omp_reduction_shuffle_and_reduce_func3, ptr @_omp_reduction_inter_warp_copy_func4)1260// CHECK-32-NEXT: [[TMP8:%.*]] = icmp eq i32 [[TMP7]], 11261// CHECK-32-NEXT: br i1 [[TMP8]], label [[DOTOMP_REDUCTION_THEN:%.*]], label [[DOTOMP_REDUCTION_DONE:%.*]]1262// CHECK-32: .omp.reduction.then:1263// CHECK-32-NEXT: [[TMP9:%.*]] = load i32, ptr [[TMP0]], align 41264// CHECK-32-NEXT: [[TMP10:%.*]] = load i32, ptr [[A1]], align 41265// CHECK-32-NEXT: [[OR5:%.*]] = or i32 [[TMP9]], [[TMP10]]1266// CHECK-32-NEXT: store i32 [[OR5]], ptr [[TMP0]], align 41267// CHECK-32-NEXT: [[TMP11:%.*]] = load i16, ptr [[TMP1]], align 21268// CHECK-32-NEXT: [[CONV6:%.*]] = sext i16 [[TMP11]] to i321269// CHECK-32-NEXT: [[TMP12:%.*]] = load i16, ptr [[B2]], align 21270// CHECK-32-NEXT: [[CONV7:%.*]] = sext i16 [[TMP12]] to i321271// CHECK-32-NEXT: [[CMP8:%.*]] = icmp sgt i32 [[CONV6]], [[CONV7]]1272// CHECK-32-NEXT: br i1 [[CMP8]], label [[COND_TRUE9:%.*]], label [[COND_FALSE10:%.*]]1273// CHECK-32: cond.true9:1274// CHECK-32-NEXT: [[TMP13:%.*]] = load i16, ptr [[TMP1]], align 21275// CHECK-32-NEXT: br label [[COND_END11:%.*]]1276// CHECK-32: cond.false10:1277// CHECK-32-NEXT: [[TMP14:%.*]] = load i16, ptr [[B2]], align 21278// CHECK-32-NEXT: br label [[COND_END11]]1279// CHECK-32: cond.end11:1280// CHECK-32-NEXT: [[COND12:%.*]] = phi i16 [ [[TMP13]], [[COND_TRUE9]] ], [ [[TMP14]], [[COND_FALSE10]] ]1281// CHECK-32-NEXT: store i16 [[COND12]], ptr [[TMP1]], align 21282// CHECK-32-NEXT: br label [[DOTOMP_REDUCTION_DONE]]1283// CHECK-32: .omp.reduction.done:1284// CHECK-32-NEXT: ret void1285//1286//1287// CHECK-32-LABEL: define {{[^@]+}}@_omp_reduction_shuffle_and_reduce_func31288// CHECK-32-SAME: (ptr noundef [[TMP0:%.*]], i16 noundef signext [[TMP1:%.*]], i16 noundef signext [[TMP2:%.*]], i16 noundef signext [[TMP3:%.*]]) #[[ATTR2]] {1289// CHECK-32-NEXT: entry:1290// CHECK-32-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 41291// CHECK-32-NEXT: [[DOTADDR1:%.*]] = alloca i16, align 21292// CHECK-32-NEXT: [[DOTADDR2:%.*]] = alloca i16, align 21293// CHECK-32-NEXT: [[DOTADDR3:%.*]] = alloca i16, align 21294// CHECK-32-NEXT: [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST:%.*]] = alloca [2 x ptr], align 41295// CHECK-32-NEXT: [[DOTOMP_REDUCTION_ELEMENT:%.*]] = alloca i32, align 41296// CHECK-32-NEXT: [[DOTOMP_REDUCTION_ELEMENT4:%.*]] = alloca i16, align 21297// CHECK-32-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 41298// CHECK-32-NEXT: store i16 [[TMP1]], ptr [[DOTADDR1]], align 21299// CHECK-32-NEXT: store i16 [[TMP2]], ptr [[DOTADDR2]], align 21300// CHECK-32-NEXT: store i16 [[TMP3]], ptr [[DOTADDR3]], align 21301// CHECK-32-NEXT: [[TMP4:%.*]] = load ptr, ptr [[DOTADDR]], align 41302// CHECK-32-NEXT: [[TMP5:%.*]] = load i16, ptr [[DOTADDR1]], align 21303// CHECK-32-NEXT: [[TMP6:%.*]] = load i16, ptr [[DOTADDR2]], align 21304// CHECK-32-NEXT: [[TMP7:%.*]] = load i16, ptr [[DOTADDR3]], align 21305// CHECK-32-NEXT: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP4]], i32 0, i32 01306// CHECK-32-NEXT: [[TMP9:%.*]] = load ptr, ptr [[TMP8]], align 41307// CHECK-32-NEXT: [[TMP10:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]], i32 0, i32 01308// CHECK-32-NEXT: [[TMP11:%.*]] = getelementptr i32, ptr [[TMP9]], i32 11309// CHECK-32-NEXT: [[TMP12:%.*]] = load i32, ptr [[TMP9]], align 41310// CHECK-32-NEXT: [[TMP13:%.*]] = call i32 @__kmpc_get_warp_size()1311// CHECK-32-NEXT: [[TMP14:%.*]] = trunc i32 [[TMP13]] to i161312// CHECK-32-NEXT: [[TMP15:%.*]] = call i32 @__kmpc_shuffle_int32(i32 [[TMP12]], i16 [[TMP6]], i16 [[TMP14]])1313// CHECK-32-NEXT: store i32 [[TMP15]], ptr [[DOTOMP_REDUCTION_ELEMENT]], align 41314// CHECK-32-NEXT: [[TMP16:%.*]] = getelementptr i32, ptr [[TMP9]], i32 11315// CHECK-32-NEXT: [[TMP17:%.*]] = getelementptr i32, ptr [[DOTOMP_REDUCTION_ELEMENT]], i32 11316// CHECK-32-NEXT: store ptr [[DOTOMP_REDUCTION_ELEMENT]], ptr [[TMP10]], align 41317// CHECK-32-NEXT: [[TMP18:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP4]], i32 0, i32 11318// CHECK-32-NEXT: [[TMP19:%.*]] = load ptr, ptr [[TMP18]], align 41319// CHECK-32-NEXT: [[TMP20:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]], i32 0, i32 11320// CHECK-32-NEXT: [[TMP21:%.*]] = getelementptr i16, ptr [[TMP19]], i32 11321// CHECK-32-NEXT: [[TMP22:%.*]] = load i16, ptr [[TMP19]], align 21322// CHECK-32-NEXT: [[TMP23:%.*]] = sext i16 [[TMP22]] to i321323// CHECK-32-NEXT: [[TMP24:%.*]] = call i32 @__kmpc_get_warp_size()1324// CHECK-32-NEXT: [[TMP25:%.*]] = trunc i32 [[TMP24]] to i161325// CHECK-32-NEXT: [[TMP26:%.*]] = call i32 @__kmpc_shuffle_int32(i32 [[TMP23]], i16 [[TMP6]], i16 [[TMP25]])1326// CHECK-32-NEXT: [[TMP27:%.*]] = trunc i32 [[TMP26]] to i161327// CHECK-32-NEXT: store i16 [[TMP27]], ptr [[DOTOMP_REDUCTION_ELEMENT4]], align 21328// CHECK-32-NEXT: [[TMP28:%.*]] = getelementptr i16, ptr [[TMP19]], i32 11329// CHECK-32-NEXT: [[TMP29:%.*]] = getelementptr i16, ptr [[DOTOMP_REDUCTION_ELEMENT4]], i32 11330// CHECK-32-NEXT: store ptr [[DOTOMP_REDUCTION_ELEMENT4]], ptr [[TMP20]], align 41331// CHECK-32-NEXT: [[TMP30:%.*]] = icmp eq i16 [[TMP7]], 01332// CHECK-32-NEXT: [[TMP31:%.*]] = icmp eq i16 [[TMP7]], 11333// CHECK-32-NEXT: [[TMP32:%.*]] = icmp ult i16 [[TMP5]], [[TMP6]]1334// CHECK-32-NEXT: [[TMP33:%.*]] = and i1 [[TMP31]], [[TMP32]]1335// CHECK-32-NEXT: [[TMP34:%.*]] = icmp eq i16 [[TMP7]], 21336// CHECK-32-NEXT: [[TMP35:%.*]] = and i16 [[TMP5]], 11337// CHECK-32-NEXT: [[TMP36:%.*]] = icmp eq i16 [[TMP35]], 01338// CHECK-32-NEXT: [[TMP37:%.*]] = and i1 [[TMP34]], [[TMP36]]1339// CHECK-32-NEXT: [[TMP38:%.*]] = icmp sgt i16 [[TMP6]], 01340// CHECK-32-NEXT: [[TMP39:%.*]] = and i1 [[TMP37]], [[TMP38]]1341// CHECK-32-NEXT: [[TMP40:%.*]] = or i1 [[TMP30]], [[TMP33]]1342// CHECK-32-NEXT: [[TMP41:%.*]] = or i1 [[TMP40]], [[TMP39]]1343// CHECK-32-NEXT: br i1 [[TMP41]], label [[THEN:%.*]], label [[ELSE:%.*]]1344// CHECK-32: then:1345// CHECK-32-NEXT: call void @"{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l35_omp_outlined_omp$reduction$reduction_func"(ptr [[TMP4]], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]]) #[[ATTR3]]1346// CHECK-32-NEXT: br label [[IFCONT:%.*]]1347// CHECK-32: else:1348// CHECK-32-NEXT: br label [[IFCONT]]1349// CHECK-32: ifcont:1350// CHECK-32-NEXT: [[TMP42:%.*]] = icmp eq i16 [[TMP7]], 11351// CHECK-32-NEXT: [[TMP43:%.*]] = icmp uge i16 [[TMP5]], [[TMP6]]1352// CHECK-32-NEXT: [[TMP44:%.*]] = and i1 [[TMP42]], [[TMP43]]1353// CHECK-32-NEXT: br i1 [[TMP44]], label [[THEN5:%.*]], label [[ELSE6:%.*]]1354// CHECK-32: then5:1355// CHECK-32-NEXT: [[TMP45:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]], i32 0, i32 01356// CHECK-32-NEXT: [[TMP46:%.*]] = load ptr, ptr [[TMP45]], align 41357// CHECK-32-NEXT: [[TMP47:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP4]], i32 0, i32 01358// CHECK-32-NEXT: [[TMP48:%.*]] = load ptr, ptr [[TMP47]], align 41359// CHECK-32-NEXT: [[TMP49:%.*]] = load i32, ptr [[TMP46]], align 41360// CHECK-32-NEXT: store i32 [[TMP49]], ptr [[TMP48]], align 41361// CHECK-32-NEXT: [[TMP50:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]], i32 0, i32 11362// CHECK-32-NEXT: [[TMP51:%.*]] = load ptr, ptr [[TMP50]], align 41363// CHECK-32-NEXT: [[TMP52:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP4]], i32 0, i32 11364// CHECK-32-NEXT: [[TMP53:%.*]] = load ptr, ptr [[TMP52]], align 41365// CHECK-32-NEXT: [[TMP54:%.*]] = load i16, ptr [[TMP51]], align 21366// CHECK-32-NEXT: store i16 [[TMP54]], ptr [[TMP53]], align 21367// CHECK-32-NEXT: br label [[IFCONT7:%.*]]1368// CHECK-32: else6:1369// CHECK-32-NEXT: br label [[IFCONT7]]1370// CHECK-32: ifcont7:1371// CHECK-32-NEXT: ret void1372//1373//1374// CHECK-32-LABEL: define {{[^@]+}}@_omp_reduction_inter_warp_copy_func41375// CHECK-32-SAME: (ptr noundef [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR2]] {1376// CHECK-32-NEXT: entry:1377// CHECK-32-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 41378// CHECK-32-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 41379// CHECK-32-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 41380// CHECK-32-NEXT: store i32 [[TMP1]], ptr [[DOTADDR1]], align 41381// CHECK-32-NEXT: [[TMP3:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()1382// CHECK-32-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()1383// CHECK-32-NEXT: [[NVPTX_LANE_ID:%.*]] = and i32 [[TMP4]], 311384// CHECK-32-NEXT: [[TMP5:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()1385// CHECK-32-NEXT: [[NVPTX_WARP_ID:%.*]] = ashr i32 [[TMP5]], 51386// CHECK-32-NEXT: [[TMP6:%.*]] = load ptr, ptr [[DOTADDR]], align 41387// CHECK-32-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1388// CHECK-32-NEXT: call void @__kmpc_barrier(ptr @[[GLOB2]], i32 [[TMP2]])1389// CHECK-32-NEXT: [[WARP_MASTER:%.*]] = icmp eq i32 [[NVPTX_LANE_ID]], 01390// CHECK-32-NEXT: br i1 [[WARP_MASTER]], label [[THEN:%.*]], label [[ELSE:%.*]]1391// CHECK-32: then:1392// CHECK-32-NEXT: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP6]], i32 0, i32 01393// CHECK-32-NEXT: [[TMP8:%.*]] = load ptr, ptr [[TMP7]], align 41394// CHECK-32-NEXT: [[TMP9:%.*]] = getelementptr inbounds [32 x i32], ptr addrspace(3) @__openmp_nvptx_data_transfer_temporary_storage, i64 0, i32 [[NVPTX_WARP_ID]]1395// CHECK-32-NEXT: [[TMP10:%.*]] = load i32, ptr [[TMP8]], align 41396// CHECK-32-NEXT: store volatile i32 [[TMP10]], ptr addrspace(3) [[TMP9]], align 41397// CHECK-32-NEXT: br label [[IFCONT:%.*]]1398// CHECK-32: else:1399// CHECK-32-NEXT: br label [[IFCONT]]1400// CHECK-32: ifcont:1401// CHECK-32-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1402// CHECK-32-NEXT: call void @__kmpc_barrier(ptr @[[GLOB2]], i32 [[TMP2]])1403// CHECK-32-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTADDR1]], align 41404// CHECK-32-NEXT: [[IS_ACTIVE_THREAD:%.*]] = icmp ult i32 [[TMP3]], [[TMP11]]1405// CHECK-32-NEXT: br i1 [[IS_ACTIVE_THREAD]], label [[THEN2:%.*]], label [[ELSE3:%.*]]1406// CHECK-32: then3:1407// CHECK-32-NEXT: [[TMP12:%.*]] = getelementptr inbounds [32 x i32], ptr addrspace(3) @__openmp_nvptx_data_transfer_temporary_storage, i64 0, i32 [[TMP3]]1408// CHECK-32-NEXT: [[TMP13:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP6]], i32 0, i32 01409// CHECK-32-NEXT: [[TMP14:%.*]] = load ptr, ptr [[TMP13]], align 41410// CHECK-32-NEXT: [[TMP15:%.*]] = load volatile i32, ptr addrspace(3) [[TMP12]], align 41411// CHECK-32-NEXT: store i32 [[TMP15]], ptr [[TMP14]], align 41412// CHECK-32-NEXT: br label [[IFCONT4:%.*]]1413// CHECK-32: else4:1414// CHECK-32-NEXT: br label [[IFCONT4]]1415// CHECK-32: ifcont5:1416// CHECK-32-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1417// CHECK-32-NEXT: call void @__kmpc_barrier(ptr @[[GLOB2]], i32 [[TMP2]])1418// CHECK-32-NEXT: [[WARP_MASTER5:%.*]] = icmp eq i32 [[NVPTX_LANE_ID]], 01419// CHECK-32-NEXT: br i1 [[WARP_MASTER5]], label [[THEN6:%.*]], label [[ELSE7:%.*]]1420// CHECK-32: then8:1421// CHECK-32-NEXT: [[TMP16:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP6]], i32 0, i32 11422// CHECK-32-NEXT: [[TMP17:%.*]] = load ptr, ptr [[TMP16]], align 41423// CHECK-32-NEXT: [[TMP18:%.*]] = getelementptr inbounds [32 x i32], ptr addrspace(3) @__openmp_nvptx_data_transfer_temporary_storage, i64 0, i32 [[NVPTX_WARP_ID]]1424// CHECK-32-NEXT: [[TMP19:%.*]] = load i16, ptr [[TMP17]], align 21425// CHECK-32-NEXT: store volatile i16 [[TMP19]], ptr addrspace(3) [[TMP18]], align 21426// CHECK-32-NEXT: br label [[IFCONT8:%.*]]1427// CHECK-32: else9:1428// CHECK-32-NEXT: br label [[IFCONT8]]1429// CHECK-32: ifcont10:1430// CHECK-32-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1431// CHECK-32-NEXT: call void @__kmpc_barrier(ptr @[[GLOB2]], i32 [[TMP2]])1432// CHECK-32-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTADDR1]], align 41433// CHECK-32-NEXT: [[IS_ACTIVE_THREAD9:%.*]] = icmp ult i32 [[TMP3]], [[TMP20]]1434// CHECK-32-NEXT: br i1 [[IS_ACTIVE_THREAD9]], label [[THEN10:%.*]], label [[ELSE11:%.*]]1435// CHECK-32: then13:1436// CHECK-32-NEXT: [[TMP21:%.*]] = getelementptr inbounds [32 x i32], ptr addrspace(3) @__openmp_nvptx_data_transfer_temporary_storage, i64 0, i32 [[TMP3]]1437// CHECK-32-NEXT: [[TMP22:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP6]], i32 0, i32 11438// CHECK-32-NEXT: [[TMP23:%.*]] = load ptr, ptr [[TMP22]], align 41439// CHECK-32-NEXT: [[TMP24:%.*]] = load volatile i16, ptr addrspace(3) [[TMP21]], align 21440// CHECK-32-NEXT: store i16 [[TMP24]], ptr [[TMP23]], align 21441// CHECK-32-NEXT: br label [[IFCONT12:%.*]]1442// CHECK-32: else14:1443// CHECK-32-NEXT: br label [[IFCONT12]]1444// CHECK-32: ifcont15:1445// CHECK-32-NEXT: ret void1446//1447//1448// CHECK-32-EX-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l241449// CHECK-32-EX-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[E:%.*]]) #[[ATTR0:[0-9]+]] {1450// CHECK-32-EX-NEXT: entry:1451// CHECK-32-EX-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 41452// CHECK-32-EX-NEXT: [[E_ADDR:%.*]] = alloca ptr, align 41453// CHECK-32-EX-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [1 x ptr], align 41454// CHECK-32-EX-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 41455// CHECK-32-EX-NEXT: store ptr [[E]], ptr [[E_ADDR]], align 41456// CHECK-32-EX-NEXT: [[TMP0:%.*]] = load ptr, ptr [[E_ADDR]], align 41457// CHECK-32-EX-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l24_kernel_environment, ptr [[DYN_PTR]])1458// CHECK-32-EX-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -11459// CHECK-32-EX-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]1460// CHECK-32-EX: user_code.entry:1461// CHECK-32-EX-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1:[0-9]+]])1462// CHECK-32-EX-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 01463// CHECK-32-EX-NEXT: store ptr [[TMP0]], ptr [[TMP3]], align 41464// CHECK-32-EX-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]+}}__Z9ftemplateIcET_i_l24_omp_outlined, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i32 1)1465// CHECK-32-EX-NEXT: call void @__kmpc_target_deinit()1466// CHECK-32-EX-NEXT: ret void1467// CHECK-32-EX: worker.exit:1468// CHECK-32-EX-NEXT: ret void1469//1470//1471// CHECK-32-EX-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l24_omp_outlined1472// CHECK-32-EX-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[E:%.*]]) #[[ATTR1:[0-9]+]] {1473// CHECK-32-EX-NEXT: entry:1474// CHECK-32-EX-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41475// CHECK-32-EX-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41476// CHECK-32-EX-NEXT: [[E_ADDR:%.*]] = alloca ptr, align 41477// CHECK-32-EX-NEXT: [[E1:%.*]] = alloca double, align 81478// CHECK-32-EX-NEXT: [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 41479// CHECK-32-EX-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41480// CHECK-32-EX-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41481// CHECK-32-EX-NEXT: store ptr [[E]], ptr [[E_ADDR]], align 41482// CHECK-32-EX-NEXT: [[TMP0:%.*]] = load ptr, ptr [[E_ADDR]], align 41483// CHECK-32-EX-NEXT: store double 0.000000e+00, ptr [[E1]], align 81484// CHECK-32-EX-NEXT: [[TMP1:%.*]] = load double, ptr [[E1]], align 81485// CHECK-32-EX-NEXT: [[ADD:%.*]] = fadd double [[TMP1]], 5.000000e+001486// CHECK-32-EX-NEXT: store double [[ADD]], ptr [[E1]], align 81487// CHECK-32-EX-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i32 0, i32 01488// CHECK-32-EX-NEXT: store ptr [[E1]], ptr [[TMP2]], align 41489// CHECK-32-EX-NEXT: [[TMP3:%.*]] = call i32 @__kmpc_nvptx_parallel_reduce_nowait_v2(ptr @[[GLOB1]], i64 8, ptr [[DOTOMP_REDUCTION_RED_LIST]], ptr @_omp_reduction_shuffle_and_reduce_func, ptr @_omp_reduction_inter_warp_copy_func)1490// CHECK-32-EX-NEXT: [[TMP4:%.*]] = icmp eq i32 [[TMP3]], 11491// CHECK-32-EX-NEXT: br i1 [[TMP4]], label [[DOTOMP_REDUCTION_THEN:%.*]], label [[DOTOMP_REDUCTION_DONE:%.*]]1492// CHECK-32-EX: .omp.reduction.then:1493// CHECK-32-EX-NEXT: [[TMP5:%.*]] = load double, ptr [[TMP0]], align 81494// CHECK-32-EX-NEXT: [[TMP6:%.*]] = load double, ptr [[E1]], align 81495// CHECK-32-EX-NEXT: [[ADD2:%.*]] = fadd double [[TMP5]], [[TMP6]]1496// CHECK-32-EX-NEXT: store double [[ADD2]], ptr [[TMP0]], align 81497// CHECK-32-EX-NEXT: br label [[DOTOMP_REDUCTION_DONE]]1498// CHECK-32-EX: .omp.reduction.done:1499// CHECK-32-EX-NEXT: ret void1500//1501//1502// CHECK-32-EX-LABEL: define {{[^@]+}}@_omp_reduction_shuffle_and_reduce_func1503// CHECK-32-EX-SAME: (ptr noundef [[TMP0:%.*]], i16 noundef signext [[TMP1:%.*]], i16 noundef signext [[TMP2:%.*]], i16 noundef signext [[TMP3:%.*]]) #[[ATTR2:[0-9]+]] {1504// CHECK-32-EX-NEXT: entry:1505// CHECK-32-EX-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 41506// CHECK-32-EX-NEXT: [[DOTADDR1:%.*]] = alloca i16, align 21507// CHECK-32-EX-NEXT: [[DOTADDR2:%.*]] = alloca i16, align 21508// CHECK-32-EX-NEXT: [[DOTADDR3:%.*]] = alloca i16, align 21509// CHECK-32-EX-NEXT: [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST:%.*]] = alloca [1 x ptr], align 41510// CHECK-32-EX-NEXT: [[DOTOMP_REDUCTION_ELEMENT:%.*]] = alloca double, align 81511// CHECK-32-EX-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 41512// CHECK-32-EX-NEXT: store i16 [[TMP1]], ptr [[DOTADDR1]], align 21513// CHECK-32-EX-NEXT: store i16 [[TMP2]], ptr [[DOTADDR2]], align 21514// CHECK-32-EX-NEXT: store i16 [[TMP3]], ptr [[DOTADDR3]], align 21515// CHECK-32-EX-NEXT: [[TMP4:%.*]] = load ptr, ptr [[DOTADDR]], align 41516// CHECK-32-EX-NEXT: [[TMP5:%.*]] = load i16, ptr [[DOTADDR1]], align 21517// CHECK-32-EX-NEXT: [[TMP6:%.*]] = load i16, ptr [[DOTADDR2]], align 21518// CHECK-32-EX-NEXT: [[TMP7:%.*]] = load i16, ptr [[DOTADDR3]], align 21519// CHECK-32-EX-NEXT: [[TMP8:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP4]], i32 0, i32 01520// CHECK-32-EX-NEXT: [[TMP9:%.*]] = load ptr, ptr [[TMP8]], align 41521// CHECK-32-EX-NEXT: [[TMP10:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]], i32 0, i32 01522// CHECK-32-EX-NEXT: [[TMP11:%.*]] = getelementptr double, ptr [[TMP9]], i32 11523// CHECK-32-EX-NEXT: [[TMP12:%.*]] = load i64, ptr [[TMP9]], align 81524// CHECK-32-EX-NEXT: [[TMP13:%.*]] = call i32 @__kmpc_get_warp_size()1525// CHECK-32-EX-NEXT: [[TMP14:%.*]] = trunc i32 [[TMP13]] to i161526// CHECK-32-EX-NEXT: [[TMP15:%.*]] = call i64 @__kmpc_shuffle_int64(i64 [[TMP12]], i16 [[TMP6]], i16 [[TMP14]])1527// CHECK-32-EX-NEXT: store i64 [[TMP15]], ptr [[DOTOMP_REDUCTION_ELEMENT]], align 81528// CHECK-32-EX-NEXT: [[TMP16:%.*]] = getelementptr i64, ptr [[TMP9]], i32 11529// CHECK-32-EX-NEXT: [[TMP17:%.*]] = getelementptr i64, ptr [[DOTOMP_REDUCTION_ELEMENT]], i32 11530// CHECK-32-EX-NEXT: store ptr [[DOTOMP_REDUCTION_ELEMENT]], ptr [[TMP10]], align 41531// CHECK-32-EX-NEXT: [[TMP18:%.*]] = icmp eq i16 [[TMP7]], 01532// CHECK-32-EX-NEXT: [[TMP19:%.*]] = icmp eq i16 [[TMP7]], 11533// CHECK-32-EX-NEXT: [[TMP20:%.*]] = icmp ult i16 [[TMP5]], [[TMP6]]1534// CHECK-32-EX-NEXT: [[TMP21:%.*]] = and i1 [[TMP19]], [[TMP20]]1535// CHECK-32-EX-NEXT: [[TMP22:%.*]] = icmp eq i16 [[TMP7]], 21536// CHECK-32-EX-NEXT: [[TMP23:%.*]] = and i16 [[TMP5]], 11537// CHECK-32-EX-NEXT: [[TMP24:%.*]] = icmp eq i16 [[TMP23]], 01538// CHECK-32-EX-NEXT: [[TMP25:%.*]] = and i1 [[TMP22]], [[TMP24]]1539// CHECK-32-EX-NEXT: [[TMP26:%.*]] = icmp sgt i16 [[TMP6]], 01540// CHECK-32-EX-NEXT: [[TMP27:%.*]] = and i1 [[TMP25]], [[TMP26]]1541// CHECK-32-EX-NEXT: [[TMP28:%.*]] = or i1 [[TMP18]], [[TMP21]]1542// CHECK-32-EX-NEXT: [[TMP29:%.*]] = or i1 [[TMP28]], [[TMP27]]1543// CHECK-32-EX-NEXT: br i1 [[TMP29]], label [[THEN:%.*]], label [[ELSE:%.*]]1544// CHECK-32-EX: then:1545// CHECK-32-EX-NEXT: call void @"{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l24_omp_outlined_omp$reduction$reduction_func"(ptr [[TMP4]], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]]) #[[ATTR3:[0-9]+]]1546// CHECK-32-EX-NEXT: br label [[IFCONT:%.*]]1547// CHECK-32-EX: else:1548// CHECK-32-EX-NEXT: br label [[IFCONT]]1549// CHECK-32-EX: ifcont:1550// CHECK-32-EX-NEXT: [[TMP30:%.*]] = icmp eq i16 [[TMP7]], 11551// CHECK-32-EX-NEXT: [[TMP31:%.*]] = icmp uge i16 [[TMP5]], [[TMP6]]1552// CHECK-32-EX-NEXT: [[TMP32:%.*]] = and i1 [[TMP30]], [[TMP31]]1553// CHECK-32-EX-NEXT: br i1 [[TMP32]], label [[THEN4:%.*]], label [[ELSE5:%.*]]1554// CHECK-32-EX: then4:1555// CHECK-32-EX-NEXT: [[TMP33:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]], i32 0, i32 01556// CHECK-32-EX-NEXT: [[TMP34:%.*]] = load ptr, ptr [[TMP33]], align 41557// CHECK-32-EX-NEXT: [[TMP35:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP4]], i32 0, i32 01558// CHECK-32-EX-NEXT: [[TMP36:%.*]] = load ptr, ptr [[TMP35]], align 41559// CHECK-32-EX-NEXT: [[TMP37:%.*]] = load double, ptr [[TMP34]], align 81560// CHECK-32-EX-NEXT: store double [[TMP37]], ptr [[TMP36]], align 81561// CHECK-32-EX-NEXT: br label [[IFCONT6:%.*]]1562// CHECK-32-EX: else5:1563// CHECK-32-EX-NEXT: br label [[IFCONT6]]1564// CHECK-32-EX: ifcont6:1565// CHECK-32-EX-NEXT: ret void1566//1567//1568// CHECK-32-EX-LABEL: define {{[^@]+}}@_omp_reduction_inter_warp_copy_func1569// CHECK-32-EX-SAME: (ptr noundef [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR2]] {1570// CHECK-32-EX-NEXT: entry:1571// CHECK-32-EX-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 41572// CHECK-32-EX-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 41573// CHECK-32-EX-NEXT: [[DOTCNT_ADDR:%.*]] = alloca i32, align 41574// CHECK-32-EX-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 41575// CHECK-32-EX-NEXT: store i32 [[TMP1]], ptr [[DOTADDR1]], align 41576// CHECK-32-EX-NEXT: [[TMP3:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()1577// CHECK-32-EX-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()1578// CHECK-32-EX-NEXT: [[NVPTX_LANE_ID:%.*]] = and i32 [[TMP4]], 311579// CHECK-32-EX-NEXT: [[TMP5:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()1580// CHECK-32-EX-NEXT: [[NVPTX_WARP_ID:%.*]] = ashr i32 [[TMP5]], 51581// CHECK-32-EX-NEXT: [[TMP6:%.*]] = load ptr, ptr [[DOTADDR]], align 41582// CHECK-32-EX-NEXT: store i32 0, ptr [[DOTCNT_ADDR]], align 41583// CHECK-32-EX-NEXT: br label [[PRECOND:%.*]]1584// CHECK-32-EX: precond:1585// CHECK-32-EX-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTCNT_ADDR]], align 41586// CHECK-32-EX-NEXT: [[TMP8:%.*]] = icmp ult i32 [[TMP7]], 21587// CHECK-32-EX-NEXT: br i1 [[TMP8]], label [[BODY:%.*]], label [[EXIT:%.*]]1588// CHECK-32-EX: body:1589// CHECK-32-EX-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1590// CHECK-32-EX-NEXT: call void @__kmpc_barrier(ptr @[[GLOB2:[0-9]+]], i32 [[TMP2]])1591// CHECK-32-EX-NEXT: [[WARP_MASTER:%.*]] = icmp eq i32 [[NVPTX_LANE_ID]], 01592// CHECK-32-EX-NEXT: br i1 [[WARP_MASTER]], label [[THEN:%.*]], label [[ELSE:%.*]]1593// CHECK-32-EX: then:1594// CHECK-32-EX-NEXT: [[TMP9:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP6]], i32 0, i32 01595// CHECK-32-EX-NEXT: [[TMP10:%.*]] = load ptr, ptr [[TMP9]], align 41596// CHECK-32-EX-NEXT: [[TMP11:%.*]] = getelementptr i32, ptr [[TMP10]], i32 [[TMP7]]1597// CHECK-32-EX-NEXT: [[TMP12:%.*]] = getelementptr inbounds [32 x i32], ptr addrspace(3) @__openmp_nvptx_data_transfer_temporary_storage, i64 0, i32 [[NVPTX_WARP_ID]]1598// CHECK-32-EX-NEXT: [[TMP13:%.*]] = load i32, ptr [[TMP11]], align 41599// CHECK-32-EX-NEXT: store volatile i32 [[TMP13]], ptr addrspace(3) [[TMP12]], align 41600// CHECK-32-EX-NEXT: br label [[IFCONT:%.*]]1601// CHECK-32-EX: else:1602// CHECK-32-EX-NEXT: br label [[IFCONT]]1603// CHECK-32-EX: ifcont:1604// CHECK-32-EX-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1605// CHECK-32-EX-NEXT: call void @__kmpc_barrier(ptr @[[GLOB2]], i32 [[TMP2]])1606// CHECK-32-EX-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTADDR1]], align 41607// CHECK-32-EX-NEXT: [[IS_ACTIVE_THREAD:%.*]] = icmp ult i32 [[TMP3]], [[TMP14]]1608// CHECK-32-EX-NEXT: br i1 [[IS_ACTIVE_THREAD]], label [[THEN2:%.*]], label [[ELSE3:%.*]]1609// CHECK-32-EX: then3:1610// CHECK-32-EX-NEXT: [[TMP15:%.*]] = getelementptr inbounds [32 x i32], ptr addrspace(3) @__openmp_nvptx_data_transfer_temporary_storage, i64 0, i32 [[TMP3]]1611// CHECK-32-EX-NEXT: [[TMP16:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP6]], i32 0, i32 01612// CHECK-32-EX-NEXT: [[TMP17:%.*]] = load ptr, ptr [[TMP16]], align 41613// CHECK-32-EX-NEXT: [[TMP18:%.*]] = getelementptr i32, ptr [[TMP17]], i32 [[TMP7]]1614// CHECK-32-EX-NEXT: [[TMP19:%.*]] = load volatile i32, ptr addrspace(3) [[TMP15]], align 41615// CHECK-32-EX-NEXT: store i32 [[TMP19]], ptr [[TMP18]], align 41616// CHECK-32-EX-NEXT: br label [[IFCONT4:%.*]]1617// CHECK-32-EX: else4:1618// CHECK-32-EX-NEXT: br label [[IFCONT4]]1619// CHECK-32-EX: ifcont5:1620// CHECK-32-EX-NEXT: [[TMP20:%.*]] = add nsw i32 [[TMP7]], 11621// CHECK-32-EX-NEXT: store i32 [[TMP20]], ptr [[DOTCNT_ADDR]], align 41622// CHECK-32-EX-NEXT: br label [[PRECOND]]1623// CHECK-32-EX: exit:1624// CHECK-32-EX-NEXT: ret void1625//1626//1627// CHECK-32-EX-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l291628// CHECK-32-EX-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef nonnull align 1 dereferenceable(1) [[C:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[D:%.*]]) #[[ATTR0]] {1629// CHECK-32-EX-NEXT: entry:1630// CHECK-32-EX-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 41631// CHECK-32-EX-NEXT: [[C_ADDR:%.*]] = alloca ptr, align 41632// CHECK-32-EX-NEXT: [[D_ADDR:%.*]] = alloca ptr, align 41633// CHECK-32-EX-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [2 x ptr], align 41634// CHECK-32-EX-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 41635// CHECK-32-EX-NEXT: store ptr [[C]], ptr [[C_ADDR]], align 41636// CHECK-32-EX-NEXT: store ptr [[D]], ptr [[D_ADDR]], align 41637// CHECK-32-EX-NEXT: [[TMP0:%.*]] = load ptr, ptr [[C_ADDR]], align 41638// CHECK-32-EX-NEXT: [[TMP1:%.*]] = load ptr, ptr [[D_ADDR]], align 41639// CHECK-32-EX-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l29_kernel_environment, ptr [[DYN_PTR]])1640// CHECK-32-EX-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP2]], -11641// CHECK-32-EX-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]1642// CHECK-32-EX: user_code.entry:1643// CHECK-32-EX-NEXT: [[TMP3:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1644// CHECK-32-EX-NEXT: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 01645// CHECK-32-EX-NEXT: store ptr [[TMP0]], ptr [[TMP4]], align 41646// CHECK-32-EX-NEXT: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 11647// CHECK-32-EX-NEXT: store ptr [[TMP1]], ptr [[TMP5]], align 41648// CHECK-32-EX-NEXT: call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP3]], i32 1, i32 -1, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l29_omp_outlined, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i32 2)1649// CHECK-32-EX-NEXT: call void @__kmpc_target_deinit()1650// CHECK-32-EX-NEXT: ret void1651// CHECK-32-EX: worker.exit:1652// CHECK-32-EX-NEXT: ret void1653//1654//1655// CHECK-32-EX-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l29_omp_outlined1656// CHECK-32-EX-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 1 dereferenceable(1) [[C:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[D:%.*]]) #[[ATTR1]] {1657// CHECK-32-EX-NEXT: entry:1658// CHECK-32-EX-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41659// CHECK-32-EX-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41660// CHECK-32-EX-NEXT: [[C_ADDR:%.*]] = alloca ptr, align 41661// CHECK-32-EX-NEXT: [[D_ADDR:%.*]] = alloca ptr, align 41662// CHECK-32-EX-NEXT: [[C1:%.*]] = alloca i8, align 11663// CHECK-32-EX-NEXT: [[D2:%.*]] = alloca float, align 41664// CHECK-32-EX-NEXT: [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [2 x ptr], align 41665// CHECK-32-EX-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41666// CHECK-32-EX-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41667// CHECK-32-EX-NEXT: store ptr [[C]], ptr [[C_ADDR]], align 41668// CHECK-32-EX-NEXT: store ptr [[D]], ptr [[D_ADDR]], align 41669// CHECK-32-EX-NEXT: [[TMP0:%.*]] = load ptr, ptr [[C_ADDR]], align 41670// CHECK-32-EX-NEXT: [[TMP1:%.*]] = load ptr, ptr [[D_ADDR]], align 41671// CHECK-32-EX-NEXT: store i8 0, ptr [[C1]], align 11672// CHECK-32-EX-NEXT: store float 1.000000e+00, ptr [[D2]], align 41673// CHECK-32-EX-NEXT: [[TMP2:%.*]] = load i8, ptr [[C1]], align 11674// CHECK-32-EX-NEXT: [[CONV:%.*]] = sext i8 [[TMP2]] to i321675// CHECK-32-EX-NEXT: [[XOR:%.*]] = xor i32 [[CONV]], 21676// CHECK-32-EX-NEXT: [[CONV3:%.*]] = trunc i32 [[XOR]] to i81677// CHECK-32-EX-NEXT: store i8 [[CONV3]], ptr [[C1]], align 11678// CHECK-32-EX-NEXT: [[TMP3:%.*]] = load float, ptr [[D2]], align 41679// CHECK-32-EX-NEXT: [[MUL:%.*]] = fmul float [[TMP3]], 3.300000e+011680// CHECK-32-EX-NEXT: store float [[MUL]], ptr [[D2]], align 41681// CHECK-32-EX-NEXT: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i32 0, i32 01682// CHECK-32-EX-NEXT: store ptr [[C1]], ptr [[TMP4]], align 41683// CHECK-32-EX-NEXT: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i32 0, i32 11684// CHECK-32-EX-NEXT: store ptr [[D2]], ptr [[TMP5]], align 41685// CHECK-32-EX-NEXT: [[TMP6:%.*]] = call i32 @__kmpc_nvptx_parallel_reduce_nowait_v2(ptr @[[GLOB1]], i64 8, ptr [[DOTOMP_REDUCTION_RED_LIST]], ptr @_omp_reduction_shuffle_and_reduce_func1, ptr @_omp_reduction_inter_warp_copy_func2)1686// CHECK-32-EX-NEXT: [[TMP7:%.*]] = icmp eq i32 [[TMP6]], 11687// CHECK-32-EX-NEXT: br i1 [[TMP7]], label [[DOTOMP_REDUCTION_THEN:%.*]], label [[DOTOMP_REDUCTION_DONE:%.*]]1688// CHECK-32-EX: .omp.reduction.then:1689// CHECK-32-EX-NEXT: [[TMP8:%.*]] = load i8, ptr [[TMP0]], align 11690// CHECK-32-EX-NEXT: [[CONV4:%.*]] = sext i8 [[TMP8]] to i321691// CHECK-32-EX-NEXT: [[TMP9:%.*]] = load i8, ptr [[C1]], align 11692// CHECK-32-EX-NEXT: [[CONV5:%.*]] = sext i8 [[TMP9]] to i321693// CHECK-32-EX-NEXT: [[XOR6:%.*]] = xor i32 [[CONV4]], [[CONV5]]1694// CHECK-32-EX-NEXT: [[CONV7:%.*]] = trunc i32 [[XOR6]] to i81695// CHECK-32-EX-NEXT: store i8 [[CONV7]], ptr [[TMP0]], align 11696// CHECK-32-EX-NEXT: [[TMP10:%.*]] = load float, ptr [[TMP1]], align 41697// CHECK-32-EX-NEXT: [[TMP11:%.*]] = load float, ptr [[D2]], align 41698// CHECK-32-EX-NEXT: [[MUL8:%.*]] = fmul float [[TMP10]], [[TMP11]]1699// CHECK-32-EX-NEXT: store float [[MUL8]], ptr [[TMP1]], align 41700// CHECK-32-EX-NEXT: br label [[DOTOMP_REDUCTION_DONE]]1701// CHECK-32-EX: .omp.reduction.done:1702// CHECK-32-EX-NEXT: ret void1703//1704//1705// CHECK-32-EX-LABEL: define {{[^@]+}}@_omp_reduction_shuffle_and_reduce_func11706// CHECK-32-EX-SAME: (ptr noundef [[TMP0:%.*]], i16 noundef signext [[TMP1:%.*]], i16 noundef signext [[TMP2:%.*]], i16 noundef signext [[TMP3:%.*]]) #[[ATTR2]] {1707// CHECK-32-EX-NEXT: entry:1708// CHECK-32-EX-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 41709// CHECK-32-EX-NEXT: [[DOTADDR1:%.*]] = alloca i16, align 21710// CHECK-32-EX-NEXT: [[DOTADDR2:%.*]] = alloca i16, align 21711// CHECK-32-EX-NEXT: [[DOTADDR3:%.*]] = alloca i16, align 21712// CHECK-32-EX-NEXT: [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST:%.*]] = alloca [2 x ptr], align 41713// CHECK-32-EX-NEXT: [[DOTOMP_REDUCTION_ELEMENT:%.*]] = alloca i8, align 11714// CHECK-32-EX-NEXT: [[DOTOMP_REDUCTION_ELEMENT4:%.*]] = alloca float, align 41715// CHECK-32-EX-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 41716// CHECK-32-EX-NEXT: store i16 [[TMP1]], ptr [[DOTADDR1]], align 21717// CHECK-32-EX-NEXT: store i16 [[TMP2]], ptr [[DOTADDR2]], align 21718// CHECK-32-EX-NEXT: store i16 [[TMP3]], ptr [[DOTADDR3]], align 21719// CHECK-32-EX-NEXT: [[TMP4:%.*]] = load ptr, ptr [[DOTADDR]], align 41720// CHECK-32-EX-NEXT: [[TMP5:%.*]] = load i16, ptr [[DOTADDR1]], align 21721// CHECK-32-EX-NEXT: [[TMP6:%.*]] = load i16, ptr [[DOTADDR2]], align 21722// CHECK-32-EX-NEXT: [[TMP7:%.*]] = load i16, ptr [[DOTADDR3]], align 21723// CHECK-32-EX-NEXT: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP4]], i32 0, i32 01724// CHECK-32-EX-NEXT: [[TMP9:%.*]] = load ptr, ptr [[TMP8]], align 41725// CHECK-32-EX-NEXT: [[TMP10:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]], i32 0, i32 01726// CHECK-32-EX-NEXT: [[TMP11:%.*]] = getelementptr i8, ptr [[TMP9]], i32 11727// CHECK-32-EX-NEXT: [[TMP12:%.*]] = load i8, ptr [[TMP9]], align 11728// CHECK-32-EX-NEXT: [[TMP13:%.*]] = sext i8 [[TMP12]] to i321729// CHECK-32-EX-NEXT: [[TMP14:%.*]] = call i32 @__kmpc_get_warp_size()1730// CHECK-32-EX-NEXT: [[TMP15:%.*]] = trunc i32 [[TMP14]] to i161731// CHECK-32-EX-NEXT: [[TMP16:%.*]] = call i32 @__kmpc_shuffle_int32(i32 [[TMP13]], i16 [[TMP6]], i16 [[TMP15]])1732// CHECK-32-EX-NEXT: [[TMP17:%.*]] = trunc i32 [[TMP16]] to i81733// CHECK-32-EX-NEXT: store i8 [[TMP17]], ptr [[DOTOMP_REDUCTION_ELEMENT]], align 11734// CHECK-32-EX-NEXT: [[TMP18:%.*]] = getelementptr i8, ptr [[TMP9]], i32 11735// CHECK-32-EX-NEXT: [[TMP19:%.*]] = getelementptr i8, ptr [[DOTOMP_REDUCTION_ELEMENT]], i32 11736// CHECK-32-EX-NEXT: store ptr [[DOTOMP_REDUCTION_ELEMENT]], ptr [[TMP10]], align 41737// CHECK-32-EX-NEXT: [[TMP20:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP4]], i32 0, i32 11738// CHECK-32-EX-NEXT: [[TMP21:%.*]] = load ptr, ptr [[TMP20]], align 41739// CHECK-32-EX-NEXT: [[TMP22:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]], i32 0, i32 11740// CHECK-32-EX-NEXT: [[TMP23:%.*]] = getelementptr float, ptr [[TMP21]], i32 11741// CHECK-32-EX-NEXT: [[TMP24:%.*]] = load i32, ptr [[TMP21]], align 41742// CHECK-32-EX-NEXT: [[TMP25:%.*]] = call i32 @__kmpc_get_warp_size()1743// CHECK-32-EX-NEXT: [[TMP26:%.*]] = trunc i32 [[TMP25]] to i161744// CHECK-32-EX-NEXT: [[TMP27:%.*]] = call i32 @__kmpc_shuffle_int32(i32 [[TMP24]], i16 [[TMP6]], i16 [[TMP26]])1745// CHECK-32-EX-NEXT: store i32 [[TMP27]], ptr [[DOTOMP_REDUCTION_ELEMENT4]], align 41746// CHECK-32-EX-NEXT: [[TMP28:%.*]] = getelementptr i32, ptr [[TMP21]], i32 11747// CHECK-32-EX-NEXT: [[TMP29:%.*]] = getelementptr i32, ptr [[DOTOMP_REDUCTION_ELEMENT4]], i32 11748// CHECK-32-EX-NEXT: store ptr [[DOTOMP_REDUCTION_ELEMENT4]], ptr [[TMP22]], align 41749// CHECK-32-EX-NEXT: [[TMP30:%.*]] = icmp eq i16 [[TMP7]], 01750// CHECK-32-EX-NEXT: [[TMP31:%.*]] = icmp eq i16 [[TMP7]], 11751// CHECK-32-EX-NEXT: [[TMP32:%.*]] = icmp ult i16 [[TMP5]], [[TMP6]]1752// CHECK-32-EX-NEXT: [[TMP33:%.*]] = and i1 [[TMP31]], [[TMP32]]1753// CHECK-32-EX-NEXT: [[TMP34:%.*]] = icmp eq i16 [[TMP7]], 21754// CHECK-32-EX-NEXT: [[TMP35:%.*]] = and i16 [[TMP5]], 11755// CHECK-32-EX-NEXT: [[TMP36:%.*]] = icmp eq i16 [[TMP35]], 01756// CHECK-32-EX-NEXT: [[TMP37:%.*]] = and i1 [[TMP34]], [[TMP36]]1757// CHECK-32-EX-NEXT: [[TMP38:%.*]] = icmp sgt i16 [[TMP6]], 01758// CHECK-32-EX-NEXT: [[TMP39:%.*]] = and i1 [[TMP37]], [[TMP38]]1759// CHECK-32-EX-NEXT: [[TMP40:%.*]] = or i1 [[TMP30]], [[TMP33]]1760// CHECK-32-EX-NEXT: [[TMP41:%.*]] = or i1 [[TMP40]], [[TMP39]]1761// CHECK-32-EX-NEXT: br i1 [[TMP41]], label [[THEN:%.*]], label [[ELSE:%.*]]1762// CHECK-32-EX: then:1763// CHECK-32-EX-NEXT: call void @"{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l29_omp_outlined_omp$reduction$reduction_func"(ptr [[TMP4]], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]]) #[[ATTR3]]1764// CHECK-32-EX-NEXT: br label [[IFCONT:%.*]]1765// CHECK-32-EX: else:1766// CHECK-32-EX-NEXT: br label [[IFCONT]]1767// CHECK-32-EX: ifcont:1768// CHECK-32-EX-NEXT: [[TMP42:%.*]] = icmp eq i16 [[TMP7]], 11769// CHECK-32-EX-NEXT: [[TMP43:%.*]] = icmp uge i16 [[TMP5]], [[TMP6]]1770// CHECK-32-EX-NEXT: [[TMP44:%.*]] = and i1 [[TMP42]], [[TMP43]]1771// CHECK-32-EX-NEXT: br i1 [[TMP44]], label [[THEN5:%.*]], label [[ELSE6:%.*]]1772// CHECK-32-EX: then5:1773// CHECK-32-EX-NEXT: [[TMP45:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]], i32 0, i32 01774// CHECK-32-EX-NEXT: [[TMP46:%.*]] = load ptr, ptr [[TMP45]], align 41775// CHECK-32-EX-NEXT: [[TMP47:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP4]], i32 0, i32 01776// CHECK-32-EX-NEXT: [[TMP48:%.*]] = load ptr, ptr [[TMP47]], align 41777// CHECK-32-EX-NEXT: [[TMP49:%.*]] = load i8, ptr [[TMP46]], align 11778// CHECK-32-EX-NEXT: store i8 [[TMP49]], ptr [[TMP48]], align 11779// CHECK-32-EX-NEXT: [[TMP50:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]], i32 0, i32 11780// CHECK-32-EX-NEXT: [[TMP51:%.*]] = load ptr, ptr [[TMP50]], align 41781// CHECK-32-EX-NEXT: [[TMP52:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP4]], i32 0, i32 11782// CHECK-32-EX-NEXT: [[TMP53:%.*]] = load ptr, ptr [[TMP52]], align 41783// CHECK-32-EX-NEXT: [[TMP54:%.*]] = load float, ptr [[TMP51]], align 41784// CHECK-32-EX-NEXT: store float [[TMP54]], ptr [[TMP53]], align 41785// CHECK-32-EX-NEXT: br label [[IFCONT7:%.*]]1786// CHECK-32-EX: else6:1787// CHECK-32-EX-NEXT: br label [[IFCONT7]]1788// CHECK-32-EX: ifcont7:1789// CHECK-32-EX-NEXT: ret void1790//1791//1792// CHECK-32-EX-LABEL: define {{[^@]+}}@_omp_reduction_inter_warp_copy_func21793// CHECK-32-EX-SAME: (ptr noundef [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR2]] {1794// CHECK-32-EX-NEXT: entry:1795// CHECK-32-EX-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 41796// CHECK-32-EX-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 41797// CHECK-32-EX-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 41798// CHECK-32-EX-NEXT: store i32 [[TMP1]], ptr [[DOTADDR1]], align 41799// CHECK-32-EX-NEXT: [[TMP3:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()1800// CHECK-32-EX-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()1801// CHECK-32-EX-NEXT: [[NVPTX_LANE_ID:%.*]] = and i32 [[TMP4]], 311802// CHECK-32-EX-NEXT: [[TMP5:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()1803// CHECK-32-EX-NEXT: [[NVPTX_WARP_ID:%.*]] = ashr i32 [[TMP5]], 51804// CHECK-32-EX-NEXT: [[TMP6:%.*]] = load ptr, ptr [[DOTADDR]], align 41805// CHECK-32-EX-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1806// CHECK-32-EX-NEXT: call void @__kmpc_barrier(ptr @[[GLOB2]], i32 [[TMP2]])1807// CHECK-32-EX-NEXT: [[WARP_MASTER:%.*]] = icmp eq i32 [[NVPTX_LANE_ID]], 01808// CHECK-32-EX-NEXT: br i1 [[WARP_MASTER]], label [[THEN:%.*]], label [[ELSE:%.*]]1809// CHECK-32-EX: then:1810// CHECK-32-EX-NEXT: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP6]], i32 0, i32 01811// CHECK-32-EX-NEXT: [[TMP8:%.*]] = load ptr, ptr [[TMP7]], align 41812// CHECK-32-EX-NEXT: [[TMP9:%.*]] = getelementptr inbounds [32 x i32], ptr addrspace(3) @__openmp_nvptx_data_transfer_temporary_storage, i64 0, i32 [[NVPTX_WARP_ID]]1813// CHECK-32-EX-NEXT: [[TMP10:%.*]] = load i8, ptr [[TMP8]], align 11814// CHECK-32-EX-NEXT: store volatile i8 [[TMP10]], ptr addrspace(3) [[TMP9]], align 11815// CHECK-32-EX-NEXT: br label [[IFCONT:%.*]]1816// CHECK-32-EX: else:1817// CHECK-32-EX-NEXT: br label [[IFCONT]]1818// CHECK-32-EX: ifcont:1819// CHECK-32-EX-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1820// CHECK-32-EX-NEXT: call void @__kmpc_barrier(ptr @[[GLOB2]], i32 [[TMP2]])1821// CHECK-32-EX-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTADDR1]], align 41822// CHECK-32-EX-NEXT: [[IS_ACTIVE_THREAD:%.*]] = icmp ult i32 [[TMP3]], [[TMP11]]1823// CHECK-32-EX-NEXT: br i1 [[IS_ACTIVE_THREAD]], label [[THEN2:%.*]], label [[ELSE3:%.*]]1824// CHECK-32-EX: then3:1825// CHECK-32-EX-NEXT: [[TMP12:%.*]] = getelementptr inbounds [32 x i32], ptr addrspace(3) @__openmp_nvptx_data_transfer_temporary_storage, i64 0, i32 [[TMP3]]1826// CHECK-32-EX-NEXT: [[TMP13:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP6]], i32 0, i32 01827// CHECK-32-EX-NEXT: [[TMP14:%.*]] = load ptr, ptr [[TMP13]], align 41828// CHECK-32-EX-NEXT: [[TMP15:%.*]] = load volatile i8, ptr addrspace(3) [[TMP12]], align 11829// CHECK-32-EX-NEXT: store i8 [[TMP15]], ptr [[TMP14]], align 11830// CHECK-32-EX-NEXT: br label [[IFCONT4:%.*]]1831// CHECK-32-EX: else4:1832// CHECK-32-EX-NEXT: br label [[IFCONT4]]1833// CHECK-32-EX: ifcont5:1834// CHECK-32-EX-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1835// CHECK-32-EX-NEXT: call void @__kmpc_barrier(ptr @[[GLOB2]], i32 [[TMP2]])1836// CHECK-32-EX-NEXT: [[WARP_MASTER5:%.*]] = icmp eq i32 [[NVPTX_LANE_ID]], 01837// CHECK-32-EX-NEXT: br i1 [[WARP_MASTER5]], label [[THEN6:%.*]], label [[ELSE7:%.*]]1838// CHECK-32-EX: then8:1839// CHECK-32-EX-NEXT: [[TMP16:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP6]], i32 0, i32 11840// CHECK-32-EX-NEXT: [[TMP17:%.*]] = load ptr, ptr [[TMP16]], align 41841// CHECK-32-EX-NEXT: [[TMP18:%.*]] = getelementptr inbounds [32 x i32], ptr addrspace(3) @__openmp_nvptx_data_transfer_temporary_storage, i64 0, i32 [[NVPTX_WARP_ID]]1842// CHECK-32-EX-NEXT: [[TMP19:%.*]] = load i32, ptr [[TMP17]], align 41843// CHECK-32-EX-NEXT: store volatile i32 [[TMP19]], ptr addrspace(3) [[TMP18]], align 41844// CHECK-32-EX-NEXT: br label [[IFCONT8:%.*]]1845// CHECK-32-EX: else9:1846// CHECK-32-EX-NEXT: br label [[IFCONT8]]1847// CHECK-32-EX: ifcont10:1848// CHECK-32-EX-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1849// CHECK-32-EX-NEXT: call void @__kmpc_barrier(ptr @[[GLOB2]], i32 [[TMP2]])1850// CHECK-32-EX-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTADDR1]], align 41851// CHECK-32-EX-NEXT: [[IS_ACTIVE_THREAD9:%.*]] = icmp ult i32 [[TMP3]], [[TMP20]]1852// CHECK-32-EX-NEXT: br i1 [[IS_ACTIVE_THREAD9]], label [[THEN10:%.*]], label [[ELSE11:%.*]]1853// CHECK-32-EX: then13:1854// CHECK-32-EX-NEXT: [[TMP21:%.*]] = getelementptr inbounds [32 x i32], ptr addrspace(3) @__openmp_nvptx_data_transfer_temporary_storage, i64 0, i32 [[TMP3]]1855// CHECK-32-EX-NEXT: [[TMP22:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP6]], i32 0, i32 11856// CHECK-32-EX-NEXT: [[TMP23:%.*]] = load ptr, ptr [[TMP22]], align 41857// CHECK-32-EX-NEXT: [[TMP24:%.*]] = load volatile i32, ptr addrspace(3) [[TMP21]], align 41858// CHECK-32-EX-NEXT: store i32 [[TMP24]], ptr [[TMP23]], align 41859// CHECK-32-EX-NEXT: br label [[IFCONT12:%.*]]1860// CHECK-32-EX: else14:1861// CHECK-32-EX-NEXT: br label [[IFCONT12]]1862// CHECK-32-EX: ifcont15:1863// CHECK-32-EX-NEXT: ret void1864//1865//1866// CHECK-32-EX-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l351867// CHECK-32-EX-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[B:%.*]]) #[[ATTR0]] {1868// CHECK-32-EX-NEXT: entry:1869// CHECK-32-EX-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 41870// CHECK-32-EX-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 41871// CHECK-32-EX-NEXT: [[B_ADDR:%.*]] = alloca ptr, align 41872// CHECK-32-EX-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [2 x ptr], align 41873// CHECK-32-EX-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 41874// CHECK-32-EX-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 41875// CHECK-32-EX-NEXT: store ptr [[B]], ptr [[B_ADDR]], align 41876// CHECK-32-EX-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 41877// CHECK-32-EX-NEXT: [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 41878// CHECK-32-EX-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l35_kernel_environment, ptr [[DYN_PTR]])1879// CHECK-32-EX-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP2]], -11880// CHECK-32-EX-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]1881// CHECK-32-EX: user_code.entry:1882// CHECK-32-EX-NEXT: [[TMP3:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])1883// CHECK-32-EX-NEXT: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 01884// CHECK-32-EX-NEXT: store ptr [[TMP0]], ptr [[TMP4]], align 41885// CHECK-32-EX-NEXT: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i32 0, i32 11886// CHECK-32-EX-NEXT: store ptr [[TMP1]], ptr [[TMP5]], align 41887// CHECK-32-EX-NEXT: call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP3]], i32 1, i32 -1, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l35_omp_outlined, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i32 2)1888// CHECK-32-EX-NEXT: call void @__kmpc_target_deinit()1889// CHECK-32-EX-NEXT: ret void1890// CHECK-32-EX: worker.exit:1891// CHECK-32-EX-NEXT: ret void1892//1893//1894// CHECK-32-EX-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l35_omp_outlined1895// CHECK-32-EX-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[A:%.*]], ptr noundef nonnull align 2 dereferenceable(2) [[B:%.*]]) #[[ATTR1]] {1896// CHECK-32-EX-NEXT: entry:1897// CHECK-32-EX-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 41898// CHECK-32-EX-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 41899// CHECK-32-EX-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 41900// CHECK-32-EX-NEXT: [[B_ADDR:%.*]] = alloca ptr, align 41901// CHECK-32-EX-NEXT: [[A1:%.*]] = alloca i32, align 41902// CHECK-32-EX-NEXT: [[B2:%.*]] = alloca i16, align 21903// CHECK-32-EX-NEXT: [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [2 x ptr], align 41904// CHECK-32-EX-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 41905// CHECK-32-EX-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 41906// CHECK-32-EX-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 41907// CHECK-32-EX-NEXT: store ptr [[B]], ptr [[B_ADDR]], align 41908// CHECK-32-EX-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 41909// CHECK-32-EX-NEXT: [[TMP1:%.*]] = load ptr, ptr [[B_ADDR]], align 41910// CHECK-32-EX-NEXT: store i32 0, ptr [[A1]], align 41911// CHECK-32-EX-NEXT: store i16 -32768, ptr [[B2]], align 21912// CHECK-32-EX-NEXT: [[TMP2:%.*]] = load i32, ptr [[A1]], align 41913// CHECK-32-EX-NEXT: [[OR:%.*]] = or i32 [[TMP2]], 11914// CHECK-32-EX-NEXT: store i32 [[OR]], ptr [[A1]], align 41915// CHECK-32-EX-NEXT: [[TMP3:%.*]] = load i16, ptr [[B2]], align 21916// CHECK-32-EX-NEXT: [[CONV:%.*]] = sext i16 [[TMP3]] to i321917// CHECK-32-EX-NEXT: [[CMP:%.*]] = icmp sgt i32 99, [[CONV]]1918// CHECK-32-EX-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]1919// CHECK-32-EX: cond.true:1920// CHECK-32-EX-NEXT: br label [[COND_END:%.*]]1921// CHECK-32-EX: cond.false:1922// CHECK-32-EX-NEXT: [[TMP4:%.*]] = load i16, ptr [[B2]], align 21923// CHECK-32-EX-NEXT: [[CONV3:%.*]] = sext i16 [[TMP4]] to i321924// CHECK-32-EX-NEXT: br label [[COND_END]]1925// CHECK-32-EX: cond.end:1926// CHECK-32-EX-NEXT: [[COND:%.*]] = phi i32 [ 99, [[COND_TRUE]] ], [ [[CONV3]], [[COND_FALSE]] ]1927// CHECK-32-EX-NEXT: [[CONV4:%.*]] = trunc i32 [[COND]] to i161928// CHECK-32-EX-NEXT: store i16 [[CONV4]], ptr [[B2]], align 21929// CHECK-32-EX-NEXT: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i32 0, i32 01930// CHECK-32-EX-NEXT: store ptr [[A1]], ptr [[TMP5]], align 41931// CHECK-32-EX-NEXT: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i32 0, i32 11932// CHECK-32-EX-NEXT: store ptr [[B2]], ptr [[TMP6]], align 41933// CHECK-32-EX-NEXT: [[TMP7:%.*]] = call i32 @__kmpc_nvptx_parallel_reduce_nowait_v2(ptr @[[GLOB1]], i64 8, ptr [[DOTOMP_REDUCTION_RED_LIST]], ptr @_omp_reduction_shuffle_and_reduce_func3, ptr @_omp_reduction_inter_warp_copy_func4)1934// CHECK-32-EX-NEXT: [[TMP8:%.*]] = icmp eq i32 [[TMP7]], 11935// CHECK-32-EX-NEXT: br i1 [[TMP8]], label [[DOTOMP_REDUCTION_THEN:%.*]], label [[DOTOMP_REDUCTION_DONE:%.*]]1936// CHECK-32-EX: .omp.reduction.then:1937// CHECK-32-EX-NEXT: [[TMP9:%.*]] = load i32, ptr [[TMP0]], align 41938// CHECK-32-EX-NEXT: [[TMP10:%.*]] = load i32, ptr [[A1]], align 41939// CHECK-32-EX-NEXT: [[OR5:%.*]] = or i32 [[TMP9]], [[TMP10]]1940// CHECK-32-EX-NEXT: store i32 [[OR5]], ptr [[TMP0]], align 41941// CHECK-32-EX-NEXT: [[TMP11:%.*]] = load i16, ptr [[TMP1]], align 21942// CHECK-32-EX-NEXT: [[CONV6:%.*]] = sext i16 [[TMP11]] to i321943// CHECK-32-EX-NEXT: [[TMP12:%.*]] = load i16, ptr [[B2]], align 21944// CHECK-32-EX-NEXT: [[CONV7:%.*]] = sext i16 [[TMP12]] to i321945// CHECK-32-EX-NEXT: [[CMP8:%.*]] = icmp sgt i32 [[CONV6]], [[CONV7]]1946// CHECK-32-EX-NEXT: br i1 [[CMP8]], label [[COND_TRUE9:%.*]], label [[COND_FALSE10:%.*]]1947// CHECK-32-EX: cond.true9:1948// CHECK-32-EX-NEXT: [[TMP13:%.*]] = load i16, ptr [[TMP1]], align 21949// CHECK-32-EX-NEXT: br label [[COND_END11:%.*]]1950// CHECK-32-EX: cond.false10:1951// CHECK-32-EX-NEXT: [[TMP14:%.*]] = load i16, ptr [[B2]], align 21952// CHECK-32-EX-NEXT: br label [[COND_END11]]1953// CHECK-32-EX: cond.end11:1954// CHECK-32-EX-NEXT: [[COND12:%.*]] = phi i16 [ [[TMP13]], [[COND_TRUE9]] ], [ [[TMP14]], [[COND_FALSE10]] ]1955// CHECK-32-EX-NEXT: store i16 [[COND12]], ptr [[TMP1]], align 21956// CHECK-32-EX-NEXT: br label [[DOTOMP_REDUCTION_DONE]]1957// CHECK-32-EX: .omp.reduction.done:1958// CHECK-32-EX-NEXT: ret void1959//1960//1961// CHECK-32-EX-LABEL: define {{[^@]+}}@_omp_reduction_shuffle_and_reduce_func31962// CHECK-32-EX-SAME: (ptr noundef [[TMP0:%.*]], i16 noundef signext [[TMP1:%.*]], i16 noundef signext [[TMP2:%.*]], i16 noundef signext [[TMP3:%.*]]) #[[ATTR2]] {1963// CHECK-32-EX-NEXT: entry:1964// CHECK-32-EX-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 41965// CHECK-32-EX-NEXT: [[DOTADDR1:%.*]] = alloca i16, align 21966// CHECK-32-EX-NEXT: [[DOTADDR2:%.*]] = alloca i16, align 21967// CHECK-32-EX-NEXT: [[DOTADDR3:%.*]] = alloca i16, align 21968// CHECK-32-EX-NEXT: [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST:%.*]] = alloca [2 x ptr], align 41969// CHECK-32-EX-NEXT: [[DOTOMP_REDUCTION_ELEMENT:%.*]] = alloca i32, align 41970// CHECK-32-EX-NEXT: [[DOTOMP_REDUCTION_ELEMENT4:%.*]] = alloca i16, align 21971// CHECK-32-EX-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 41972// CHECK-32-EX-NEXT: store i16 [[TMP1]], ptr [[DOTADDR1]], align 21973// CHECK-32-EX-NEXT: store i16 [[TMP2]], ptr [[DOTADDR2]], align 21974// CHECK-32-EX-NEXT: store i16 [[TMP3]], ptr [[DOTADDR3]], align 21975// CHECK-32-EX-NEXT: [[TMP4:%.*]] = load ptr, ptr [[DOTADDR]], align 41976// CHECK-32-EX-NEXT: [[TMP5:%.*]] = load i16, ptr [[DOTADDR1]], align 21977// CHECK-32-EX-NEXT: [[TMP6:%.*]] = load i16, ptr [[DOTADDR2]], align 21978// CHECK-32-EX-NEXT: [[TMP7:%.*]] = load i16, ptr [[DOTADDR3]], align 21979// CHECK-32-EX-NEXT: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP4]], i32 0, i32 01980// CHECK-32-EX-NEXT: [[TMP9:%.*]] = load ptr, ptr [[TMP8]], align 41981// CHECK-32-EX-NEXT: [[TMP10:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]], i32 0, i32 01982// CHECK-32-EX-NEXT: [[TMP11:%.*]] = getelementptr i32, ptr [[TMP9]], i32 11983// CHECK-32-EX-NEXT: [[TMP12:%.*]] = load i32, ptr [[TMP9]], align 41984// CHECK-32-EX-NEXT: [[TMP13:%.*]] = call i32 @__kmpc_get_warp_size()1985// CHECK-32-EX-NEXT: [[TMP14:%.*]] = trunc i32 [[TMP13]] to i161986// CHECK-32-EX-NEXT: [[TMP15:%.*]] = call i32 @__kmpc_shuffle_int32(i32 [[TMP12]], i16 [[TMP6]], i16 [[TMP14]])1987// CHECK-32-EX-NEXT: store i32 [[TMP15]], ptr [[DOTOMP_REDUCTION_ELEMENT]], align 41988// CHECK-32-EX-NEXT: [[TMP16:%.*]] = getelementptr i32, ptr [[TMP9]], i32 11989// CHECK-32-EX-NEXT: [[TMP17:%.*]] = getelementptr i32, ptr [[DOTOMP_REDUCTION_ELEMENT]], i32 11990// CHECK-32-EX-NEXT: store ptr [[DOTOMP_REDUCTION_ELEMENT]], ptr [[TMP10]], align 41991// CHECK-32-EX-NEXT: [[TMP18:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP4]], i32 0, i32 11992// CHECK-32-EX-NEXT: [[TMP19:%.*]] = load ptr, ptr [[TMP18]], align 41993// CHECK-32-EX-NEXT: [[TMP20:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]], i32 0, i32 11994// CHECK-32-EX-NEXT: [[TMP21:%.*]] = getelementptr i16, ptr [[TMP19]], i32 11995// CHECK-32-EX-NEXT: [[TMP22:%.*]] = load i16, ptr [[TMP19]], align 21996// CHECK-32-EX-NEXT: [[TMP23:%.*]] = sext i16 [[TMP22]] to i321997// CHECK-32-EX-NEXT: [[TMP24:%.*]] = call i32 @__kmpc_get_warp_size()1998// CHECK-32-EX-NEXT: [[TMP25:%.*]] = trunc i32 [[TMP24]] to i161999// CHECK-32-EX-NEXT: [[TMP26:%.*]] = call i32 @__kmpc_shuffle_int32(i32 [[TMP23]], i16 [[TMP6]], i16 [[TMP25]])2000// CHECK-32-EX-NEXT: [[TMP27:%.*]] = trunc i32 [[TMP26]] to i162001// CHECK-32-EX-NEXT: store i16 [[TMP27]], ptr [[DOTOMP_REDUCTION_ELEMENT4]], align 22002// CHECK-32-EX-NEXT: [[TMP28:%.*]] = getelementptr i16, ptr [[TMP19]], i32 12003// CHECK-32-EX-NEXT: [[TMP29:%.*]] = getelementptr i16, ptr [[DOTOMP_REDUCTION_ELEMENT4]], i32 12004// CHECK-32-EX-NEXT: store ptr [[DOTOMP_REDUCTION_ELEMENT4]], ptr [[TMP20]], align 42005// CHECK-32-EX-NEXT: [[TMP30:%.*]] = icmp eq i16 [[TMP7]], 02006// CHECK-32-EX-NEXT: [[TMP31:%.*]] = icmp eq i16 [[TMP7]], 12007// CHECK-32-EX-NEXT: [[TMP32:%.*]] = icmp ult i16 [[TMP5]], [[TMP6]]2008// CHECK-32-EX-NEXT: [[TMP33:%.*]] = and i1 [[TMP31]], [[TMP32]]2009// CHECK-32-EX-NEXT: [[TMP34:%.*]] = icmp eq i16 [[TMP7]], 22010// CHECK-32-EX-NEXT: [[TMP35:%.*]] = and i16 [[TMP5]], 12011// CHECK-32-EX-NEXT: [[TMP36:%.*]] = icmp eq i16 [[TMP35]], 02012// CHECK-32-EX-NEXT: [[TMP37:%.*]] = and i1 [[TMP34]], [[TMP36]]2013// CHECK-32-EX-NEXT: [[TMP38:%.*]] = icmp sgt i16 [[TMP6]], 02014// CHECK-32-EX-NEXT: [[TMP39:%.*]] = and i1 [[TMP37]], [[TMP38]]2015// CHECK-32-EX-NEXT: [[TMP40:%.*]] = or i1 [[TMP30]], [[TMP33]]2016// CHECK-32-EX-NEXT: [[TMP41:%.*]] = or i1 [[TMP40]], [[TMP39]]2017// CHECK-32-EX-NEXT: br i1 [[TMP41]], label [[THEN:%.*]], label [[ELSE:%.*]]2018// CHECK-32-EX: then:2019// CHECK-32-EX-NEXT: call void @"{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIcET_i_l35_omp_outlined_omp$reduction$reduction_func"(ptr [[TMP4]], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]]) #[[ATTR3]]2020// CHECK-32-EX-NEXT: br label [[IFCONT:%.*]]2021// CHECK-32-EX: else:2022// CHECK-32-EX-NEXT: br label [[IFCONT]]2023// CHECK-32-EX: ifcont:2024// CHECK-32-EX-NEXT: [[TMP42:%.*]] = icmp eq i16 [[TMP7]], 12025// CHECK-32-EX-NEXT: [[TMP43:%.*]] = icmp uge i16 [[TMP5]], [[TMP6]]2026// CHECK-32-EX-NEXT: [[TMP44:%.*]] = and i1 [[TMP42]], [[TMP43]]2027// CHECK-32-EX-NEXT: br i1 [[TMP44]], label [[THEN5:%.*]], label [[ELSE6:%.*]]2028// CHECK-32-EX: then5:2029// CHECK-32-EX-NEXT: [[TMP45:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]], i32 0, i32 02030// CHECK-32-EX-NEXT: [[TMP46:%.*]] = load ptr, ptr [[TMP45]], align 42031// CHECK-32-EX-NEXT: [[TMP47:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP4]], i32 0, i32 02032// CHECK-32-EX-NEXT: [[TMP48:%.*]] = load ptr, ptr [[TMP47]], align 42033// CHECK-32-EX-NEXT: [[TMP49:%.*]] = load i32, ptr [[TMP46]], align 42034// CHECK-32-EX-NEXT: store i32 [[TMP49]], ptr [[TMP48]], align 42035// CHECK-32-EX-NEXT: [[TMP50:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]], i32 0, i32 12036// CHECK-32-EX-NEXT: [[TMP51:%.*]] = load ptr, ptr [[TMP50]], align 42037// CHECK-32-EX-NEXT: [[TMP52:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP4]], i32 0, i32 12038// CHECK-32-EX-NEXT: [[TMP53:%.*]] = load ptr, ptr [[TMP52]], align 42039// CHECK-32-EX-NEXT: [[TMP54:%.*]] = load i16, ptr [[TMP51]], align 22040// CHECK-32-EX-NEXT: store i16 [[TMP54]], ptr [[TMP53]], align 22041// CHECK-32-EX-NEXT: br label [[IFCONT7:%.*]]2042// CHECK-32-EX: else6:2043// CHECK-32-EX-NEXT: br label [[IFCONT7]]2044// CHECK-32-EX: ifcont7:2045// CHECK-32-EX-NEXT: ret void2046//2047//2048// CHECK-32-EX-LABEL: define {{[^@]+}}@_omp_reduction_inter_warp_copy_func42049// CHECK-32-EX-SAME: (ptr noundef [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR2]] {2050// CHECK-32-EX-NEXT: entry:2051// CHECK-32-EX-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 42052// CHECK-32-EX-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 42053// CHECK-32-EX-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 42054// CHECK-32-EX-NEXT: store i32 [[TMP1]], ptr [[DOTADDR1]], align 42055// CHECK-32-EX-NEXT: [[TMP3:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()2056// CHECK-32-EX-NEXT: [[TMP4:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()2057// CHECK-32-EX-NEXT: [[NVPTX_LANE_ID:%.*]] = and i32 [[TMP4]], 312058// CHECK-32-EX-NEXT: [[TMP5:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()2059// CHECK-32-EX-NEXT: [[NVPTX_WARP_ID:%.*]] = ashr i32 [[TMP5]], 52060// CHECK-32-EX-NEXT: [[TMP6:%.*]] = load ptr, ptr [[DOTADDR]], align 42061// CHECK-32-EX-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])2062// CHECK-32-EX-NEXT: call void @__kmpc_barrier(ptr @[[GLOB2]], i32 [[TMP2]])2063// CHECK-32-EX-NEXT: [[WARP_MASTER:%.*]] = icmp eq i32 [[NVPTX_LANE_ID]], 02064// CHECK-32-EX-NEXT: br i1 [[WARP_MASTER]], label [[THEN:%.*]], label [[ELSE:%.*]]2065// CHECK-32-EX: then:2066// CHECK-32-EX-NEXT: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP6]], i32 0, i32 02067// CHECK-32-EX-NEXT: [[TMP8:%.*]] = load ptr, ptr [[TMP7]], align 42068// CHECK-32-EX-NEXT: [[TMP9:%.*]] = getelementptr inbounds [32 x i32], ptr addrspace(3) @__openmp_nvptx_data_transfer_temporary_storage, i64 0, i32 [[NVPTX_WARP_ID]]2069// CHECK-32-EX-NEXT: [[TMP10:%.*]] = load i32, ptr [[TMP8]], align 42070// CHECK-32-EX-NEXT: store volatile i32 [[TMP10]], ptr addrspace(3) [[TMP9]], align 42071// CHECK-32-EX-NEXT: br label [[IFCONT:%.*]]2072// CHECK-32-EX: else:2073// CHECK-32-EX-NEXT: br label [[IFCONT]]2074// CHECK-32-EX: ifcont:2075// CHECK-32-EX-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])2076// CHECK-32-EX-NEXT: call void @__kmpc_barrier(ptr @[[GLOB2]], i32 [[TMP2]])2077// CHECK-32-EX-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTADDR1]], align 42078// CHECK-32-EX-NEXT: [[IS_ACTIVE_THREAD:%.*]] = icmp ult i32 [[TMP3]], [[TMP11]]2079// CHECK-32-EX-NEXT: br i1 [[IS_ACTIVE_THREAD]], label [[THEN2:%.*]], label [[ELSE3:%.*]]2080// CHECK-32-EX: then3:2081// CHECK-32-EX-NEXT: [[TMP12:%.*]] = getelementptr inbounds [32 x i32], ptr addrspace(3) @__openmp_nvptx_data_transfer_temporary_storage, i64 0, i32 [[TMP3]]2082// CHECK-32-EX-NEXT: [[TMP13:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP6]], i32 0, i32 02083// CHECK-32-EX-NEXT: [[TMP14:%.*]] = load ptr, ptr [[TMP13]], align 42084// CHECK-32-EX-NEXT: [[TMP15:%.*]] = load volatile i32, ptr addrspace(3) [[TMP12]], align 42085// CHECK-32-EX-NEXT: store i32 [[TMP15]], ptr [[TMP14]], align 42086// CHECK-32-EX-NEXT: br label [[IFCONT4:%.*]]2087// CHECK-32-EX: else4:2088// CHECK-32-EX-NEXT: br label [[IFCONT4]]2089// CHECK-32-EX: ifcont5:2090// CHECK-32-EX-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])2091// CHECK-32-EX-NEXT: call void @__kmpc_barrier(ptr @[[GLOB2]], i32 [[TMP2]])2092// CHECK-32-EX-NEXT: [[WARP_MASTER5:%.*]] = icmp eq i32 [[NVPTX_LANE_ID]], 02093// CHECK-32-EX-NEXT: br i1 [[WARP_MASTER5]], label [[THEN6:%.*]], label [[ELSE7:%.*]]2094// CHECK-32-EX: then8:2095// CHECK-32-EX-NEXT: [[TMP16:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP6]], i32 0, i32 12096// CHECK-32-EX-NEXT: [[TMP17:%.*]] = load ptr, ptr [[TMP16]], align 42097// CHECK-32-EX-NEXT: [[TMP18:%.*]] = getelementptr inbounds [32 x i32], ptr addrspace(3) @__openmp_nvptx_data_transfer_temporary_storage, i64 0, i32 [[NVPTX_WARP_ID]]2098// CHECK-32-EX-NEXT: [[TMP19:%.*]] = load i16, ptr [[TMP17]], align 22099// CHECK-32-EX-NEXT: store volatile i16 [[TMP19]], ptr addrspace(3) [[TMP18]], align 22100// CHECK-32-EX-NEXT: br label [[IFCONT8:%.*]]2101// CHECK-32-EX: else9:2102// CHECK-32-EX-NEXT: br label [[IFCONT8]]2103// CHECK-32-EX: ifcont10:2104// CHECK-32-EX-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])2105// CHECK-32-EX-NEXT: call void @__kmpc_barrier(ptr @[[GLOB2]], i32 [[TMP2]])2106// CHECK-32-EX-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTADDR1]], align 42107// CHECK-32-EX-NEXT: [[IS_ACTIVE_THREAD9:%.*]] = icmp ult i32 [[TMP3]], [[TMP20]]2108// CHECK-32-EX-NEXT: br i1 [[IS_ACTIVE_THREAD9]], label [[THEN10:%.*]], label [[ELSE11:%.*]]2109// CHECK-32-EX: then13:2110// CHECK-32-EX-NEXT: [[TMP21:%.*]] = getelementptr inbounds [32 x i32], ptr addrspace(3) @__openmp_nvptx_data_transfer_temporary_storage, i64 0, i32 [[TMP3]]2111// CHECK-32-EX-NEXT: [[TMP22:%.*]] = getelementptr inbounds [2 x ptr], ptr [[TMP6]], i32 0, i32 12112// CHECK-32-EX-NEXT: [[TMP23:%.*]] = load ptr, ptr [[TMP22]], align 42113// CHECK-32-EX-NEXT: [[TMP24:%.*]] = load volatile i16, ptr addrspace(3) [[TMP21]], align 22114// CHECK-32-EX-NEXT: store i16 [[TMP24]], ptr [[TMP23]], align 22115// CHECK-32-EX-NEXT: br label [[IFCONT12:%.*]]2116// CHECK-32-EX: else14:2117// CHECK-32-EX-NEXT: br label [[IFCONT12]]2118// CHECK-32-EX: ifcont15:2119// CHECK-32-EX-NEXT: ret void2120//2121