brintos / llvm-project-archived public Read only

1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _ --version 62// RUN: %clang_cc1 -no-enable-noundef-analysis -x c++ -O1 -disable-llvm-optzns -verify -fopenmp -internal-isystem %S/../Headers/Inputs/include -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc3// RUN: %clang_cc1 -no-enable-noundef-analysis -x c++ -O1 -disable-llvm-optzns -verify -fopenmp -internal-isystem %S/../Headers/Inputs/include -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown  -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK14// RUN: %clang_cc1 -no-enable-noundef-analysis -x c++ -O1 -disable-llvm-optzns -verify -fopenmp -internal-isystem %S/../Headers/Inputs/include -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc5// RUN: %clang_cc1 -no-enable-noundef-analysis -x c++ -O1 -disable-llvm-optzns -verify -fopenmp -internal-isystem %S/../Headers/Inputs/include -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix=CHECK16// RUN: %clang_cc1 -no-enable-noundef-analysis -x c++ -O1 -disable-llvm-optzns -verify -fopenmp -internal-isystem %S/../Headers/Inputs/include -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -fexceptions -fcxx-exceptions -aux-triple powerpc64le-unknown-unknown -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix=CHECK17// expected-no-diagnostics8#ifndef HEADER9#define HEADER10 11#include <complex>12 13// Verify we do not add tbaa metadata to type punned memory operations:14template <typename T>15void complex_reduction() {16#pragma omp target teams distribute17  for (int ib = 0; ib < 100; ib++) {18    std::complex<T> partial_sum;19    const int istart = ib * 4;20    const int iend = (ib + 1) * 4;21#pragma omp parallel for reduction(+ \22                                   : partial_sum)23    for (int i = istart; i < iend; i++)24      partial_sum += std::complex<T>(i, i);25  }26}27 28void test() {29  complex_reduction<float>();30  complex_reduction<double>();31}32#endif33// CHECK1-LABEL: define weak_odr protected ptx_kernel void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z17complex_reductionIfEvv_l16(34// CHECK1-SAME: ptr noalias [[DYN_PTR:%.*]]) #[[ATTR0:[0-9]+]] {35// CHECK1-NEXT:  [[ENTRY:.*:]]36// CHECK1-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 837// CHECK1-NEXT:    [[DOTZERO_ADDR:%.*]] = alloca i32, align 438// CHECK1-NEXT:    [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 439// CHECK1-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8, !tbaa [[ANYPTR_TBAA10:![0-9]+]]40// CHECK1-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z17complex_reductionIfEvv_l16_kernel_environment, ptr [[DYN_PTR]])41// CHECK1-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -142// CHECK1-NEXT:    br i1 [[EXEC_USER_CODE]], label %[[USER_CODE_ENTRY:.*]], label %[[WORKER_EXIT:.*]]43// CHECK1:       [[USER_CODE_ENTRY]]:44// CHECK1-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1:[0-9]+]])45// CHECK1-NEXT:    store i32 0, ptr [[DOTZERO_ADDR]], align 446// CHECK1-NEXT:    store i32 [[TMP1]], ptr [[DOTTHREADID_TEMP_]], align 4, !tbaa [[INT_TBAA6:![0-9]+]]47// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z17complex_reductionIfEvv_l16_omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTZERO_ADDR]]) #[[ATTR4:[0-9]+]]48// CHECK1-NEXT:    call void @__kmpc_target_deinit()49// CHECK1-NEXT:    ret void50// CHECK1:       [[WORKER_EXIT]]:51// CHECK1-NEXT:    ret void52//53//54// CHECK1-LABEL: define internal void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z17complex_reductionIfEvv_l16_omp_outlined(55// CHECK1-SAME: ptr noalias [[DOTGLOBAL_TID_:%.*]], ptr noalias [[DOTBOUND_TID_:%.*]]) #[[ATTR1:[0-9]+]] {56// CHECK1-NEXT:  [[ENTRY:.*:]]57// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 858// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 859// CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 460// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 461// CHECK1-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 462// CHECK1-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 463// CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 464// CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 465// CHECK1-NEXT:    [[IB:%.*]] = alloca i32, align 466// CHECK1-NEXT:    [[REF_TMP:%.*]] = alloca float, align 467// CHECK1-NEXT:    [[REF_TMP2:%.*]] = alloca float, align 468// CHECK1-NEXT:    [[CAPTURED_VARS_ADDRS:%.*]] = alloca [3 x ptr], align 869// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8, !tbaa [[INTPTR_TBAA12:![0-9]+]]70// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8, !tbaa [[INTPTR_TBAA12]]71// CHECK1-NEXT:    [[ISTART:%.*]] = call align 16 ptr @__kmpc_alloc_shared(i64 4)72// CHECK1-NEXT:    [[IEND:%.*]] = call align 16 ptr @__kmpc_alloc_shared(i64 4)73// CHECK1-NEXT:    [[PARTIAL_SUM:%.*]] = call align 16 ptr @__kmpc_alloc_shared(i64 8)74// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[DOTOMP_IV]]) #[[ATTR4]]75// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[DOTOMP_LB]]) #[[ATTR4]]76// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4, !tbaa [[INT_TBAA6]]77// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[DOTOMP_UB]]) #[[ATTR4]]78// CHECK1-NEXT:    store i32 99, ptr [[DOTOMP_UB]], align 4, !tbaa [[INT_TBAA6]]79// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[DOTOMP_STRIDE]]) #[[ATTR4]]80// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4, !tbaa [[INT_TBAA6]]81// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[DOTOMP_IS_LAST]]) #[[ATTR4]]82// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4, !tbaa [[INT_TBAA6]]83// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[IB]]) #[[ATTR4]]84// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 885// CHECK1-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4, !tbaa [[INT_TBAA6]]86// CHECK1-NEXT:    call void @__kmpc_distribute_static_init_4(ptr @[[GLOB2:[0-9]+]], i32 [[TMP1]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)87// CHECK1-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !tbaa [[INT_TBAA6]]88// CHECK1-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 9989// CHECK1-NEXT:    br i1 [[CMP]], label %[[COND_TRUE:.*]], label %[[COND_FALSE:.*]]90// CHECK1:       [[COND_TRUE]]:91// CHECK1-NEXT:    br label %[[COND_END:.*]]92// CHECK1:       [[COND_FALSE]]:93// CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !tbaa [[INT_TBAA6]]94// CHECK1-NEXT:    br label %[[COND_END]]95// CHECK1:       [[COND_END]]:96// CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ 99, %[[COND_TRUE]] ], [ [[TMP3]], %[[COND_FALSE]] ]97// CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4, !tbaa [[INT_TBAA6]]98// CHECK1-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4, !tbaa [[INT_TBAA6]]99// CHECK1-NEXT:    store i32 [[TMP4]], ptr [[DOTOMP_IV]], align 4, !tbaa [[INT_TBAA6]]100// CHECK1-NEXT:    br label %[[OMP_INNER_FOR_COND:.*]]101// CHECK1:       [[OMP_INNER_FOR_COND]]:102// CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !tbaa [[INT_TBAA6]]103// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !tbaa [[INT_TBAA6]]104// CHECK1-NEXT:    [[CMP1:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]105// CHECK1-NEXT:    br i1 [[CMP1]], label %[[OMP_INNER_FOR_BODY:.*]], label %[[OMP_INNER_FOR_COND_CLEANUP:.*]]106// CHECK1:       [[OMP_INNER_FOR_COND_CLEANUP]]:107// CHECK1-NEXT:    br label %[[OMP_INNER_FOR_END:.*]]108// CHECK1:       [[OMP_INNER_FOR_BODY]]:109// CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !tbaa [[INT_TBAA6]]110// CHECK1-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1111// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]112// CHECK1-NEXT:    store i32 [[ADD]], ptr [[IB]], align 4, !tbaa [[INT_TBAA6]]113// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[REF_TMP]]) #[[ATTR4]]114// CHECK1-NEXT:    store float 0.000000e+00, ptr [[REF_TMP]], align 4, !tbaa [[FLOAT_TBAA14:![0-9]+]]115// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[REF_TMP2]]) #[[ATTR4]]116// CHECK1-NEXT:    store float 0.000000e+00, ptr [[REF_TMP2]], align 4, !tbaa [[FLOAT_TBAA14]]117// CHECK1-NEXT:    call void @_ZNSt7complexIfEC1ERKfS2_(ptr nonnull align 4 dereferenceable(8) [[PARTIAL_SUM]], ptr nonnull align 4 dereferenceable(4) [[REF_TMP]], ptr nonnull align 4 dereferenceable(4) [[REF_TMP2]]) #[[ATTR11:[0-9]+]]118// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[REF_TMP2]]) #[[ATTR4]]119// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[REF_TMP]]) #[[ATTR4]]120// CHECK1-NEXT:    [[TMP8:%.*]] = load i32, ptr [[IB]], align 4, !tbaa [[INT_TBAA6]]121// CHECK1-NEXT:    [[MUL3:%.*]] = mul nsw i32 [[TMP8]], 4122// CHECK1-NEXT:    store i32 [[MUL3]], ptr [[ISTART]], align 4, !tbaa [[INT_TBAA6]]123// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[IB]], align 4, !tbaa [[INT_TBAA6]]124// CHECK1-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP9]], 1125// CHECK1-NEXT:    [[MUL5:%.*]] = mul nsw i32 [[ADD4]], 4126// CHECK1-NEXT:    store i32 [[MUL5]], ptr [[IEND]], align 4, !tbaa [[INT_TBAA6]]127// CHECK1-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [3 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 0128// CHECK1-NEXT:    store ptr [[ISTART]], ptr [[TMP10]], align 8, !tbaa [[ANYPTR_TBAA16:![0-9]+]]129// CHECK1-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 1130// CHECK1-NEXT:    store ptr [[IEND]], ptr [[TMP11]], align 8, !tbaa [[ANYPTR_TBAA16]]131// CHECK1-NEXT:    [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 2132// CHECK1-NEXT:    store ptr [[PARTIAL_SUM]], ptr [[TMP12]], align 8, !tbaa [[ANYPTR_TBAA16]]133// CHECK1-NEXT:    call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z17complex_reductionIfEvv_l16_omp_outlined_omp_outlined, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z17complex_reductionIfEvv_l16_omp_outlined_omp_outlined_wrapper, ptr [[CAPTURED_VARS_ADDRS]], i64 3)134// CHECK1-NEXT:    br label %[[OMP_BODY_CONTINUE:.*]]135// CHECK1:       [[OMP_BODY_CONTINUE]]:136// CHECK1-NEXT:    br label %[[OMP_INNER_FOR_INC:.*]]137// CHECK1:       [[OMP_INNER_FOR_INC]]:138// CHECK1-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !tbaa [[INT_TBAA6]]139// CHECK1-NEXT:    [[ADD6:%.*]] = add nsw i32 [[TMP13]], 1140// CHECK1-NEXT:    store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 4, !tbaa [[INT_TBAA6]]141// CHECK1-NEXT:    br label %[[OMP_INNER_FOR_COND]]142// CHECK1:       [[OMP_INNER_FOR_END]]:143// CHECK1-NEXT:    br label %[[OMP_LOOP_EXIT:.*]]144// CHECK1:       [[OMP_LOOP_EXIT]]:145// CHECK1-NEXT:    call void @__kmpc_distribute_static_fini(ptr @[[GLOB2]], i32 [[TMP1]])146// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[IB]]) #[[ATTR4]]147// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[DOTOMP_IS_LAST]]) #[[ATTR4]]148// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[DOTOMP_STRIDE]]) #[[ATTR4]]149// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[DOTOMP_UB]]) #[[ATTR4]]150// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[DOTOMP_LB]]) #[[ATTR4]]151// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[DOTOMP_IV]]) #[[ATTR4]]152// CHECK1-NEXT:    call void @__kmpc_free_shared(ptr [[PARTIAL_SUM]], i64 8)153// CHECK1-NEXT:    call void @__kmpc_free_shared(ptr [[IEND]], i64 4)154// CHECK1-NEXT:    call void @__kmpc_free_shared(ptr [[ISTART]], i64 4)155// CHECK1-NEXT:    ret void156//157//158// CHECK1-LABEL: define linkonce_odr hidden void @_ZNSt7complexIfEC1ERKfS2_(159// CHECK1-SAME: ptr nonnull align 4 dereferenceable(8) [[THIS:%.*]], ptr nonnull align 4 dereferenceable(4) [[__RE:%.*]], ptr nonnull align 4 dereferenceable(4) [[__IM:%.*]]) unnamed_addr #[[ATTR5:[0-9]+]] comdat align 2 {160// CHECK1-NEXT:  [[ENTRY:.*:]]161// CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8162// CHECK1-NEXT:    [[__RE_ADDR:%.*]] = alloca ptr, align 8163// CHECK1-NEXT:    [[__IM_ADDR:%.*]] = alloca ptr, align 8164// CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8, !tbaa [[_ZTSST7COMPLEXIFEPTR_TBAA18:![0-9]+]]165// CHECK1-NEXT:    store ptr [[__RE]], ptr [[__RE_ADDR]], align 8, !tbaa [[FLOATPTR_TBAA20:![0-9]+]]166// CHECK1-NEXT:    store ptr [[__IM]], ptr [[__IM_ADDR]], align 8, !tbaa [[FLOATPTR_TBAA20]]167// CHECK1-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8168// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[__RE_ADDR]], align 8169// CHECK1-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[__IM_ADDR]], align 8170// CHECK1-NEXT:    call void @_ZNSt7complexIfEC2ERKfS2_(ptr nonnull align 4 dereferenceable(8) [[THIS1]], ptr nonnull align 4 dereferenceable(4) [[TMP0]], ptr nonnull align 4 dereferenceable(4) [[TMP1]]) #[[ATTR11]]171// CHECK1-NEXT:    ret void172//173//174// CHECK1-LABEL: define internal void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z17complex_reductionIfEvv_l16_omp_outlined_omp_outlined(175// CHECK1-SAME: ptr noalias [[DOTGLOBAL_TID_:%.*]], ptr noalias [[DOTBOUND_TID_:%.*]], ptr nonnull align 4 dereferenceable(4) [[ISTART:%.*]], ptr nonnull align 4 dereferenceable(4) [[IEND:%.*]], ptr nonnull align 4 dereferenceable(8) [[PARTIAL_SUM:%.*]]) #[[ATTR1]] {176// CHECK1-NEXT:  [[ENTRY:.*:]]177// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8178// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8179// CHECK1-NEXT:    [[ISTART_ADDR:%.*]] = alloca ptr, align 8180// CHECK1-NEXT:    [[IEND_ADDR:%.*]] = alloca ptr, align 8181// CHECK1-NEXT:    [[PARTIAL_SUM_ADDR:%.*]] = alloca ptr, align 8182// CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4183// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4184// CHECK1-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4185// CHECK1-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4186// CHECK1-NEXT:    [[DOTCAPTURE_EXPR_2:%.*]] = alloca i32, align 4187// CHECK1-NEXT:    [[I:%.*]] = alloca i32, align 4188// CHECK1-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4189// CHECK1-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4190// CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4191// CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4192// CHECK1-NEXT:    [[PARTIAL_SUM5:%.*]] = alloca %"class.std::complex", align 4193// CHECK1-NEXT:    [[REF_TMP:%.*]] = alloca float, align 4194// CHECK1-NEXT:    [[REF_TMP6:%.*]] = alloca float, align 4195// CHECK1-NEXT:    [[I7:%.*]] = alloca i32, align 4196// CHECK1-NEXT:    [[REF_TMP14:%.*]] = alloca %"class.std::complex", align 4197// CHECK1-NEXT:    [[REF_TMP15:%.*]] = alloca float, align 4198// CHECK1-NEXT:    [[REF_TMP16:%.*]] = alloca float, align 4199// CHECK1-NEXT:    [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 8200// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8, !tbaa [[INTPTR_TBAA12]]201// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8, !tbaa [[INTPTR_TBAA12]]202// CHECK1-NEXT:    store ptr [[ISTART]], ptr [[ISTART_ADDR]], align 8, !tbaa [[INTPTR_TBAA12]]203// CHECK1-NEXT:    store ptr [[IEND]], ptr [[IEND_ADDR]], align 8, !tbaa [[INTPTR_TBAA12]]204// CHECK1-NEXT:    store ptr [[PARTIAL_SUM]], ptr [[PARTIAL_SUM_ADDR]], align 8, !tbaa [[_ZTSST7COMPLEXIFEPTR_TBAA18]]205// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[ISTART_ADDR]], align 8, !tbaa [[INTPTR_TBAA12]], !nonnull [[META22:![0-9]+]], !align [[META23:![0-9]+]]206// CHECK1-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[IEND_ADDR]], align 8, !tbaa [[INTPTR_TBAA12]], !nonnull [[META22]], !align [[META23]]207// CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[PARTIAL_SUM_ADDR]], align 8, !tbaa [[_ZTSST7COMPLEXIFEPTR_TBAA18]], !nonnull [[META22]], !align [[META23]]208// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[DOTOMP_IV]]) #[[ATTR4]]209// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[DOTCAPTURE_EXPR_]]) #[[ATTR4]]210// CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[TMP0]], align 4, !tbaa [[INT_TBAA6]]211// CHECK1-NEXT:    store i32 [[TMP3]], ptr [[DOTCAPTURE_EXPR_]], align 4, !tbaa [[INT_TBAA6]]212// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[DOTCAPTURE_EXPR_1]]) #[[ATTR4]]213// CHECK1-NEXT:    [[TMP4:%.*]] = load i32, ptr [[TMP1]], align 4, !tbaa [[INT_TBAA6]]214// CHECK1-NEXT:    store i32 [[TMP4]], ptr [[DOTCAPTURE_EXPR_1]], align 4, !tbaa [[INT_TBAA6]]215// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[DOTCAPTURE_EXPR_2]]) #[[ATTR4]]216// CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4, !tbaa [[INT_TBAA6]]217// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4, !tbaa [[INT_TBAA6]]218// CHECK1-NEXT:    [[SUB:%.*]] = sub i32 [[TMP5]], [[TMP6]]219// CHECK1-NEXT:    [[SUB3:%.*]] = sub i32 [[SUB]], 1220// CHECK1-NEXT:    [[ADD:%.*]] = add i32 [[SUB3]], 1221// CHECK1-NEXT:    [[DIV:%.*]] = udiv i32 [[ADD]], 1222// CHECK1-NEXT:    [[SUB4:%.*]] = sub i32 [[DIV]], 1223// CHECK1-NEXT:    store i32 [[SUB4]], ptr [[DOTCAPTURE_EXPR_2]], align 4, !tbaa [[INT_TBAA6]]224// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[I]]) #[[ATTR4]]225// CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4, !tbaa [[INT_TBAA6]]226// CHECK1-NEXT:    store i32 [[TMP7]], ptr [[I]], align 4, !tbaa [[INT_TBAA6]]227// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[I]]) #[[ATTR4]]228// CHECK1-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4, !tbaa [[INT_TBAA6]]229// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4, !tbaa [[INT_TBAA6]]230// CHECK1-NEXT:    [[CMP:%.*]] = icmp slt i32 [[TMP8]], [[TMP9]]231// CHECK1-NEXT:    br i1 [[CMP]], label %[[OMP_PRECOND_THEN:.*]], label %[[OMP_PRECOND_END:.*]]232// CHECK1:       [[OMP_PRECOND_THEN]]:233// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[DOTOMP_LB]]) #[[ATTR4]]234// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4, !tbaa [[INT_TBAA6]]235// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[DOTOMP_UB]]) #[[ATTR4]]236// CHECK1-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_2]], align 4, !tbaa [[INT_TBAA6]]237// CHECK1-NEXT:    store i32 [[TMP10]], ptr [[DOTOMP_UB]], align 4, !tbaa [[INT_TBAA6]]238// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[DOTOMP_STRIDE]]) #[[ATTR4]]239// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4, !tbaa [[INT_TBAA6]]240// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[DOTOMP_IS_LAST]]) #[[ATTR4]]241// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4, !tbaa [[INT_TBAA6]]242// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[PARTIAL_SUM5]]) #[[ATTR4]]243// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[REF_TMP]]) #[[ATTR4]]244// CHECK1-NEXT:    store float 0.000000e+00, ptr [[REF_TMP]], align 4, !tbaa [[FLOAT_TBAA14]]245// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[REF_TMP6]]) #[[ATTR4]]246// CHECK1-NEXT:    store float 0.000000e+00, ptr [[REF_TMP6]], align 4, !tbaa [[FLOAT_TBAA14]]247// CHECK1-NEXT:    call void @_ZNSt7complexIfEC1ERKfS2_(ptr nonnull align 4 dereferenceable(8) [[PARTIAL_SUM5]], ptr nonnull align 4 dereferenceable(4) [[REF_TMP]], ptr nonnull align 4 dereferenceable(4) [[REF_TMP6]]) #[[ATTR11]]248// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[REF_TMP6]]) #[[ATTR4]]249// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[REF_TMP]]) #[[ATTR4]]250// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[I7]]) #[[ATTR4]]251// CHECK1-NEXT:    [[TMP11:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8252// CHECK1-NEXT:    [[TMP12:%.*]] = load i32, ptr [[TMP11]], align 4, !tbaa [[INT_TBAA6]]253// CHECK1-NEXT:    call void @__kmpc_for_static_init_4u(ptr @[[GLOB3:[0-9]+]], i32 [[TMP12]], i32 33, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)254// CHECK1-NEXT:    br label %[[OMP_DISPATCH_COND:.*]]255// CHECK1:       [[OMP_DISPATCH_COND]]:256// CHECK1-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !tbaa [[INT_TBAA6]]257// CHECK1-NEXT:    [[TMP14:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_2]], align 4, !tbaa [[INT_TBAA6]]258// CHECK1-NEXT:    [[CMP8:%.*]] = icmp ugt i32 [[TMP13]], [[TMP14]]259// CHECK1-NEXT:    br i1 [[CMP8]], label %[[COND_TRUE:.*]], label %[[COND_FALSE:.*]]260// CHECK1:       [[COND_TRUE]]:261// CHECK1-NEXT:    [[TMP15:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_2]], align 4, !tbaa [[INT_TBAA6]]262// CHECK1-NEXT:    br label %[[COND_END:.*]]263// CHECK1:       [[COND_FALSE]]:264// CHECK1-NEXT:    [[TMP16:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !tbaa [[INT_TBAA6]]265// CHECK1-NEXT:    br label %[[COND_END]]266// CHECK1:       [[COND_END]]:267// CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ [[TMP15]], %[[COND_TRUE]] ], [ [[TMP16]], %[[COND_FALSE]] ]268// CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4, !tbaa [[INT_TBAA6]]269// CHECK1-NEXT:    [[TMP17:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4, !tbaa [[INT_TBAA6]]270// CHECK1-NEXT:    store i32 [[TMP17]], ptr [[DOTOMP_IV]], align 4, !tbaa [[INT_TBAA6]]271// CHECK1-NEXT:    [[TMP18:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !tbaa [[INT_TBAA6]]272// CHECK1-NEXT:    [[TMP19:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !tbaa [[INT_TBAA6]]273// CHECK1-NEXT:    [[ADD9:%.*]] = add i32 [[TMP19]], 1274// CHECK1-NEXT:    [[CMP10:%.*]] = icmp ult i32 [[TMP18]], [[ADD9]]275// CHECK1-NEXT:    br i1 [[CMP10]], label %[[OMP_DISPATCH_BODY:.*]], label %[[OMP_DISPATCH_CLEANUP:.*]]276// CHECK1:       [[OMP_DISPATCH_CLEANUP]]:277// CHECK1-NEXT:    br label %[[OMP_DISPATCH_END:.*]]278// CHECK1:       [[OMP_DISPATCH_BODY]]:279// CHECK1-NEXT:    br label %[[OMP_INNER_FOR_COND:.*]]280// CHECK1:       [[OMP_INNER_FOR_COND]]:281// CHECK1-NEXT:    [[TMP20:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !tbaa [[INT_TBAA6]]282// CHECK1-NEXT:    [[TMP21:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !tbaa [[INT_TBAA6]]283// CHECK1-NEXT:    [[ADD11:%.*]] = add i32 [[TMP21]], 1284// CHECK1-NEXT:    [[CMP12:%.*]] = icmp ult i32 [[TMP20]], [[ADD11]]285// CHECK1-NEXT:    br i1 [[CMP12]], label %[[OMP_INNER_FOR_BODY:.*]], label %[[OMP_INNER_FOR_COND_CLEANUP:.*]]286// CHECK1:       [[OMP_INNER_FOR_COND_CLEANUP]]:287// CHECK1-NEXT:    br label %[[OMP_INNER_FOR_END:.*]]288// CHECK1:       [[OMP_INNER_FOR_BODY]]:289// CHECK1-NEXT:    [[TMP22:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4, !tbaa [[INT_TBAA6]]290// CHECK1-NEXT:    [[TMP23:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !tbaa [[INT_TBAA6]]291// CHECK1-NEXT:    [[MUL:%.*]] = mul i32 [[TMP23]], 1292// CHECK1-NEXT:    [[ADD13:%.*]] = add i32 [[TMP22]], [[MUL]]293// CHECK1-NEXT:    store i32 [[ADD13]], ptr [[I7]], align 4, !tbaa [[INT_TBAA6]]294// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[REF_TMP14]]) #[[ATTR4]]295// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[REF_TMP15]]) #[[ATTR4]]296// CHECK1-NEXT:    [[TMP24:%.*]] = load i32, ptr [[I7]], align 4, !tbaa [[INT_TBAA6]]297// CHECK1-NEXT:    [[CONV:%.*]] = sitofp i32 [[TMP24]] to float298// CHECK1-NEXT:    store float [[CONV]], ptr [[REF_TMP15]], align 4, !tbaa [[FLOAT_TBAA14]]299// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[REF_TMP16]]) #[[ATTR4]]300// CHECK1-NEXT:    [[TMP25:%.*]] = load i32, ptr [[I7]], align 4, !tbaa [[INT_TBAA6]]301// CHECK1-NEXT:    [[CONV17:%.*]] = sitofp i32 [[TMP25]] to float302// CHECK1-NEXT:    store float [[CONV17]], ptr [[REF_TMP16]], align 4, !tbaa [[FLOAT_TBAA14]]303// CHECK1-NEXT:    call void @_ZNSt7complexIfEC1ERKfS2_(ptr nonnull align 4 dereferenceable(8) [[REF_TMP14]], ptr nonnull align 4 dereferenceable(4) [[REF_TMP15]], ptr nonnull align 4 dereferenceable(4) [[REF_TMP16]]) #[[ATTR11]]304// CHECK1-NEXT:    [[CALL:%.*]] = call nonnull align 4 dereferenceable(8) ptr @_ZNSt7complexIfEpLIfEERS0_RKS_IT_E(ptr nonnull align 4 dereferenceable(8) [[PARTIAL_SUM5]], ptr nonnull align 4 dereferenceable(8) [[REF_TMP14]]) #[[ATTR11]]305// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[REF_TMP16]]) #[[ATTR4]]306// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[REF_TMP15]]) #[[ATTR4]]307// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[REF_TMP14]]) #[[ATTR4]]308// CHECK1-NEXT:    br label %[[OMP_BODY_CONTINUE:.*]]309// CHECK1:       [[OMP_BODY_CONTINUE]]:310// CHECK1-NEXT:    br label %[[OMP_INNER_FOR_INC:.*]]311// CHECK1:       [[OMP_INNER_FOR_INC]]:312// CHECK1-NEXT:    [[TMP26:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !tbaa [[INT_TBAA6]]313// CHECK1-NEXT:    [[ADD18:%.*]] = add i32 [[TMP26]], 1314// CHECK1-NEXT:    store i32 [[ADD18]], ptr [[DOTOMP_IV]], align 4, !tbaa [[INT_TBAA6]]315// CHECK1-NEXT:    br label %[[OMP_INNER_FOR_COND]]316// CHECK1:       [[OMP_INNER_FOR_END]]:317// CHECK1-NEXT:    br label %[[OMP_DISPATCH_INC:.*]]318// CHECK1:       [[OMP_DISPATCH_INC]]:319// CHECK1-NEXT:    [[TMP27:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4, !tbaa [[INT_TBAA6]]320// CHECK1-NEXT:    [[TMP28:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4, !tbaa [[INT_TBAA6]]321// CHECK1-NEXT:    [[ADD19:%.*]] = add i32 [[TMP27]], [[TMP28]]322// CHECK1-NEXT:    store i32 [[ADD19]], ptr [[DOTOMP_LB]], align 4, !tbaa [[INT_TBAA6]]323// CHECK1-NEXT:    [[TMP29:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !tbaa [[INT_TBAA6]]324// CHECK1-NEXT:    [[TMP30:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4, !tbaa [[INT_TBAA6]]325// CHECK1-NEXT:    [[ADD20:%.*]] = add i32 [[TMP29]], [[TMP30]]326// CHECK1-NEXT:    store i32 [[ADD20]], ptr [[DOTOMP_UB]], align 4, !tbaa [[INT_TBAA6]]327// CHECK1-NEXT:    br label %[[OMP_DISPATCH_COND]]328// CHECK1:       [[OMP_DISPATCH_END]]:329// CHECK1-NEXT:    [[TMP31:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8330// CHECK1-NEXT:    [[TMP32:%.*]] = load i32, ptr [[TMP31]], align 4, !tbaa [[INT_TBAA6]]331// CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB3]], i32 [[TMP32]])332// CHECK1-NEXT:    [[TMP33:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i64 0, i64 0333// CHECK1-NEXT:    store ptr [[PARTIAL_SUM5]], ptr [[TMP33]], align 8334// CHECK1-NEXT:    [[TMP34:%.*]] = 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)335// CHECK1-NEXT:    [[TMP35:%.*]] = icmp eq i32 [[TMP34]], 1336// CHECK1-NEXT:    br i1 [[TMP35]], [[DOTOMP_REDUCTION_THEN:label %.*]], [[DOTOMP_REDUCTION_DONE:label %.*]]337// CHECK1:       [[_OMP_REDUCTION_THEN:.*:]]338// CHECK1-NEXT:    [[CALL21:%.*]] = call nonnull align 4 dereferenceable(8) ptr @_ZNSt7complexIfEpLIfEERS0_RKS_IT_E(ptr nonnull align 4 dereferenceable(8) [[TMP2]], ptr nonnull align 4 dereferenceable(8) [[PARTIAL_SUM5]]) #[[ATTR11]]339// CHECK1-NEXT:    br [[DOTOMP_REDUCTION_DONE]]340// CHECK1:       [[_OMP_REDUCTION_DONE:.*:]]341// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[I7]]) #[[ATTR4]]342// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[PARTIAL_SUM5]]) #[[ATTR4]]343// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[DOTOMP_IS_LAST]]) #[[ATTR4]]344// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[DOTOMP_STRIDE]]) #[[ATTR4]]345// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[DOTOMP_UB]]) #[[ATTR4]]346// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[DOTOMP_LB]]) #[[ATTR4]]347// CHECK1-NEXT:    br label %[[OMP_PRECOND_END]]348// CHECK1:       [[OMP_PRECOND_END]]:349// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[DOTCAPTURE_EXPR_2]]) #[[ATTR4]]350// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[DOTCAPTURE_EXPR_1]]) #[[ATTR4]]351// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[DOTCAPTURE_EXPR_]]) #[[ATTR4]]352// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[DOTOMP_IV]]) #[[ATTR4]]353// CHECK1-NEXT:    ret void354//355//356// CHECK1-LABEL: define linkonce_odr hidden nonnull align 4 dereferenceable(8) ptr @_ZNSt7complexIfEpLIfEERS0_RKS_IT_E(357// CHECK1-SAME: ptr nonnull align 4 dereferenceable(8) [[THIS:%.*]], ptr nonnull align 4 dereferenceable(8) [[__C:%.*]]) #[[ATTR5]] comdat align 2 {358// CHECK1-NEXT:  [[ENTRY:.*:]]359// CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8360// CHECK1-NEXT:    [[__C_ADDR:%.*]] = alloca ptr, align 8361// CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8, !tbaa [[_ZTSST7COMPLEXIFEPTR_TBAA18]]362// CHECK1-NEXT:    store ptr [[__C]], ptr [[__C_ADDR]], align 8, !tbaa [[_ZTSST7COMPLEXIFEPTR_TBAA18]]363// CHECK1-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8364// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[__C_ADDR]], align 8, !tbaa [[_ZTSST7COMPLEXIFEPTR_TBAA18]], !nonnull [[META22]], !align [[META23]]365// CHECK1-NEXT:    [[CALL:%.*]] = call float @_ZNKSt7complexIfE4realEv(ptr nonnull align 4 dereferenceable(8) [[TMP0]]) #[[ATTR11]]366// CHECK1-NEXT:    [[__RE_:%.*]] = getelementptr inbounds nuw %"class.std::complex", ptr [[THIS1]], i32 0, i32 0367// CHECK1-NEXT:    [[TMP1:%.*]] = load float, ptr [[__RE_]], align 4, !tbaa [[FLOAT_TBAA24:![0-9]+]]368// CHECK1-NEXT:    [[ADD:%.*]] = fadd float [[TMP1]], [[CALL]]369// CHECK1-NEXT:    store float [[ADD]], ptr [[__RE_]], align 4, !tbaa [[FLOAT_TBAA24]]370// CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[__C_ADDR]], align 8, !tbaa [[_ZTSST7COMPLEXIFEPTR_TBAA18]], !nonnull [[META22]], !align [[META23]]371// CHECK1-NEXT:    [[CALL2:%.*]] = call float @_ZNKSt7complexIfE4imagEv(ptr nonnull align 4 dereferenceable(8) [[TMP2]]) #[[ATTR11]]372// CHECK1-NEXT:    [[__IM_:%.*]] = getelementptr inbounds nuw %"class.std::complex", ptr [[THIS1]], i32 0, i32 1373// CHECK1-NEXT:    [[TMP3:%.*]] = load float, ptr [[__IM_]], align 4, !tbaa [[FLOAT_TBAA26:![0-9]+]]374// CHECK1-NEXT:    [[ADD3:%.*]] = fadd float [[TMP3]], [[CALL2]]375// CHECK1-NEXT:    store float [[ADD3]], ptr [[__IM_]], align 4, !tbaa [[FLOAT_TBAA26]]376// CHECK1-NEXT:    ret ptr [[THIS1]]377//378//379// CHECK1-LABEL: define internal void @_omp_reduction_shuffle_and_reduce_func(380// CHECK1-SAME: ptr noundef [[TMP0:%.*]], i16 noundef signext [[TMP1:%.*]], i16 noundef signext [[TMP2:%.*]], i16 noundef signext [[TMP3:%.*]]) #[[ATTR1]] {381// CHECK1-NEXT:  [[ENTRY:.*:]]382// CHECK1-NEXT:    [[DOTADDR:%.*]] = alloca ptr, align 8383// CHECK1-NEXT:    [[DOTADDR1:%.*]] = alloca i16, align 2384// CHECK1-NEXT:    [[DOTADDR2:%.*]] = alloca i16, align 2385// CHECK1-NEXT:    [[DOTADDR3:%.*]] = alloca i16, align 2386// CHECK1-NEXT:    [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST:%.*]] = alloca [1 x ptr], align 8387// CHECK1-NEXT:    [[DOTOMP_REDUCTION_ELEMENT:%.*]] = alloca %"class.std::complex", align 8388// CHECK1-NEXT:    store ptr [[TMP0]], ptr [[DOTADDR]], align 8389// CHECK1-NEXT:    store i16 [[TMP1]], ptr [[DOTADDR1]], align 2390// CHECK1-NEXT:    store i16 [[TMP2]], ptr [[DOTADDR2]], align 2391// CHECK1-NEXT:    store i16 [[TMP3]], ptr [[DOTADDR3]], align 2392// CHECK1-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[DOTADDR]], align 8393// CHECK1-NEXT:    [[TMP5:%.*]] = load i16, ptr [[DOTADDR1]], align 2394// CHECK1-NEXT:    [[TMP6:%.*]] = load i16, ptr [[DOTADDR2]], align 2395// CHECK1-NEXT:    [[TMP7:%.*]] = load i16, ptr [[DOTADDR3]], align 2396// CHECK1-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP4]], i64 0, i64 0397// CHECK1-NEXT:    [[TMP9:%.*]] = load ptr, ptr [[TMP8]], align 8398// CHECK1-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]], i64 0, i64 0399// CHECK1-NEXT:    [[TMP11:%.*]] = getelementptr %"class.std::complex", ptr [[TMP9]], i64 1400// CHECK1-NEXT:    [[TMP12:%.*]] = load i64, ptr [[TMP9]], align 8401// CHECK1-NEXT:    [[TMP13:%.*]] = call i32 @__kmpc_get_warp_size()402// CHECK1-NEXT:    [[TMP14:%.*]] = trunc i32 [[TMP13]] to i16403// CHECK1-NEXT:    [[TMP15:%.*]] = call i64 @__kmpc_shuffle_int64(i64 [[TMP12]], i16 [[TMP6]], i16 [[TMP14]])404// CHECK1-NEXT:    store i64 [[TMP15]], ptr [[DOTOMP_REDUCTION_ELEMENT]], align 8405// CHECK1-NEXT:    [[TMP16:%.*]] = getelementptr i64, ptr [[TMP9]], i64 1406// CHECK1-NEXT:    [[TMP17:%.*]] = getelementptr i64, ptr [[DOTOMP_REDUCTION_ELEMENT]], i64 1407// CHECK1-NEXT:    store ptr [[DOTOMP_REDUCTION_ELEMENT]], ptr [[TMP10]], align 8408// CHECK1-NEXT:    [[TMP18:%.*]] = icmp eq i16 [[TMP7]], 0409// CHECK1-NEXT:    [[TMP19:%.*]] = icmp eq i16 [[TMP7]], 1410// CHECK1-NEXT:    [[TMP20:%.*]] = icmp ult i16 [[TMP5]], [[TMP6]]411// CHECK1-NEXT:    [[TMP21:%.*]] = and i1 [[TMP19]], [[TMP20]]412// CHECK1-NEXT:    [[TMP22:%.*]] = icmp eq i16 [[TMP7]], 2413// CHECK1-NEXT:    [[TMP23:%.*]] = and i16 [[TMP5]], 1414// CHECK1-NEXT:    [[TMP24:%.*]] = icmp eq i16 [[TMP23]], 0415// CHECK1-NEXT:    [[TMP25:%.*]] = and i1 [[TMP22]], [[TMP24]]416// CHECK1-NEXT:    [[TMP26:%.*]] = icmp sgt i16 [[TMP6]], 0417// CHECK1-NEXT:    [[TMP27:%.*]] = and i1 [[TMP25]], [[TMP26]]418// CHECK1-NEXT:    [[TMP28:%.*]] = or i1 [[TMP18]], [[TMP21]]419// CHECK1-NEXT:    [[TMP29:%.*]] = or i1 [[TMP28]], [[TMP27]]420// CHECK1-NEXT:    br i1 [[TMP29]], label %[[THEN:.*]], label %[[ELSE:.*]]421// CHECK1:       [[THEN]]:422// CHECK1-NEXT:    call void @"{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z17complex_reductionIfEvv_l16_omp_outlined_omp_outlined_omp$reduction$reduction_func"(ptr [[TMP4]], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]]) #[[ATTR4]]423// CHECK1-NEXT:    br label %[[IFCONT:.*]]424// CHECK1:       [[ELSE]]:425// CHECK1-NEXT:    br label %[[IFCONT]]426// CHECK1:       [[IFCONT]]:427// CHECK1-NEXT:    [[TMP30:%.*]] = icmp eq i16 [[TMP7]], 1428// CHECK1-NEXT:    [[TMP31:%.*]] = icmp uge i16 [[TMP5]], [[TMP6]]429// CHECK1-NEXT:    [[TMP32:%.*]] = and i1 [[TMP30]], [[TMP31]]430// CHECK1-NEXT:    br i1 [[TMP32]], label %[[THEN4:.*]], label %[[ELSE5:.*]]431// CHECK1:       [[THEN4]]:432// CHECK1-NEXT:    [[TMP33:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]], i64 0, i64 0433// CHECK1-NEXT:    [[TMP34:%.*]] = load ptr, ptr [[TMP33]], align 8434// CHECK1-NEXT:    [[TMP35:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP4]], i64 0, i64 0435// CHECK1-NEXT:    [[TMP36:%.*]] = load ptr, ptr [[TMP35]], align 8436// CHECK1-NEXT:    call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[TMP36]], ptr align 8 [[TMP34]], i64 8, i1 false)437// CHECK1-NEXT:    br label %[[IFCONT6:.*]]438// CHECK1:       [[ELSE5]]:439// CHECK1-NEXT:    br label %[[IFCONT6]]440// CHECK1:       [[IFCONT6]]:441// CHECK1-NEXT:    ret void442//443//444// CHECK1-LABEL: define internal void @_omp_reduction_inter_warp_copy_func(445// CHECK1-SAME: ptr noundef [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR1]] {446// CHECK1-NEXT:  [[ENTRY:.*:]]447// CHECK1-NEXT:    [[DOTADDR:%.*]] = alloca ptr, align 8448// CHECK1-NEXT:    [[DOTADDR1:%.*]] = alloca i32, align 4449// CHECK1-NEXT:    [[DOTCNT_ADDR:%.*]] = alloca i32, align 4450// CHECK1-NEXT:    store ptr [[TMP0]], ptr [[DOTADDR]], align 8451// CHECK1-NEXT:    store i32 [[TMP1]], ptr [[DOTADDR1]], align 4452// CHECK1-NEXT:    [[TMP2:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()453// CHECK1-NEXT:    [[TMP3:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()454// CHECK1-NEXT:    [[NVPTX_LANE_ID:%.*]] = and i32 [[TMP3]], 31455// CHECK1-NEXT:    [[TMP4:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()456// CHECK1-NEXT:    [[NVPTX_WARP_ID:%.*]] = ashr i32 [[TMP4]], 5457// CHECK1-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[DOTADDR]], align 8458// CHECK1-NEXT:    store i32 0, ptr [[DOTCNT_ADDR]], align 4459// CHECK1-NEXT:    br label %[[PRECOND:.*]]460// CHECK1:       [[PRECOND]]:461// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCNT_ADDR]], align 4462// CHECK1-NEXT:    [[TMP7:%.*]] = icmp ult i32 [[TMP6]], 2463// CHECK1-NEXT:    br i1 [[TMP7]], label %[[BODY:.*]], label %[[EXIT:.*]]464// CHECK1:       [[BODY]]:465// CHECK1-NEXT:    [[OMP_GLOBAL_THREAD_NUM:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])466// CHECK1-NEXT:    call void @__kmpc_barrier(ptr @[[GLOB4:[0-9]+]], i32 [[OMP_GLOBAL_THREAD_NUM]])467// CHECK1-NEXT:    [[WARP_MASTER:%.*]] = icmp eq i32 [[NVPTX_LANE_ID]], 0468// CHECK1-NEXT:    br i1 [[WARP_MASTER]], label %[[THEN:.*]], label %[[ELSE:.*]]469// CHECK1:       [[THEN]]:470// CHECK1-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP5]], i64 0, i64 0471// CHECK1-NEXT:    [[TMP9:%.*]] = load ptr, ptr [[TMP8]], align 8472// CHECK1-NEXT:    [[TMP10:%.*]] = getelementptr i32, ptr [[TMP9]], i32 [[TMP6]]473// CHECK1-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [32 x i32], ptr addrspace(3) @__openmp_nvptx_data_transfer_temporary_storage, i64 0, i32 [[NVPTX_WARP_ID]]474// CHECK1-NEXT:    [[TMP12:%.*]] = load i32, ptr [[TMP10]], align 4475// CHECK1-NEXT:    store volatile i32 [[TMP12]], ptr addrspace(3) [[TMP11]], align 4476// CHECK1-NEXT:    br label %[[IFCONT:.*]]477// CHECK1:       [[ELSE]]:478// CHECK1-NEXT:    br label %[[IFCONT]]479// CHECK1:       [[IFCONT]]:480// CHECK1-NEXT:    [[OMP_GLOBAL_THREAD_NUM2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])481// CHECK1-NEXT:    call void @__kmpc_barrier(ptr @[[GLOB4]], i32 [[OMP_GLOBAL_THREAD_NUM2]])482// CHECK1-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTADDR1]], align 4483// CHECK1-NEXT:    [[IS_ACTIVE_THREAD:%.*]] = icmp ult i32 [[TMP2]], [[TMP13]]484// CHECK1-NEXT:    br i1 [[IS_ACTIVE_THREAD]], label %[[THEN3:.*]], label %[[ELSE4:.*]]485// CHECK1:       [[THEN3]]:486// CHECK1-NEXT:    [[TMP14:%.*]] = getelementptr inbounds [32 x i32], ptr addrspace(3) @__openmp_nvptx_data_transfer_temporary_storage, i64 0, i32 [[TMP2]]487// CHECK1-NEXT:    [[TMP15:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP5]], i64 0, i64 0488// CHECK1-NEXT:    [[TMP16:%.*]] = load ptr, ptr [[TMP15]], align 8489// CHECK1-NEXT:    [[TMP17:%.*]] = getelementptr i32, ptr [[TMP16]], i32 [[TMP6]]490// CHECK1-NEXT:    [[TMP18:%.*]] = load volatile i32, ptr addrspace(3) [[TMP14]], align 4491// CHECK1-NEXT:    store i32 [[TMP18]], ptr [[TMP17]], align 4492// CHECK1-NEXT:    br label %[[IFCONT5:.*]]493// CHECK1:       [[ELSE4]]:494// CHECK1-NEXT:    br label %[[IFCONT5]]495// CHECK1:       [[IFCONT5]]:496// CHECK1-NEXT:    [[TMP19:%.*]] = add nsw i32 [[TMP6]], 1497// CHECK1-NEXT:    store i32 [[TMP19]], ptr [[DOTCNT_ADDR]], align 4498// CHECK1-NEXT:    br label %[[PRECOND]]499// CHECK1:       [[EXIT]]:500// CHECK1-NEXT:    ret void501//502//503// CHECK1-LABEL: define internal void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z17complex_reductionIfEvv_l16_omp_outlined_omp_outlined_wrapper(504// CHECK1-SAME: i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR8:[0-9]+]] {505// CHECK1-NEXT:  [[ENTRY:.*:]]506// CHECK1-NEXT:    [[DOTADDR:%.*]] = alloca i16, align 2507// CHECK1-NEXT:    [[DOTADDR1:%.*]] = alloca i32, align 4508// CHECK1-NEXT:    [[DOTZERO_ADDR:%.*]] = alloca i32, align 4509// CHECK1-NEXT:    [[GLOBAL_ARGS:%.*]] = alloca ptr, align 8510// CHECK1-NEXT:    store i16 [[TMP0]], ptr [[DOTADDR]], align 2, !tbaa [[SHORT_TBAA27:![0-9]+]]511// CHECK1-NEXT:    store i32 [[TMP1]], ptr [[DOTADDR1]], align 4, !tbaa [[INT_TBAA6]]512// CHECK1-NEXT:    store i32 0, ptr [[DOTZERO_ADDR]], align 4513// CHECK1-NEXT:    call void @__kmpc_get_shared_variables(ptr [[GLOBAL_ARGS]])514// CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[GLOBAL_ARGS]], align 8515// CHECK1-NEXT:    [[TMP3:%.*]] = getelementptr inbounds ptr, ptr [[TMP2]], i64 0516// CHECK1-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[TMP3]], align 8, !tbaa [[ANYPTR_TBAA10]]517// CHECK1-NEXT:    [[TMP5:%.*]] = getelementptr inbounds ptr, ptr [[TMP2]], i64 1518// CHECK1-NEXT:    [[TMP6:%.*]] = load ptr, ptr [[TMP5]], align 8, !tbaa [[ANYPTR_TBAA10]]519// CHECK1-NEXT:    [[TMP7:%.*]] = getelementptr inbounds ptr, ptr [[TMP2]], i64 2520// CHECK1-NEXT:    [[TMP8:%.*]] = load ptr, ptr [[TMP7]], align 8, !tbaa [[ANYPTR_TBAA10]]521// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z17complex_reductionIfEvv_l16_omp_outlined_omp_outlined(ptr [[DOTADDR1]], ptr [[DOTZERO_ADDR]], ptr [[TMP4]], ptr [[TMP6]], ptr [[TMP8]]) #[[ATTR4]]522// CHECK1-NEXT:    ret void523//524//525// CHECK1-LABEL: define weak_odr protected ptx_kernel void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z17complex_reductionIdEvv_l16(526// CHECK1-SAME: ptr noalias [[DYN_PTR:%.*]]) #[[ATTR0]] {527// CHECK1-NEXT:  [[ENTRY:.*:]]528// CHECK1-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8529// CHECK1-NEXT:    [[DOTZERO_ADDR:%.*]] = alloca i32, align 4530// CHECK1-NEXT:    [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4531// CHECK1-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8, !tbaa [[ANYPTR_TBAA10]]532// CHECK1-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z17complex_reductionIdEvv_l16_kernel_environment, ptr [[DYN_PTR]])533// CHECK1-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1534// CHECK1-NEXT:    br i1 [[EXEC_USER_CODE]], label %[[USER_CODE_ENTRY:.*]], label %[[WORKER_EXIT:.*]]535// CHECK1:       [[USER_CODE_ENTRY]]:536// CHECK1-NEXT:    [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])537// CHECK1-NEXT:    store i32 0, ptr [[DOTZERO_ADDR]], align 4538// CHECK1-NEXT:    store i32 [[TMP1]], ptr [[DOTTHREADID_TEMP_]], align 4, !tbaa [[INT_TBAA6]]539// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z17complex_reductionIdEvv_l16_omp_outlined(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTZERO_ADDR]]) #[[ATTR4]]540// CHECK1-NEXT:    call void @__kmpc_target_deinit()541// CHECK1-NEXT:    ret void542// CHECK1:       [[WORKER_EXIT]]:543// CHECK1-NEXT:    ret void544//545//546// CHECK1-LABEL: define internal void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z17complex_reductionIdEvv_l16_omp_outlined(547// CHECK1-SAME: ptr noalias [[DOTGLOBAL_TID_:%.*]], ptr noalias [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {548// CHECK1-NEXT:  [[ENTRY:.*:]]549// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8550// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8551// CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4552// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4553// CHECK1-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4554// CHECK1-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4555// CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4556// CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4557// CHECK1-NEXT:    [[IB:%.*]] = alloca i32, align 4558// CHECK1-NEXT:    [[REF_TMP:%.*]] = alloca double, align 8559// CHECK1-NEXT:    [[REF_TMP2:%.*]] = alloca double, align 8560// CHECK1-NEXT:    [[CAPTURED_VARS_ADDRS:%.*]] = alloca [3 x ptr], align 8561// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8, !tbaa [[INTPTR_TBAA12]]562// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8, !tbaa [[INTPTR_TBAA12]]563// CHECK1-NEXT:    [[ISTART:%.*]] = call align 16 ptr @__kmpc_alloc_shared(i64 4)564// CHECK1-NEXT:    [[IEND:%.*]] = call align 16 ptr @__kmpc_alloc_shared(i64 4)565// CHECK1-NEXT:    [[PARTIAL_SUM:%.*]] = call align 16 ptr @__kmpc_alloc_shared(i64 16)566// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[DOTOMP_IV]]) #[[ATTR4]]567// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[DOTOMP_LB]]) #[[ATTR4]]568// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4, !tbaa [[INT_TBAA6]]569// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[DOTOMP_UB]]) #[[ATTR4]]570// CHECK1-NEXT:    store i32 99, ptr [[DOTOMP_UB]], align 4, !tbaa [[INT_TBAA6]]571// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[DOTOMP_STRIDE]]) #[[ATTR4]]572// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4, !tbaa [[INT_TBAA6]]573// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[DOTOMP_IS_LAST]]) #[[ATTR4]]574// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4, !tbaa [[INT_TBAA6]]575// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[IB]]) #[[ATTR4]]576// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8577// CHECK1-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4, !tbaa [[INT_TBAA6]]578// CHECK1-NEXT:    call void @__kmpc_distribute_static_init_4(ptr @[[GLOB2]], i32 [[TMP1]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)579// CHECK1-NEXT:    [[TMP2:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !tbaa [[INT_TBAA6]]580// CHECK1-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 99581// CHECK1-NEXT:    br i1 [[CMP]], label %[[COND_TRUE:.*]], label %[[COND_FALSE:.*]]582// CHECK1:       [[COND_TRUE]]:583// CHECK1-NEXT:    br label %[[COND_END:.*]]584// CHECK1:       [[COND_FALSE]]:585// CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !tbaa [[INT_TBAA6]]586// CHECK1-NEXT:    br label %[[COND_END]]587// CHECK1:       [[COND_END]]:588// CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ 99, %[[COND_TRUE]] ], [ [[TMP3]], %[[COND_FALSE]] ]589// CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4, !tbaa [[INT_TBAA6]]590// CHECK1-NEXT:    [[TMP4:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4, !tbaa [[INT_TBAA6]]591// CHECK1-NEXT:    store i32 [[TMP4]], ptr [[DOTOMP_IV]], align 4, !tbaa [[INT_TBAA6]]592// CHECK1-NEXT:    br label %[[OMP_INNER_FOR_COND:.*]]593// CHECK1:       [[OMP_INNER_FOR_COND]]:594// CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !tbaa [[INT_TBAA6]]595// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !tbaa [[INT_TBAA6]]596// CHECK1-NEXT:    [[CMP1:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]597// CHECK1-NEXT:    br i1 [[CMP1]], label %[[OMP_INNER_FOR_BODY:.*]], label %[[OMP_INNER_FOR_COND_CLEANUP:.*]]598// CHECK1:       [[OMP_INNER_FOR_COND_CLEANUP]]:599// CHECK1-NEXT:    br label %[[OMP_INNER_FOR_END:.*]]600// CHECK1:       [[OMP_INNER_FOR_BODY]]:601// CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !tbaa [[INT_TBAA6]]602// CHECK1-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP7]], 1603// CHECK1-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]604// CHECK1-NEXT:    store i32 [[ADD]], ptr [[IB]], align 4, !tbaa [[INT_TBAA6]]605// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[REF_TMP]]) #[[ATTR4]]606// CHECK1-NEXT:    store double 0.000000e+00, ptr [[REF_TMP]], align 8, !tbaa [[DOUBLE_TBAA29:![0-9]+]]607// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[REF_TMP2]]) #[[ATTR4]]608// CHECK1-NEXT:    store double 0.000000e+00, ptr [[REF_TMP2]], align 8, !tbaa [[DOUBLE_TBAA29]]609// CHECK1-NEXT:    call void @_ZNSt7complexIdEC1ERKdS2_(ptr nonnull align 8 dereferenceable(16) [[PARTIAL_SUM]], ptr nonnull align 8 dereferenceable(8) [[REF_TMP]], ptr nonnull align 8 dereferenceable(8) [[REF_TMP2]]) #[[ATTR11]]610// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[REF_TMP2]]) #[[ATTR4]]611// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[REF_TMP]]) #[[ATTR4]]612// CHECK1-NEXT:    [[TMP8:%.*]] = load i32, ptr [[IB]], align 4, !tbaa [[INT_TBAA6]]613// CHECK1-NEXT:    [[MUL3:%.*]] = mul nsw i32 [[TMP8]], 4614// CHECK1-NEXT:    store i32 [[MUL3]], ptr [[ISTART]], align 4, !tbaa [[INT_TBAA6]]615// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[IB]], align 4, !tbaa [[INT_TBAA6]]616// CHECK1-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP9]], 1617// CHECK1-NEXT:    [[MUL5:%.*]] = mul nsw i32 [[ADD4]], 4618// CHECK1-NEXT:    store i32 [[MUL5]], ptr [[IEND]], align 4, !tbaa [[INT_TBAA6]]619// CHECK1-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [3 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 0620// CHECK1-NEXT:    store ptr [[ISTART]], ptr [[TMP10]], align 8, !tbaa [[ANYPTR_TBAA16]]621// CHECK1-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [3 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 1622// CHECK1-NEXT:    store ptr [[IEND]], ptr [[TMP11]], align 8, !tbaa [[ANYPTR_TBAA16]]623// CHECK1-NEXT:    [[TMP12:%.*]] = getelementptr inbounds [3 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 2624// CHECK1-NEXT:    store ptr [[PARTIAL_SUM]], ptr [[TMP12]], align 8, !tbaa [[ANYPTR_TBAA16]]625// CHECK1-NEXT:    call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z17complex_reductionIdEvv_l16_omp_outlined_omp_outlined, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z17complex_reductionIdEvv_l16_omp_outlined_omp_outlined_wrapper, ptr [[CAPTURED_VARS_ADDRS]], i64 3)626// CHECK1-NEXT:    br label %[[OMP_BODY_CONTINUE:.*]]627// CHECK1:       [[OMP_BODY_CONTINUE]]:628// CHECK1-NEXT:    br label %[[OMP_INNER_FOR_INC:.*]]629// CHECK1:       [[OMP_INNER_FOR_INC]]:630// CHECK1-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !tbaa [[INT_TBAA6]]631// CHECK1-NEXT:    [[ADD6:%.*]] = add nsw i32 [[TMP13]], 1632// CHECK1-NEXT:    store i32 [[ADD6]], ptr [[DOTOMP_IV]], align 4, !tbaa [[INT_TBAA6]]633// CHECK1-NEXT:    br label %[[OMP_INNER_FOR_COND]]634// CHECK1:       [[OMP_INNER_FOR_END]]:635// CHECK1-NEXT:    br label %[[OMP_LOOP_EXIT:.*]]636// CHECK1:       [[OMP_LOOP_EXIT]]:637// CHECK1-NEXT:    call void @__kmpc_distribute_static_fini(ptr @[[GLOB2]], i32 [[TMP1]])638// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[IB]]) #[[ATTR4]]639// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[DOTOMP_IS_LAST]]) #[[ATTR4]]640// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[DOTOMP_STRIDE]]) #[[ATTR4]]641// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[DOTOMP_UB]]) #[[ATTR4]]642// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[DOTOMP_LB]]) #[[ATTR4]]643// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[DOTOMP_IV]]) #[[ATTR4]]644// CHECK1-NEXT:    call void @__kmpc_free_shared(ptr [[PARTIAL_SUM]], i64 16)645// CHECK1-NEXT:    call void @__kmpc_free_shared(ptr [[IEND]], i64 4)646// CHECK1-NEXT:    call void @__kmpc_free_shared(ptr [[ISTART]], i64 4)647// CHECK1-NEXT:    ret void648//649//650// CHECK1-LABEL: define linkonce_odr hidden void @_ZNSt7complexIdEC1ERKdS2_(651// CHECK1-SAME: ptr nonnull align 8 dereferenceable(16) [[THIS:%.*]], ptr nonnull align 8 dereferenceable(8) [[__RE:%.*]], ptr nonnull align 8 dereferenceable(8) [[__IM:%.*]]) unnamed_addr #[[ATTR5]] comdat align 2 {652// CHECK1-NEXT:  [[ENTRY:.*:]]653// CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8654// CHECK1-NEXT:    [[__RE_ADDR:%.*]] = alloca ptr, align 8655// CHECK1-NEXT:    [[__IM_ADDR:%.*]] = alloca ptr, align 8656// CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8, !tbaa [[_ZTSST7COMPLEXIDEPTR_TBAA31:![0-9]+]]657// CHECK1-NEXT:    store ptr [[__RE]], ptr [[__RE_ADDR]], align 8, !tbaa [[DOUBLEPTR_TBAA33:![0-9]+]]658// CHECK1-NEXT:    store ptr [[__IM]], ptr [[__IM_ADDR]], align 8, !tbaa [[DOUBLEPTR_TBAA33]]659// CHECK1-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8660// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[__RE_ADDR]], align 8661// CHECK1-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[__IM_ADDR]], align 8662// CHECK1-NEXT:    call void @_ZNSt7complexIdEC2ERKdS2_(ptr nonnull align 8 dereferenceable(16) [[THIS1]], ptr nonnull align 8 dereferenceable(8) [[TMP0]], ptr nonnull align 8 dereferenceable(8) [[TMP1]]) #[[ATTR11]]663// CHECK1-NEXT:    ret void664//665//666// CHECK1-LABEL: define internal void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z17complex_reductionIdEvv_l16_omp_outlined_omp_outlined(667// CHECK1-SAME: ptr noalias [[DOTGLOBAL_TID_:%.*]], ptr noalias [[DOTBOUND_TID_:%.*]], ptr nonnull align 4 dereferenceable(4) [[ISTART:%.*]], ptr nonnull align 4 dereferenceable(4) [[IEND:%.*]], ptr nonnull align 8 dereferenceable(16) [[PARTIAL_SUM:%.*]]) #[[ATTR1]] {668// CHECK1-NEXT:  [[ENTRY:.*:]]669// CHECK1-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8670// CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8671// CHECK1-NEXT:    [[ISTART_ADDR:%.*]] = alloca ptr, align 8672// CHECK1-NEXT:    [[IEND_ADDR:%.*]] = alloca ptr, align 8673// CHECK1-NEXT:    [[PARTIAL_SUM_ADDR:%.*]] = alloca ptr, align 8674// CHECK1-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4675// CHECK1-NEXT:    [[TMP:%.*]] = alloca i32, align 4676// CHECK1-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4677// CHECK1-NEXT:    [[DOTCAPTURE_EXPR_1:%.*]] = alloca i32, align 4678// CHECK1-NEXT:    [[DOTCAPTURE_EXPR_2:%.*]] = alloca i32, align 4679// CHECK1-NEXT:    [[I:%.*]] = alloca i32, align 4680// CHECK1-NEXT:    [[DOTOMP_LB:%.*]] = alloca i32, align 4681// CHECK1-NEXT:    [[DOTOMP_UB:%.*]] = alloca i32, align 4682// CHECK1-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4683// CHECK1-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4684// CHECK1-NEXT:    [[PARTIAL_SUM5:%.*]] = alloca %"class.std::complex.0", align 8685// CHECK1-NEXT:    [[REF_TMP:%.*]] = alloca double, align 8686// CHECK1-NEXT:    [[REF_TMP6:%.*]] = alloca double, align 8687// CHECK1-NEXT:    [[I7:%.*]] = alloca i32, align 4688// CHECK1-NEXT:    [[REF_TMP14:%.*]] = alloca %"class.std::complex.0", align 8689// CHECK1-NEXT:    [[REF_TMP15:%.*]] = alloca double, align 8690// CHECK1-NEXT:    [[REF_TMP16:%.*]] = alloca double, align 8691// CHECK1-NEXT:    [[DOTOMP_REDUCTION_RED_LIST:%.*]] = alloca [1 x ptr], align 8692// CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8, !tbaa [[INTPTR_TBAA12]]693// CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8, !tbaa [[INTPTR_TBAA12]]694// CHECK1-NEXT:    store ptr [[ISTART]], ptr [[ISTART_ADDR]], align 8, !tbaa [[INTPTR_TBAA12]]695// CHECK1-NEXT:    store ptr [[IEND]], ptr [[IEND_ADDR]], align 8, !tbaa [[INTPTR_TBAA12]]696// CHECK1-NEXT:    store ptr [[PARTIAL_SUM]], ptr [[PARTIAL_SUM_ADDR]], align 8, !tbaa [[_ZTSST7COMPLEXIDEPTR_TBAA31]]697// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[ISTART_ADDR]], align 8, !tbaa [[INTPTR_TBAA12]], !nonnull [[META22]], !align [[META23]]698// CHECK1-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[IEND_ADDR]], align 8, !tbaa [[INTPTR_TBAA12]], !nonnull [[META22]], !align [[META23]]699// CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[PARTIAL_SUM_ADDR]], align 8, !tbaa [[_ZTSST7COMPLEXIDEPTR_TBAA31]], !nonnull [[META22]], !align [[META35:![0-9]+]]700// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[DOTOMP_IV]]) #[[ATTR4]]701// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[DOTCAPTURE_EXPR_]]) #[[ATTR4]]702// CHECK1-NEXT:    [[TMP3:%.*]] = load i32, ptr [[TMP0]], align 4, !tbaa [[INT_TBAA6]]703// CHECK1-NEXT:    store i32 [[TMP3]], ptr [[DOTCAPTURE_EXPR_]], align 4, !tbaa [[INT_TBAA6]]704// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[DOTCAPTURE_EXPR_1]]) #[[ATTR4]]705// CHECK1-NEXT:    [[TMP4:%.*]] = load i32, ptr [[TMP1]], align 4, !tbaa [[INT_TBAA6]]706// CHECK1-NEXT:    store i32 [[TMP4]], ptr [[DOTCAPTURE_EXPR_1]], align 4, !tbaa [[INT_TBAA6]]707// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[DOTCAPTURE_EXPR_2]]) #[[ATTR4]]708// CHECK1-NEXT:    [[TMP5:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4, !tbaa [[INT_TBAA6]]709// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4, !tbaa [[INT_TBAA6]]710// CHECK1-NEXT:    [[SUB:%.*]] = sub i32 [[TMP5]], [[TMP6]]711// CHECK1-NEXT:    [[SUB3:%.*]] = sub i32 [[SUB]], 1712// CHECK1-NEXT:    [[ADD:%.*]] = add i32 [[SUB3]], 1713// CHECK1-NEXT:    [[DIV:%.*]] = udiv i32 [[ADD]], 1714// CHECK1-NEXT:    [[SUB4:%.*]] = sub i32 [[DIV]], 1715// CHECK1-NEXT:    store i32 [[SUB4]], ptr [[DOTCAPTURE_EXPR_2]], align 4, !tbaa [[INT_TBAA6]]716// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[I]]) #[[ATTR4]]717// CHECK1-NEXT:    [[TMP7:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4, !tbaa [[INT_TBAA6]]718// CHECK1-NEXT:    store i32 [[TMP7]], ptr [[I]], align 4, !tbaa [[INT_TBAA6]]719// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[I]]) #[[ATTR4]]720// CHECK1-NEXT:    [[TMP8:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4, !tbaa [[INT_TBAA6]]721// CHECK1-NEXT:    [[TMP9:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_1]], align 4, !tbaa [[INT_TBAA6]]722// CHECK1-NEXT:    [[CMP:%.*]] = icmp slt i32 [[TMP8]], [[TMP9]]723// CHECK1-NEXT:    br i1 [[CMP]], label %[[OMP_PRECOND_THEN:.*]], label %[[OMP_PRECOND_END:.*]]724// CHECK1:       [[OMP_PRECOND_THEN]]:725// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[DOTOMP_LB]]) #[[ATTR4]]726// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_LB]], align 4, !tbaa [[INT_TBAA6]]727// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[DOTOMP_UB]]) #[[ATTR4]]728// CHECK1-NEXT:    [[TMP10:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_2]], align 4, !tbaa [[INT_TBAA6]]729// CHECK1-NEXT:    store i32 [[TMP10]], ptr [[DOTOMP_UB]], align 4, !tbaa [[INT_TBAA6]]730// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[DOTOMP_STRIDE]]) #[[ATTR4]]731// CHECK1-NEXT:    store i32 1, ptr [[DOTOMP_STRIDE]], align 4, !tbaa [[INT_TBAA6]]732// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[DOTOMP_IS_LAST]]) #[[ATTR4]]733// CHECK1-NEXT:    store i32 0, ptr [[DOTOMP_IS_LAST]], align 4, !tbaa [[INT_TBAA6]]734// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[PARTIAL_SUM5]]) #[[ATTR4]]735// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[REF_TMP]]) #[[ATTR4]]736// CHECK1-NEXT:    store double 0.000000e+00, ptr [[REF_TMP]], align 8, !tbaa [[DOUBLE_TBAA29]]737// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[REF_TMP6]]) #[[ATTR4]]738// CHECK1-NEXT:    store double 0.000000e+00, ptr [[REF_TMP6]], align 8, !tbaa [[DOUBLE_TBAA29]]739// CHECK1-NEXT:    call void @_ZNSt7complexIdEC1ERKdS2_(ptr nonnull align 8 dereferenceable(16) [[PARTIAL_SUM5]], ptr nonnull align 8 dereferenceable(8) [[REF_TMP]], ptr nonnull align 8 dereferenceable(8) [[REF_TMP6]]) #[[ATTR11]]740// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[REF_TMP6]]) #[[ATTR4]]741// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[REF_TMP]]) #[[ATTR4]]742// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[I7]]) #[[ATTR4]]743// CHECK1-NEXT:    [[TMP11:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8744// CHECK1-NEXT:    [[TMP12:%.*]] = load i32, ptr [[TMP11]], align 4, !tbaa [[INT_TBAA6]]745// CHECK1-NEXT:    call void @__kmpc_for_static_init_4u(ptr @[[GLOB3]], i32 [[TMP12]], i32 33, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)746// CHECK1-NEXT:    br label %[[OMP_DISPATCH_COND:.*]]747// CHECK1:       [[OMP_DISPATCH_COND]]:748// CHECK1-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !tbaa [[INT_TBAA6]]749// CHECK1-NEXT:    [[TMP14:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_2]], align 4, !tbaa [[INT_TBAA6]]750// CHECK1-NEXT:    [[CMP8:%.*]] = icmp ugt i32 [[TMP13]], [[TMP14]]751// CHECK1-NEXT:    br i1 [[CMP8]], label %[[COND_TRUE:.*]], label %[[COND_FALSE:.*]]752// CHECK1:       [[COND_TRUE]]:753// CHECK1-NEXT:    [[TMP15:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_2]], align 4, !tbaa [[INT_TBAA6]]754// CHECK1-NEXT:    br label %[[COND_END:.*]]755// CHECK1:       [[COND_FALSE]]:756// CHECK1-NEXT:    [[TMP16:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !tbaa [[INT_TBAA6]]757// CHECK1-NEXT:    br label %[[COND_END]]758// CHECK1:       [[COND_END]]:759// CHECK1-NEXT:    [[COND:%.*]] = phi i32 [ [[TMP15]], %[[COND_TRUE]] ], [ [[TMP16]], %[[COND_FALSE]] ]760// CHECK1-NEXT:    store i32 [[COND]], ptr [[DOTOMP_UB]], align 4, !tbaa [[INT_TBAA6]]761// CHECK1-NEXT:    [[TMP17:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4, !tbaa [[INT_TBAA6]]762// CHECK1-NEXT:    store i32 [[TMP17]], ptr [[DOTOMP_IV]], align 4, !tbaa [[INT_TBAA6]]763// CHECK1-NEXT:    [[TMP18:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !tbaa [[INT_TBAA6]]764// CHECK1-NEXT:    [[TMP19:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !tbaa [[INT_TBAA6]]765// CHECK1-NEXT:    [[ADD9:%.*]] = add i32 [[TMP19]], 1766// CHECK1-NEXT:    [[CMP10:%.*]] = icmp ult i32 [[TMP18]], [[ADD9]]767// CHECK1-NEXT:    br i1 [[CMP10]], label %[[OMP_DISPATCH_BODY:.*]], label %[[OMP_DISPATCH_CLEANUP:.*]]768// CHECK1:       [[OMP_DISPATCH_CLEANUP]]:769// CHECK1-NEXT:    br label %[[OMP_DISPATCH_END:.*]]770// CHECK1:       [[OMP_DISPATCH_BODY]]:771// CHECK1-NEXT:    br label %[[OMP_INNER_FOR_COND:.*]]772// CHECK1:       [[OMP_INNER_FOR_COND]]:773// CHECK1-NEXT:    [[TMP20:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !tbaa [[INT_TBAA6]]774// CHECK1-NEXT:    [[TMP21:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !tbaa [[INT_TBAA6]]775// CHECK1-NEXT:    [[ADD11:%.*]] = add i32 [[TMP21]], 1776// CHECK1-NEXT:    [[CMP12:%.*]] = icmp ult i32 [[TMP20]], [[ADD11]]777// CHECK1-NEXT:    br i1 [[CMP12]], label %[[OMP_INNER_FOR_BODY:.*]], label %[[OMP_INNER_FOR_COND_CLEANUP:.*]]778// CHECK1:       [[OMP_INNER_FOR_COND_CLEANUP]]:779// CHECK1-NEXT:    br label %[[OMP_INNER_FOR_END:.*]]780// CHECK1:       [[OMP_INNER_FOR_BODY]]:781// CHECK1-NEXT:    [[TMP22:%.*]] = load i32, ptr [[DOTCAPTURE_EXPR_]], align 4, !tbaa [[INT_TBAA6]]782// CHECK1-NEXT:    [[TMP23:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !tbaa [[INT_TBAA6]]783// CHECK1-NEXT:    [[MUL:%.*]] = mul i32 [[TMP23]], 1784// CHECK1-NEXT:    [[ADD13:%.*]] = add i32 [[TMP22]], [[MUL]]785// CHECK1-NEXT:    store i32 [[ADD13]], ptr [[I7]], align 4, !tbaa [[INT_TBAA6]]786// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[REF_TMP14]]) #[[ATTR4]]787// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[REF_TMP15]]) #[[ATTR4]]788// CHECK1-NEXT:    [[TMP24:%.*]] = load i32, ptr [[I7]], align 4, !tbaa [[INT_TBAA6]]789// CHECK1-NEXT:    [[CONV:%.*]] = sitofp i32 [[TMP24]] to double790// CHECK1-NEXT:    store double [[CONV]], ptr [[REF_TMP15]], align 8, !tbaa [[DOUBLE_TBAA29]]791// CHECK1-NEXT:    call void @llvm.lifetime.start.p0(ptr [[REF_TMP16]]) #[[ATTR4]]792// CHECK1-NEXT:    [[TMP25:%.*]] = load i32, ptr [[I7]], align 4, !tbaa [[INT_TBAA6]]793// CHECK1-NEXT:    [[CONV17:%.*]] = sitofp i32 [[TMP25]] to double794// CHECK1-NEXT:    store double [[CONV17]], ptr [[REF_TMP16]], align 8, !tbaa [[DOUBLE_TBAA29]]795// CHECK1-NEXT:    call void @_ZNSt7complexIdEC1ERKdS2_(ptr nonnull align 8 dereferenceable(16) [[REF_TMP14]], ptr nonnull align 8 dereferenceable(8) [[REF_TMP15]], ptr nonnull align 8 dereferenceable(8) [[REF_TMP16]]) #[[ATTR11]]796// CHECK1-NEXT:    [[CALL:%.*]] = call nonnull align 8 dereferenceable(16) ptr @_ZNSt7complexIdEpLIdEERS0_RKS_IT_E(ptr nonnull align 8 dereferenceable(16) [[PARTIAL_SUM5]], ptr nonnull align 8 dereferenceable(16) [[REF_TMP14]]) #[[ATTR11]]797// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[REF_TMP16]]) #[[ATTR4]]798// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[REF_TMP15]]) #[[ATTR4]]799// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[REF_TMP14]]) #[[ATTR4]]800// CHECK1-NEXT:    br label %[[OMP_BODY_CONTINUE:.*]]801// CHECK1:       [[OMP_BODY_CONTINUE]]:802// CHECK1-NEXT:    br label %[[OMP_INNER_FOR_INC:.*]]803// CHECK1:       [[OMP_INNER_FOR_INC]]:804// CHECK1-NEXT:    [[TMP26:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4, !tbaa [[INT_TBAA6]]805// CHECK1-NEXT:    [[ADD18:%.*]] = add i32 [[TMP26]], 1806// CHECK1-NEXT:    store i32 [[ADD18]], ptr [[DOTOMP_IV]], align 4, !tbaa [[INT_TBAA6]]807// CHECK1-NEXT:    br label %[[OMP_INNER_FOR_COND]]808// CHECK1:       [[OMP_INNER_FOR_END]]:809// CHECK1-NEXT:    br label %[[OMP_DISPATCH_INC:.*]]810// CHECK1:       [[OMP_DISPATCH_INC]]:811// CHECK1-NEXT:    [[TMP27:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4, !tbaa [[INT_TBAA6]]812// CHECK1-NEXT:    [[TMP28:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4, !tbaa [[INT_TBAA6]]813// CHECK1-NEXT:    [[ADD19:%.*]] = add i32 [[TMP27]], [[TMP28]]814// CHECK1-NEXT:    store i32 [[ADD19]], ptr [[DOTOMP_LB]], align 4, !tbaa [[INT_TBAA6]]815// CHECK1-NEXT:    [[TMP29:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4, !tbaa [[INT_TBAA6]]816// CHECK1-NEXT:    [[TMP30:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4, !tbaa [[INT_TBAA6]]817// CHECK1-NEXT:    [[ADD20:%.*]] = add i32 [[TMP29]], [[TMP30]]818// CHECK1-NEXT:    store i32 [[ADD20]], ptr [[DOTOMP_UB]], align 4, !tbaa [[INT_TBAA6]]819// CHECK1-NEXT:    br label %[[OMP_DISPATCH_COND]]820// CHECK1:       [[OMP_DISPATCH_END]]:821// CHECK1-NEXT:    [[TMP31:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8822// CHECK1-NEXT:    [[TMP32:%.*]] = load i32, ptr [[TMP31]], align 4, !tbaa [[INT_TBAA6]]823// CHECK1-NEXT:    call void @__kmpc_for_static_fini(ptr @[[GLOB3]], i32 [[TMP32]])824// CHECK1-NEXT:    [[TMP33:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_RED_LIST]], i64 0, i64 0825// CHECK1-NEXT:    store ptr [[PARTIAL_SUM5]], ptr [[TMP33]], align 8826// CHECK1-NEXT:    [[TMP34:%.*]] = call i32 @__kmpc_nvptx_parallel_reduce_nowait_v2(ptr @[[GLOB1]], i64 16, ptr [[DOTOMP_REDUCTION_RED_LIST]], ptr @_omp_reduction_shuffle_and_reduce_func1, ptr @_omp_reduction_inter_warp_copy_func2)827// CHECK1-NEXT:    [[TMP35:%.*]] = icmp eq i32 [[TMP34]], 1828// CHECK1-NEXT:    br i1 [[TMP35]], [[DOTOMP_REDUCTION_THEN:label %.*]], [[DOTOMP_REDUCTION_DONE:label %.*]]829// CHECK1:       [[_OMP_REDUCTION_THEN:.*:]]830// CHECK1-NEXT:    [[CALL21:%.*]] = call nonnull align 8 dereferenceable(16) ptr @_ZNSt7complexIdEpLIdEERS0_RKS_IT_E(ptr nonnull align 8 dereferenceable(16) [[TMP2]], ptr nonnull align 8 dereferenceable(16) [[PARTIAL_SUM5]]) #[[ATTR11]]831// CHECK1-NEXT:    br [[DOTOMP_REDUCTION_DONE]]832// CHECK1:       [[_OMP_REDUCTION_DONE:.*:]]833// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[I7]]) #[[ATTR4]]834// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[PARTIAL_SUM5]]) #[[ATTR4]]835// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[DOTOMP_IS_LAST]]) #[[ATTR4]]836// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[DOTOMP_STRIDE]]) #[[ATTR4]]837// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[DOTOMP_UB]]) #[[ATTR4]]838// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[DOTOMP_LB]]) #[[ATTR4]]839// CHECK1-NEXT:    br label %[[OMP_PRECOND_END]]840// CHECK1:       [[OMP_PRECOND_END]]:841// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[DOTCAPTURE_EXPR_2]]) #[[ATTR4]]842// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[DOTCAPTURE_EXPR_1]]) #[[ATTR4]]843// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[DOTCAPTURE_EXPR_]]) #[[ATTR4]]844// CHECK1-NEXT:    call void @llvm.lifetime.end.p0(ptr [[DOTOMP_IV]]) #[[ATTR4]]845// CHECK1-NEXT:    ret void846//847//848// CHECK1-LABEL: define linkonce_odr hidden nonnull align 8 dereferenceable(16) ptr @_ZNSt7complexIdEpLIdEERS0_RKS_IT_E(849// CHECK1-SAME: ptr nonnull align 8 dereferenceable(16) [[THIS:%.*]], ptr nonnull align 8 dereferenceable(16) [[__C:%.*]]) #[[ATTR5]] comdat align 2 {850// CHECK1-NEXT:  [[ENTRY:.*:]]851// CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 8852// CHECK1-NEXT:    [[__C_ADDR:%.*]] = alloca ptr, align 8853// CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8, !tbaa [[_ZTSST7COMPLEXIDEPTR_TBAA31]]854// CHECK1-NEXT:    store ptr [[__C]], ptr [[__C_ADDR]], align 8, !tbaa [[_ZTSST7COMPLEXIDEPTR_TBAA31]]855// CHECK1-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8856// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[__C_ADDR]], align 8, !tbaa [[_ZTSST7COMPLEXIDEPTR_TBAA31]], !nonnull [[META22]], !align [[META35]]857// CHECK1-NEXT:    [[CALL:%.*]] = call double @_ZNKSt7complexIdE4realEv(ptr nonnull align 8 dereferenceable(16) [[TMP0]]) #[[ATTR11]]858// CHECK1-NEXT:    [[__RE_:%.*]] = getelementptr inbounds nuw %"class.std::complex.0", ptr [[THIS1]], i32 0, i32 0859// CHECK1-NEXT:    [[TMP1:%.*]] = load double, ptr [[__RE_]], align 8, !tbaa [[DOUBLE_TBAA36:![0-9]+]]860// CHECK1-NEXT:    [[ADD:%.*]] = fadd double [[TMP1]], [[CALL]]861// CHECK1-NEXT:    store double [[ADD]], ptr [[__RE_]], align 8, !tbaa [[DOUBLE_TBAA36]]862// CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[__C_ADDR]], align 8, !tbaa [[_ZTSST7COMPLEXIDEPTR_TBAA31]], !nonnull [[META22]], !align [[META35]]863// CHECK1-NEXT:    [[CALL2:%.*]] = call double @_ZNKSt7complexIdE4imagEv(ptr nonnull align 8 dereferenceable(16) [[TMP2]]) #[[ATTR11]]864// CHECK1-NEXT:    [[__IM_:%.*]] = getelementptr inbounds nuw %"class.std::complex.0", ptr [[THIS1]], i32 0, i32 1865// CHECK1-NEXT:    [[TMP3:%.*]] = load double, ptr [[__IM_]], align 8, !tbaa [[DOUBLE_TBAA38:![0-9]+]]866// CHECK1-NEXT:    [[ADD3:%.*]] = fadd double [[TMP3]], [[CALL2]]867// CHECK1-NEXT:    store double [[ADD3]], ptr [[__IM_]], align 8, !tbaa [[DOUBLE_TBAA38]]868// CHECK1-NEXT:    ret ptr [[THIS1]]869//870//871// CHECK1-LABEL: define internal void @_omp_reduction_shuffle_and_reduce_func1(872// CHECK1-SAME: ptr noundef [[TMP0:%.*]], i16 noundef signext [[TMP1:%.*]], i16 noundef signext [[TMP2:%.*]], i16 noundef signext [[TMP3:%.*]]) #[[ATTR1]] {873// CHECK1-NEXT:  [[ENTRY:.*]]:874// CHECK1-NEXT:    [[DOTADDR:%.*]] = alloca ptr, align 8875// CHECK1-NEXT:    [[DOTADDR1:%.*]] = alloca i16, align 2876// CHECK1-NEXT:    [[DOTADDR2:%.*]] = alloca i16, align 2877// CHECK1-NEXT:    [[DOTADDR3:%.*]] = alloca i16, align 2878// CHECK1-NEXT:    [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST:%.*]] = alloca [1 x ptr], align 8879// CHECK1-NEXT:    [[DOTOMP_REDUCTION_ELEMENT:%.*]] = alloca %"class.std::complex.0", align 8880// CHECK1-NEXT:    store ptr [[TMP0]], ptr [[DOTADDR]], align 8881// CHECK1-NEXT:    store i16 [[TMP1]], ptr [[DOTADDR1]], align 2882// CHECK1-NEXT:    store i16 [[TMP2]], ptr [[DOTADDR2]], align 2883// CHECK1-NEXT:    store i16 [[TMP3]], ptr [[DOTADDR3]], align 2884// CHECK1-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[DOTADDR]], align 8885// CHECK1-NEXT:    [[TMP5:%.*]] = load i16, ptr [[DOTADDR1]], align 2886// CHECK1-NEXT:    [[TMP6:%.*]] = load i16, ptr [[DOTADDR2]], align 2887// CHECK1-NEXT:    [[TMP7:%.*]] = load i16, ptr [[DOTADDR3]], align 2888// CHECK1-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP4]], i64 0, i64 0889// CHECK1-NEXT:    [[TMP9:%.*]] = load ptr, ptr [[TMP8]], align 8890// CHECK1-NEXT:    [[TMP10:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]], i64 0, i64 0891// CHECK1-NEXT:    [[TMP11:%.*]] = getelementptr %"class.std::complex.0", ptr [[TMP9]], i64 1892// CHECK1-NEXT:    br [[DOTSHUFFLE_PRE_COND:label %.*]]893// CHECK1:       [[_SHUFFLE_PRE_COND:.*:]]894// CHECK1-NEXT:    [[TMP12:%.*]] = phi ptr [ [[TMP9]], %[[ENTRY]] ], [ [[TMP23:%.*]], %[[DOTSHUFFLE_THEN:.*]] ]895// CHECK1-NEXT:    [[TMP13:%.*]] = phi ptr [ [[DOTOMP_REDUCTION_ELEMENT]], %[[ENTRY]] ], [ [[TMP24:%.*]], %[[DOTSHUFFLE_THEN]] ]896// CHECK1-NEXT:    [[TMP14:%.*]] = ptrtoint ptr [[TMP11]] to i64897// CHECK1-NEXT:    [[TMP15:%.*]] = ptrtoint ptr [[TMP12]] to i64898// CHECK1-NEXT:    [[TMP16:%.*]] = sub i64 [[TMP14]], [[TMP15]]899// CHECK1-NEXT:    [[TMP17:%.*]] = sdiv exact i64 [[TMP16]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64)900// CHECK1-NEXT:    [[TMP18:%.*]] = icmp sgt i64 [[TMP17]], 7901// CHECK1-NEXT:    br i1 [[TMP18]], label %[[DOTSHUFFLE_THEN]], [[DOTSHUFFLE_EXIT:label %.*]]902// CHECK1:       [[_SHUFFLE_THEN:.*:]]903// CHECK1-NEXT:    [[TMP19:%.*]] = load i64, ptr [[TMP12]], align 8904// CHECK1-NEXT:    [[TMP20:%.*]] = call i32 @__kmpc_get_warp_size()905// CHECK1-NEXT:    [[TMP21:%.*]] = trunc i32 [[TMP20]] to i16906// CHECK1-NEXT:    [[TMP22:%.*]] = call i64 @__kmpc_shuffle_int64(i64 [[TMP19]], i16 [[TMP6]], i16 [[TMP21]])907// CHECK1-NEXT:    store i64 [[TMP22]], ptr [[TMP13]], align 8908// CHECK1-NEXT:    [[TMP23]] = getelementptr i64, ptr [[TMP12]], i64 1909// CHECK1-NEXT:    [[TMP24]] = getelementptr i64, ptr [[TMP13]], i64 1910// CHECK1-NEXT:    br [[DOTSHUFFLE_PRE_COND]]911// CHECK1:       [[_SHUFFLE_EXIT:.*:]]912// CHECK1-NEXT:    store ptr [[DOTOMP_REDUCTION_ELEMENT]], ptr [[TMP10]], align 8913// CHECK1-NEXT:    [[TMP25:%.*]] = icmp eq i16 [[TMP7]], 0914// CHECK1-NEXT:    [[TMP26:%.*]] = icmp eq i16 [[TMP7]], 1915// CHECK1-NEXT:    [[TMP27:%.*]] = icmp ult i16 [[TMP5]], [[TMP6]]916// CHECK1-NEXT:    [[TMP28:%.*]] = and i1 [[TMP26]], [[TMP27]]917// CHECK1-NEXT:    [[TMP29:%.*]] = icmp eq i16 [[TMP7]], 2918// CHECK1-NEXT:    [[TMP30:%.*]] = and i16 [[TMP5]], 1919// CHECK1-NEXT:    [[TMP31:%.*]] = icmp eq i16 [[TMP30]], 0920// CHECK1-NEXT:    [[TMP32:%.*]] = and i1 [[TMP29]], [[TMP31]]921// CHECK1-NEXT:    [[TMP33:%.*]] = icmp sgt i16 [[TMP6]], 0922// CHECK1-NEXT:    [[TMP34:%.*]] = and i1 [[TMP32]], [[TMP33]]923// CHECK1-NEXT:    [[TMP35:%.*]] = or i1 [[TMP25]], [[TMP28]]924// CHECK1-NEXT:    [[TMP36:%.*]] = or i1 [[TMP35]], [[TMP34]]925// CHECK1-NEXT:    br i1 [[TMP36]], label %[[THEN:.*]], label %[[ELSE:.*]]926// CHECK1:       [[THEN]]:927// CHECK1-NEXT:    call void @"{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z17complex_reductionIdEvv_l16_omp_outlined_omp_outlined_omp$reduction$reduction_func"(ptr [[TMP4]], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]]) #[[ATTR4]]928// CHECK1-NEXT:    br label %[[IFCONT:.*]]929// CHECK1:       [[ELSE]]:930// CHECK1-NEXT:    br label %[[IFCONT]]931// CHECK1:       [[IFCONT]]:932// CHECK1-NEXT:    [[TMP37:%.*]] = icmp eq i16 [[TMP7]], 1933// CHECK1-NEXT:    [[TMP38:%.*]] = icmp uge i16 [[TMP5]], [[TMP6]]934// CHECK1-NEXT:    [[TMP39:%.*]] = and i1 [[TMP37]], [[TMP38]]935// CHECK1-NEXT:    br i1 [[TMP39]], label %[[THEN4:.*]], label %[[ELSE5:.*]]936// CHECK1:       [[THEN4]]:937// CHECK1-NEXT:    [[TMP40:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOMP_REDUCTION_REMOTE_REDUCE_LIST]], i64 0, i64 0938// CHECK1-NEXT:    [[TMP41:%.*]] = load ptr, ptr [[TMP40]], align 8939// CHECK1-NEXT:    [[TMP42:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP4]], i64 0, i64 0940// CHECK1-NEXT:    [[TMP43:%.*]] = load ptr, ptr [[TMP42]], align 8941// CHECK1-NEXT:    call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[TMP43]], ptr align 8 [[TMP41]], i64 16, i1 false)942// CHECK1-NEXT:    br label %[[IFCONT6:.*]]943// CHECK1:       [[ELSE5]]:944// CHECK1-NEXT:    br label %[[IFCONT6]]945// CHECK1:       [[IFCONT6]]:946// CHECK1-NEXT:    ret void947//948//949// CHECK1-LABEL: define internal void @_omp_reduction_inter_warp_copy_func2(950// CHECK1-SAME: ptr noundef [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR1]] {951// CHECK1-NEXT:  [[ENTRY:.*:]]952// CHECK1-NEXT:    [[DOTADDR:%.*]] = alloca ptr, align 8953// CHECK1-NEXT:    [[DOTADDR1:%.*]] = alloca i32, align 4954// CHECK1-NEXT:    [[DOTCNT_ADDR:%.*]] = alloca i32, align 4955// CHECK1-NEXT:    store ptr [[TMP0]], ptr [[DOTADDR]], align 8956// CHECK1-NEXT:    store i32 [[TMP1]], ptr [[DOTADDR1]], align 4957// CHECK1-NEXT:    [[TMP2:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()958// CHECK1-NEXT:    [[TMP3:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()959// CHECK1-NEXT:    [[NVPTX_LANE_ID:%.*]] = and i32 [[TMP3]], 31960// CHECK1-NEXT:    [[TMP4:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()961// CHECK1-NEXT:    [[NVPTX_WARP_ID:%.*]] = ashr i32 [[TMP4]], 5962// CHECK1-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[DOTADDR]], align 8963// CHECK1-NEXT:    store i32 0, ptr [[DOTCNT_ADDR]], align 4964// CHECK1-NEXT:    br label %[[PRECOND:.*]]965// CHECK1:       [[PRECOND]]:966// CHECK1-NEXT:    [[TMP6:%.*]] = load i32, ptr [[DOTCNT_ADDR]], align 4967// CHECK1-NEXT:    [[TMP7:%.*]] = icmp ult i32 [[TMP6]], 4968// CHECK1-NEXT:    br i1 [[TMP7]], label %[[BODY:.*]], label %[[EXIT:.*]]969// CHECK1:       [[BODY]]:970// CHECK1-NEXT:    [[OMP_GLOBAL_THREAD_NUM:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])971// CHECK1-NEXT:    call void @__kmpc_barrier(ptr @[[GLOB4]], i32 [[OMP_GLOBAL_THREAD_NUM]])972// CHECK1-NEXT:    [[WARP_MASTER:%.*]] = icmp eq i32 [[NVPTX_LANE_ID]], 0973// CHECK1-NEXT:    br i1 [[WARP_MASTER]], label %[[THEN:.*]], label %[[ELSE:.*]]974// CHECK1:       [[THEN]]:975// CHECK1-NEXT:    [[TMP8:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP5]], i64 0, i64 0976// CHECK1-NEXT:    [[TMP9:%.*]] = load ptr, ptr [[TMP8]], align 8977// CHECK1-NEXT:    [[TMP10:%.*]] = getelementptr i32, ptr [[TMP9]], i32 [[TMP6]]978// CHECK1-NEXT:    [[TMP11:%.*]] = getelementptr inbounds [32 x i32], ptr addrspace(3) @__openmp_nvptx_data_transfer_temporary_storage, i64 0, i32 [[NVPTX_WARP_ID]]979// CHECK1-NEXT:    [[TMP12:%.*]] = load i32, ptr [[TMP10]], align 4980// CHECK1-NEXT:    store volatile i32 [[TMP12]], ptr addrspace(3) [[TMP11]], align 4981// CHECK1-NEXT:    br label %[[IFCONT:.*]]982// CHECK1:       [[ELSE]]:983// CHECK1-NEXT:    br label %[[IFCONT]]984// CHECK1:       [[IFCONT]]:985// CHECK1-NEXT:    [[OMP_GLOBAL_THREAD_NUM2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])986// CHECK1-NEXT:    call void @__kmpc_barrier(ptr @[[GLOB4]], i32 [[OMP_GLOBAL_THREAD_NUM2]])987// CHECK1-NEXT:    [[TMP13:%.*]] = load i32, ptr [[DOTADDR1]], align 4988// CHECK1-NEXT:    [[IS_ACTIVE_THREAD:%.*]] = icmp ult i32 [[TMP2]], [[TMP13]]989// CHECK1-NEXT:    br i1 [[IS_ACTIVE_THREAD]], label %[[THEN3:.*]], label %[[ELSE4:.*]]990// CHECK1:       [[THEN3]]:991// CHECK1-NEXT:    [[TMP14:%.*]] = getelementptr inbounds [32 x i32], ptr addrspace(3) @__openmp_nvptx_data_transfer_temporary_storage, i64 0, i32 [[TMP2]]992// CHECK1-NEXT:    [[TMP15:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP5]], i64 0, i64 0993// CHECK1-NEXT:    [[TMP16:%.*]] = load ptr, ptr [[TMP15]], align 8994// CHECK1-NEXT:    [[TMP17:%.*]] = getelementptr i32, ptr [[TMP16]], i32 [[TMP6]]995// CHECK1-NEXT:    [[TMP18:%.*]] = load volatile i32, ptr addrspace(3) [[TMP14]], align 4996// CHECK1-NEXT:    store i32 [[TMP18]], ptr [[TMP17]], align 4997// CHECK1-NEXT:    br label %[[IFCONT5:.*]]998// CHECK1:       [[ELSE4]]:999// CHECK1-NEXT:    br label %[[IFCONT5]]1000// CHECK1:       [[IFCONT5]]:1001// CHECK1-NEXT:    [[TMP19:%.*]] = add nsw i32 [[TMP6]], 11002// CHECK1-NEXT:    store i32 [[TMP19]], ptr [[DOTCNT_ADDR]], align 41003// CHECK1-NEXT:    br label %[[PRECOND]]1004// CHECK1:       [[EXIT]]:1005// CHECK1-NEXT:    ret void1006//1007//1008// CHECK1-LABEL: define internal void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z17complex_reductionIdEvv_l16_omp_outlined_omp_outlined_wrapper(1009// CHECK1-SAME: i16 zeroext [[TMP0:%.*]], i32 [[TMP1:%.*]]) #[[ATTR8]] {1010// CHECK1-NEXT:  [[ENTRY:.*:]]1011// CHECK1-NEXT:    [[DOTADDR:%.*]] = alloca i16, align 21012// CHECK1-NEXT:    [[DOTADDR1:%.*]] = alloca i32, align 41013// CHECK1-NEXT:    [[DOTZERO_ADDR:%.*]] = alloca i32, align 41014// CHECK1-NEXT:    [[GLOBAL_ARGS:%.*]] = alloca ptr, align 81015// CHECK1-NEXT:    store i16 [[TMP0]], ptr [[DOTADDR]], align 2, !tbaa [[SHORT_TBAA27]]1016// CHECK1-NEXT:    store i32 [[TMP1]], ptr [[DOTADDR1]], align 4, !tbaa [[INT_TBAA6]]1017// CHECK1-NEXT:    store i32 0, ptr [[DOTZERO_ADDR]], align 41018// CHECK1-NEXT:    call void @__kmpc_get_shared_variables(ptr [[GLOBAL_ARGS]])1019// CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[GLOBAL_ARGS]], align 81020// CHECK1-NEXT:    [[TMP3:%.*]] = getelementptr inbounds ptr, ptr [[TMP2]], i64 01021// CHECK1-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[TMP3]], align 8, !tbaa [[ANYPTR_TBAA10]]1022// CHECK1-NEXT:    [[TMP5:%.*]] = getelementptr inbounds ptr, ptr [[TMP2]], i64 11023// CHECK1-NEXT:    [[TMP6:%.*]] = load ptr, ptr [[TMP5]], align 8, !tbaa [[ANYPTR_TBAA10]]1024// CHECK1-NEXT:    [[TMP7:%.*]] = getelementptr inbounds ptr, ptr [[TMP2]], i64 21025// CHECK1-NEXT:    [[TMP8:%.*]] = load ptr, ptr [[TMP7]], align 8, !tbaa [[ANYPTR_TBAA10]]1026// CHECK1-NEXT:    call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z17complex_reductionIdEvv_l16_omp_outlined_omp_outlined(ptr [[DOTADDR1]], ptr [[DOTZERO_ADDR]], ptr [[TMP4]], ptr [[TMP6]], ptr [[TMP8]]) #[[ATTR4]]1027// CHECK1-NEXT:    ret void1028//1029//1030// CHECK1-LABEL: define linkonce_odr hidden void @_ZNSt7complexIfEC2ERKfS2_(1031// CHECK1-SAME: ptr nonnull align 4 dereferenceable(8) [[THIS:%.*]], ptr nonnull align 4 dereferenceable(4) [[__RE:%.*]], ptr nonnull align 4 dereferenceable(4) [[__IM:%.*]]) unnamed_addr #[[ATTR5]] comdat align 2 {1032// CHECK1-NEXT:  [[ENTRY:.*:]]1033// CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 81034// CHECK1-NEXT:    [[__RE_ADDR:%.*]] = alloca ptr, align 81035// CHECK1-NEXT:    [[__IM_ADDR:%.*]] = alloca ptr, align 81036// CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8, !tbaa [[_ZTSST7COMPLEXIFEPTR_TBAA18]]1037// CHECK1-NEXT:    store ptr [[__RE]], ptr [[__RE_ADDR]], align 8, !tbaa [[FLOATPTR_TBAA20]]1038// CHECK1-NEXT:    store ptr [[__IM]], ptr [[__IM_ADDR]], align 8, !tbaa [[FLOATPTR_TBAA20]]1039// CHECK1-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81040// CHECK1-NEXT:    [[__RE_:%.*]] = getelementptr inbounds nuw %"class.std::complex", ptr [[THIS1]], i32 0, i32 01041// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[__RE_ADDR]], align 8, !tbaa [[FLOATPTR_TBAA20]], !nonnull [[META22]], !align [[META23]]1042// CHECK1-NEXT:    [[TMP1:%.*]] = load float, ptr [[TMP0]], align 4, !tbaa [[FLOAT_TBAA14]]1043// CHECK1-NEXT:    store float [[TMP1]], ptr [[__RE_]], align 4, !tbaa [[FLOAT_TBAA24]]1044// CHECK1-NEXT:    [[__IM_:%.*]] = getelementptr inbounds nuw %"class.std::complex", ptr [[THIS1]], i32 0, i32 11045// CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[__IM_ADDR]], align 8, !tbaa [[FLOATPTR_TBAA20]], !nonnull [[META22]], !align [[META23]]1046// CHECK1-NEXT:    [[TMP3:%.*]] = load float, ptr [[TMP2]], align 4, !tbaa [[FLOAT_TBAA14]]1047// CHECK1-NEXT:    store float [[TMP3]], ptr [[__IM_]], align 4, !tbaa [[FLOAT_TBAA26]]1048// CHECK1-NEXT:    ret void1049//1050//1051// CHECK1-LABEL: define linkonce_odr hidden float @_ZNKSt7complexIfE4realEv(1052// CHECK1-SAME: ptr nonnull align 4 dereferenceable(8) [[THIS:%.*]]) #[[ATTR5]] comdat align 2 {1053// CHECK1-NEXT:  [[ENTRY:.*:]]1054// CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 81055// CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8, !tbaa [[_ZTSST7COMPLEXIFEPTR_TBAA18]]1056// CHECK1-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81057// CHECK1-NEXT:    [[__RE_:%.*]] = getelementptr inbounds nuw %"class.std::complex", ptr [[THIS1]], i32 0, i32 01058// CHECK1-NEXT:    [[TMP0:%.*]] = load float, ptr [[__RE_]], align 4, !tbaa [[FLOAT_TBAA24]]1059// CHECK1-NEXT:    ret float [[TMP0]]1060//1061//1062// CHECK1-LABEL: define linkonce_odr hidden float @_ZNKSt7complexIfE4imagEv(1063// CHECK1-SAME: ptr nonnull align 4 dereferenceable(8) [[THIS:%.*]]) #[[ATTR5]] comdat align 2 {1064// CHECK1-NEXT:  [[ENTRY:.*:]]1065// CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 81066// CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8, !tbaa [[_ZTSST7COMPLEXIFEPTR_TBAA18]]1067// CHECK1-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81068// CHECK1-NEXT:    [[__IM_:%.*]] = getelementptr inbounds nuw %"class.std::complex", ptr [[THIS1]], i32 0, i32 11069// CHECK1-NEXT:    [[TMP0:%.*]] = load float, ptr [[__IM_]], align 4, !tbaa [[FLOAT_TBAA26]]1070// CHECK1-NEXT:    ret float [[TMP0]]1071//1072//1073// CHECK1-LABEL: define linkonce_odr hidden void @_ZNSt7complexIdEC2ERKdS2_(1074// CHECK1-SAME: ptr nonnull align 8 dereferenceable(16) [[THIS:%.*]], ptr nonnull align 8 dereferenceable(8) [[__RE:%.*]], ptr nonnull align 8 dereferenceable(8) [[__IM:%.*]]) unnamed_addr #[[ATTR5]] comdat align 2 {1075// CHECK1-NEXT:  [[ENTRY:.*:]]1076// CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 81077// CHECK1-NEXT:    [[__RE_ADDR:%.*]] = alloca ptr, align 81078// CHECK1-NEXT:    [[__IM_ADDR:%.*]] = alloca ptr, align 81079// CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8, !tbaa [[_ZTSST7COMPLEXIDEPTR_TBAA31]]1080// CHECK1-NEXT:    store ptr [[__RE]], ptr [[__RE_ADDR]], align 8, !tbaa [[DOUBLEPTR_TBAA33]]1081// CHECK1-NEXT:    store ptr [[__IM]], ptr [[__IM_ADDR]], align 8, !tbaa [[DOUBLEPTR_TBAA33]]1082// CHECK1-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81083// CHECK1-NEXT:    [[__RE_:%.*]] = getelementptr inbounds nuw %"class.std::complex.0", ptr [[THIS1]], i32 0, i32 01084// CHECK1-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[__RE_ADDR]], align 8, !tbaa [[DOUBLEPTR_TBAA33]], !nonnull [[META22]], !align [[META35]]1085// CHECK1-NEXT:    [[TMP1:%.*]] = load double, ptr [[TMP0]], align 8, !tbaa [[DOUBLE_TBAA29]]1086// CHECK1-NEXT:    store double [[TMP1]], ptr [[__RE_]], align 8, !tbaa [[DOUBLE_TBAA36]]1087// CHECK1-NEXT:    [[__IM_:%.*]] = getelementptr inbounds nuw %"class.std::complex.0", ptr [[THIS1]], i32 0, i32 11088// CHECK1-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[__IM_ADDR]], align 8, !tbaa [[DOUBLEPTR_TBAA33]], !nonnull [[META22]], !align [[META35]]1089// CHECK1-NEXT:    [[TMP3:%.*]] = load double, ptr [[TMP2]], align 8, !tbaa [[DOUBLE_TBAA29]]1090// CHECK1-NEXT:    store double [[TMP3]], ptr [[__IM_]], align 8, !tbaa [[DOUBLE_TBAA38]]1091// CHECK1-NEXT:    ret void1092//1093//1094// CHECK1-LABEL: define linkonce_odr hidden double @_ZNKSt7complexIdE4realEv(1095// CHECK1-SAME: ptr nonnull align 8 dereferenceable(16) [[THIS:%.*]]) #[[ATTR5]] comdat align 2 {1096// CHECK1-NEXT:  [[ENTRY:.*:]]1097// CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 81098// CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8, !tbaa [[_ZTSST7COMPLEXIDEPTR_TBAA31]]1099// CHECK1-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81100// CHECK1-NEXT:    [[__RE_:%.*]] = getelementptr inbounds nuw %"class.std::complex.0", ptr [[THIS1]], i32 0, i32 01101// CHECK1-NEXT:    [[TMP0:%.*]] = load double, ptr [[__RE_]], align 8, !tbaa [[DOUBLE_TBAA36]]1102// CHECK1-NEXT:    ret double [[TMP0]]1103//1104//1105// CHECK1-LABEL: define linkonce_odr hidden double @_ZNKSt7complexIdE4imagEv(1106// CHECK1-SAME: ptr nonnull align 8 dereferenceable(16) [[THIS:%.*]]) #[[ATTR5]] comdat align 2 {1107// CHECK1-NEXT:  [[ENTRY:.*:]]1108// CHECK1-NEXT:    [[THIS_ADDR:%.*]] = alloca ptr, align 81109// CHECK1-NEXT:    store ptr [[THIS]], ptr [[THIS_ADDR]], align 8, !tbaa [[_ZTSST7COMPLEXIDEPTR_TBAA31]]1110// CHECK1-NEXT:    [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 81111// CHECK1-NEXT:    [[__IM_:%.*]] = getelementptr inbounds nuw %"class.std::complex.0", ptr [[THIS1]], i32 0, i32 11112// CHECK1-NEXT:    [[TMP0:%.*]] = load double, ptr [[__IM_]], align 8, !tbaa [[DOUBLE_TBAA38]]1113// CHECK1-NEXT:    ret double [[TMP0]]1114//1115//.1116// CHECK1: [[INT_TBAA6]] = !{[[META7:![0-9]+]], [[META7]], i64 0}1117// CHECK1: [[META7]] = !{!"int", [[META8:![0-9]+]], i64 0}1118// CHECK1: [[META8]] = !{!"omnipotent char", [[META9:![0-9]+]], i64 0}1119// CHECK1: [[META9]] = !{!"Simple C++ TBAA"}1120// CHECK1: [[ANYPTR_TBAA10]] = !{[[META11:![0-9]+]], [[META11]], i64 0}1121// CHECK1: [[META11]] = !{!"any pointer", [[META8]], i64 0}1122// CHECK1: [[INTPTR_TBAA12]] = !{[[META13:![0-9]+]], [[META13]], i64 0}1123// CHECK1: [[META13]] = !{!"p1 int", [[META11]], i64 0}1124// CHECK1: [[FLOAT_TBAA14]] = !{[[META15:![0-9]+]], [[META15]], i64 0}1125// CHECK1: [[META15]] = !{!"float", [[META8]], i64 0}1126// CHECK1: [[ANYPTR_TBAA16]] = !{[[META17:![0-9]+]], [[META17]], i64 0}1127// CHECK1: [[META17]] = !{!"any p2 pointer", [[META11]], i64 0}1128// CHECK1: [[_ZTSST7COMPLEXIFEPTR_TBAA18]] = !{[[META19:![0-9]+]], [[META19]], i64 0}1129// CHECK1: [[META19]] = !{!"p1 _ZTSSt7complexIfE", [[META11]], i64 0}1130// CHECK1: [[FLOATPTR_TBAA20]] = !{[[META21:![0-9]+]], [[META21]], i64 0}1131// CHECK1: [[META21]] = !{!"p1 float", [[META11]], i64 0}1132// CHECK1: [[META22]] = !{}1133// CHECK1: [[META23]] = !{i64 4}1134// CHECK1: [[FLOAT_TBAA24]] = !{[[META25:![0-9]+]], [[META15]], i64 0}1135// CHECK1: [[META25]] = !{!"_ZTSSt7complexIfE", [[META15]], i64 0, [[META15]], i64 4}1136// CHECK1: [[FLOAT_TBAA26]] = !{[[META25]], [[META15]], i64 4}1137// CHECK1: [[SHORT_TBAA27]] = !{[[META28:![0-9]+]], [[META28]], i64 0}1138// CHECK1: [[META28]] = !{!"short", [[META8]], i64 0}1139// CHECK1: [[DOUBLE_TBAA29]] = !{[[META30:![0-9]+]], [[META30]], i64 0}1140// CHECK1: [[META30]] = !{!"double", [[META8]], i64 0}1141// CHECK1: [[_ZTSST7COMPLEXIDEPTR_TBAA31]] = !{[[META32:![0-9]+]], [[META32]], i64 0}1142// CHECK1: [[META32]] = !{!"p1 _ZTSSt7complexIdE", [[META11]], i64 0}1143// CHECK1: [[DOUBLEPTR_TBAA33]] = !{[[META34:![0-9]+]], [[META34]], i64 0}1144// CHECK1: [[META34]] = !{!"p1 double", [[META11]], i64 0}1145// CHECK1: [[META35]] = !{i64 8}1146// CHECK1: [[DOUBLE_TBAA36]] = !{[[META37:![0-9]+]], [[META30]], i64 0}1147// CHECK1: [[META37]] = !{!"_ZTSSt7complexIdE", [[META30]], i64 0, [[META30]], i64 8}1148// CHECK1: [[DOUBLE_TBAA38]] = !{[[META37]], [[META30]], i64 8}1149//.1150