391 lines · cpp
1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _2// REQUIRES: amdgpu-registered-target3 4// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-ppc-host.bc5// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple amdgcn-amd-amdhsa -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=IR-GPU6 7// RUN: %clang_cc1 -verify -triple x86_64-pc-linux-gnu -fopenmp -emit-llvm %s -o - | FileCheck %s --check-prefix=IR8 9// Check same results after serialization round-trip10// RUN: %clang_cc1 -verify -triple x86_64-pc-linux-gnu -fopenmp -emit-pch -o %t %s11// RUN: %clang_cc1 -verify -triple x86_64-pc-linux-gnu -fopenmp -include-pch %t -emit-llvm %s -o - | FileCheck %s --check-prefix=IR-PCH12 13// expected-no-diagnostics14 15#ifndef HEADER16#define HEADER17 18typedef void **omp_allocator_handle_t;19extern const omp_allocator_handle_t omp_null_allocator;20extern const omp_allocator_handle_t omp_default_mem_alloc;21extern const omp_allocator_handle_t omp_large_cap_mem_alloc;22extern const omp_allocator_handle_t omp_const_mem_alloc;23extern const omp_allocator_handle_t omp_high_bw_mem_alloc;24extern const omp_allocator_handle_t omp_low_lat_mem_alloc;25extern const omp_allocator_handle_t omp_cgroup_mem_alloc;26extern const omp_allocator_handle_t omp_pteam_mem_alloc;27extern const omp_allocator_handle_t omp_thread_mem_alloc;28 29extern int omp_get_thread_num(void);30 31#define N 6432 33int main() {34 int x = 0;35 int device_result[N] = {0};36 37 #pragma omp target parallel loop num_threads(N) uses_allocators(omp_pteam_mem_alloc) allocate(omp_pteam_mem_alloc: x) private(x) map(from: device_result)38 for (int i = 0; i < N; i++) {39 x = omp_get_thread_num();40 device_result[i] = i + x;41 }42}43#endif44// IR-GPU-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l3745// IR-GPU-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], ptr noundef nonnull align 4 dereferenceable(256) [[DEVICE_RESULT:%.*]], ptr noundef [[OMP_PTEAM_MEM_ALLOC:%.*]]) #[[ATTR0:[0-9]+]] {46// IR-GPU-NEXT: entry:47// IR-GPU-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)48// IR-GPU-NEXT: [[DEVICE_RESULT_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)49// IR-GPU-NEXT: [[OMP_PTEAM_MEM_ALLOC_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)50// IR-GPU-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [2 x ptr], align 8, addrspace(5)51// IR-GPU-NEXT: [[DYN_PTR_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DYN_PTR_ADDR]] to ptr52// IR-GPU-NEXT: [[DEVICE_RESULT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DEVICE_RESULT_ADDR]] to ptr53// IR-GPU-NEXT: [[OMP_PTEAM_MEM_ALLOC_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OMP_PTEAM_MEM_ALLOC_ADDR]] to ptr54// IR-GPU-NEXT: [[CAPTURED_VARS_ADDRS_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[CAPTURED_VARS_ADDRS]] to ptr55// IR-GPU-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR_ASCAST]], align 856// IR-GPU-NEXT: store ptr [[DEVICE_RESULT]], ptr [[DEVICE_RESULT_ADDR_ASCAST]], align 857// IR-GPU-NEXT: store ptr [[OMP_PTEAM_MEM_ALLOC]], ptr [[OMP_PTEAM_MEM_ALLOC_ADDR_ASCAST]], align 858// IR-GPU-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DEVICE_RESULT_ADDR_ASCAST]], align 859// IR-GPU-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr addrspacecast (ptr addrspace(1) @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l37_kernel_environment to ptr), ptr [[DYN_PTR]])60// IR-GPU-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -161// IR-GPU-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]62// IR-GPU: user_code.entry:63// IR-GPU-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr addrspacecast (ptr addrspace(1) @[[GLOB1:[0-9]+]] to ptr))64// IR-GPU-NEXT: [[TMP3:%.*]] = load ptr, ptr [[OMP_PTEAM_MEM_ALLOC_ADDR_ASCAST]], align 865// IR-GPU-NEXT: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[CAPTURED_VARS_ADDRS_ASCAST]], i64 0, i64 066// IR-GPU-NEXT: store ptr [[TMP0]], ptr [[TMP4]], align 867// IR-GPU-NEXT: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[CAPTURED_VARS_ADDRS_ASCAST]], i64 0, i64 168// IR-GPU-NEXT: store ptr [[TMP3]], ptr [[TMP5]], align 869// IR-GPU-NEXT: call void @__kmpc_parallel_51(ptr addrspacecast (ptr addrspace(1) @[[GLOB1]] to ptr), i32 [[TMP2]], i32 1, i32 64, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l37_omp_outlined, ptr null, ptr [[CAPTURED_VARS_ADDRS_ASCAST]], i64 2)70// IR-GPU-NEXT: call void @__kmpc_target_deinit()71// IR-GPU-NEXT: ret void72// IR-GPU: worker.exit:73// IR-GPU-NEXT: ret void74//75//76// IR-GPU-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l37_omp_outlined77// IR-GPU-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(256) [[DEVICE_RESULT:%.*]], ptr noundef [[OMP_PTEAM_MEM_ALLOC:%.*]]) #[[ATTR1:[0-9]+]] {78// IR-GPU-NEXT: entry:79// IR-GPU-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8, addrspace(5)80// IR-GPU-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8, addrspace(5)81// IR-GPU-NEXT: [[DEVICE_RESULT_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)82// IR-GPU-NEXT: [[OMP_PTEAM_MEM_ALLOC_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)83// IR-GPU-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4, addrspace(5)84// IR-GPU-NEXT: [[TMP:%.*]] = alloca i32, align 4, addrspace(5)85// IR-GPU-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4, addrspace(5)86// IR-GPU-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4, addrspace(5)87// IR-GPU-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4, addrspace(5)88// IR-GPU-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4, addrspace(5)89// IR-GPU-NEXT: [[I:%.*]] = alloca i32, align 4, addrspace(5)90// IR-GPU-NEXT: [[DOTGLOBAL_TID__ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTGLOBAL_TID__ADDR]] to ptr91// IR-GPU-NEXT: [[DOTBOUND_TID__ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTBOUND_TID__ADDR]] to ptr92// IR-GPU-NEXT: [[DEVICE_RESULT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DEVICE_RESULT_ADDR]] to ptr93// IR-GPU-NEXT: [[OMP_PTEAM_MEM_ALLOC_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OMP_PTEAM_MEM_ALLOC_ADDR]] to ptr94// IR-GPU-NEXT: [[DOTOMP_IV_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_IV]] to ptr95// IR-GPU-NEXT: [[TMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[TMP]] to ptr96// IR-GPU-NEXT: [[DOTOMP_LB_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_LB]] to ptr97// IR-GPU-NEXT: [[DOTOMP_UB_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_UB]] to ptr98// IR-GPU-NEXT: [[DOTOMP_STRIDE_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_STRIDE]] to ptr99// IR-GPU-NEXT: [[DOTOMP_IS_LAST_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_IS_LAST]] to ptr100// IR-GPU-NEXT: [[I_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[I]] to ptr101// IR-GPU-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR_ASCAST]], align 8102// IR-GPU-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR_ASCAST]], align 8103// IR-GPU-NEXT: store ptr [[DEVICE_RESULT]], ptr [[DEVICE_RESULT_ADDR_ASCAST]], align 8104// IR-GPU-NEXT: store ptr [[OMP_PTEAM_MEM_ALLOC]], ptr [[OMP_PTEAM_MEM_ALLOC_ADDR_ASCAST]], align 8105// IR-GPU-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DEVICE_RESULT_ADDR_ASCAST]], align 8106// IR-GPU-NEXT: store i32 0, ptr [[DOTOMP_LB_ASCAST]], align 4107// IR-GPU-NEXT: store i32 63, ptr [[DOTOMP_UB_ASCAST]], align 4108// IR-GPU-NEXT: store i32 1, ptr [[DOTOMP_STRIDE_ASCAST]], align 4109// IR-GPU-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST_ASCAST]], align 4110// IR-GPU-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR_ASCAST]], align 8111// IR-GPU-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4112// IR-GPU-NEXT: call void @__kmpc_for_static_init_4(ptr addrspacecast (ptr addrspace(1) @[[GLOB2:[0-9]+]] to ptr), i32 [[TMP2]], i32 33, ptr [[DOTOMP_IS_LAST_ASCAST]], ptr [[DOTOMP_LB_ASCAST]], ptr [[DOTOMP_UB_ASCAST]], ptr [[DOTOMP_STRIDE_ASCAST]], i32 1, i32 1)113// IR-GPU-NEXT: br label [[OMP_DISPATCH_COND:%.*]]114// IR-GPU: omp.dispatch.cond:115// IR-GPU-NEXT: [[TMP3:%.*]] = load i32, ptr [[DOTOMP_UB_ASCAST]], align 4116// IR-GPU-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP3]], 63117// IR-GPU-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]118// IR-GPU: cond.true:119// IR-GPU-NEXT: br label [[COND_END:%.*]]120// IR-GPU: cond.false:121// IR-GPU-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB_ASCAST]], align 4122// IR-GPU-NEXT: br label [[COND_END]]123// IR-GPU: cond.end:124// IR-GPU-NEXT: [[COND:%.*]] = phi i32 [ 63, [[COND_TRUE]] ], [ [[TMP4]], [[COND_FALSE]] ]125// IR-GPU-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB_ASCAST]], align 4126// IR-GPU-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_LB_ASCAST]], align 4127// IR-GPU-NEXT: store i32 [[TMP5]], ptr [[DOTOMP_IV_ASCAST]], align 4128// IR-GPU-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 4129// IR-GPU-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB_ASCAST]], align 4130// IR-GPU-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP6]], [[TMP7]]131// IR-GPU-NEXT: br i1 [[CMP1]], label [[OMP_DISPATCH_BODY:%.*]], label [[OMP_DISPATCH_END:%.*]]132// IR-GPU: omp.dispatch.body:133// IR-GPU-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]134// IR-GPU: omp.inner.for.cond:135// IR-GPU-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 4136// IR-GPU-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_UB_ASCAST]], align 4137// IR-GPU-NEXT: [[CMP2:%.*]] = icmp sle i32 [[TMP8]], [[TMP9]]138// IR-GPU-NEXT: br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]139// IR-GPU: omp.inner.for.body:140// IR-GPU-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 4141// IR-GPU-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP10]], 1142// IR-GPU-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]143// IR-GPU-NEXT: store i32 [[ADD]], ptr [[I_ASCAST]], align 4144// IR-GPU-NEXT: [[CALL:%.*]] = call noundef i32 @_Z18omp_get_thread_numv() #[[ATTR5:[0-9]+]]145// IR-GPU-NEXT: store i32 [[CALL]], ptr addrspacecast (ptr addrspace(3) @x to ptr), align 4146// IR-GPU-NEXT: [[TMP11:%.*]] = load i32, ptr [[I_ASCAST]], align 4147// IR-GPU-NEXT: [[TMP12:%.*]] = load i32, ptr addrspacecast (ptr addrspace(3) @x to ptr), align 4148// IR-GPU-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP11]], [[TMP12]]149// IR-GPU-NEXT: [[TMP13:%.*]] = load i32, ptr [[I_ASCAST]], align 4150// IR-GPU-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP13]] to i64151// IR-GPU-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [64 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]152// IR-GPU-NEXT: store i32 [[ADD3]], ptr [[ARRAYIDX]], align 4153// IR-GPU-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]154// IR-GPU: omp.body.continue:155// IR-GPU-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]156// IR-GPU: omp.inner.for.inc:157// IR-GPU-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 4158// IR-GPU-NEXT: [[ADD4:%.*]] = add nsw i32 [[TMP14]], 1159// IR-GPU-NEXT: store i32 [[ADD4]], ptr [[DOTOMP_IV_ASCAST]], align 4160// IR-GPU-NEXT: br label [[OMP_INNER_FOR_COND]]161// IR-GPU: omp.inner.for.end:162// IR-GPU-NEXT: br label [[OMP_DISPATCH_INC:%.*]]163// IR-GPU: omp.dispatch.inc:164// IR-GPU-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_LB_ASCAST]], align 4165// IR-GPU-NEXT: [[TMP16:%.*]] = load i32, ptr [[DOTOMP_STRIDE_ASCAST]], align 4166// IR-GPU-NEXT: [[ADD5:%.*]] = add nsw i32 [[TMP15]], [[TMP16]]167// IR-GPU-NEXT: store i32 [[ADD5]], ptr [[DOTOMP_LB_ASCAST]], align 4168// IR-GPU-NEXT: [[TMP17:%.*]] = load i32, ptr [[DOTOMP_UB_ASCAST]], align 4169// IR-GPU-NEXT: [[TMP18:%.*]] = load i32, ptr [[DOTOMP_STRIDE_ASCAST]], align 4170// IR-GPU-NEXT: [[ADD6:%.*]] = add nsw i32 [[TMP17]], [[TMP18]]171// IR-GPU-NEXT: store i32 [[ADD6]], ptr [[DOTOMP_UB_ASCAST]], align 4172// IR-GPU-NEXT: br label [[OMP_DISPATCH_COND]]173// IR-GPU: omp.dispatch.end:174// IR-GPU-NEXT: call void @__kmpc_for_static_fini(ptr addrspacecast (ptr addrspace(1) @[[GLOB2]] to ptr), i32 [[TMP2]])175// IR-GPU-NEXT: ret void176//177//178// IR-LABEL: define {{[^@]+}}@main179// IR-SAME: () #[[ATTR0:[0-9]+]] {180// IR-NEXT: entry:181// IR-NEXT: [[X:%.*]] = alloca i32, align 4182// IR-NEXT: [[DEVICE_RESULT:%.*]] = alloca [64 x i32], align 16183// IR-NEXT: store i32 0, ptr [[X]], align 4184// IR-NEXT: call void @llvm.memset.p0.i64(ptr align 16 [[DEVICE_RESULT]], i8 0, i64 256, i1 false)185// IR-NEXT: [[TMP0:%.*]] = load ptr, ptr @omp_pteam_mem_alloc, align 8186// IR-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l37(ptr [[DEVICE_RESULT]], ptr [[TMP0]]) #[[ATTR3:[0-9]+]]187// IR-NEXT: ret i32 0188//189//190// IR-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l37191// IR-SAME: (ptr noundef nonnull align 4 dereferenceable(256) [[DEVICE_RESULT:%.*]], ptr noundef [[OMP_PTEAM_MEM_ALLOC:%.*]]) #[[ATTR2:[0-9]+]] {192// IR-NEXT: entry:193// IR-NEXT: [[DEVICE_RESULT_ADDR:%.*]] = alloca ptr, align 8194// IR-NEXT: [[OMP_PTEAM_MEM_ALLOC_ADDR:%.*]] = alloca ptr, align 8195// IR-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB2:[0-9]+]])196// IR-NEXT: store ptr [[DEVICE_RESULT]], ptr [[DEVICE_RESULT_ADDR]], align 8197// IR-NEXT: store ptr [[OMP_PTEAM_MEM_ALLOC]], ptr [[OMP_PTEAM_MEM_ALLOC_ADDR]], align 8198// IR-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DEVICE_RESULT_ADDR]], align 8199// IR-NEXT: call void @__kmpc_push_num_threads(ptr @[[GLOB2]], i32 [[TMP0]], i32 64)200// IR-NEXT: [[TMP2:%.*]] = load ptr, ptr [[OMP_PTEAM_MEM_ALLOC_ADDR]], align 8201// IR-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB2]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l37.omp_outlined, ptr [[TMP1]], ptr [[TMP2]])202// IR-NEXT: ret void203//204//205// IR-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l37.omp_outlined206// IR-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(256) [[DEVICE_RESULT:%.*]], ptr noundef [[OMP_PTEAM_MEM_ALLOC:%.*]]) #[[ATTR2]] {207// IR-NEXT: entry:208// IR-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8209// IR-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8210// IR-NEXT: [[DEVICE_RESULT_ADDR:%.*]] = alloca ptr, align 8211// IR-NEXT: [[OMP_PTEAM_MEM_ALLOC_ADDR:%.*]] = alloca ptr, align 8212// IR-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4213// IR-NEXT: [[TMP:%.*]] = alloca i32, align 4214// IR-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4215// IR-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4216// IR-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4217// IR-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4218// IR-NEXT: [[I:%.*]] = alloca i32, align 4219// IR-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8220// IR-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8221// IR-NEXT: store ptr [[DEVICE_RESULT]], ptr [[DEVICE_RESULT_ADDR]], align 8222// IR-NEXT: store ptr [[OMP_PTEAM_MEM_ALLOC]], ptr [[OMP_PTEAM_MEM_ALLOC_ADDR]], align 8223// IR-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DEVICE_RESULT_ADDR]], align 8224// IR-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4225// IR-NEXT: store i32 63, ptr [[DOTOMP_UB]], align 4226// IR-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4227// IR-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4228// IR-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8229// IR-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4230// IR-NEXT: [[TMP3:%.*]] = load ptr, ptr @omp_pteam_mem_alloc, align 8231// IR-NEXT: [[DOTX__VOID_ADDR:%.*]] = call ptr @__kmpc_alloc(i32 [[TMP2]], i64 4, ptr [[TMP3]])232// IR-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP2]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)233// IR-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4234// IR-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP4]], 63235// IR-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]236// IR: cond.true:237// IR-NEXT: br label [[COND_END:%.*]]238// IR: cond.false:239// IR-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4240// IR-NEXT: br label [[COND_END]]241// IR: cond.end:242// IR-NEXT: [[COND:%.*]] = phi i32 [ 63, [[COND_TRUE]] ], [ [[TMP5]], [[COND_FALSE]] ]243// IR-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 4244// IR-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4245// IR-NEXT: store i32 [[TMP6]], ptr [[DOTOMP_IV]], align 4246// IR-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]247// IR: omp.inner.for.cond:248// IR-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4249// IR-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4250// IR-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP7]], [[TMP8]]251// IR-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_COND_CLEANUP:%.*]]252// IR: omp.inner.for.cond.cleanup:253// IR-NEXT: br label [[OMP_INNER_FOR_END:%.*]]254// IR: omp.inner.for.body:255// IR-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4256// IR-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP9]], 1257// IR-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]258// IR-NEXT: store i32 [[ADD]], ptr [[I]], align 4259// IR-NEXT: [[CALL:%.*]] = call noundef i32 @_Z18omp_get_thread_numv()260// IR-NEXT: store i32 [[CALL]], ptr [[DOTX__VOID_ADDR]], align 4261// IR-NEXT: [[TMP10:%.*]] = load i32, ptr [[I]], align 4262// IR-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTX__VOID_ADDR]], align 4263// IR-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP10]], [[TMP11]]264// IR-NEXT: [[TMP12:%.*]] = load i32, ptr [[I]], align 4265// IR-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP12]] to i64266// IR-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [64 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]267// IR-NEXT: store i32 [[ADD2]], ptr [[ARRAYIDX]], align 4268// IR-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]269// IR: omp.body.continue:270// IR-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]271// IR: omp.inner.for.inc:272// IR-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4273// IR-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP13]], 1274// IR-NEXT: store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4275// IR-NEXT: br label [[OMP_INNER_FOR_COND]]276// IR: omp.inner.for.end:277// IR-NEXT: br label [[OMP_LOOP_EXIT:%.*]]278// IR: omp.loop.exit:279// IR-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])280// IR-NEXT: [[TMP14:%.*]] = load ptr, ptr @omp_pteam_mem_alloc, align 8281// IR-NEXT: call void @__kmpc_free(i32 [[TMP2]], ptr [[DOTX__VOID_ADDR]], ptr [[TMP14]])282// IR-NEXT: ret void283//284//285// IR-PCH-LABEL: define {{[^@]+}}@main286// IR-PCH-SAME: () #[[ATTR0:[0-9]+]] {287// IR-PCH-NEXT: entry:288// IR-PCH-NEXT: [[X:%.*]] = alloca i32, align 4289// IR-PCH-NEXT: [[DEVICE_RESULT:%.*]] = alloca [64 x i32], align 16290// IR-PCH-NEXT: store i32 0, ptr [[X]], align 4291// IR-PCH-NEXT: call void @llvm.memset.p0.i64(ptr align 16 [[DEVICE_RESULT]], i8 0, i64 256, i1 false)292// IR-PCH-NEXT: [[TMP0:%.*]] = load ptr, ptr @omp_pteam_mem_alloc, align 8293// IR-PCH-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l37(ptr [[DEVICE_RESULT]], ptr [[TMP0]]) #[[ATTR3:[0-9]+]]294// IR-PCH-NEXT: ret i32 0295//296//297// IR-PCH-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l37298// IR-PCH-SAME: (ptr noundef nonnull align 4 dereferenceable(256) [[DEVICE_RESULT:%.*]], ptr noundef [[OMP_PTEAM_MEM_ALLOC:%.*]]) #[[ATTR2:[0-9]+]] {299// IR-PCH-NEXT: entry:300// IR-PCH-NEXT: [[DEVICE_RESULT_ADDR:%.*]] = alloca ptr, align 8301// IR-PCH-NEXT: [[OMP_PTEAM_MEM_ALLOC_ADDR:%.*]] = alloca ptr, align 8302// IR-PCH-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB2:[0-9]+]])303// IR-PCH-NEXT: store ptr [[DEVICE_RESULT]], ptr [[DEVICE_RESULT_ADDR]], align 8304// IR-PCH-NEXT: store ptr [[OMP_PTEAM_MEM_ALLOC]], ptr [[OMP_PTEAM_MEM_ALLOC_ADDR]], align 8305// IR-PCH-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DEVICE_RESULT_ADDR]], align 8306// IR-PCH-NEXT: call void @__kmpc_push_num_threads(ptr @[[GLOB2]], i32 [[TMP0]], i32 64)307// IR-PCH-NEXT: [[TMP2:%.*]] = load ptr, ptr [[OMP_PTEAM_MEM_ALLOC_ADDR]], align 8308// IR-PCH-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB2]], i32 2, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l37.omp_outlined, ptr [[TMP1]], ptr [[TMP2]])309// IR-PCH-NEXT: ret void310//311//312// IR-PCH-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l37.omp_outlined313// IR-PCH-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(256) [[DEVICE_RESULT:%.*]], ptr noundef [[OMP_PTEAM_MEM_ALLOC:%.*]]) #[[ATTR2]] {314// IR-PCH-NEXT: entry:315// IR-PCH-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8316// IR-PCH-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8317// IR-PCH-NEXT: [[DEVICE_RESULT_ADDR:%.*]] = alloca ptr, align 8318// IR-PCH-NEXT: [[OMP_PTEAM_MEM_ALLOC_ADDR:%.*]] = alloca ptr, align 8319// IR-PCH-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4320// IR-PCH-NEXT: [[TMP:%.*]] = alloca i32, align 4321// IR-PCH-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4322// IR-PCH-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4323// IR-PCH-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4324// IR-PCH-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4325// IR-PCH-NEXT: [[I:%.*]] = alloca i32, align 4326// IR-PCH-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8327// IR-PCH-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8328// IR-PCH-NEXT: store ptr [[DEVICE_RESULT]], ptr [[DEVICE_RESULT_ADDR]], align 8329// IR-PCH-NEXT: store ptr [[OMP_PTEAM_MEM_ALLOC]], ptr [[OMP_PTEAM_MEM_ALLOC_ADDR]], align 8330// IR-PCH-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DEVICE_RESULT_ADDR]], align 8331// IR-PCH-NEXT: store i32 0, ptr [[DOTOMP_LB]], align 4332// IR-PCH-NEXT: store i32 63, ptr [[DOTOMP_UB]], align 4333// IR-PCH-NEXT: store i32 1, ptr [[DOTOMP_STRIDE]], align 4334// IR-PCH-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST]], align 4335// IR-PCH-NEXT: [[TMP1:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8336// IR-PCH-NEXT: [[TMP2:%.*]] = load i32, ptr [[TMP1]], align 4337// IR-PCH-NEXT: [[TMP3:%.*]] = load ptr, ptr @omp_pteam_mem_alloc, align 8338// IR-PCH-NEXT: [[DOTX__VOID_ADDR:%.*]] = call ptr @__kmpc_alloc(i32 [[TMP2]], i64 4, ptr [[TMP3]])339// IR-PCH-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP2]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)340// IR-PCH-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4341// IR-PCH-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP4]], 63342// IR-PCH-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]343// IR-PCH: cond.true:344// IR-PCH-NEXT: br label [[COND_END:%.*]]345// IR-PCH: cond.false:346// IR-PCH-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4347// IR-PCH-NEXT: br label [[COND_END]]348// IR-PCH: cond.end:349// IR-PCH-NEXT: [[COND:%.*]] = phi i32 [ 63, [[COND_TRUE]] ], [ [[TMP5]], [[COND_FALSE]] ]350// IR-PCH-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB]], align 4351// IR-PCH-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_LB]], align 4352// IR-PCH-NEXT: store i32 [[TMP6]], ptr [[DOTOMP_IV]], align 4353// IR-PCH-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]354// IR-PCH: omp.inner.for.cond:355// IR-PCH-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4356// IR-PCH-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB]], align 4357// IR-PCH-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP7]], [[TMP8]]358// IR-PCH-NEXT: br i1 [[CMP1]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_COND_CLEANUP:%.*]]359// IR-PCH: omp.inner.for.cond.cleanup:360// IR-PCH-NEXT: br label [[OMP_INNER_FOR_END:%.*]]361// IR-PCH: omp.inner.for.body:362// IR-PCH-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4363// IR-PCH-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP9]], 1364// IR-PCH-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]365// IR-PCH-NEXT: store i32 [[ADD]], ptr [[I]], align 4366// IR-PCH-NEXT: [[CALL:%.*]] = call noundef i32 @_Z18omp_get_thread_numv()367// IR-PCH-NEXT: store i32 [[CALL]], ptr [[DOTX__VOID_ADDR]], align 4368// IR-PCH-NEXT: [[TMP10:%.*]] = load i32, ptr [[I]], align 4369// IR-PCH-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTX__VOID_ADDR]], align 4370// IR-PCH-NEXT: [[ADD2:%.*]] = add nsw i32 [[TMP10]], [[TMP11]]371// IR-PCH-NEXT: [[TMP12:%.*]] = load i32, ptr [[I]], align 4372// IR-PCH-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP12]] to i64373// IR-PCH-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [64 x i32], ptr [[TMP0]], i64 0, i64 [[IDXPROM]]374// IR-PCH-NEXT: store i32 [[ADD2]], ptr [[ARRAYIDX]], align 4375// IR-PCH-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]376// IR-PCH: omp.body.continue:377// IR-PCH-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]378// IR-PCH: omp.inner.for.inc:379// IR-PCH-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV]], align 4380// IR-PCH-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP13]], 1381// IR-PCH-NEXT: store i32 [[ADD3]], ptr [[DOTOMP_IV]], align 4382// IR-PCH-NEXT: br label [[OMP_INNER_FOR_COND]]383// IR-PCH: omp.inner.for.end:384// IR-PCH-NEXT: br label [[OMP_LOOP_EXIT:%.*]]385// IR-PCH: omp.loop.exit:386// IR-PCH-NEXT: call void @__kmpc_for_static_fini(ptr @[[GLOB1]], i32 [[TMP2]])387// IR-PCH-NEXT: [[TMP14:%.*]] = load ptr, ptr @omp_pteam_mem_alloc, align 8388// IR-PCH-NEXT: call void @__kmpc_free(i32 [[TMP2]], ptr [[DOTX__VOID_ADDR]], ptr [[TMP14]])389// IR-PCH-NEXT: ret void390//391