This is an archive of the discontinued LLVM Phabricator instance.

Differential D138614

[Clang][OpenMP][AMDGPU] Fix capture of variably modified type alias in teams distribute
Needs ReviewPublic

Authored by doru1004 on Nov 23 2022, 2:40 PM.

Details

Reviewers
carlo.bertolli
ronl
ABataev
jhuber6
jdoerfert
Summary

This patch fixes an issue with failing to properly capture a variably modified type alias of an already mapped flat array. The test showcases a minimal example of the failing scenario.

Diff Detail

Event Timeline

doru1004 created this revision.Nov 23 2022, 2:40 PM
Herald added a project: Restricted Project. · View Herald TranscriptNov 23 2022, 2:40 PM
Herald added subscribers: kosarev, jeroen.dobbelaere, guansong and 4 others. · View Herald Transcript
doru1004 requested review of this revision.Nov 23 2022, 2:40 PM
Herald added a reviewer: jdoerfert. · View Herald TranscriptNov 23 2022, 2:40 PM
Herald added a project: Restricted Project. · View Herald Transcript
Herald added subscribers: cfe-commits, sstefan1, wdng. · View Herald Transcript
Harbormaster completed remote builds in B199294: Diff 477610.Nov 23 2022, 4:13 PM
carlo.bertolli added a comment.Nov 29 2022, 9:38 AM

LGTM but @ABataev should review/approve this

ABataev added inline comments.Nov 29 2022, 9:52 AM
clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
243–244

Is it for pointers only? Or for other types too?

doru1004 added inline comments.Nov 29 2022, 10:29 AM
clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
243–244

I am not sure about that distinction; it allows for the same types as before except that when the directive is not a combined parallel for directive it returns immediately because the variable will then have to be captured and shared.

doru1004 added inline comments.Nov 29 2022, 10:34 AM
clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
243–244

Correction: it allows for VMTs to pass through in cases in which you don't use a combined parallel for directive.

ABataev added inline comments.Nov 29 2022, 10:49 AM
clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
243–244

IsVMTTy?

243–244

How can it be captured by value? Why does it happen? My first question relates to the assertion you changed. I assume, it applies only to the pointers. Otherwise, please add more context how we can capture mariably modifiable type by value.

doru1004 added inline comments.Nov 29 2022, 12:26 PM
clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
243–244

Yes it should be called IsVMTTy

doru1004 updated this revision to Diff 478693.Nov 29 2022, 12:54 PM
ABataev added inline comments.Nov 29 2022, 1:19 PM
clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
243–244

Still naming, IsVMTTy.

doru1004 updated this revision to Diff 478707.Nov 29 2022, 1:25 PM
Harbormaster completed remote builds in B200103: Diff 478707.Nov 29 2022, 4:19 PM
doru1004 added inline comments.Nov 30 2022, 7:13 AM
clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
243–244

It only applies to pointers, it's the pointer that needs to be used correctly.

doru1004 marked an inline comment as done.Dec 5 2022, 7:04 AM
doru1004 added a comment.Dec 7 2022, 7:45 AM

ping

doru1004 added a comment.Dec 13 2022, 6:02 AM

ping

ABataev added inline comments.Dec 13 2022, 6:12 AM
clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
243–244

Can you change the assert check then to something like this:

