diff --git a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp --- a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp @@ -445,14 +445,7 @@ void VisitImplicitCastExpr(const ImplicitCastExpr *E) { if (!E) return; - if (E->getCastKind() == CK_ArrayToPointerDecay) { - const bool SavedAllEscaped = AllEscaped; - AllEscaped = true; - Visit(E->getSubExpr()); - AllEscaped = SavedAllEscaped; - } else { - Visit(E->getSubExpr()); - } + Visit(E->getSubExpr()); } void VisitExpr(const Expr *E) { if (!E) diff --git a/clang/test/OpenMP/target_alloc_shared_emission.cpp b/clang/test/OpenMP/target_alloc_shared_emission.cpp new file mode 100644 --- /dev/null +++ b/clang/test/OpenMP/target_alloc_shared_emission.cpp @@ -0,0 +1,827 @@ +// 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 _ +// REQUIRES: amdgpu-registered-target + + +// Test target codegen - host bc file has to be created first. +// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-ppc-host-amd.bc +// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple amdgcn-amd-amdhsa -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-target-debug -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host-amd.bc -o - | FileCheck %s --check-prefix=CHECK-AMD + +// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host-nvidia.bc +// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-unknown-unknown -emit-llvm %s -fopenmp-target-debug -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host-nvidia.bc -o - | FileCheck %s --check-prefix=CHECK-NVIDIA + +// expected-no-diagnostics + +#ifndef HEADER +#define HEADER + +void foo(int *stack); + +void emits_alloc_shared(const int *localPadding , int *res) +{ + int stack[64]; + int stackptr = 0; + stack[stackptr++] = -1; + *res = 0; + + do + { + if(localPadding[0] > 0) + stack[stackptr++] = 0; + *res = stack[--stackptr]; + foo(&stack[2]); + } while (*res > 0); +} + +void does_not_emit_alloc_shared(const int *localPadding , int *res) +{ + int stack[64]; + int stackptr = 0; + stack[stackptr++] = -1; + *res = 0; + + do + { + if(localPadding[0] > 0) + stack[stackptr++] = 0; + *res = stack[--stackptr]; + } while (*res > 0); +} + +#define N 1000 + +int main() { + const int maz = 1; + const int may = 2; + const int max = 3; + int res; + int localPadding[N]; +#pragma omp target teams distribute parallel for map(tofrom: localPadding[:N],maz, may, max) + + for (int pi = 0; pi < N; pi++) + { + for (int hz = 0; hz <= maz; hz++) + for (int hy = 0; hy <= may; hy++) + for (int hx = 0; hx <= max; hx++) { + emits_alloc_shared(localPadding, &res); + does_not_emit_alloc_shared(localPadding, &res); + } + localPadding[pi] = res; + } + return 0; +} + +#endif +// CHECK-AMD-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l58 +// CHECK-AMD-SAME: (ptr noundef nonnull align 4 dereferenceable(4000) [[LOCALPADDING:%.*]], i64 noundef [[RES:%.*]]) #[[ATTR0:[0-9]+]] { +// CHECK-AMD-NEXT: entry: +// CHECK-AMD-NEXT: [[LOCALPADDING_ADDR:%.*]] = alloca ptr, align 8, addrspace(5) +// CHECK-AMD-NEXT: [[RES_ADDR:%.*]] = alloca i64, align 8, addrspace(5) +// CHECK-AMD-NEXT: [[RES_CASTED:%.*]] = alloca i64, align 8, addrspace(5) +// CHECK-AMD-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4, addrspace(5) +// CHECK-AMD-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4, addrspace(5) +// CHECK-AMD-NEXT: [[LOCALPADDING_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[LOCALPADDING_ADDR]] to ptr +// CHECK-AMD-NEXT: [[RES_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RES_ADDR]] to ptr +// CHECK-AMD-NEXT: [[RES_CASTED_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RES_CASTED]] to ptr +// CHECK-AMD-NEXT: [[DOTZERO_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTZERO_ADDR]] to ptr +// CHECK-AMD-NEXT: [[DOTTHREADID_TEMP__ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTTHREADID_TEMP_]] to ptr +// CHECK-AMD-NEXT: store ptr [[LOCALPADDING]], ptr [[LOCALPADDING_ADDR_ASCAST]], align 8 +// CHECK-AMD-NEXT: store i64 [[RES]], ptr [[RES_ADDR_ASCAST]], align 8 +// CHECK-AMD-NEXT: [[TMP0:%.*]] = load ptr, ptr [[LOCALPADDING_ADDR_ASCAST]], align 8 +// CHECK-AMD-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr addrspacecast (ptr addrspace(1) @[[GLOB1:[0-9]+]] to ptr), i8 2, i1 false) +// CHECK-AMD-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1 +// CHECK-AMD-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] +// CHECK-AMD: user_code.entry: +// CHECK-AMD-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr addrspacecast (ptr addrspace(1) @[[GLOB1]] to ptr)) +// CHECK-AMD-NEXT: [[TMP3:%.*]] = load i32, ptr [[RES_ADDR_ASCAST]], align 4 +// CHECK-AMD-NEXT: store i32 [[TMP3]], ptr [[RES_CASTED_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[TMP4:%.*]] = load i64, ptr [[RES_CASTED_ASCAST]], align 8 +// CHECK-AMD-NEXT: store i32 0, ptr [[DOTZERO_ADDR_ASCAST]], align 4 +// CHECK-AMD-NEXT: store i32 [[TMP2]], ptr [[DOTTHREADID_TEMP__ASCAST]], align 4 +// CHECK-AMD-NEXT: call void @__omp_outlined__(ptr [[DOTTHREADID_TEMP__ASCAST]], ptr [[DOTZERO_ADDR_ASCAST]], ptr [[TMP0]], i64 [[TMP4]]) #[[ATTR3:[0-9]+]] +// CHECK-AMD-NEXT: call void @__kmpc_target_deinit(ptr addrspacecast (ptr addrspace(1) @[[GLOB1]] to ptr), i8 2) +// CHECK-AMD-NEXT: ret void +// CHECK-AMD: worker.exit: +// CHECK-AMD-NEXT: ret void +// +// +// CHECK-AMD-LABEL: define {{[^@]+}}@__omp_outlined__ +// CHECK-AMD-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[LOCALPADDING:%.*]], i64 noundef [[RES:%.*]]) #[[ATTR1:[0-9]+]] { +// CHECK-AMD-NEXT: entry: +// CHECK-AMD-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8, addrspace(5) +// CHECK-AMD-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8, addrspace(5) +// CHECK-AMD-NEXT: [[LOCALPADDING_ADDR:%.*]] = alloca ptr, align 8, addrspace(5) +// CHECK-AMD-NEXT: [[RES_ADDR:%.*]] = alloca i64, align 8, addrspace(5) +// CHECK-AMD-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4, addrspace(5) +// CHECK-AMD-NEXT: [[TMP:%.*]] = alloca i32, align 4, addrspace(5) +// CHECK-AMD-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4, addrspace(5) +// CHECK-AMD-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4, addrspace(5) +// CHECK-AMD-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4, addrspace(5) +// CHECK-AMD-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4, addrspace(5) +// CHECK-AMD-NEXT: [[PI:%.*]] = alloca i32, align 4, addrspace(5) +// CHECK-AMD-NEXT: [[RES_CASTED:%.*]] = alloca i64, align 8, addrspace(5) +// CHECK-AMD-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [4 x ptr], align 8, addrspace(5) +// CHECK-AMD-NEXT: [[DOTGLOBAL_TID__ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTGLOBAL_TID__ADDR]] to ptr +// CHECK-AMD-NEXT: [[DOTBOUND_TID__ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTBOUND_TID__ADDR]] to ptr +// CHECK-AMD-NEXT: [[LOCALPADDING_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[LOCALPADDING_ADDR]] to ptr +// CHECK-AMD-NEXT: [[RES_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RES_ADDR]] to ptr +// CHECK-AMD-NEXT: [[DOTOMP_IV_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_IV]] to ptr +// CHECK-AMD-NEXT: [[TMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[TMP]] to ptr +// CHECK-AMD-NEXT: [[DOTOMP_COMB_LB_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_COMB_LB]] to ptr +// CHECK-AMD-NEXT: [[DOTOMP_COMB_UB_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_COMB_UB]] to ptr +// CHECK-AMD-NEXT: [[DOTOMP_STRIDE_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_STRIDE]] to ptr +// CHECK-AMD-NEXT: [[DOTOMP_IS_LAST_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_IS_LAST]] to ptr +// CHECK-AMD-NEXT: [[PI_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[PI]] to ptr +// CHECK-AMD-NEXT: [[RES_CASTED_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RES_CASTED]] to ptr +// CHECK-AMD-NEXT: [[CAPTURED_VARS_ADDRS_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[CAPTURED_VARS_ADDRS]] to ptr +// CHECK-AMD-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR_ASCAST]], align 8 +// CHECK-AMD-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR_ASCAST]], align 8 +// CHECK-AMD-NEXT: store ptr [[LOCALPADDING]], ptr [[LOCALPADDING_ADDR_ASCAST]], align 8 +// CHECK-AMD-NEXT: store i64 [[RES]], ptr [[RES_ADDR_ASCAST]], align 8 +// CHECK-AMD-NEXT: [[TMP0:%.*]] = load ptr, ptr [[LOCALPADDING_ADDR_ASCAST]], align 8 +// CHECK-AMD-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB_ASCAST]], align 4 +// CHECK-AMD-NEXT: store i32 999, ptr [[DOTOMP_COMB_UB_ASCAST]], align 4 +// CHECK-AMD-NEXT: store i32 1, ptr [[DOTOMP_STRIDE_ASCAST]], align 4 +// CHECK-AMD-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @__kmpc_get_hardware_num_threads_in_block() +// CHECK-AMD-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR_ASCAST]], align 8 +// CHECK-AMD-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4 +// CHECK-AMD-NEXT: call void @__kmpc_distribute_static_init_4(ptr addrspacecast (ptr addrspace(1) @[[GLOB2:[0-9]+]] to ptr), i32 [[TMP2]], i32 91, ptr [[DOTOMP_IS_LAST_ASCAST]], ptr [[DOTOMP_COMB_LB_ASCAST]], ptr [[DOTOMP_COMB_UB_ASCAST]], ptr [[DOTOMP_STRIDE_ASCAST]], i32 1, i32 [[NVPTX_NUM_THREADS]]) +// CHECK-AMD-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 999 +// CHECK-AMD-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]] +// CHECK-AMD: cond.true: +// CHECK-AMD-NEXT: br label [[COND_END:%.*]] +// CHECK-AMD: cond.false: +// CHECK-AMD-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_UB_ASCAST]], align 4 +// CHECK-AMD-NEXT: br label [[COND_END]] +// CHECK-AMD: cond.end: +// CHECK-AMD-NEXT: [[COND:%.*]] = phi i32 [ 999, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ] +// CHECK-AMD-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_COMB_LB_ASCAST]], align 4 +// CHECK-AMD-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV_ASCAST]], align 4 +// CHECK-AMD-NEXT: br label [[OMP_INNER_FOR_COND:%.*]] +// CHECK-AMD: omp.inner.for.cond: +// CHECK-AMD-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[CMP1:%.*]] = icmp slt i32 [[TMP6]], 1000 +// CHECK-AMD-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]] +// CHECK-AMD: omp.inner.for.body: +// CHECK-AMD-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_LB_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[TMP8:%.*]] = zext i32 [[TMP7]] to i64 +// CHECK-AMD-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_COMB_UB_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[TMP10:%.*]] = zext i32 [[TMP9]] to i64 +// CHECK-AMD-NEXT: [[TMP11:%.*]] = load i32, ptr [[RES_ADDR_ASCAST]], align 4 +// CHECK-AMD-NEXT: store i32 [[TMP11]], ptr [[RES_CASTED_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[TMP12:%.*]] = load i64, ptr [[RES_CASTED_ASCAST]], align 8 +// CHECK-AMD-NEXT: [[TMP13:%.*]] = getelementptr inbounds [4 x ptr], ptr [[CAPTURED_VARS_ADDRS_ASCAST]], i64 0, i64 0 +// CHECK-AMD-NEXT: [[TMP14:%.*]] = inttoptr i64 [[TMP8]] to ptr +// CHECK-AMD-NEXT: store ptr [[TMP14]], ptr [[TMP13]], align 8 +// CHECK-AMD-NEXT: [[TMP15:%.*]] = getelementptr inbounds [4 x ptr], ptr [[CAPTURED_VARS_ADDRS_ASCAST]], i64 0, i64 1 +// CHECK-AMD-NEXT: [[TMP16:%.*]] = inttoptr i64 [[TMP10]] to ptr +// CHECK-AMD-NEXT: store ptr [[TMP16]], ptr [[TMP15]], align 8 +// CHECK-AMD-NEXT: [[TMP17:%.*]] = getelementptr inbounds [4 x ptr], ptr [[CAPTURED_VARS_ADDRS_ASCAST]], i64 0, i64 2 +// CHECK-AMD-NEXT: store ptr [[TMP0]], ptr [[TMP17]], align 8 +// CHECK-AMD-NEXT: [[TMP18:%.*]] = getelementptr inbounds [4 x ptr], ptr [[CAPTURED_VARS_ADDRS_ASCAST]], i64 0, i64 3 +// CHECK-AMD-NEXT: [[TMP19:%.*]] = inttoptr i64 [[TMP12]] to ptr +// CHECK-AMD-NEXT: store ptr [[TMP19]], ptr [[TMP18]], align 8 +// CHECK-AMD-NEXT: call void @__kmpc_parallel_51(ptr addrspacecast (ptr addrspace(1) @[[GLOB1]] to ptr), i32 [[TMP2]], i32 1, i32 -1, i32 -1, ptr @__omp_outlined__.1, ptr null, ptr [[CAPTURED_VARS_ADDRS_ASCAST]], i64 4) +// CHECK-AMD-NEXT: br label [[OMP_INNER_FOR_INC:%.*]] +// CHECK-AMD: omp.inner.for.inc: +// CHECK-AMD-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[TMP21:%.*]] = load i32, ptr [[DOTOMP_STRIDE_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP20]], [[TMP21]] +// CHECK-AMD-NEXT: store i32 [[ADD]], ptr [[DOTOMP_IV_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[TMP22:%.*]] = load i32, ptr [[DOTOMP_COMB_LB_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[TMP23:%.*]] = load i32, ptr [[DOTOMP_STRIDE_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP22]], [[TMP23]] +// CHECK-AMD-NEXT: store i32 [[ADD2]], ptr [[DOTOMP_COMB_LB_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[TMP24:%.*]] = load i32, ptr [[DOTOMP_COMB_UB_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[TMP25:%.*]] = load i32, ptr [[DOTOMP_STRIDE_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP24]], [[TMP25]] +// CHECK-AMD-NEXT: store i32 [[ADD3]], ptr [[DOTOMP_COMB_UB_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[TMP26:%.*]] = load i32, ptr [[DOTOMP_COMB_UB_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP26]], 999 +// CHECK-AMD-NEXT: br i1 [[CMP4]], label [[COND_TRUE5:%.*]], label [[COND_FALSE6:%.*]] +// CHECK-AMD: cond.true5: +// CHECK-AMD-NEXT: br label [[COND_END7:%.*]] +// CHECK-AMD: cond.false6: +// CHECK-AMD-NEXT: [[TMP27:%.*]] = load i32, ptr [[DOTOMP_COMB_UB_ASCAST]], align 4 +// CHECK-AMD-NEXT: br label [[COND_END7]] +// CHECK-AMD: cond.end7: +// CHECK-AMD-NEXT: [[COND8:%.*]] = phi i32 [ 999, [[COND_TRUE5]] ], [ [[TMP27]], [[COND_FALSE6]] ] +// CHECK-AMD-NEXT: store i32 [[COND8]], ptr [[DOTOMP_COMB_UB_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[TMP28:%.*]] = load i32, ptr [[DOTOMP_COMB_LB_ASCAST]], align 4 +// CHECK-AMD-NEXT: store i32 [[TMP28]], ptr [[DOTOMP_IV_ASCAST]], align 4 +// CHECK-AMD-NEXT: br label [[OMP_INNER_FOR_COND]] +// CHECK-AMD: omp.inner.for.end: +// CHECK-AMD-NEXT: br label [[OMP_LOOP_EXIT:%.*]] +// CHECK-AMD: omp.loop.exit: +// CHECK-AMD-NEXT: call void @__kmpc_distribute_static_fini(ptr addrspacecast (ptr addrspace(1) @[[GLOB2]] to ptr), i32 [[TMP2]]) +// CHECK-AMD-NEXT: ret void +// +// +// CHECK-AMD-LABEL: define {{[^@]+}}@__omp_outlined__.1 +// CHECK-AMD-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[LOCALPADDING:%.*]], i64 noundef [[RES:%.*]]) #[[ATTR1]] { +// CHECK-AMD-NEXT: entry: +// CHECK-AMD-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8, addrspace(5) +// CHECK-AMD-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8, addrspace(5) +// CHECK-AMD-NEXT: [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 8, addrspace(5) +// CHECK-AMD-NEXT: [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 8, addrspace(5) +// CHECK-AMD-NEXT: [[LOCALPADDING_ADDR:%.*]] = alloca ptr, align 8, addrspace(5) +// CHECK-AMD-NEXT: [[RES_ADDR:%.*]] = alloca i64, align 8, addrspace(5) +// CHECK-AMD-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4, addrspace(5) +// CHECK-AMD-NEXT: [[TMP:%.*]] = alloca i32, align 4, addrspace(5) +// CHECK-AMD-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4, addrspace(5) +// CHECK-AMD-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4, addrspace(5) +// CHECK-AMD-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4, addrspace(5) +// CHECK-AMD-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4, addrspace(5) +// CHECK-AMD-NEXT: [[PI:%.*]] = alloca i32, align 4, addrspace(5) +// CHECK-AMD-NEXT: [[HZ:%.*]] = alloca i32, align 4, addrspace(5) +// CHECK-AMD-NEXT: [[HY:%.*]] = alloca i32, align 4, addrspace(5) +// CHECK-AMD-NEXT: [[HX:%.*]] = alloca i32, align 4, addrspace(5) +// CHECK-AMD-NEXT: [[DOTGLOBAL_TID__ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTGLOBAL_TID__ADDR]] to ptr +// CHECK-AMD-NEXT: [[DOTBOUND_TID__ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTBOUND_TID__ADDR]] to ptr +// CHECK-AMD-NEXT: [[DOTPREVIOUS_LB__ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTPREVIOUS_LB__ADDR]] to ptr +// CHECK-AMD-NEXT: [[DOTPREVIOUS_UB__ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTPREVIOUS_UB__ADDR]] to ptr +// CHECK-AMD-NEXT: [[LOCALPADDING_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[LOCALPADDING_ADDR]] to ptr +// CHECK-AMD-NEXT: [[RES_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RES_ADDR]] to ptr +// CHECK-AMD-NEXT: [[DOTOMP_IV_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_IV]] to ptr +// CHECK-AMD-NEXT: [[TMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[TMP]] to ptr +// CHECK-AMD-NEXT: [[DOTOMP_LB_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_LB]] to ptr +// CHECK-AMD-NEXT: [[DOTOMP_UB_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_UB]] to ptr +// CHECK-AMD-NEXT: [[DOTOMP_STRIDE_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_STRIDE]] to ptr +// CHECK-AMD-NEXT: [[DOTOMP_IS_LAST_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_IS_LAST]] to ptr +// CHECK-AMD-NEXT: [[PI_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[PI]] to ptr +// CHECK-AMD-NEXT: [[HZ_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[HZ]] to ptr +// CHECK-AMD-NEXT: [[HY_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[HY]] to ptr +// CHECK-AMD-NEXT: [[HX_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[HX]] to ptr +// CHECK-AMD-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR_ASCAST]], align 8 +// CHECK-AMD-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR_ASCAST]], align 8 +// CHECK-AMD-NEXT: store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR_ASCAST]], align 8 +// CHECK-AMD-NEXT: store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR_ASCAST]], align 8 +// CHECK-AMD-NEXT: store ptr [[LOCALPADDING]], ptr [[LOCALPADDING_ADDR_ASCAST]], align 8 +// CHECK-AMD-NEXT: store i64 [[RES]], ptr [[RES_ADDR_ASCAST]], align 8 +// CHECK-AMD-NEXT: [[TMP0:%.*]] = load ptr, ptr [[LOCALPADDING_ADDR_ASCAST]], align 8 +// CHECK-AMD-NEXT: store i32 0, ptr [[DOTOMP_LB_ASCAST]], align 4 +// CHECK-AMD-NEXT: store i32 999, ptr [[DOTOMP_UB_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[TMP1:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR_ASCAST]], align 8 +// CHECK-AMD-NEXT: [[CONV:%.*]] = trunc i64 [[TMP1]] to i32 +// CHECK-AMD-NEXT: [[TMP2:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR_ASCAST]], align 8 +// CHECK-AMD-NEXT: [[CONV1:%.*]] = trunc i64 [[TMP2]] to i32 +// CHECK-AMD-NEXT: store i32 [[CONV]], ptr [[DOTOMP_LB_ASCAST]], align 4 +// CHECK-AMD-NEXT: store i32 [[CONV1]], ptr [[DOTOMP_UB_ASCAST]], align 4 +// CHECK-AMD-NEXT: store i32 1, ptr [[DOTOMP_STRIDE_ASCAST]], align 4 +// CHECK-AMD-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[TMP3:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR_ASCAST]], align 8 +// CHECK-AMD-NEXT: [[TMP4:%.*]] = load i32, ptr [[TMP3]], align 4 +// CHECK-AMD-NEXT: call void @__kmpc_for_static_init_4(ptr addrspacecast (ptr addrspace(1) @[[GLOB3:[0-9]+]] to ptr), i32 [[TMP4]], i32 33, ptr [[DOTOMP_IS_LAST_ASCAST]], ptr [[DOTOMP_LB_ASCAST]], ptr [[DOTOMP_UB_ASCAST]], ptr [[DOTOMP_STRIDE_ASCAST]], i32 1, i32 1) +// CHECK-AMD-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB_ASCAST]], align 4 +// CHECK-AMD-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV_ASCAST]], align 4 +// CHECK-AMD-NEXT: br label [[OMP_INNER_FOR_COND:%.*]] +// CHECK-AMD: omp.inner.for.cond: +// CHECK-AMD-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[CONV2:%.*]] = sext i32 [[TMP6]] to i64 +// CHECK-AMD-NEXT: [[TMP7:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR_ASCAST]], align 8 +// CHECK-AMD-NEXT: [[CMP:%.*]] = icmp ule i64 [[CONV2]], [[TMP7]] +// CHECK-AMD-NEXT: br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]] +// CHECK-AMD: omp.inner.for.body: +// CHECK-AMD-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP8]], 1 +// CHECK-AMD-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]] +// CHECK-AMD-NEXT: store i32 [[ADD]], ptr [[PI_ASCAST]], align 4 +// CHECK-AMD-NEXT: store i32 0, ptr [[HZ_ASCAST]], align 4 +// CHECK-AMD-NEXT: br label [[FOR_COND:%.*]] +// CHECK-AMD: for.cond: +// CHECK-AMD-NEXT: [[TMP9:%.*]] = load i32, ptr [[HZ_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[CMP3:%.*]] = icmp sle i32 [[TMP9]], 1 +// CHECK-AMD-NEXT: br i1 [[CMP3]], label [[FOR_BODY:%.*]], label [[FOR_END16:%.*]] +// CHECK-AMD: for.body: +// CHECK-AMD-NEXT: store i32 0, ptr [[HY_ASCAST]], align 4 +// CHECK-AMD-NEXT: br label [[FOR_COND4:%.*]] +// CHECK-AMD: for.cond4: +// CHECK-AMD-NEXT: [[TMP10:%.*]] = load i32, ptr [[HY_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP10]], 2 +// CHECK-AMD-NEXT: br i1 [[CMP5]], label [[FOR_BODY6:%.*]], label [[FOR_END13:%.*]] +// CHECK-AMD: for.body6: +// CHECK-AMD-NEXT: store i32 0, ptr [[HX_ASCAST]], align 4 +// CHECK-AMD-NEXT: br label [[FOR_COND7:%.*]] +// CHECK-AMD: for.cond7: +// CHECK-AMD-NEXT: [[TMP11:%.*]] = load i32, ptr [[HX_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[CMP8:%.*]] = icmp sle i32 [[TMP11]], 3 +// CHECK-AMD-NEXT: br i1 [[CMP8]], label [[FOR_BODY9:%.*]], label [[FOR_END:%.*]] +// CHECK-AMD: for.body9: +// CHECK-AMD-NEXT: [[ARRAYDECAY:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i64 0, i64 0 +// CHECK-AMD-NEXT: call void @_Z18emits_alloc_sharedPKiPi(ptr noundef [[ARRAYDECAY]], ptr noundef [[RES_ADDR_ASCAST]]) #[[ATTR8:[0-9]+]] +// CHECK-AMD-NEXT: [[ARRAYDECAY10:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i64 0, i64 0 +// CHECK-AMD-NEXT: call void @_Z26does_not_emit_alloc_sharedPKiPi(ptr noundef [[ARRAYDECAY10]], ptr noundef [[RES_ADDR_ASCAST]]) #[[ATTR8]] +// CHECK-AMD-NEXT: br label [[FOR_INC:%.*]] +// CHECK-AMD: for.inc: +// CHECK-AMD-NEXT: [[TMP12:%.*]] = load i32, ptr [[HX_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[INC:%.*]] = add nsw i32 [[TMP12]], 1 +// CHECK-AMD-NEXT: store i32 [[INC]], ptr [[HX_ASCAST]], align 4 +// CHECK-AMD-NEXT: br label [[FOR_COND7]], !llvm.loop [[LOOP7:![0-9]+]] +// CHECK-AMD: for.end: +// CHECK-AMD-NEXT: br label [[FOR_INC11:%.*]] +// CHECK-AMD: for.inc11: +// CHECK-AMD-NEXT: [[TMP13:%.*]] = load i32, ptr [[HY_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[INC12:%.*]] = add nsw i32 [[TMP13]], 1 +// CHECK-AMD-NEXT: store i32 [[INC12]], ptr [[HY_ASCAST]], align 4 +// CHECK-AMD-NEXT: br label [[FOR_COND4]], !llvm.loop [[LOOP9:![0-9]+]] +// CHECK-AMD: for.end13: +// CHECK-AMD-NEXT: br label [[FOR_INC14:%.*]] +// CHECK-AMD: for.inc14: +// CHECK-AMD-NEXT: [[TMP14:%.*]] = load i32, ptr [[HZ_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[INC15:%.*]] = add nsw i32 [[TMP14]], 1 +// CHECK-AMD-NEXT: store i32 [[INC15]], ptr [[HZ_ASCAST]], align 4 +// CHECK-AMD-NEXT: br label [[FOR_COND]], !llvm.loop [[LOOP10:![0-9]+]] +// CHECK-AMD: for.end16: +// CHECK-AMD-NEXT: [[TMP15:%.*]] = load i32, ptr [[RES_ADDR_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[TMP16:%.*]] = load i32, ptr [[PI_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP16]] to i64 +// CHECK-AMD-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]] +// CHECK-AMD-NEXT: store i32 [[TMP15]], ptr [[ARRAYIDX]], align 4 +// CHECK-AMD-NEXT: br label [[OMP_BODY_CONTINUE:%.*]] +// CHECK-AMD: omp.body.continue: +// CHECK-AMD-NEXT: br label [[OMP_INNER_FOR_INC:%.*]] +// CHECK-AMD: omp.inner.for.inc: +// CHECK-AMD-NEXT: [[TMP17:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[TMP18:%.*]] = load i32, ptr [[DOTOMP_STRIDE_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[ADD17:%.*]] = add nsw i32 [[TMP17]], [[TMP18]] +// CHECK-AMD-NEXT: store i32 [[ADD17]], ptr [[DOTOMP_IV_ASCAST]], align 4 +// CHECK-AMD-NEXT: br label [[OMP_INNER_FOR_COND]] +// CHECK-AMD: omp.inner.for.end: +// CHECK-AMD-NEXT: br label [[OMP_LOOP_EXIT:%.*]] +// CHECK-AMD: omp.loop.exit: +// CHECK-AMD-NEXT: call void @__kmpc_distribute_static_fini(ptr addrspacecast (ptr addrspace(1) @[[GLOB2]] to ptr), i32 [[TMP4]]) +// CHECK-AMD-NEXT: ret void +// +// +// CHECK-AMD-LABEL: define {{[^@]+}}@_Z18emits_alloc_sharedPKiPi +// CHECK-AMD-SAME: (ptr noundef [[LOCALPADDING:%.*]], ptr noundef [[RES:%.*]]) #[[ATTR2:[0-9]+]] { +// CHECK-AMD-NEXT: entry: +// CHECK-AMD-NEXT: [[LOCALPADDING_ADDR:%.*]] = alloca ptr, align 8, addrspace(5) +// CHECK-AMD-NEXT: [[RES_ADDR:%.*]] = alloca ptr, align 8, addrspace(5) +// CHECK-AMD-NEXT: [[STACKPTR:%.*]] = alloca i32, align 4, addrspace(5) +// CHECK-AMD-NEXT: [[LOCALPADDING_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[LOCALPADDING_ADDR]] to ptr +// CHECK-AMD-NEXT: [[RES_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RES_ADDR]] to ptr +// CHECK-AMD-NEXT: [[STACKPTR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[STACKPTR]] to ptr +// CHECK-AMD-NEXT: [[STACK:%.*]] = call align 8 ptr @__kmpc_alloc_shared(i64 256) +// CHECK-AMD-NEXT: store ptr [[LOCALPADDING]], ptr [[LOCALPADDING_ADDR_ASCAST]], align 8 +// CHECK-AMD-NEXT: store ptr [[RES]], ptr [[RES_ADDR_ASCAST]], align 8 +// CHECK-AMD-NEXT: store i32 0, ptr [[STACKPTR_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[TMP0:%.*]] = load i32, ptr [[STACKPTR_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[INC:%.*]] = add nsw i32 [[TMP0]], 1 +// CHECK-AMD-NEXT: store i32 [[INC]], ptr [[STACKPTR_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP0]] to i64 +// CHECK-AMD-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [64 x i32], ptr [[STACK]], i64 0, i64 [[IDXPROM]] +// CHECK-AMD-NEXT: store i32 -1, ptr [[ARRAYIDX]], align 4 +// CHECK-AMD-NEXT: [[TMP1:%.*]] = load ptr, ptr [[RES_ADDR_ASCAST]], align 8 +// CHECK-AMD-NEXT: store i32 0, ptr [[TMP1]], align 4 +// CHECK-AMD-NEXT: br label [[DO_BODY:%.*]] +// CHECK-AMD: do.body: +// CHECK-AMD-NEXT: [[TMP2:%.*]] = load ptr, ptr [[LOCALPADDING_ADDR_ASCAST]], align 8 +// CHECK-AMD-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds i32, ptr [[TMP2]], i64 0 +// CHECK-AMD-NEXT: [[TMP3:%.*]] = load i32, ptr [[ARRAYIDX1]], align 4 +// CHECK-AMD-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 0 +// CHECK-AMD-NEXT: br i1 [[CMP]], label [[IF_THEN:%.*]], label [[IF_END:%.*]] +// CHECK-AMD: if.then: +// CHECK-AMD-NEXT: [[TMP4:%.*]] = load i32, ptr [[STACKPTR_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[INC2:%.*]] = add nsw i32 [[TMP4]], 1 +// CHECK-AMD-NEXT: store i32 [[INC2]], ptr [[STACKPTR_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[IDXPROM3:%.*]] = sext i32 [[TMP4]] to i64 +// CHECK-AMD-NEXT: [[ARRAYIDX4:%.*]] = getelementptr inbounds [64 x i32], ptr [[STACK]], i64 0, i64 [[IDXPROM3]] +// CHECK-AMD-NEXT: store i32 0, ptr [[ARRAYIDX4]], align 4 +// CHECK-AMD-NEXT: br label [[IF_END]] +// CHECK-AMD: if.end: +// CHECK-AMD-NEXT: [[TMP5:%.*]] = load i32, ptr [[STACKPTR_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[DEC:%.*]] = add nsw i32 [[TMP5]], -1 +// CHECK-AMD-NEXT: store i32 [[DEC]], ptr [[STACKPTR_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[IDXPROM5:%.*]] = sext i32 [[DEC]] to i64 +// CHECK-AMD-NEXT: [[ARRAYIDX6:%.*]] = getelementptr inbounds [64 x i32], ptr [[STACK]], i64 0, i64 [[IDXPROM5]] +// CHECK-AMD-NEXT: [[TMP6:%.*]] = load i32, ptr [[ARRAYIDX6]], align 4 +// CHECK-AMD-NEXT: [[TMP7:%.*]] = load ptr, ptr [[RES_ADDR_ASCAST]], align 8 +// CHECK-AMD-NEXT: store i32 [[TMP6]], ptr [[TMP7]], align 4 +// CHECK-AMD-NEXT: [[ARRAYIDX7:%.*]] = getelementptr inbounds [64 x i32], ptr [[STACK]], i64 0, i64 2 +// CHECK-AMD-NEXT: call void @_Z3fooPi(ptr noundef [[ARRAYIDX7]]) #[[ATTR8]] +// CHECK-AMD-NEXT: br label [[DO_COND:%.*]] +// CHECK-AMD: do.cond: +// CHECK-AMD-NEXT: [[TMP8:%.*]] = load ptr, ptr [[RES_ADDR_ASCAST]], align 8 +// CHECK-AMD-NEXT: [[TMP9:%.*]] = load i32, ptr [[TMP8]], align 4 +// CHECK-AMD-NEXT: [[CMP8:%.*]] = icmp sgt i32 [[TMP9]], 0 +// CHECK-AMD-NEXT: br i1 [[CMP8]], label [[DO_BODY]], label [[DO_END:%.*]], !llvm.loop [[LOOP11:![0-9]+]] +// CHECK-AMD: do.end: +// CHECK-AMD-NEXT: call void @__kmpc_free_shared(ptr [[STACK]], i64 256) +// CHECK-AMD-NEXT: ret void +// +// +// CHECK-AMD-LABEL: define {{[^@]+}}@_Z26does_not_emit_alloc_sharedPKiPi +// CHECK-AMD-SAME: (ptr noundef [[LOCALPADDING:%.*]], ptr noundef [[RES:%.*]]) #[[ATTR2]] { +// CHECK-AMD-NEXT: entry: +// CHECK-AMD-NEXT: [[LOCALPADDING_ADDR:%.*]] = alloca ptr, align 8, addrspace(5) +// CHECK-AMD-NEXT: [[RES_ADDR:%.*]] = alloca ptr, align 8, addrspace(5) +// CHECK-AMD-NEXT: [[STACK:%.*]] = alloca [64 x i32], align 4, addrspace(5) +// CHECK-AMD-NEXT: [[STACKPTR:%.*]] = alloca i32, align 4, addrspace(5) +// CHECK-AMD-NEXT: [[LOCALPADDING_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[LOCALPADDING_ADDR]] to ptr +// CHECK-AMD-NEXT: [[RES_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RES_ADDR]] to ptr +// CHECK-AMD-NEXT: [[STACK_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[STACK]] to ptr +// CHECK-AMD-NEXT: [[STACKPTR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[STACKPTR]] to ptr +// CHECK-AMD-NEXT: store ptr [[LOCALPADDING]], ptr [[LOCALPADDING_ADDR_ASCAST]], align 8 +// CHECK-AMD-NEXT: store ptr [[RES]], ptr [[RES_ADDR_ASCAST]], align 8 +// CHECK-AMD-NEXT: store i32 0, ptr [[STACKPTR_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[TMP0:%.*]] = load i32, ptr [[STACKPTR_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[INC:%.*]] = add nsw i32 [[TMP0]], 1 +// CHECK-AMD-NEXT: store i32 [[INC]], ptr [[STACKPTR_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP0]] to i64 +// CHECK-AMD-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [64 x i32], ptr [[STACK_ASCAST]], i64 0, i64 [[IDXPROM]] +// CHECK-AMD-NEXT: store i32 -1, ptr [[ARRAYIDX]], align 4 +// CHECK-AMD-NEXT: [[TMP1:%.*]] = load ptr, ptr [[RES_ADDR_ASCAST]], align 8 +// CHECK-AMD-NEXT: store i32 0, ptr [[TMP1]], align 4 +// CHECK-AMD-NEXT: br label [[DO_BODY:%.*]] +// CHECK-AMD: do.body: +// CHECK-AMD-NEXT: [[TMP2:%.*]] = load ptr, ptr [[LOCALPADDING_ADDR_ASCAST]], align 8 +// CHECK-AMD-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds i32, ptr [[TMP2]], i64 0 +// CHECK-AMD-NEXT: [[TMP3:%.*]] = load i32, ptr [[ARRAYIDX1]], align 4 +// CHECK-AMD-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 0 +// CHECK-AMD-NEXT: br i1 [[CMP]], label [[IF_THEN:%.*]], label [[IF_END:%.*]] +// CHECK-AMD: if.then: +// CHECK-AMD-NEXT: [[TMP4:%.*]] = load i32, ptr [[STACKPTR_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[INC2:%.*]] = add nsw i32 [[TMP4]], 1 +// CHECK-AMD-NEXT: store i32 [[INC2]], ptr [[STACKPTR_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[IDXPROM3:%.*]] = sext i32 [[TMP4]] to i64 +// CHECK-AMD-NEXT: [[ARRAYIDX4:%.*]] = getelementptr inbounds [64 x i32], ptr [[STACK_ASCAST]], i64 0, i64 [[IDXPROM3]] +// CHECK-AMD-NEXT: store i32 0, ptr [[ARRAYIDX4]], align 4 +// CHECK-AMD-NEXT: br label [[IF_END]] +// CHECK-AMD: if.end: +// CHECK-AMD-NEXT: [[TMP5:%.*]] = load i32, ptr [[STACKPTR_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[DEC:%.*]] = add nsw i32 [[TMP5]], -1 +// CHECK-AMD-NEXT: store i32 [[DEC]], ptr [[STACKPTR_ASCAST]], align 4 +// CHECK-AMD-NEXT: [[IDXPROM5:%.*]] = sext i32 [[DEC]] to i64 +// CHECK-AMD-NEXT: [[ARRAYIDX6:%.*]] = getelementptr inbounds [64 x i32], ptr [[STACK_ASCAST]], i64 0, i64 [[IDXPROM5]] +// CHECK-AMD-NEXT: [[TMP6:%.*]] = load i32, ptr [[ARRAYIDX6]], align 4 +// CHECK-AMD-NEXT: [[TMP7:%.*]] = load ptr, ptr [[RES_ADDR_ASCAST]], align 8 +// CHECK-AMD-NEXT: store i32 [[TMP6]], ptr [[TMP7]], align 4 +// CHECK-AMD-NEXT: br label [[DO_COND:%.*]] +// CHECK-AMD: do.cond: +// CHECK-AMD-NEXT: [[TMP8:%.*]] = load ptr, ptr [[RES_ADDR_ASCAST]], align 8 +// CHECK-AMD-NEXT: [[TMP9:%.*]] = load i32, ptr [[TMP8]], align 4 +// CHECK-AMD-NEXT: [[CMP7:%.*]] = icmp sgt i32 [[TMP9]], 0 +// CHECK-AMD-NEXT: br i1 [[CMP7]], label [[DO_BODY]], label [[DO_END:%.*]], !llvm.loop [[LOOP12:![0-9]+]] +// CHECK-AMD: do.end: +// CHECK-AMD-NEXT: ret void +// +// +// CHECK-NVIDIA-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l58 +// CHECK-NVIDIA-SAME: (ptr noundef nonnull align 4 dereferenceable(4000) [[LOCALPADDING:%.*]], i64 noundef [[RES:%.*]]) #[[ATTR0:[0-9]+]] { +// CHECK-NVIDIA-NEXT: entry: +// CHECK-NVIDIA-NEXT: [[LOCALPADDING_ADDR:%.*]] = alloca ptr, align 8 +// CHECK-NVIDIA-NEXT: [[RES_ADDR:%.*]] = alloca i64, align 8 +// CHECK-NVIDIA-NEXT: [[RES_CASTED:%.*]] = alloca i64, align 8 +// CHECK-NVIDIA-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4 +// CHECK-NVIDIA-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4 +// CHECK-NVIDIA-NEXT: store ptr [[LOCALPADDING]], ptr [[LOCALPADDING_ADDR]], align 8 +// CHECK-NVIDIA-NEXT: store i64 [[RES]], ptr [[RES_ADDR]], align 8 +// CHECK-NVIDIA-NEXT: [[TMP0:%.*]] = load ptr, ptr [[LOCALPADDING_ADDR]], align 8 +// CHECK-NVIDIA-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr @[[GLOB1:[0-9]+]], i8 2, i1 false) +// CHECK-NVIDIA-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1 +// CHECK-NVIDIA-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] +// CHECK-NVIDIA: user_code.entry: +// CHECK-NVIDIA-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]]) +// CHECK-NVIDIA-NEXT: [[TMP3:%.*]] = load i32, ptr [[RES_ADDR]], align 4 +// CHECK-NVIDIA-NEXT: store i32 [[TMP3]], ptr [[RES_CASTED]], align 4 +// CHECK-NVIDIA-NEXT: [[TMP4:%.*]] = load i64, ptr [[RES_CASTED]], align 8 +// CHECK-NVIDIA-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 4 +// CHECK-NVIDIA-NEXT: store i32 [[TMP2]], ptr [[DOTTHREADID_TEMP_]], align 4 +// CHECK-NVIDIA-NEXT: call void @__omp_outlined__(ptr [[DOTTHREADID_TEMP_]], ptr [[DOTZERO_ADDR]], ptr [[TMP0]], i64 [[TMP4]]) #[[ATTR3:[0-9]+]] +// CHECK-NVIDIA-NEXT: call void @__kmpc_target_deinit(ptr @[[GLOB1]], i8 2) +// CHECK-NVIDIA-NEXT: ret void +// CHECK-NVIDIA: worker.exit: +// CHECK-NVIDIA-NEXT: ret void +// +// +// CHECK-NVIDIA-LABEL: define {{[^@]+}}@__omp_outlined__ +// CHECK-NVIDIA-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[LOCALPADDING:%.*]], i64 noundef [[RES:%.*]]) #[[ATTR1:[0-9]+]] { +// CHECK-NVIDIA-NEXT: entry: +// CHECK-NVIDIA-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8 +// CHECK-NVIDIA-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8 +// CHECK-NVIDIA-NEXT: [[LOCALPADDING_ADDR:%.*]] = alloca ptr, align 8 +// CHECK-NVIDIA-NEXT: [[RES_ADDR:%.*]] = alloca i64, align 8 +// CHECK-NVIDIA-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4 +// CHECK-NVIDIA-NEXT: [[TMP:%.*]] = alloca i32, align 4 +// CHECK-NVIDIA-NEXT: [[DOTOMP_COMB_LB:%.*]] = alloca i32, align 4 +// CHECK-NVIDIA-NEXT: [[DOTOMP_COMB_UB:%.*]] = alloca i32, align 4 +// CHECK-NVIDIA-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4 +// CHECK-NVIDIA-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4 +// CHECK-NVIDIA-NEXT: [[PI:%.*]] = alloca i32, align 4 +// CHECK-NVIDIA-NEXT: [[RES_CASTED:%.*]] = alloca i64, align 8 +// CHECK-NVIDIA-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [4 x ptr], align 8 +// CHECK-NVIDIA-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8 +// CHECK-NVIDIA-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8 +// CHECK-NVIDIA-NEXT: store ptr [[LOCALPADDING]], ptr [[LOCALPADDING_ADDR]], align 8 +// CHECK-NVIDIA-NEXT: store i64 [[RES]], ptr [[RES_ADDR]], align 8 +// CHECK-NVIDIA-NEXT: [[TMP0:%.*]] = load ptr, ptr [[LOCALPADDING_ADDR]], align 8 +// CHECK-NVIDIA-NEXT: store i32 0, ptr [[DOTOMP_COMB_LB]], align 4 +// CHECK-NVIDIA-NEXT: store i32 999, ptr [[DOTOMP_COMB_UB]], align 4 +// CHECK-NVIDIA-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4 +// CHECK-NVIDIA-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4 +// CHECK-NVIDIA-NEXT: [[NVPTX_NUM_THREADS:%.*]] = call i32 @__kmpc_get_hardware_num_threads_in_block() +// CHECK-NVIDIA-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8 +// CHECK-NVIDIA-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4 +// CHECK-NVIDIA-NEXT: call void @__kmpc_distribute_static_init_4(ptr @[[GLOB2:[0-9]+]], i32 [[TMP2]], i32 91, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 [[NVPTX_NUM_THREADS]]) +// CHECK-NVIDIA-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4 +// CHECK-NVIDIA-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 999 +// CHECK-NVIDIA-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]] +// CHECK-NVIDIA: cond.true: +// CHECK-NVIDIA-NEXT: br label [[COND_END:%.*]] +// CHECK-NVIDIA: cond.false: +// CHECK-NVIDIA-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4 +// CHECK-NVIDIA-NEXT: br label [[COND_END]] +// CHECK-NVIDIA: cond.end: +// CHECK-NVIDIA-NEXT: [[COND:%.*]] = phi i32 [ 999, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ] +// CHECK-NVIDIA-NEXT: store i32 [[COND]], ptr [[DOTOMP_COMB_UB]], align 4 +// CHECK-NVIDIA-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4 +// CHECK-NVIDIA-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4 +// CHECK-NVIDIA-NEXT: br label [[OMP_INNER_FOR_COND:%.*]] +// CHECK-NVIDIA: omp.inner.for.cond: +// CHECK-NVIDIA-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4 +// CHECK-NVIDIA-NEXT: [[CMP1:%.*]] = icmp slt i32 [[TMP6]], 1000 +// CHECK-NVIDIA-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]] +// CHECK-NVIDIA: omp.inner.for.body: +// CHECK-NVIDIA-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4 +// CHECK-NVIDIA-NEXT: [[TMP8:%.*]] = zext i32 [[TMP7]] to i64 +// CHECK-NVIDIA-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4 +// CHECK-NVIDIA-NEXT: [[TMP10:%.*]] = zext i32 [[TMP9]] to i64 +// CHECK-NVIDIA-NEXT: [[TMP11:%.*]] = load i32, ptr [[RES_ADDR]], align 4 +// CHECK-NVIDIA-NEXT: store i32 [[TMP11]], ptr [[RES_CASTED]], align 4 +// CHECK-NVIDIA-NEXT: [[TMP12:%.*]] = load i64, ptr [[RES_CASTED]], align 8 +// CHECK-NVIDIA-NEXT: [[TMP13:%.*]] = getelementptr inbounds [4 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 0 +// CHECK-NVIDIA-NEXT: [[TMP14:%.*]] = inttoptr i64 [[TMP8]] to ptr +// CHECK-NVIDIA-NEXT: store ptr [[TMP14]], ptr [[TMP13]], align 8 +// CHECK-NVIDIA-NEXT: [[TMP15:%.*]] = getelementptr inbounds [4 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 1 +// CHECK-NVIDIA-NEXT: [[TMP16:%.*]] = inttoptr i64 [[TMP10]] to ptr +// CHECK-NVIDIA-NEXT: store ptr [[TMP16]], ptr [[TMP15]], align 8 +// CHECK-NVIDIA-NEXT: [[TMP17:%.*]] = getelementptr inbounds [4 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 2 +// CHECK-NVIDIA-NEXT: store ptr [[TMP0]], ptr [[TMP17]], align 8 +// CHECK-NVIDIA-NEXT: [[TMP18:%.*]] = getelementptr inbounds [4 x ptr], ptr [[CAPTURED_VARS_ADDRS]], i64 0, i64 3 +// CHECK-NVIDIA-NEXT: [[TMP19:%.*]] = inttoptr i64 [[TMP12]] to ptr +// CHECK-NVIDIA-NEXT: store ptr [[TMP19]], ptr [[TMP18]], align 8 +// CHECK-NVIDIA-NEXT: call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP2]], i32 1, i32 -1, i32 -1, ptr @__omp_outlined__1, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i64 4) +// CHECK-NVIDIA-NEXT: br label [[OMP_INNER_FOR_INC:%.*]] +// CHECK-NVIDIA: omp.inner.for.inc: +// CHECK-NVIDIA-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4 +// CHECK-NVIDIA-NEXT: [[TMP21:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4 +// CHECK-NVIDIA-NEXT: [[ADD:%.*]] = add nsw i32 [[TMP20]], [[TMP21]] +// CHECK-NVIDIA-NEXT: store i32 [[ADD]], ptr [[DOTOMP_IV]], align 4 +// CHECK-NVIDIA-NEXT: [[TMP22:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4 +// CHECK-NVIDIA-NEXT: [[TMP23:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4 +// CHECK-NVIDIA-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP22]], [[TMP23]] +// CHECK-NVIDIA-NEXT: store i32 [[ADD2]], ptr [[DOTOMP_COMB_LB]], align 4 +// CHECK-NVIDIA-NEXT: [[TMP24:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4 +// CHECK-NVIDIA-NEXT: [[TMP25:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4 +// CHECK-NVIDIA-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP24]], [[TMP25]] +// CHECK-NVIDIA-NEXT: store i32 [[ADD3]], ptr [[DOTOMP_COMB_UB]], align 4 +// CHECK-NVIDIA-NEXT: [[TMP26:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4 +// CHECK-NVIDIA-NEXT: [[CMP4:%.*]] = icmp sgt i32 [[TMP26]], 999 +// CHECK-NVIDIA-NEXT: br i1 [[CMP4]], label [[COND_TRUE5:%.*]], label [[COND_FALSE6:%.*]] +// CHECK-NVIDIA: cond.true5: +// CHECK-NVIDIA-NEXT: br label [[COND_END7:%.*]] +// CHECK-NVIDIA: cond.false6: +// CHECK-NVIDIA-NEXT: [[TMP27:%.*]] = load i32, ptr [[DOTOMP_COMB_UB]], align 4 +// CHECK-NVIDIA-NEXT: br label [[COND_END7]] +// CHECK-NVIDIA: cond.end7: +// CHECK-NVIDIA-NEXT: [[COND8:%.*]] = phi i32 [ 999, [[COND_TRUE5]] ], [ [[TMP27]], [[COND_FALSE6]] ] +// CHECK-NVIDIA-NEXT: store i32 [[COND8]], ptr [[DOTOMP_COMB_UB]], align 4 +// CHECK-NVIDIA-NEXT: [[TMP28:%.*]] = load i32, ptr [[DOTOMP_COMB_LB]], align 4 +// CHECK-NVIDIA-NEXT: store i32 [[TMP28]], ptr [[DOTOMP_IV]], align 4 +// CHECK-NVIDIA-NEXT: br label [[OMP_INNER_FOR_COND]] +// CHECK-NVIDIA: omp.inner.for.end: +// CHECK-NVIDIA-NEXT: br label [[OMP_LOOP_EXIT:%.*]] +// CHECK-NVIDIA: omp.loop.exit: +// CHECK-NVIDIA-NEXT: call void @__kmpc_distribute_static_fini(ptr @[[GLOB2]], i32 [[TMP2]]) +// CHECK-NVIDIA-NEXT: ret void +// +// +// CHECK-NVIDIA-LABEL: define {{[^@]+}}@__omp_outlined__1 +// CHECK-NVIDIA-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[DOTPREVIOUS_LB_:%.*]], i64 noundef [[DOTPREVIOUS_UB_:%.*]], ptr noundef nonnull align 4 dereferenceable(4000) [[LOCALPADDING:%.*]], i64 noundef [[RES:%.*]]) #[[ATTR1]] { +// CHECK-NVIDIA-NEXT: entry: +// CHECK-NVIDIA-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8 +// CHECK-NVIDIA-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8 +// CHECK-NVIDIA-NEXT: [[DOTPREVIOUS_LB__ADDR:%.*]] = alloca i64, align 8 +// CHECK-NVIDIA-NEXT: [[DOTPREVIOUS_UB__ADDR:%.*]] = alloca i64, align 8 +// CHECK-NVIDIA-NEXT: [[LOCALPADDING_ADDR:%.*]] = alloca ptr, align 8 +// CHECK-NVIDIA-NEXT: [[RES_ADDR:%.*]] = alloca i64, align 8 +// CHECK-NVIDIA-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4 +// CHECK-NVIDIA-NEXT: [[TMP:%.*]] = alloca i32, align 4 +// CHECK-NVIDIA-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4 +// CHECK-NVIDIA-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4 +// CHECK-NVIDIA-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4 +// CHECK-NVIDIA-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4 +// CHECK-NVIDIA-NEXT: [[PI:%.*]] = alloca i32, align 4 +// CHECK-NVIDIA-NEXT: [[HZ:%.*]] = alloca i32, align 4 +// CHECK-NVIDIA-NEXT: [[HY:%.*]] = alloca i32, align 4 +// CHECK-NVIDIA-NEXT: [[HX:%.*]] = alloca i32, align 4 +// CHECK-NVIDIA-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8 +// CHECK-NVIDIA-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8 +// CHECK-NVIDIA-NEXT: store i64 [[DOTPREVIOUS_LB_]], ptr [[DOTPREVIOUS_LB__ADDR]], align 8 +// CHECK-NVIDIA-NEXT: store i64 [[DOTPREVIOUS_UB_]], ptr [[DOTPREVIOUS_UB__ADDR]], align 8 +// CHECK-NVIDIA-NEXT: store ptr [[LOCALPADDING]], ptr [[LOCALPADDING_ADDR]], align 8 +// CHECK-NVIDIA-NEXT: store i64 [[RES]], ptr [[RES_ADDR]], align 8 +// CHECK-NVIDIA-NEXT: [[TMP0:%.*]] = load ptr, ptr [[LOCALPADDING_ADDR]], align 8 +// CHECK-NVIDIA-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4 +// CHECK-NVIDIA-NEXT: store i32 999, ptr [[DOTOMP_UB]], align 4 +// CHECK-NVIDIA-NEXT: [[TMP1:%.*]] = load i64, ptr [[DOTPREVIOUS_LB__ADDR]], align 8 +// CHECK-NVIDIA-NEXT: [[CONV:%.*]] = trunc i64 [[TMP1]] to i32 +// CHECK-NVIDIA-NEXT: [[TMP2:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 8 +// CHECK-NVIDIA-NEXT: [[CONV1:%.*]] = trunc i64 [[TMP2]] to i32 +// CHECK-NVIDIA-NEXT: store i32 [[CONV]], ptr [[DOTOMP_LB]], align 4 +// CHECK-NVIDIA-NEXT: store i32 [[CONV1]], ptr [[DOTOMP_UB]], align 4 +// CHECK-NVIDIA-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4 +// CHECK-NVIDIA-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4 +// CHECK-NVIDIA-NEXT: [[TMP3:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8 +// CHECK-NVIDIA-NEXT: [[TMP4:%.*]] = load i32, ptr [[TMP3]], align 4 +// CHECK-NVIDIA-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB3:[0-9]+]], i32 [[TMP4]], i32 33, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1) +// CHECK-NVIDIA-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4 +// CHECK-NVIDIA-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV]], align 4 +// CHECK-NVIDIA-NEXT: br label [[OMP_INNER_FOR_COND:%.*]] +// CHECK-NVIDIA: omp.inner.for.cond: +// CHECK-NVIDIA-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4 +// CHECK-NVIDIA-NEXT: [[CONV2:%.*]] = sext i32 [[TMP6]] to i64 +// CHECK-NVIDIA-NEXT: [[TMP7:%.*]] = load i64, ptr [[DOTPREVIOUS_UB__ADDR]], align 8 +// CHECK-NVIDIA-NEXT: [[CMP:%.*]] = icmp ule i64 [[CONV2]], [[TMP7]] +// CHECK-NVIDIA-NEXT: br i1 [[CMP]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]] +// CHECK-NVIDIA: omp.inner.for.body: +// CHECK-NVIDIA-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4 +// CHECK-NVIDIA-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP8]], 1 +// CHECK-NVIDIA-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]] +// CHECK-NVIDIA-NEXT: store i32 [[ADD]], ptr [[PI]], align 4 +// CHECK-NVIDIA-NEXT: store i32 0, ptr [[HZ]], align 4 +// CHECK-NVIDIA-NEXT: br label [[FOR_COND:%.*]] +// CHECK-NVIDIA: for.cond: +// CHECK-NVIDIA-NEXT: [[TMP9:%.*]] = load i32, ptr [[HZ]], align 4 +// CHECK-NVIDIA-NEXT: [[CMP3:%.*]] = icmp sle i32 [[TMP9]], 1 +// CHECK-NVIDIA-NEXT: br i1 [[CMP3]], label [[FOR_BODY:%.*]], label [[FOR_END16:%.*]] +// CHECK-NVIDIA: for.body: +// CHECK-NVIDIA-NEXT: store i32 0, ptr [[HY]], align 4 +// CHECK-NVIDIA-NEXT: br label [[FOR_COND4:%.*]] +// CHECK-NVIDIA: for.cond4: +// CHECK-NVIDIA-NEXT: [[TMP10:%.*]] = load i32, ptr [[HY]], align 4 +// CHECK-NVIDIA-NEXT: [[CMP5:%.*]] = icmp sle i32 [[TMP10]], 2 +// CHECK-NVIDIA-NEXT: br i1 [[CMP5]], label [[FOR_BODY6:%.*]], label [[FOR_END13:%.*]] +// CHECK-NVIDIA: for.body6: +// CHECK-NVIDIA-NEXT: store i32 0, ptr [[HX]], align 4 +// CHECK-NVIDIA-NEXT: br label [[FOR_COND7:%.*]] +// CHECK-NVIDIA: for.cond7: +// CHECK-NVIDIA-NEXT: [[TMP11:%.*]] = load i32, ptr [[HX]], align 4 +// CHECK-NVIDIA-NEXT: [[CMP8:%.*]] = icmp sle i32 [[TMP11]], 3 +// CHECK-NVIDIA-NEXT: br i1 [[CMP8]], label [[FOR_BODY9:%.*]], label [[FOR_END:%.*]] +// CHECK-NVIDIA: for.body9: +// CHECK-NVIDIA-NEXT: [[ARRAYDECAY:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i64 0, i64 0 +// CHECK-NVIDIA-NEXT: call void @_Z18emits_alloc_sharedPKiPi(ptr noundef [[ARRAYDECAY]], ptr noundef [[RES_ADDR]]) #[[ATTR8:[0-9]+]] +// CHECK-NVIDIA-NEXT: [[ARRAYDECAY10:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i64 0, i64 0 +// CHECK-NVIDIA-NEXT: call void @_Z26does_not_emit_alloc_sharedPKiPi(ptr noundef [[ARRAYDECAY10]], ptr noundef [[RES_ADDR]]) #[[ATTR8]] +// CHECK-NVIDIA-NEXT: br label [[FOR_INC:%.*]] +// CHECK-NVIDIA: for.inc: +// CHECK-NVIDIA-NEXT: [[TMP12:%.*]] = load i32, ptr [[HX]], align 4 +// CHECK-NVIDIA-NEXT: [[INC:%.*]] = add nsw i32 [[TMP12]], 1 +// CHECK-NVIDIA-NEXT: store i32 [[INC]], ptr [[HX]], align 4 +// CHECK-NVIDIA-NEXT: br label [[FOR_COND7]], !llvm.loop [[LOOP6:![0-9]+]] +// CHECK-NVIDIA: for.end: +// CHECK-NVIDIA-NEXT: br label [[FOR_INC11:%.*]] +// CHECK-NVIDIA: for.inc11: +// CHECK-NVIDIA-NEXT: [[TMP13:%.*]] = load i32, ptr [[HY]], align 4 +// CHECK-NVIDIA-NEXT: [[INC12:%.*]] = add nsw i32 [[TMP13]], 1 +// CHECK-NVIDIA-NEXT: store i32 [[INC12]], ptr [[HY]], align 4 +// CHECK-NVIDIA-NEXT: br label [[FOR_COND4]], !llvm.loop [[LOOP8:![0-9]+]] +// CHECK-NVIDIA: for.end13: +// CHECK-NVIDIA-NEXT: br label [[FOR_INC14:%.*]] +// CHECK-NVIDIA: for.inc14: +// CHECK-NVIDIA-NEXT: [[TMP14:%.*]] = load i32, ptr [[HZ]], align 4 +// CHECK-NVIDIA-NEXT: [[INC15:%.*]] = add nsw i32 [[TMP14]], 1 +// CHECK-NVIDIA-NEXT: store i32 [[INC15]], ptr [[HZ]], align 4 +// CHECK-NVIDIA-NEXT: br label [[FOR_COND]], !llvm.loop [[LOOP9:![0-9]+]] +// CHECK-NVIDIA: for.end16: +// CHECK-NVIDIA-NEXT: [[TMP15:%.*]] = load i32, ptr [[RES_ADDR]], align 4 +// CHECK-NVIDIA-NEXT: [[TMP16:%.*]] = load i32, ptr [[PI]], align 4 +// CHECK-NVIDIA-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP16]] to i64 +// CHECK-NVIDIA-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [1000 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]] +// CHECK-NVIDIA-NEXT: store i32 [[TMP15]], ptr [[ARRAYIDX]], align 4 +// CHECK-NVIDIA-NEXT: br label [[OMP_BODY_CONTINUE:%.*]] +// CHECK-NVIDIA: omp.body.continue: +// CHECK-NVIDIA-NEXT: br label [[OMP_INNER_FOR_INC:%.*]] +// CHECK-NVIDIA: omp.inner.for.inc: +// CHECK-NVIDIA-NEXT: [[TMP17:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4 +// CHECK-NVIDIA-NEXT: [[TMP18:%.*]] = load i32, ptr [[DOTOMP_STRIDE]], align 4 +// CHECK-NVIDIA-NEXT: [[ADD17:%.*]] = add nsw i32 [[TMP17]], [[TMP18]] +// CHECK-NVIDIA-NEXT: store i32 [[ADD17]], ptr [[DOTOMP_IV]], align 4 +// CHECK-NVIDIA-NEXT: br label [[OMP_INNER_FOR_COND]] +// CHECK-NVIDIA: omp.inner.for.end: +// CHECK-NVIDIA-NEXT: br label [[OMP_LOOP_EXIT:%.*]] +// CHECK-NVIDIA: omp.loop.exit: +// CHECK-NVIDIA-NEXT: call void @__kmpc_distribute_static_fini(ptr @[[GLOB2]], i32 [[TMP4]]) +// CHECK-NVIDIA-NEXT: ret void +// +// +// CHECK-NVIDIA-LABEL: define {{[^@]+}}@_Z18emits_alloc_sharedPKiPi +// CHECK-NVIDIA-SAME: (ptr noundef [[LOCALPADDING:%.*]], ptr noundef [[RES:%.*]]) #[[ATTR2:[0-9]+]] { +// CHECK-NVIDIA-NEXT: entry: +// CHECK-NVIDIA-NEXT: [[LOCALPADDING_ADDR:%.*]] = alloca ptr, align 8 +// CHECK-NVIDIA-NEXT: [[RES_ADDR:%.*]] = alloca ptr, align 8 +// CHECK-NVIDIA-NEXT: [[STACKPTR:%.*]] = alloca i32, align 4 +// CHECK-NVIDIA-NEXT: [[STACK:%.*]] = call align 8 ptr @__kmpc_alloc_shared(i64 256) +// CHECK-NVIDIA-NEXT: store ptr [[LOCALPADDING]], ptr [[LOCALPADDING_ADDR]], align 8 +// CHECK-NVIDIA-NEXT: store ptr [[RES]], ptr [[RES_ADDR]], align 8 +// CHECK-NVIDIA-NEXT: store i32 0, ptr [[STACKPTR]], align 4 +// CHECK-NVIDIA-NEXT: [[TMP0:%.*]] = load i32, ptr [[STACKPTR]], align 4 +// CHECK-NVIDIA-NEXT: [[INC:%.*]] = add nsw i32 [[TMP0]], 1 +// CHECK-NVIDIA-NEXT: store i32 [[INC]], ptr [[STACKPTR]], align 4 +// CHECK-NVIDIA-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP0]] to i64 +// CHECK-NVIDIA-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [64 x i32], ptr [[STACK]], i64 0, i64 [[IDXPROM]] +// CHECK-NVIDIA-NEXT: store i32 -1, ptr [[ARRAYIDX]], align 4 +// CHECK-NVIDIA-NEXT: [[TMP1:%.*]] = load ptr, ptr [[RES_ADDR]], align 8 +// CHECK-NVIDIA-NEXT: store i32 0, ptr [[TMP1]], align 4 +// CHECK-NVIDIA-NEXT: br label [[DO_BODY:%.*]] +// CHECK-NVIDIA: do.body: +// CHECK-NVIDIA-NEXT: [[TMP2:%.*]] = load ptr, ptr [[LOCALPADDING_ADDR]], align 8 +// CHECK-NVIDIA-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds i32, ptr [[TMP2]], i64 0 +// CHECK-NVIDIA-NEXT: [[TMP3:%.*]] = load i32, ptr [[ARRAYIDX1]], align 4 +// CHECK-NVIDIA-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 0 +// CHECK-NVIDIA-NEXT: br i1 [[CMP]], label [[IF_THEN:%.*]], label [[IF_END:%.*]] +// CHECK-NVIDIA: if.then: +// CHECK-NVIDIA-NEXT: [[TMP4:%.*]] = load i32, ptr [[STACKPTR]], align 4 +// CHECK-NVIDIA-NEXT: [[INC2:%.*]] = add nsw i32 [[TMP4]], 1 +// CHECK-NVIDIA-NEXT: store i32 [[INC2]], ptr [[STACKPTR]], align 4 +// CHECK-NVIDIA-NEXT: [[IDXPROM3:%.*]] = sext i32 [[TMP4]] to i64 +// CHECK-NVIDIA-NEXT: [[ARRAYIDX4:%.*]] = getelementptr inbounds [64 x i32], ptr [[STACK]], i64 0, i64 [[IDXPROM3]] +// CHECK-NVIDIA-NEXT: store i32 0, ptr [[ARRAYIDX4]], align 4 +// CHECK-NVIDIA-NEXT: br label [[IF_END]] +// CHECK-NVIDIA: if.end: +// CHECK-NVIDIA-NEXT: [[TMP5:%.*]] = load i32, ptr [[STACKPTR]], align 4 +// CHECK-NVIDIA-NEXT: [[DEC:%.*]] = add nsw i32 [[TMP5]], -1 +// CHECK-NVIDIA-NEXT: store i32 [[DEC]], ptr [[STACKPTR]], align 4 +// CHECK-NVIDIA-NEXT: [[IDXPROM5:%.*]] = sext i32 [[DEC]] to i64 +// CHECK-NVIDIA-NEXT: [[ARRAYIDX6:%.*]] = getelementptr inbounds [64 x i32], ptr [[STACK]], i64 0, i64 [[IDXPROM5]] +// CHECK-NVIDIA-NEXT: [[TMP6:%.*]] = load i32, ptr [[ARRAYIDX6]], align 4 +// CHECK-NVIDIA-NEXT: [[TMP7:%.*]] = load ptr, ptr [[RES_ADDR]], align 8 +// CHECK-NVIDIA-NEXT: store i32 [[TMP6]], ptr [[TMP7]], align 4 +// CHECK-NVIDIA-NEXT: [[ARRAYIDX7:%.*]] = getelementptr inbounds [64 x i32], ptr [[STACK]], i64 0, i64 2 +// CHECK-NVIDIA-NEXT: call void @_Z3fooPi(ptr noundef [[ARRAYIDX7]]) #[[ATTR8]] +// CHECK-NVIDIA-NEXT: br label [[DO_COND:%.*]] +// CHECK-NVIDIA: do.cond: +// CHECK-NVIDIA-NEXT: [[TMP8:%.*]] = load ptr, ptr [[RES_ADDR]], align 8 +// CHECK-NVIDIA-NEXT: [[TMP9:%.*]] = load i32, ptr [[TMP8]], align 4 +// CHECK-NVIDIA-NEXT: [[CMP8:%.*]] = icmp sgt i32 [[TMP9]], 0 +// CHECK-NVIDIA-NEXT: br i1 [[CMP8]], label [[DO_BODY]], label [[DO_END:%.*]], !llvm.loop [[LOOP10:![0-9]+]] +// CHECK-NVIDIA: do.end: +// CHECK-NVIDIA-NEXT: call void @__kmpc_free_shared(ptr [[STACK]], i64 256) +// CHECK-NVIDIA-NEXT: ret void +// +// +// CHECK-NVIDIA-LABEL: define {{[^@]+}}@_Z26does_not_emit_alloc_sharedPKiPi +// CHECK-NVIDIA-SAME: (ptr noundef [[LOCALPADDING:%.*]], ptr noundef [[RES:%.*]]) #[[ATTR2]] { +// CHECK-NVIDIA-NEXT: entry: +// CHECK-NVIDIA-NEXT: [[LOCALPADDING_ADDR:%.*]] = alloca ptr, align 8 +// CHECK-NVIDIA-NEXT: [[RES_ADDR:%.*]] = alloca ptr, align 8 +// CHECK-NVIDIA-NEXT: [[STACK:%.*]] = alloca [64 x i32], align 4 +// CHECK-NVIDIA-NEXT: [[STACKPTR:%.*]] = alloca i32, align 4 +// CHECK-NVIDIA-NEXT: store ptr [[LOCALPADDING]], ptr [[LOCALPADDING_ADDR]], align 8 +// CHECK-NVIDIA-NEXT: store ptr [[RES]], ptr [[RES_ADDR]], align 8 +// CHECK-NVIDIA-NEXT: store i32 0, ptr [[STACKPTR]], align 4 +// CHECK-NVIDIA-NEXT: [[TMP0:%.*]] = load i32, ptr [[STACKPTR]], align 4 +// CHECK-NVIDIA-NEXT: [[INC:%.*]] = add nsw i32 [[TMP0]], 1 +// CHECK-NVIDIA-NEXT: store i32 [[INC]], ptr [[STACKPTR]], align 4 +// CHECK-NVIDIA-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP0]] to i64 +// CHECK-NVIDIA-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [64 x i32], ptr [[STACK]], i64 0, i64 [[IDXPROM]] +// CHECK-NVIDIA-NEXT: store i32 -1, ptr [[ARRAYIDX]], align 4 +// CHECK-NVIDIA-NEXT: [[TMP1:%.*]] = load ptr, ptr [[RES_ADDR]], align 8 +// CHECK-NVIDIA-NEXT: store i32 0, ptr [[TMP1]], align 4 +// CHECK-NVIDIA-NEXT: br label [[DO_BODY:%.*]] +// CHECK-NVIDIA: do.body: +// CHECK-NVIDIA-NEXT: [[TMP2:%.*]] = load ptr, ptr [[LOCALPADDING_ADDR]], align 8 +// CHECK-NVIDIA-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds i32, ptr [[TMP2]], i64 0 +// CHECK-NVIDIA-NEXT: [[TMP3:%.*]] = load i32, ptr [[ARRAYIDX1]], align 4 +// CHECK-NVIDIA-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 0 +// CHECK-NVIDIA-NEXT: br i1 [[CMP]], label [[IF_THEN:%.*]], label [[IF_END:%.*]] +// CHECK-NVIDIA: if.then: +// CHECK-NVIDIA-NEXT: [[TMP4:%.*]] = load i32, ptr [[STACKPTR]], align 4 +// CHECK-NVIDIA-NEXT: [[INC2:%.*]] = add nsw i32 [[TMP4]], 1 +// CHECK-NVIDIA-NEXT: store i32 [[INC2]], ptr [[STACKPTR]], align 4 +// CHECK-NVIDIA-NEXT: [[IDXPROM3:%.*]] = sext i32 [[TMP4]] to i64 +// CHECK-NVIDIA-NEXT: [[ARRAYIDX4:%.*]] = getelementptr inbounds [64 x i32], ptr [[STACK]], i64 0, i64 [[IDXPROM3]] +// CHECK-NVIDIA-NEXT: store i32 0, ptr [[ARRAYIDX4]], align 4 +// CHECK-NVIDIA-NEXT: br label [[IF_END]] +// CHECK-NVIDIA: if.end: +// CHECK-NVIDIA-NEXT: [[TMP5:%.*]] = load i32, ptr [[STACKPTR]], align 4 +// CHECK-NVIDIA-NEXT: [[DEC:%.*]] = add nsw i32 [[TMP5]], -1 +// CHECK-NVIDIA-NEXT: store i32 [[DEC]], ptr [[STACKPTR]], align 4 +// CHECK-NVIDIA-NEXT: [[IDXPROM5:%.*]] = sext i32 [[DEC]] to i64 +// CHECK-NVIDIA-NEXT: [[ARRAYIDX6:%.*]] = getelementptr inbounds [64 x i32], ptr [[STACK]], i64 0, i64 [[IDXPROM5]] +// CHECK-NVIDIA-NEXT: [[TMP6:%.*]] = load i32, ptr [[ARRAYIDX6]], align 4 +// CHECK-NVIDIA-NEXT: [[TMP7:%.*]] = load ptr, ptr [[RES_ADDR]], align 8 +// CHECK-NVIDIA-NEXT: store i32 [[TMP6]], ptr [[TMP7]], align 4 +// CHECK-NVIDIA-NEXT: br label [[DO_COND:%.*]] +// CHECK-NVIDIA: do.cond: +// CHECK-NVIDIA-NEXT: [[TMP8:%.*]] = load ptr, ptr [[RES_ADDR]], align 8 +// CHECK-NVIDIA-NEXT: [[TMP9:%.*]] = load i32, ptr [[TMP8]], align 4 +// CHECK-NVIDIA-NEXT: [[CMP7:%.*]] = icmp sgt i32 [[TMP9]], 0 +// CHECK-NVIDIA-NEXT: br i1 [[CMP7]], label [[DO_BODY]], label [[DO_END:%.*]], !llvm.loop [[LOOP11:![0-9]+]] +// CHECK-NVIDIA: do.end: +// CHECK-NVIDIA-NEXT: ret void +//