assert((VD->getType()->isPointerType() || !VD->getType()->isVariablyModifiedType()) &&
246–249

Also, can you make it specific to pinter types?

Revision Contents

PathSize
clang/
lib/
CodeGen/
8 lines
test/
OpenMP/
308 lines

Diff 478707

clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp

Show First 20 Lines • Show All 234 Lines • ▼ Show 20 Lines if (auto *CSI = CGF.CapturedStmtInfo) {
return; return;
if (((Attr->getCaptureKind() != OMPC_map) && if (((Attr->getCaptureKind() != OMPC_map) &&
!isOpenMPPrivate(Attr->getCaptureKind())) || !isOpenMPPrivate(Attr->getCaptureKind())) ||
((Attr->getCaptureKind() == OMPC_map) && ((Attr->getCaptureKind() == OMPC_map) &&
!FD->getType()->isAnyPointerType())) !FD->getType()->isAnyPointerType()))
return; return;
} }
if (!FD->getType()->isReferenceType()) { if (!FD->getType()->isReferenceType()) {
assert(!VD->getType()->isVariablyModifiedType() && bool IsVMTTy = VD->getType()->isVariablyModifiedType();
assert((!IsVMTTy || (IsVMTTy && !IsForCombinedParallelRegion)) &&
ABataevUnsubmitted
Not Done
ReplyInline Actions

Is it for pointers only? Or for other types too?

ABataev: Is it for pointers only? Or for other types too?
doru1004AuthorUnsubmitted
Done
ReplyInline Actions

I am not sure about that distinction; it allows for the same types as before except that when the directive is not a combined parallel for directive it returns immediately because the variable will then have to be captured and shared.

doru1004: I am not sure about that distinction; it allows for the same types as before except that when…
doru1004AuthorUnsubmitted
Done
ReplyInline Actions

Correction: it allows for VMTs to pass through in cases in which you don't use a combined parallel for directive.

doru1004: Correction: it allows for VMTs to pass through in cases in which you don't use a combined…
ABataevUnsubmitted
Not Done
ReplyInline Actions

How can it be captured by value? Why does it happen? My first question relates to the assertion you changed. I assume, it applies only to the pointers. Otherwise, please add more context how we can capture mariably modifiable type by value.

ABataev: How can it be captured by value? Why does it happen? My first question relates to the assertion…
doru1004AuthorUnsubmitted
Done
ReplyInline Actions

It only applies to pointers, it's the pointer that needs to be used correctly.

doru1004: It only applies to pointers, it's the pointer that needs to be used correctly.
ABataevUnsubmitted
Not Done
ReplyInline Actions

Can you change the assert check then to something like this:

assert((VD->getType()->isPointerType() || !VD->getType()->isVariablyModifiedType()) &&
ABataev: Can you change the assert check then to something like this: ``` assert((VD->getType()…
ABataevUnsubmitted
Not Done
ReplyInline Actions

IsVMTTy?

ABataev: `IsVMTTy`?
doru1004AuthorUnsubmitted
Done
ReplyInline Actions

Yes it should be called IsVMTTy

doru1004: Yes it should be called `IsVMTTy`
ABataevUnsubmitted
Done
ReplyInline Actions

Still naming, IsVMTTy.

ABataev: Still naming, `IsVMTTy`.
"Parameter captured by value with variably modified type"); "Parameter captured by value with variably modified type");
if (!IsVMTTy)
EscapedParameters.insert(VD); EscapedParameters.insert(VD);
else if (!IsForCombinedParallelRegion)
return;
ABataevUnsubmitted
Not Done
ReplyInline Actions

Also, can you make it specific to pinter types?

ABataev: Also, can you make it specific to pinter types?
} else if (!IsForCombinedParallelRegion) { } else if (!IsForCombinedParallelRegion) {
return; return;
} }
} }
} }
if ((!CGF.CapturedStmtInfo || if ((!CGF.CapturedStmtInfo ||
(IsForCombinedParallelRegion && CGF.CapturedStmtInfo)) && (IsForCombinedParallelRegion && CGF.CapturedStmtInfo)) &&
VD->getType()->isReferenceType()) VD->getType()->isReferenceType())
▲ Show 20 Lines • Show All 3,473 Lines • Show Last 20 Lines

clang/test/OpenMP/target_teams_distribute_with_alias.c

  • This file was added.
// 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 _
// RUN: %clang_cc1 -fopenmp -x c -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-ppc-host.bc
// RUN: %clang_cc1 -verify -triple amdgcn-amd-amdhsa -fopenmp -x c -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK1
// expected-no-diagnostics
void foo(void) {
const int m = 800;
const int n = 800;
int size = m * n;
float *out;
#pragma omp target data map(from: out[0:size])
{
float(*__restrict out_)[n] = (float(*__restrict)[n])out;
#pragma omp target teams distribute
for (int x = 0; x < m; x += 1) {
#pragma omp parallel for
for (int y = 0; y < n; y += 1) {
out_[x][y] = 1;
}
}
}
}
// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_foo_l18
// CHECK1-SAME: (i64 noundef [[M:%.*]], i64 noundef [[N:%.*]], i64 noundef [[VLA:%.*]], ptr noundef [[OUT_:%.*]]) #[[ATTR0:[0-9]+]] {
// CHECK1-NEXT: entry:
// CHECK1-NEXT: [[M_ADDR:%.*]] = alloca i64, align 8, addrspace(5)
// CHECK1-NEXT: [[N_ADDR:%.*]] = alloca i64, align 8, addrspace(5)
// CHECK1-NEXT: [[VLA_ADDR:%.*]] = alloca i64, align 8, addrspace(5)
// CHECK1-NEXT: [[OUT__ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
// CHECK1-NEXT: [[M_CASTED:%.*]] = alloca i64, align 8, addrspace(5)
// CHECK1-NEXT: [[N_CASTED:%.*]] = alloca i64, align 8, addrspace(5)
// CHECK1-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
// CHECK1-NEXT: [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4, addrspace(5)
// CHECK1-NEXT: [[M_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[M_ADDR]] to ptr
// CHECK1-NEXT: [[N_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[N_ADDR]] to ptr
// CHECK1-NEXT: [[VLA_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[VLA_ADDR]] to ptr
// CHECK1-NEXT: [[OUT__ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT__ADDR]] to ptr
// CHECK1-NEXT: [[M_CASTED_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[M_CASTED]] to ptr
// CHECK1-NEXT: [[N_CASTED_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[N_CASTED]] to ptr
// CHECK1-NEXT: [[DOTZERO_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTZERO_ADDR]] to ptr
// CHECK1-NEXT: [[DOTTHREADID_TEMP__ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTTHREADID_TEMP_]] to ptr
// CHECK1-NEXT: store i64 [[M]], ptr [[M_ADDR_ASCAST]], align 8
// CHECK1-NEXT: store i64 [[N]], ptr [[N_ADDR_ASCAST]], align 8
// CHECK1-NEXT: store i64 [[VLA]], ptr [[VLA_ADDR_ASCAST]], align 8
// CHECK1-NEXT: store ptr [[OUT_]], ptr [[OUT__ADDR_ASCAST]], align 8
// CHECK1-NEXT: [[TMP0:%.*]] = load i64, ptr [[VLA_ADDR_ASCAST]], align 8
// CHECK1-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr addrspacecast (ptr addrspace(1) @[[GLOB1:[0-9]+]] to ptr), i8 1, i1 true, i1 true)
// CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
// CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
// CHECK1: user_code.entry:
// CHECK1-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr addrspacecast (ptr addrspace(1) @[[GLOB1]] to ptr))
// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[M_ADDR_ASCAST]], align 4
// CHECK1-NEXT: store i32 [[TMP3]], ptr [[M_CASTED_ASCAST]], align 4
// CHECK1-NEXT: [[TMP4:%.*]] = load i64, ptr [[M_CASTED_ASCAST]], align 8
// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[N_ADDR_ASCAST]], align 4
// CHECK1-NEXT: store i32 [[TMP5]], ptr [[N_CASTED_ASCAST]], align 4
// CHECK1-NEXT: [[TMP6:%.*]] = load i64, ptr [[N_CASTED_ASCAST]], align 8
// CHECK1-NEXT: [[TMP7:%.*]] = load ptr, ptr [[OUT__ADDR_ASCAST]], align 8
// CHECK1-NEXT: store i32 0, ptr [[DOTZERO_ADDR_ASCAST]], align 4
// CHECK1-NEXT: store i32 [[TMP2]], ptr [[DOTTHREADID_TEMP__ASCAST]], align 4
// CHECK1-NEXT: call void @__omp_outlined__(ptr [[DOTTHREADID_TEMP__ASCAST]], ptr [[DOTZERO_ADDR_ASCAST]], i64 [[TMP4]], i64 [[TMP6]], i64 [[TMP0]], ptr [[TMP7]]) #[[ATTR3:[0-9]+]]
// CHECK1-NEXT: call void @__kmpc_target_deinit(ptr addrspacecast (ptr addrspace(1) @[[GLOB1]] to ptr), i8 1, i1 true)
// CHECK1-NEXT: ret void
// CHECK1: worker.exit:
// CHECK1-NEXT: ret void
//
//
// CHECK1-LABEL: define {{[^@]+}}@__omp_outlined__
// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], i64 noundef [[M:%.*]], i64 noundef [[N:%.*]], i64 noundef [[VLA:%.*]], ptr noundef [[OUT_:%.*]]) #[[ATTR1:[0-9]+]] {
// CHECK1-NEXT: entry:
// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
// CHECK1-NEXT: [[M_ADDR:%.*]] = alloca i64, align 8, addrspace(5)
// CHECK1-NEXT: [[N_ADDR:%.*]] = alloca i64, align 8, addrspace(5)
// CHECK1-NEXT: [[VLA_ADDR:%.*]] = alloca i64, align 8, addrspace(5)
// CHECK1-NEXT: [[OUT__ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4, addrspace(5)
// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4, addrspace(5)
// CHECK1-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4, addrspace(5)
// CHECK1-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4, addrspace(5)
// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4, addrspace(5)
// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4, addrspace(5)
// CHECK1-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [4 x ptr], align 8, addrspace(5)
// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTGLOBAL_TID__ADDR]] to ptr
// CHECK1-NEXT: [[DOTBOUND_TID__ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTBOUND_TID__ADDR]] to ptr
// CHECK1-NEXT: [[M_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[M_ADDR]] to ptr
// CHECK1-NEXT: [[N_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[N_ADDR]] to ptr
// CHECK1-NEXT: [[VLA_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[VLA_ADDR]] to ptr
// CHECK1-NEXT: [[OUT__ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT__ADDR]] to ptr
// CHECK1-NEXT: [[DOTOMP_IV_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_IV]] to ptr
// CHECK1-NEXT: [[TMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[TMP]] to ptr
// CHECK1-NEXT: [[DOTOMP_LB_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_LB]] to ptr
// CHECK1-NEXT: [[DOTOMP_UB_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_UB]] to ptr
// CHECK1-NEXT: [[DOTOMP_STRIDE_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_STRIDE]] to ptr
// CHECK1-NEXT: [[DOTOMP_IS_LAST_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_IS_LAST]] to ptr
// CHECK1-NEXT: [[CAPTURED_VARS_ADDRS_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[CAPTURED_VARS_ADDRS]] to ptr
// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR_ASCAST]], align 8
// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR_ASCAST]], align 8
// CHECK1-NEXT: store i64 [[M]], ptr [[M_ADDR_ASCAST]], align 8
// CHECK1-NEXT: store i64 [[N]], ptr [[N_ADDR_ASCAST]], align 8
// CHECK1-NEXT: store i64 [[VLA]], ptr [[VLA_ADDR_ASCAST]], align 8
// CHECK1-NEXT: store ptr [[OUT_]], ptr [[OUT__ADDR_ASCAST]], align 8
// CHECK1-NEXT: [[TMP0:%.*]] = load i64, ptr [[VLA_ADDR_ASCAST]], align 8
// CHECK1-NEXT: [[TMP1:%.*]] = load i32, ptr [[N_ADDR_ASCAST]], align 4
// CHECK1-NEXT: [[N1:%.*]] = call align 8 ptr @__kmpc_alloc_shared(i64 4)
// CHECK1-NEXT: store i32 [[TMP1]], ptr [[N1]], align 4
// CHECK1-NEXT: [[X:%.*]] = call align 8 ptr @__kmpc_alloc_shared(i64 4)
// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_LB_ASCAST]], align 4
// CHECK1-NEXT: store i32 799, ptr [[DOTOMP_UB_ASCAST]], align 4
// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE_ASCAST]], align 4
// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST_ASCAST]], align 4
// CHECK1-NEXT: [[TMP2:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR_ASCAST]], align 8
// CHECK1-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 4
// CHECK1-NEXT: call void @__kmpc_distribute_static_init_4(ptr addrspacecast (ptr addrspace(1) @[[GLOB2:[0-9]+]] to ptr), i32 [[TMP3]], i32 92, ptr [[DOTOMP_IS_LAST_ASCAST]], ptr [[DOTOMP_LB_ASCAST]], ptr [[DOTOMP_UB_ASCAST]], ptr [[DOTOMP_STRIDE_ASCAST]], i32 1, i32 1)
// CHECK1-NEXT: [[TMP4:%.*]] = load i32, ptr [[DOTOMP_UB_ASCAST]], align 4
// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP4]], 799
// CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]
// CHECK1: cond.true:
// CHECK1-NEXT: br label [[COND_END:%.*]]
// CHECK1: cond.false:
// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[DOTOMP_UB_ASCAST]], align 4
// CHECK1-NEXT: br label [[COND_END]]
// CHECK1: cond.end:
// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 799, [[COND_TRUE]] ], [ [[TMP5]], [[COND_FALSE]] ]
// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB_ASCAST]], align 4
// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_LB_ASCAST]], align 4
// CHECK1-NEXT: store i32 [[TMP6]], ptr [[DOTOMP_IV_ASCAST]], align 4
// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]
// CHECK1: omp.inner.for.cond:
// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 4
// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_UB_ASCAST]], align 4
// CHECK1-NEXT: [[CMP2:%.*]] = icmp sle i32 [[TMP7]], [[TMP8]]
// CHECK1-NEXT: br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
// CHECK1: omp.inner.for.body:
// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 4
// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP9]], 1
// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]
// CHECK1-NEXT: store i32 [[ADD]], ptr [[X]], align 4
// CHECK1-NEXT: [[TMP10:%.*]] = getelementptr inbounds [4 x ptr], ptr [[CAPTURED_VARS_ADDRS_ASCAST]], i64 0, i64 0
// CHECK1-NEXT: store ptr [[N1]], ptr [[TMP10]], align 8
// CHECK1-NEXT: [[TMP11:%.*]] = getelementptr inbounds [4 x ptr], ptr [[CAPTURED_VARS_ADDRS_ASCAST]], i64 0, i64 1
// CHECK1-NEXT: [[TMP12:%.*]] = inttoptr i64 [[TMP0]] to ptr
// CHECK1-NEXT: store ptr [[TMP12]], ptr [[TMP11]], align 8
// CHECK1-NEXT: [[TMP13:%.*]] = getelementptr inbounds [4 x ptr], ptr [[CAPTURED_VARS_ADDRS_ASCAST]], i64 0, i64 2
// CHECK1-NEXT: store ptr [[OUT__ADDR_ASCAST]], ptr [[TMP13]], align 8
// CHECK1-NEXT: [[TMP14:%.*]] = getelementptr inbounds [4 x ptr], ptr [[CAPTURED_VARS_ADDRS_ASCAST]], i64 0, i64 3
// CHECK1-NEXT: store ptr [[X]], ptr [[TMP14]], align 8
// CHECK1-NEXT: call void @__kmpc_parallel_51(ptr addrspacecast (ptr addrspace(1) @[[GLOB1]] to ptr), i32 [[TMP3]], i32 1, i32 -1, i32 -1, ptr @__omp_outlined__.1, ptr @__omp_outlined__.1_wrapper, ptr [[CAPTURED_VARS_ADDRS_ASCAST]], i64 4)
// CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]
// CHECK1: omp.body.continue:
// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]
// CHECK1: omp.inner.for.inc:
// CHECK1-NEXT: [[TMP15:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 4
// CHECK1-NEXT: [[ADD3:%.*]] = add nsw i32 [[TMP15]], 1
// CHECK1-NEXT: store i32 [[ADD3]], ptr [[DOTOMP_IV_ASCAST]], align 4
// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]
// CHECK1: omp.inner.for.end:
// CHECK1-NEXT: br label [[OMP_LOOP_EXIT:%.*]]
// CHECK1: omp.loop.exit:
// CHECK1-NEXT: call void @__kmpc_distribute_static_fini(ptr addrspacecast (ptr addrspace(1) @[[GLOB2]] to ptr), i32 [[TMP3]])
// CHECK1-NEXT: call void @__kmpc_free_shared(ptr [[X]], i64 4)
// CHECK1-NEXT: call void @__kmpc_free_shared(ptr [[N1]], i64 4)
// CHECK1-NEXT: ret void
//
//
// CHECK1-LABEL: define {{[^@]+}}@__omp_outlined__.1
// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[N:%.*]], i64 noundef [[VLA:%.*]], ptr noundef nonnull align 8 dereferenceable(8) [[OUT_:%.*]], ptr noundef nonnull align 4 dereferenceable(4) [[X:%.*]]) #[[ATTR1]] {
// CHECK1-NEXT: entry:
// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
// CHECK1-NEXT: [[N_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
// CHECK1-NEXT: [[VLA_ADDR:%.*]] = alloca i64, align 8, addrspace(5)
// CHECK1-NEXT: [[OUT__ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
// CHECK1-NEXT: [[X_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
// CHECK1-NEXT: [[DOTOMP_IV:%.*]] = alloca i32, align 4, addrspace(5)
// CHECK1-NEXT: [[TMP:%.*]] = alloca i32, align 4, addrspace(5)
// CHECK1-NEXT: [[DOTOMP_LB:%.*]] = alloca i32, align 4, addrspace(5)
// CHECK1-NEXT: [[DOTOMP_UB:%.*]] = alloca i32, align 4, addrspace(5)
// CHECK1-NEXT: [[DOTOMP_STRIDE:%.*]] = alloca i32, align 4, addrspace(5)
// CHECK1-NEXT: [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4, addrspace(5)
// CHECK1-NEXT: [[Y:%.*]] = alloca i32, align 4, addrspace(5)
// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTGLOBAL_TID__ADDR]] to ptr
// CHECK1-NEXT: [[DOTBOUND_TID__ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTBOUND_TID__ADDR]] to ptr
// CHECK1-NEXT: [[N_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[N_ADDR]] to ptr
// CHECK1-NEXT: [[VLA_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[VLA_ADDR]] to ptr
// CHECK1-NEXT: [[OUT__ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT__ADDR]] to ptr
// CHECK1-NEXT: [[X_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X_ADDR]] to ptr
// CHECK1-NEXT: [[DOTOMP_IV_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_IV]] to ptr
// CHECK1-NEXT: [[TMP_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[TMP]] to ptr
// CHECK1-NEXT: [[DOTOMP_LB_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_LB]] to ptr
// CHECK1-NEXT: [[DOTOMP_UB_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_UB]] to ptr
// CHECK1-NEXT: [[DOTOMP_STRIDE_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_STRIDE]] to ptr
// CHECK1-NEXT: [[DOTOMP_IS_LAST_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTOMP_IS_LAST]] to ptr
// CHECK1-NEXT: [[Y_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[Y]] to ptr
// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR_ASCAST]], align 8
// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR_ASCAST]], align 8
// CHECK1-NEXT: store ptr [[N]], ptr [[N_ADDR_ASCAST]], align 8
// CHECK1-NEXT: store i64 [[VLA]], ptr [[VLA_ADDR_ASCAST]], align 8
// CHECK1-NEXT: store ptr [[OUT_]], ptr [[OUT__ADDR_ASCAST]], align 8
// CHECK1-NEXT: store ptr [[X]], ptr [[X_ADDR_ASCAST]], align 8
// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[N_ADDR_ASCAST]], align 8
// CHECK1-NEXT: [[TMP1:%.*]] = load i64, ptr [[VLA_ADDR_ASCAST]], align 8
// CHECK1-NEXT: [[TMP2:%.*]] = load ptr, ptr [[OUT__ADDR_ASCAST]], align 8
// CHECK1-NEXT: [[TMP3:%.*]] = load ptr, ptr [[X_ADDR_ASCAST]], align 8
// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_LB_ASCAST]], align 4
// CHECK1-NEXT: store i32 799, ptr [[DOTOMP_UB_ASCAST]], align 4
// CHECK1-NEXT: store i32 1, ptr [[DOTOMP_STRIDE_ASCAST]], align 4
// CHECK1-NEXT: store i32 0, ptr [[DOTOMP_IS_LAST_ASCAST]], align 4
// CHECK1-NEXT: [[TMP4:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR_ASCAST]], align 8
// CHECK1-NEXT: [[TMP5:%.*]] = load i32, ptr [[TMP4]], align 4
// CHECK1-NEXT: call void @__kmpc_for_static_init_4(ptr addrspacecast (ptr addrspace(1) @[[GLOB3:[0-9]+]] to ptr), i32 [[TMP5]], i32 33, ptr [[DOTOMP_IS_LAST_ASCAST]], ptr [[DOTOMP_LB_ASCAST]], ptr [[DOTOMP_UB_ASCAST]], ptr [[DOTOMP_STRIDE_ASCAST]], i32 1, i32 1)
// CHECK1-NEXT: br label [[OMP_DISPATCH_COND:%.*]]
// CHECK1: omp.dispatch.cond:
// CHECK1-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTOMP_UB_ASCAST]], align 4
// CHECK1-NEXT: [[CMP:%.*]] = icmp sgt i32 [[TMP6]], 799
// CHECK1-NEXT: br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]
// CHECK1: cond.true:
// CHECK1-NEXT: br label [[COND_END:%.*]]
// CHECK1: cond.false:
// CHECK1-NEXT: [[TMP7:%.*]] = load i32, ptr [[DOTOMP_UB_ASCAST]], align 4
// CHECK1-NEXT: br label [[COND_END]]
// CHECK1: cond.end:
// CHECK1-NEXT: [[COND:%.*]] = phi i32 [ 799, [[COND_TRUE]] ], [ [[TMP7]], [[COND_FALSE]] ]
// CHECK1-NEXT: store i32 [[COND]], ptr [[DOTOMP_UB_ASCAST]], align 4
// CHECK1-NEXT: [[TMP8:%.*]] = load i32, ptr [[DOTOMP_LB_ASCAST]], align 4
// CHECK1-NEXT: store i32 [[TMP8]], ptr [[DOTOMP_IV_ASCAST]], align 4
// CHECK1-NEXT: [[TMP9:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 4
// CHECK1-NEXT: [[TMP10:%.*]] = load i32, ptr [[DOTOMP_UB_ASCAST]], align 4
// CHECK1-NEXT: [[CMP1:%.*]] = icmp sle i32 [[TMP9]], [[TMP10]]
// CHECK1-NEXT: br i1 [[CMP1]], label [[OMP_DISPATCH_BODY:%.*]], label [[OMP_DISPATCH_END:%.*]]
// CHECK1: omp.dispatch.body:
// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND:%.*]]
// CHECK1: omp.inner.for.cond:
// CHECK1-NEXT: [[TMP11:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 4
// CHECK1-NEXT: [[TMP12:%.*]] = load i32, ptr [[DOTOMP_UB_ASCAST]], align 4
// CHECK1-NEXT: [[CMP2:%.*]] = icmp sle i32 [[TMP11]], [[TMP12]]
// CHECK1-NEXT: br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
// CHECK1: omp.inner.for.body:
// CHECK1-NEXT: [[TMP13:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 4
// CHECK1-NEXT: [[MUL:%.*]] = mul nsw i32 [[TMP13]], 1
// CHECK1-NEXT: [[ADD:%.*]] = add nsw i32 0, [[MUL]]
// CHECK1-NEXT: store i32 [[ADD]], ptr [[Y_ASCAST]], align 4
// CHECK1-NEXT: [[TMP14:%.*]] = load ptr, ptr [[TMP2]], align 8
// CHECK1-NEXT: [[TMP15:%.*]] = load i32, ptr [[TMP3]], align 4
// CHECK1-NEXT: [[IDXPROM:%.*]] = sext i32 [[TMP15]] to i64
// CHECK1-NEXT: [[TMP16:%.*]] = mul nsw i64 [[IDXPROM]], [[TMP1]]
// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds float, ptr [[TMP14]], i64 [[TMP16]]
// CHECK1-NEXT: [[TMP17:%.*]] = load i32, ptr [[Y_ASCAST]], align 4
// CHECK1-NEXT: [[IDXPROM3:%.*]] = sext i32 [[TMP17]] to i64
// CHECK1-NEXT: [[ARRAYIDX4:%.*]] = getelementptr inbounds float, ptr [[ARRAYIDX]], i64 [[IDXPROM3]]
// CHECK1-NEXT: store float 1.000000e+00, ptr [[ARRAYIDX4]], align 4
// CHECK1-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]
// CHECK1: omp.body.continue:
// CHECK1-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]
// CHECK1: omp.inner.for.inc:
// CHECK1-NEXT: [[TMP18:%.*]] = load i32, ptr [[DOTOMP_IV_ASCAST]], align 4
// CHECK1-NEXT: [[ADD5:%.*]] = add nsw i32 [[TMP18]], 1
// CHECK1-NEXT: store i32 [[ADD5]], ptr [[DOTOMP_IV_ASCAST]], align 4
// CHECK1-NEXT: br label [[OMP_INNER_FOR_COND]]
// CHECK1: omp.inner.for.end:
// CHECK1-NEXT: br label [[OMP_DISPATCH_INC:%.*]]
// CHECK1: omp.dispatch.inc:
// CHECK1-NEXT: [[TMP19:%.*]] = load i32, ptr [[DOTOMP_LB_ASCAST]], align 4
// CHECK1-NEXT: [[TMP20:%.*]] = load i32, ptr [[DOTOMP_STRIDE_ASCAST]], align 4
// CHECK1-NEXT: [[ADD6:%.*]] = add nsw i32 [[TMP19]], [[TMP20]]
// CHECK1-NEXT: store i32 [[ADD6]], ptr [[DOTOMP_LB_ASCAST]], align 4
// CHECK1-NEXT: [[TMP21:%.*]] = load i32, ptr [[DOTOMP_UB_ASCAST]], align 4
// CHECK1-NEXT: [[TMP22:%.*]] = load i32, ptr [[DOTOMP_STRIDE_ASCAST]], align 4
// CHECK1-NEXT: [[ADD7:%.*]] = add nsw i32 [[TMP21]], [[TMP22]]
// CHECK1-NEXT: store i32 [[ADD7]], ptr [[DOTOMP_UB_ASCAST]], align 4
// CHECK1-NEXT: br label [[OMP_DISPATCH_COND]]
// CHECK1: omp.dispatch.end:
// CHECK1-NEXT: call void @__kmpc_for_static_fini(ptr addrspacecast (ptr addrspace(1) @[[GLOB3]] to ptr), i32 [[TMP5]])
// CHECK1-NEXT: ret void
//
//
// CHECK1-LABEL: define {{[^@]+}}@__omp_outlined__.1_wrapper
// CHECK1-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR4:[0-9]+]] {
// CHECK1-NEXT: entry:
// CHECK1-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2, addrspace(5)
// CHECK1-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4, addrspace(5)
// CHECK1-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
// CHECK1-NEXT: [[GLOBAL_ARGS:%.*]] = alloca ptr, align 8, addrspace(5)
// CHECK1-NEXT: [[DOTADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTADDR]] to ptr
// CHECK1-NEXT: [[DOTADDR1_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTADDR1]] to ptr
// CHECK1-NEXT: [[DOTZERO_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DOTZERO_ADDR]] to ptr
// CHECK1-NEXT: [[GLOBAL_ARGS_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[GLOBAL_ARGS]] to ptr
// CHECK1-NEXT: store i16 [[TMP0]], ptr [[DOTADDR_ASCAST]], align 2
// CHECK1-NEXT: store i32 [[TMP1]], ptr [[DOTADDR1_ASCAST]], align 4
// CHECK1-NEXT: store i32 0, ptr [[DOTZERO_ADDR_ASCAST]], align 4
// CHECK1-NEXT: call void @__kmpc_get_shared_variables(ptr [[GLOBAL_ARGS_ASCAST]])
// CHECK1-NEXT: [[TMP2:%.*]] = load ptr, ptr [[GLOBAL_ARGS_ASCAST]], align 8
// CHECK1-NEXT: [[TMP3:%.*]] = getelementptr inbounds ptr, ptr [[TMP2]], i64 0
// CHECK1-NEXT: [[TMP4:%.*]] = load ptr, ptr [[TMP3]], align 8
// CHECK1-NEXT: [[TMP5:%.*]] = getelementptr inbounds ptr, ptr [[TMP2]], i64 1
// CHECK1-NEXT: [[TMP6:%.*]] = load i64, ptr [[TMP5]], align 8
// CHECK1-NEXT: [[TMP7:%.*]] = getelementptr inbounds ptr, ptr [[TMP2]], i64 2
// CHECK1-NEXT: [[TMP8:%.*]] = load ptr, ptr [[TMP7]], align 8
// CHECK1-NEXT: [[TMP9:%.*]] = getelementptr inbounds ptr, ptr [[TMP2]], i64 3
// CHECK1-NEXT: [[TMP10:%.*]] = load ptr, ptr [[TMP9]], align 8
// CHECK1-NEXT: call void @__omp_outlined__.1(ptr [[DOTADDR1_ASCAST]], ptr [[DOTZERO_ADDR_ASCAST]], ptr [[TMP4]], i64 [[TMP6]], ptr [[TMP8]], ptr [[TMP10]]) #[[ATTR3]]
// CHECK1-NEXT: ret void
//