This is an archive of the discontinued LLVM Phabricator instance.

Differential D106345

[OpenMP] Run rewriteDeviceCodeStateMachine in the Module not CGSCC pass
ClosedPublic

Authored by jdoerfert on Jul 20 2021, 12:16 AM.

Details

Reviewers
jhuber6
tianshilei1992
ggeorgakoudis
ye-luo
Commits
rGbe2b56964698: [OpenMP] Run rewriteDeviceCodeStateMachine in the Module not CGSCC pass
Summary

While rewriteDeviceCodeStateMachine should probably be folded into
buildCustomStateMachine, we at least need the optimization to happen.
This was not reliably the case in the CGSCC pass but in the Module pass
it seems to work reliably.

This also ports a test to the new kernel encoding (target_init/deinit),
and makes sure we cannot run the kernel in SPMD mode.

Diff Detail

Event Timeline

jdoerfert created this revision.Jul 20 2021, 12:16 AM
Herald added subscribers: ormris, guansong, bollu and 2 others. · View Herald TranscriptJul 20 2021, 12:16 AM
jdoerfert requested review of this revision.Jul 20 2021, 12:16 AM
Herald added a project: Restricted Project. · View Herald TranscriptJul 20 2021, 12:16 AM
Herald added subscribers: llvm-commits, sstefan1. · View Herald Transcript
Harbormaster completed remote builds in B115033: Diff 360030.Jul 20 2021, 12:58 AM
jhuber6 accepted this revision.Jul 20 2021, 4:54 AM

LGTM

This revision is now accepted and ready to land.Jul 20 2021, 4:54 AM
This revision was landed with ongoing or failed builds.Jul 26 2021, 7:26 PM
Closed by commit rGbe2b56964698: [OpenMP] Run rewriteDeviceCodeStateMachine in the Module not CGSCC pass (authored by jdoerfert). · Explain Why
This revision was automatically updated to reflect the committed changes.
jdoerfert added a commit: rGbe2b56964698: [OpenMP] Run rewriteDeviceCodeStateMachine in the Module not CGSCC pass.

Revision Contents

PathSize
llvm/
lib/
Transforms/
IPO/
4 lines
test/
Transforms/
OpenMP/
24 lines
199 lines

Diff 361886

llvm/lib/Transforms/IPO/OpenMPOpt.cpp

Show First 20 Lines • Show All 713 Lines • ▼ Show 20 Lines LLVM_DEBUG(dbgs() << TAG << "Run on SCC with " << SCC.size()
<< OMPInfoCache.ModuleSlice.size() << " functions\n"); << OMPInfoCache.ModuleSlice.size() << " functions\n");
if (IsModulePass) { if (IsModulePass) {
Changed |= runAttributor(IsModulePass); Changed |= runAttributor(IsModulePass);
// Recollect uses, in case Attributor deleted any. // Recollect uses, in case Attributor deleted any.
OMPInfoCache.recollectUses(); OMPInfoCache.recollectUses();
// TODO: This should be folded into buildCustomStateMachine.
Changed |= rewriteDeviceCodeStateMachine();
if (remarksEnabled()) if (remarksEnabled())
analysisGlobalization(); analysisGlobalization();
} else { } else {
if (PrintICVValues) if (PrintICVValues)
printICVs(); printICVs();
if (PrintOpenMPKernels) if (PrintOpenMPKernels)
printKernels(); printKernels();
Changed |= runAttributor(IsModulePass); Changed |= runAttributor(IsModulePass);
// Recollect uses, in case Attributor deleted any. // Recollect uses, in case Attributor deleted any.
OMPInfoCache.recollectUses(); OMPInfoCache.recollectUses();
Changed |= deleteParallelRegions(); Changed |= deleteParallelRegions();
Changed |= rewriteDeviceCodeStateMachine();
if (HideMemoryTransferLatency) if (HideMemoryTransferLatency)
Changed |= hideMemTransfersLatency(); Changed |= hideMemTransfersLatency();
Changed |= deduplicateRuntimeCalls(); Changed |= deduplicateRuntimeCalls();
if (EnableParallelRegionMerging) { if (EnableParallelRegionMerging) {
if (mergeParallelRegions()) { if (mergeParallelRegions()) {
deduplicateRuntimeCalls(); deduplicateRuntimeCalls();
Changed = true; Changed = true;
▲ Show 20 Lines • Show All 3,492 Lines • Show Last 20 Lines

llvm/test/Transforms/OpenMP/custom_state_machines.ll

Show First 20 Lines • Show All 997 Lines • ▼ Show 20 Lines
; CHECK-NEXT: [[WORKER_WORK_FN_ADDR_CAST:%.*]] = bitcast i8* [[WORKER_WORK_FN]] to void (i16, i32)* ; CHECK-NEXT: [[WORKER_WORK_FN_ADDR_CAST:%.*]] = bitcast i8* [[WORKER_WORK_FN]] to void (i16, i32)*
; CHECK-NEXT: [[WORKER_IS_DONE:%.*]] = icmp eq i8* [[WORKER_WORK_FN]], null ; CHECK-NEXT: [[WORKER_IS_DONE:%.*]] = icmp eq i8* [[WORKER_WORK_FN]], null
; CHECK-NEXT: br i1 [[WORKER_IS_DONE]], label [[WORKER_STATE_MACHINE_FINISHED:%.*]], label [[WORKER_STATE_MACHINE_IS_ACTIVE_CHECK:%.*]] ; CHECK-NEXT: br i1 [[WORKER_IS_DONE]], label [[WORKER_STATE_MACHINE_FINISHED:%.*]], label [[WORKER_STATE_MACHINE_IS_ACTIVE_CHECK:%.*]]
; CHECK: worker_state_machine.finished: ; CHECK: worker_state_machine.finished:
; CHECK-NEXT: ret void ; CHECK-NEXT: ret void
; CHECK: worker_state_machine.is_active.check: ; CHECK: worker_state_machine.is_active.check:
; CHECK-NEXT: br i1 [[WORKER_IS_ACTIVE]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK:%.*]], label [[WORKER_STATE_MACHINE_DONE_BARRIER:%.*]] ; CHECK-NEXT: br i1 [[WORKER_IS_ACTIVE]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK:%.*]], label [[WORKER_STATE_MACHINE_DONE_BARRIER:%.*]]
; CHECK: worker_state_machine.parallel_region.check: ; CHECK: worker_state_machine.parallel_region.check:
; CHECK-NEXT: [[WORKER_CHECK_PARALLEL_REGION:%.*]] = icmp eq void (i16, i32)* [[WORKER_WORK_FN_ADDR_CAST]], @__omp_outlined__2_wrapper ; CHECK-NEXT: [[WORKER_CHECK_PARALLEL_REGION:%.*]] = icmp eq void (i16, i32)* [[WORKER_WORK_FN_ADDR_CAST]], bitcast (i8* @__omp_outlined__2_wrapper.ID to void (i16, i32)*)
; CHECK-NEXT: br i1 [[WORKER_CHECK_PARALLEL_REGION]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK1:%.*]] ; CHECK-NEXT: br i1 [[WORKER_CHECK_PARALLEL_REGION]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK1:%.*]]
; CHECK: worker_state_machine.parallel_region.execute: ; CHECK: worker_state_machine.parallel_region.execute:
; CHECK-NEXT: call void @__omp_outlined__2_wrapper(i16 0, i32 [[TMP0]]) ; CHECK-NEXT: call void @__omp_outlined__2_wrapper(i16 0, i32 [[TMP0]])
; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END:%.*]] ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END:%.*]]
; CHECK: worker_state_machine.parallel_region.check1: ; CHECK: worker_state_machine.parallel_region.check1:
; CHECK-NEXT: br i1 true, label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE2:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK3:%.*]] ; CHECK-NEXT: br i1 true, label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE2:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK3:%.*]]
; CHECK: worker_state_machine.parallel_region.execute2: ; CHECK: worker_state_machine.parallel_region.execute2:
; CHECK-NEXT: call void @__omp_outlined__3_wrapper(i16 0, i32 [[TMP0]]) ; CHECK-NEXT: call void @__omp_outlined__3_wrapper(i16 0, i32 [[TMP0]])
Show All 26 Lines
; CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 ; CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
; CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 ; CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
; CHECK-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8 ; CHECK-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8
; CHECK-NEXT: [[CAPTURED_VARS_ADDRS1:%.*]] = alloca [0 x i8*], align 8 ; CHECK-NEXT: [[CAPTURED_VARS_ADDRS1:%.*]] = alloca [0 x i8*], align 8
; CHECK-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 ; CHECK-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8
; CHECK-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 ; CHECK-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8
; CHECK-NEXT: [[TMP0:%.*]] = load i32, i32* [[DOTGLOBAL_TID_]], align 4 ; CHECK-NEXT: [[TMP0:%.*]] = load i32, i32* [[DOTGLOBAL_TID_]], align 4
; CHECK-NEXT: [[TMP1:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** ; CHECK-NEXT: [[TMP1:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
; CHECK-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__2 to i8*), i8* noundef bitcast (void (i16, i32)* @__omp_outlined__2_wrapper to i8*), i8** noundef [[TMP1]], i64 noundef 0) ; CHECK-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__2 to i8*), i8* noundef @__omp_outlined__2_wrapper.ID, i8** noundef [[TMP1]], i64 noundef 0)
; CHECK-NEXT: call void @no_parallel_region_in_here.internalized() #[[ATTR7]] ; CHECK-NEXT: call void @no_parallel_region_in_here.internalized() #[[ATTR7]]
; CHECK-NEXT: [[TMP2:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS1]] to i8** ; CHECK-NEXT: [[TMP2:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS1]] to i8**
; CHECK-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__3 to i8*), i8* noundef bitcast (void (i16, i32)* @__omp_outlined__3_wrapper to i8*), i8** noundef [[TMP2]], i64 noundef 0) ; CHECK-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__3 to i8*), i8* noundef @__omp_outlined__3_wrapper.ID, i8** noundef [[TMP2]], i64 noundef 0)
; CHECK-NEXT: ret void ; CHECK-NEXT: ret void
; ;
; ;
; CHECK: Function Attrs: convergent norecurse nounwind ; CHECK: Function Attrs: convergent norecurse nounwind
; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__2 ; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__2
; CHECK-SAME: (i32* noalias nofree [[DOTGLOBAL_TID_:%.*]], i32* noalias nofree [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] { ; CHECK-SAME: (i32* noalias nofree [[DOTGLOBAL_TID_:%.*]], i32* noalias nofree [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
; CHECK-NEXT: entry: ; CHECK-NEXT: entry:
; CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 ; CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
▲ Show 20 Lines • Show All 72 Lines • ▼ Show 20 Lines
; CHECK-NEXT: br i1 [[WORKER_IS_ACTIVE]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK:%.*]], label [[WORKER_STATE_MACHINE_DONE_BARRIER:%.*]] ; CHECK-NEXT: br i1 [[WORKER_IS_ACTIVE]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK:%.*]], label [[WORKER_STATE_MACHINE_DONE_BARRIER:%.*]]
; CHECK: worker_state_machine.parallel_region.check: ; CHECK: worker_state_machine.parallel_region.check:
; CHECK-NEXT: [[WORKER_CHECK_PARALLEL_REGION:%.*]] = icmp eq void (i16, i32)* [[WORKER_WORK_FN_ADDR_CAST]], @__omp_outlined__17_wrapper ; CHECK-NEXT: [[WORKER_CHECK_PARALLEL_REGION:%.*]] = icmp eq void (i16, i32)* [[WORKER_WORK_FN_ADDR_CAST]], @__omp_outlined__17_wrapper
; CHECK-NEXT: br i1 [[WORKER_CHECK_PARALLEL_REGION]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK1:%.*]] ; CHECK-NEXT: br i1 [[WORKER_CHECK_PARALLEL_REGION]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK1:%.*]]
; CHECK: worker_state_machine.parallel_region.execute: ; CHECK: worker_state_machine.parallel_region.execute:
; CHECK-NEXT: call void @__omp_outlined__17_wrapper(i16 0, i32 [[TMP0]]) ; CHECK-NEXT: call void @__omp_outlined__17_wrapper(i16 0, i32 [[TMP0]])
; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END:%.*]] ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END:%.*]]
; CHECK: worker_state_machine.parallel_region.check1: ; CHECK: worker_state_machine.parallel_region.check1:
; CHECK-NEXT: [[WORKER_CHECK_PARALLEL_REGION4:%.*]] = icmp eq void (i16, i32)* [[WORKER_WORK_FN_ADDR_CAST]], @__omp_outlined__5_wrapper ; CHECK-NEXT: [[WORKER_CHECK_PARALLEL_REGION4:%.*]] = icmp eq void (i16, i32)* [[WORKER_WORK_FN_ADDR_CAST]], bitcast (i8* @__omp_outlined__5_wrapper.ID to void (i16, i32)*)
; CHECK-NEXT: br i1 [[WORKER_CHECK_PARALLEL_REGION4]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE2:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK3:%.*]] ; CHECK-NEXT: br i1 [[WORKER_CHECK_PARALLEL_REGION4]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE2:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK3:%.*]]
; CHECK: worker_state_machine.parallel_region.execute2: ; CHECK: worker_state_machine.parallel_region.execute2:
; CHECK-NEXT: call void @__omp_outlined__5_wrapper(i16 0, i32 [[TMP0]]) ; CHECK-NEXT: call void @__omp_outlined__5_wrapper(i16 0, i32 [[TMP0]])
; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END]] ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END]]
; CHECK: worker_state_machine.parallel_region.check3: ; CHECK: worker_state_machine.parallel_region.check3:
; CHECK-NEXT: br i1 true, label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE5:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK6:%.*]] ; CHECK-NEXT: br i1 true, label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE5:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK6:%.*]]
; CHECK: worker_state_machine.parallel_region.execute5: ; CHECK: worker_state_machine.parallel_region.execute5:
; CHECK-NEXT: call void @__omp_outlined__18_wrapper(i16 0, i32 [[TMP0]]) ; CHECK-NEXT: call void @__omp_outlined__18_wrapper(i16 0, i32 [[TMP0]])
Show All 27 Lines
; CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 ; CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
; CHECK-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8 ; CHECK-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8
; CHECK-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 ; CHECK-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8
; CHECK-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 ; CHECK-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8
; CHECK-NEXT: call void @simple_state_machine_interprocedural_before.internalized() #[[ATTR7]] ; CHECK-NEXT: call void @simple_state_machine_interprocedural_before.internalized() #[[ATTR7]]
; CHECK-NEXT: call void @no_parallel_region_in_here.internalized() #[[ATTR7]] ; CHECK-NEXT: call void @no_parallel_region_in_here.internalized() #[[ATTR7]]
; CHECK-NEXT: [[TMP0:%.*]] = load i32, i32* [[DOTGLOBAL_TID_]], align 4 ; CHECK-NEXT: [[TMP0:%.*]] = load i32, i32* [[DOTGLOBAL_TID_]], align 4
; CHECK-NEXT: [[TMP1:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** ; CHECK-NEXT: [[TMP1:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
; CHECK-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__5 to i8*), i8* noundef bitcast (void (i16, i32)* @__omp_outlined__5_wrapper to i8*), i8** noundef [[TMP1]], i64 noundef 0) ; CHECK-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__5 to i8*), i8* noundef @__omp_outlined__5_wrapper.ID, i8** noundef [[TMP1]], i64 noundef 0)
; CHECK-NEXT: call void @simple_state_machine_interprocedural_after.internalized() #[[ATTR7]] ; CHECK-NEXT: call void @simple_state_machine_interprocedural_after.internalized() #[[ATTR7]]
; CHECK-NEXT: ret void ; CHECK-NEXT: ret void
; ;
; ;
; CHECK: Function Attrs: convergent nounwind ; CHECK: Function Attrs: convergent nounwind
; CHECK-LABEL: define {{[^@]+}}@simple_state_machine_interprocedural_before.internalized ; CHECK-LABEL: define {{[^@]+}}@simple_state_machine_interprocedural_before.internalized
; CHECK-SAME: () #[[ATTR1]] { ; CHECK-SAME: () #[[ATTR1]] {
; CHECK-NEXT: entry: ; CHECK-NEXT: entry:
▲ Show 20 Lines • Show All 83 Lines • ▼ Show 20 Lines
; CHECK-NEXT: [[WORKER_WORK_FN_ADDR_CAST:%.*]] = bitcast i8* [[WORKER_WORK_FN]] to void (i16, i32)* ; CHECK-NEXT: [[WORKER_WORK_FN_ADDR_CAST:%.*]] = bitcast i8* [[WORKER_WORK_FN]] to void (i16, i32)*
; CHECK-NEXT: [[WORKER_IS_DONE:%.*]] = icmp eq i8* [[WORKER_WORK_FN]], null ; CHECK-NEXT: [[WORKER_IS_DONE:%.*]] = icmp eq i8* [[WORKER_WORK_FN]], null
; CHECK-NEXT: br i1 [[WORKER_IS_DONE]], label [[WORKER_STATE_MACHINE_FINISHED:%.*]], label [[WORKER_STATE_MACHINE_IS_ACTIVE_CHECK:%.*]] ; CHECK-NEXT: br i1 [[WORKER_IS_DONE]], label [[WORKER_STATE_MACHINE_FINISHED:%.*]], label [[WORKER_STATE_MACHINE_IS_ACTIVE_CHECK:%.*]]
; CHECK: worker_state_machine.finished: ; CHECK: worker_state_machine.finished:
; CHECK-NEXT: ret void ; CHECK-NEXT: ret void
; CHECK: worker_state_machine.is_active.check: ; CHECK: worker_state_machine.is_active.check:
; CHECK-NEXT: br i1 [[WORKER_IS_ACTIVE]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK:%.*]], label [[WORKER_STATE_MACHINE_DONE_BARRIER:%.*]] ; CHECK-NEXT: br i1 [[WORKER_IS_ACTIVE]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK:%.*]], label [[WORKER_STATE_MACHINE_DONE_BARRIER:%.*]]
; CHECK: worker_state_machine.parallel_region.check: ; CHECK: worker_state_machine.parallel_region.check:
; CHECK-NEXT: [[WORKER_CHECK_PARALLEL_REGION:%.*]] = icmp eq void (i16, i32)* [[WORKER_WORK_FN_ADDR_CAST]], @__omp_outlined__7_wrapper ; CHECK-NEXT: [[WORKER_CHECK_PARALLEL_REGION:%.*]] = icmp eq void (i16, i32)* [[WORKER_WORK_FN_ADDR_CAST]], bitcast (i8* @__omp_outlined__7_wrapper.ID to void (i16, i32)*)
; CHECK-NEXT: br i1 [[WORKER_CHECK_PARALLEL_REGION]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK1:%.*]] ; CHECK-NEXT: br i1 [[WORKER_CHECK_PARALLEL_REGION]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK1:%.*]]
; CHECK: worker_state_machine.parallel_region.execute: ; CHECK: worker_state_machine.parallel_region.execute:
; CHECK-NEXT: call void @__omp_outlined__7_wrapper(i16 0, i32 [[TMP0]]) ; CHECK-NEXT: call void @__omp_outlined__7_wrapper(i16 0, i32 [[TMP0]])
; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END:%.*]] ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END:%.*]]
; CHECK: worker_state_machine.parallel_region.check1: ; CHECK: worker_state_machine.parallel_region.check1:
; CHECK-NEXT: [[WORKER_CHECK_PARALLEL_REGION4:%.*]] = icmp eq void (i16, i32)* [[WORKER_WORK_FN_ADDR_CAST]], @__omp_outlined__8_wrapper ; CHECK-NEXT: [[WORKER_CHECK_PARALLEL_REGION4:%.*]] = icmp eq void (i16, i32)* [[WORKER_WORK_FN_ADDR_CAST]], bitcast (i8* @__omp_outlined__8_wrapper.ID to void (i16, i32)*)
; CHECK-NEXT: br i1 [[WORKER_CHECK_PARALLEL_REGION4]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE2:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_FALLBACK_EXECUTE:%.*]] ; CHECK-NEXT: br i1 [[WORKER_CHECK_PARALLEL_REGION4]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE2:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_FALLBACK_EXECUTE:%.*]]
; CHECK: worker_state_machine.parallel_region.execute2: ; CHECK: worker_state_machine.parallel_region.execute2:
; CHECK-NEXT: call void @__omp_outlined__8_wrapper(i16 0, i32 [[TMP0]]) ; CHECK-NEXT: call void @__omp_outlined__8_wrapper(i16 0, i32 [[TMP0]])
; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END]] ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END]]
; CHECK: worker_state_machine.parallel_region.fallback.execute: ; CHECK: worker_state_machine.parallel_region.fallback.execute:
; CHECK-NEXT: call void [[WORKER_WORK_FN_ADDR_CAST]](i16 0, i32 [[TMP0]]) ; CHECK-NEXT: call void [[WORKER_WORK_FN_ADDR_CAST]](i16 0, i32 [[TMP0]])
; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END]] ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END]]
; CHECK: worker_state_machine.parallel_region.end: ; CHECK: worker_state_machine.parallel_region.end:
Show All 22 Lines
; CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 ; CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
; CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 ; CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
; CHECK-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8 ; CHECK-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8
; CHECK-NEXT: [[CAPTURED_VARS_ADDRS1:%.*]] = alloca [0 x i8*], align 8 ; CHECK-NEXT: [[CAPTURED_VARS_ADDRS1:%.*]] = alloca [0 x i8*], align 8
; CHECK-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 ; CHECK-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8
; CHECK-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 ; CHECK-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8
; CHECK-NEXT: [[TMP0:%.*]] = load i32, i32* [[DOTGLOBAL_TID_]], align 4 ; CHECK-NEXT: [[TMP0:%.*]] = load i32, i32* [[DOTGLOBAL_TID_]], align 4
; CHECK-NEXT: [[TMP1:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** ; CHECK-NEXT: [[TMP1:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
; CHECK-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__7 to i8*), i8* noundef bitcast (void (i16, i32)* @__omp_outlined__7_wrapper to i8*), i8** noundef [[TMP1]], i64 noundef 0) ; CHECK-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__7 to i8*), i8* noundef @__omp_outlined__7_wrapper.ID, i8** noundef [[TMP1]], i64 noundef 0)
; CHECK-NEXT: [[TMP2:%.*]] = call i32 @unknown() #[[ATTR8]] ; CHECK-NEXT: [[TMP2:%.*]] = call i32 @unknown() #[[ATTR8]]
; CHECK-NEXT: [[TMP3:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS1]] to i8** ; CHECK-NEXT: [[TMP3:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS1]] to i8**
; CHECK-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__8 to i8*), i8* noundef bitcast (void (i16, i32)* @__omp_outlined__8_wrapper to i8*), i8** noundef [[TMP3]], i64 noundef 0) ; CHECK-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__8 to i8*), i8* noundef @__omp_outlined__8_wrapper.ID, i8** noundef [[TMP3]], i64 noundef 0)
; CHECK-NEXT: ret void ; CHECK-NEXT: ret void
; ;
; ;
; CHECK: Function Attrs: convergent norecurse nounwind ; CHECK: Function Attrs: convergent norecurse nounwind
; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__7 ; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__7
; CHECK-SAME: (i32* noalias nofree [[DOTGLOBAL_TID_:%.*]], i32* noalias nofree [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] { ; CHECK-SAME: (i32* noalias nofree [[DOTGLOBAL_TID_:%.*]], i32* noalias nofree [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
; CHECK-NEXT: entry: ; CHECK-NEXT: entry:
; CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 ; CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
▲ Show 20 Lines • Show All 66 Lines • ▼ Show 20 Lines
; CHECK-NEXT: [[WORKER_WORK_FN_ADDR_CAST:%.*]] = bitcast i8* [[WORKER_WORK_FN]] to void (i16, i32)* ; CHECK-NEXT: [[WORKER_WORK_FN_ADDR_CAST:%.*]] = bitcast i8* [[WORKER_WORK_FN]] to void (i16, i32)*
; CHECK-NEXT: [[WORKER_IS_DONE:%.*]] = icmp eq i8* [[WORKER_WORK_FN]], null ; CHECK-NEXT: [[WORKER_IS_DONE:%.*]] = icmp eq i8* [[WORKER_WORK_FN]], null
; CHECK-NEXT: br i1 [[WORKER_IS_DONE]], label [[WORKER_STATE_MACHINE_FINISHED:%.*]], label [[WORKER_STATE_MACHINE_IS_ACTIVE_CHECK:%.*]] ; CHECK-NEXT: br i1 [[WORKER_IS_DONE]], label [[WORKER_STATE_MACHINE_FINISHED:%.*]], label [[WORKER_STATE_MACHINE_IS_ACTIVE_CHECK:%.*]]
; CHECK: worker_state_machine.finished: ; CHECK: worker_state_machine.finished:
; CHECK-NEXT: ret void ; CHECK-NEXT: ret void
; CHECK: worker_state_machine.is_active.check: ; CHECK: worker_state_machine.is_active.check:
; CHECK-NEXT: br i1 [[WORKER_IS_ACTIVE]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK:%.*]], label [[WORKER_STATE_MACHINE_DONE_BARRIER:%.*]] ; CHECK-NEXT: br i1 [[WORKER_IS_ACTIVE]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK:%.*]], label [[WORKER_STATE_MACHINE_DONE_BARRIER:%.*]]
; CHECK: worker_state_machine.parallel_region.check: ; CHECK: worker_state_machine.parallel_region.check:
; CHECK-NEXT: [[WORKER_CHECK_PARALLEL_REGION:%.*]] = icmp eq void (i16, i32)* [[WORKER_WORK_FN_ADDR_CAST]], @__omp_outlined__10_wrapper ; CHECK-NEXT: [[WORKER_CHECK_PARALLEL_REGION:%.*]] = icmp eq void (i16, i32)* [[WORKER_WORK_FN_ADDR_CAST]], bitcast (i8* @__omp_outlined__10_wrapper.ID to void (i16, i32)*)
; CHECK-NEXT: br i1 [[WORKER_CHECK_PARALLEL_REGION]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK1:%.*]] ; CHECK-NEXT: br i1 [[WORKER_CHECK_PARALLEL_REGION]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK1:%.*]]
; CHECK: worker_state_machine.parallel_region.execute: ; CHECK: worker_state_machine.parallel_region.execute:
; CHECK-NEXT: call void @__omp_outlined__10_wrapper(i16 0, i32 [[TMP0]]) ; CHECK-NEXT: call void @__omp_outlined__10_wrapper(i16 0, i32 [[TMP0]])
; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END:%.*]] ; CHECK-NEXT: br label [[WORKER_STATE_MACHINE_PARALLEL_REGION_END:%.*]]
; CHECK: worker_state_machine.parallel_region.check1: ; CHECK: worker_state_machine.parallel_region.check1:
; CHECK-NEXT: br i1 true, label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE2:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK3:%.*]] ; CHECK-NEXT: br i1 true, label [[WORKER_STATE_MACHINE_PARALLEL_REGION_EXECUTE2:%.*]], label [[WORKER_STATE_MACHINE_PARALLEL_REGION_CHECK3:%.*]]
; CHECK: worker_state_machine.parallel_region.execute2: ; CHECK: worker_state_machine.parallel_region.execute2:
; CHECK-NEXT: call void @__omp_outlined__11_wrapper(i16 0, i32 [[TMP0]]) ; CHECK-NEXT: call void @__omp_outlined__11_wrapper(i16 0, i32 [[TMP0]])
Show All 26 Lines
; CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 ; CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
; CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8 ; CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
; CHECK-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8 ; CHECK-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x i8*], align 8
; CHECK-NEXT: [[CAPTURED_VARS_ADDRS1:%.*]] = alloca [0 x i8*], align 8 ; CHECK-NEXT: [[CAPTURED_VARS_ADDRS1:%.*]] = alloca [0 x i8*], align 8
; CHECK-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8 ; CHECK-NEXT: store i32* [[DOTGLOBAL_TID_]], i32** [[DOTGLOBAL_TID__ADDR]], align 8
; CHECK-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8 ; CHECK-NEXT: store i32* [[DOTBOUND_TID_]], i32** [[DOTBOUND_TID__ADDR]], align 8
; CHECK-NEXT: [[TMP0:%.*]] = load i32, i32* [[DOTGLOBAL_TID_]], align 4 ; CHECK-NEXT: [[TMP0:%.*]] = load i32, i32* [[DOTGLOBAL_TID_]], align 4
; CHECK-NEXT: [[TMP1:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8** ; CHECK-NEXT: [[TMP1:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS]] to i8**
; CHECK-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__10 to i8*), i8* noundef bitcast (void (i16, i32)* @__omp_outlined__10_wrapper to i8*), i8** noundef [[TMP1]], i64 noundef 0) ; CHECK-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__10 to i8*), i8* noundef @__omp_outlined__10_wrapper.ID, i8** noundef [[TMP1]], i64 noundef 0)
; CHECK-NEXT: call void @unknown_no_openmp() #[[ATTR9:[0-9]+]] ; CHECK-NEXT: call void @unknown_no_openmp() #[[ATTR9:[0-9]+]]
; CHECK-NEXT: [[TMP2:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS1]] to i8** ; CHECK-NEXT: [[TMP2:%.*]] = bitcast [0 x i8*]* [[CAPTURED_VARS_ADDRS1]] to i8**
; CHECK-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__11 to i8*), i8* noundef bitcast (void (i16, i32)* @__omp_outlined__11_wrapper to i8*), i8** noundef [[TMP2]], i64 noundef 0) ; CHECK-NEXT: call void @__kmpc_parallel_51(%struct.ident_t* noundef @[[GLOB1]], i32 [[TMP0]], i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__11 to i8*), i8* noundef @__omp_outlined__11_wrapper.ID, i8** noundef [[TMP2]], i64 noundef 0)
; CHECK-NEXT: ret void ; CHECK-NEXT: ret void
; ;
; ;
; CHECK: Function Attrs: convergent norecurse nounwind ; CHECK: Function Attrs: convergent norecurse nounwind
; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__10 ; CHECK-LABEL: define {{[^@]+}}@__omp_outlined__10
; CHECK-SAME: (i32* noalias nofree [[DOTGLOBAL_TID_:%.*]], i32* noalias nofree [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] { ; CHECK-SAME: (i32* noalias nofree [[DOTGLOBAL_TID_:%.*]], i32* noalias nofree [[DOTBOUND_TID_:%.*]]) #[[ATTR0]] {
; CHECK-NEXT: entry: ; CHECK-NEXT: entry:
; CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8 ; CHECK-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
▲ Show 20 Lines • Show All 414 Lines • Show Last 20 Lines

llvm/test/Transforms/OpenMP/gpu_state_machine_function_ptr_replacement.ll

; RUN: opt -S -passes=openmp-opt-cgscc -openmp-ir-builder-optimistic-attributes -pass-remarks=openmp-opt -openmp-print-gpu-kernels < %s | FileCheck %s ; RUN: opt -S -passes=openmp-opt -openmp-ir-builder-optimistic-attributes -pass-remarks=openmp-opt -openmp-print-gpu-kernels < %s | FileCheck %s
; RUN: opt -S -passes=openmp-opt-cgscc -pass-remarks=openmp-opt -openmp-print-gpu-kernels < %s | FileCheck %s ; RUN: opt -S -passes=openmp-opt -pass-remarks=openmp-opt -openmp-print-gpu-kernels < %s | FileCheck %s
; RUN: opt -S -openmp-opt-cgscc -pass-remarks=openmp-opt -openmp-print-gpu-kernels < %s | FileCheck %s
; C input used for this test: ; C input used for this test:
; void bar(void) { ; void bar(void) {
; #pragma omp parallel ; #pragma omp parallel
; { } ; { }
; } ; }
; void foo(void) { ; void foo(void) {
; #pragma omp target teams ; #pragma omp target teams
; { ; {
; #pragma omp parallel ; #pragma omp parallel
; {} ; {}
; bar(); ; bar();
; unknown();
; #pragma omp parallel ; #pragma omp parallel
; {} ; {}
; } ; }
; } ; }
; Verify we replace the function pointer uses for the first and last outlined ; Verify we replace the function pointer uses for the first and last outlined
; region (1 and 3) but not for the middle one (2) because it could be called from ; region (1 and 3) but not for the middle one (2) because it could be called from
; another kernel. ; another kernel.
; CHECK-DAG: @__omp_outlined__1_wrapper.ID = private constant i8 undef ; CHECK-DAG: @__omp_outlined__1_wrapper.ID = private constant i8 undef
; CHECK-DAG: @__omp_outlined__3_wrapper.ID = private constant i8 undef ; CHECK-DAG: @__omp_outlined__2_wrapper.ID = private constant i8 undef
; CHECK-DAG: icmp eq i8* %5, @__omp_outlined__1_wrapper.ID ; CHECK-DAG: icmp eq void (i16, i32)* %worker.work_fn.addr_cast, bitcast (i8* @__omp_outlined__1_wrapper.ID to void (i16, i32)*)
; CHECK-DAG: icmp eq i8* %7, @__omp_outlined__3_wrapper.ID ; CHECK-DAG: icmp eq void (i16, i32)* %worker.work_fn.addr_cast, bitcast (i8* @__omp_outlined__2_wrapper.ID to void (i16, i32)*)
; CHECK-DAG: call void @__kmpc_parallel_51(%struct.ident_t* noundef @1, i32 %1, i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__1 to i8*), i8* noundef @__omp_outlined__1_wrapper.ID, i8** noundef %2, i64 noundef 0)
; CHECK-DAG: call void @__kmpc_parallel_51(%struct.ident_t* @1, i32 %0, i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*)* @__omp_outlined__2 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__2_wrapper to i8*), i8** %1, i64 0) ; CHECK-DAG: call void @__kmpc_parallel_51(%struct.ident_t* noundef @1, i32 %{{.*}}, i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__1 to i8*), i8* noundef @__omp_outlined__1_wrapper.ID, i8** noundef %{{.*}}, i64 noundef 0)
; CHECK-DAG: call void @__kmpc_parallel_51(%struct.ident_t* noundef @1, i32 %1, i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__3 to i8*), i8* noundef @__omp_outlined__3_wrapper.ID, i8** noundef %3, i64 noundef 0) ; CHECK-DAG: call void @__kmpc_parallel_51(%struct.ident_t* noundef @1, i32 %{{.*}}, i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__2 to i8*), i8* noundef @__omp_outlined__2_wrapper.ID, i8** noundef %{{.*}}, i64 noundef 0)
; CHECK-DAG: call void @__kmpc_parallel_51(%struct.ident_t* noundef @2, i32 %{{.*}}, i32 noundef 1, i32 noundef -1, i32 noundef -1, i8* noundef bitcast (void (i32*, i32*)* @__omp_outlined__3 to i8*), i8* noundef bitcast (void (i16, i32)* @__omp_outlined__3_wrapper to i8*), i8** noundef %{{.*}}, i64 noundef 0)
%struct.ident_t = type { i32, i32, i32, i32, i8* } %struct.ident_t = type { i32, i32, i32, i32, i8* }
@0 = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00", align 1 @0 = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00", align 1
@1 = private unnamed_addr constant %struct.ident_t { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* @0, i32 0, i32 0) }, align 8 @1 = private unnamed_addr constant %struct.ident_t { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* @0, i32 0, i32 0) }, align 8
@__omp_offloading_10301_87b2c_foo_l7_exec_mode = weak constant i8 1
@2 = private unnamed_addr constant %struct.ident_t { i32 0, i32 2, i32 2, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* @0, i32 0, i32 0) }, align 8
@llvm.compiler.used = appending global [1 x i8*] [i8* @__omp_offloading_10301_87b2c_foo_l7_exec_mode], section "llvm.metadata"
define internal void @__omp_offloading_50_6dfa0f01_foo_l6_worker() { define weak void @__omp_offloading_10301_87b2c_foo_l7() {
entry:
%work_fn = alloca i8*, align 8
%exec_status = alloca i8, align 1
store i8* null, i8** %work_fn, align 8
store i8 0, i8* %exec_status, align 1
br label %.await.work
.await.work: ; preds = %.barrier.parallel, %entry
call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
%0 = call i1 @__kmpc_kernel_parallel(i8** %work_fn)
%1 = zext i1 %0 to i8
store i8 %1, i8* %exec_status, align 1
%2 = load i8*, i8** %work_fn, align 8
%should_terminate = icmp eq i8* %2, null
br i1 %should_terminate, label %.exit, label %.select.workers
.select.workers: ; preds = %.await.work
%3 = load i8, i8* %exec_status, align 1
%is_active = icmp ne i8 %3, 0
br i1 %is_active, label %.execute.parallel, label %.barrier.parallel
.execute.parallel: ; preds = %.select.workers
%4 = call i32 @__kmpc_global_thread_num(%struct.ident_t* @1)
%5 = load i8*, i8** %work_fn, align 8
%work_match = icmp eq i8* %5, bitcast (void (i16, i32)* @__omp_outlined__1_wrapper to i8*)
br i1 %work_match, label %.execute.fn, label %.check.next
.execute.fn: ; preds = %.execute.parallel
call void @__omp_outlined__1_wrapper(i16 zeroext 0, i32 %4)
br label %.terminate.parallel
.check.next: ; preds = %.execute.parallel
%6 = load i8*, i8** %work_fn, align 8
%work_match1 = icmp eq i8* %6, bitcast (void (i16, i32)* @__omp_outlined__2_wrapper to i8*)
br i1 %work_match1, label %.execute.fn2, label %.check.next3
.execute.fn2: ; preds = %.check.next
call void @__omp_outlined__2_wrapper(i16 zeroext 0, i32 %4)
br label %.terminate.parallel
.check.next3: ; preds = %.check.next
%7 = load i8*, i8** %work_fn, align 8
%work_match4 = icmp eq i8* %7, bitcast (void (i16, i32)* @__omp_outlined__3_wrapper to i8*)
br i1 %work_match4, label %.execute.fn5, label %.check.next6
.execute.fn5: ; preds = %.check.next3
call void @__omp_outlined__3_wrapper(i16 zeroext 0, i32 %4)
br label %.terminate.parallel
.check.next6: ; preds = %.check.next3
%8 = bitcast i8* %2 to void (i16, i32)*
call void %8(i16 0, i32 %4)
br label %.terminate.parallel
.terminate.parallel: ; preds = %.check.next6, %.execute.fn5, %.execute.fn2, %.execute.fn
call void @__kmpc_kernel_end_parallel()
br label %.barrier.parallel
.barrier.parallel: ; preds = %.terminate.parallel, %.select.workers
call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
br label %.await.work
.exit: ; preds = %.await.work
ret void
}
define weak void @__omp_offloading_50_6dfa0f01_foo_l6() {
entry: entry:
%.zero.addr = alloca i32, align 4 %.zero.addr = alloca i32, align 4
%.threadid_temp. = alloca i32, align 4 %.threadid_temp. = alloca i32, align 4
store i32 0, i32* %.zero.addr, align 4 store i32 0, i32* %.zero.addr, align 4
%nvptx_tid = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() %0 = call i32 @__kmpc_target_init(%struct.ident_t* @1, i1 false, i1 true, i1 true)
%nvptx_num_threads = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x() %exec_user_code = icmp eq i32 %0, -1
%nvptx_warp_size = call i32 @llvm.nvvm.read.ptx.sreg.warpsize() br i1 %exec_user_code, label %user_code.entry, label %worker.exit
%thread_limit = sub nuw i32 %nvptx_num_threads, %nvptx_warp_size
%0 = icmp ult i32 %nvptx_tid, %thread_limit user_code.entry: ; preds = %entry
br i1 %0, label %.worker, label %.mastercheck %1 = call i32 @__kmpc_global_thread_num(%struct.ident_t* @1)
store i32 %1, i32* %.threadid_temp., align 4
.worker: ; preds = %entry
call void @__omp_offloading_50_6dfa0f01_foo_l6_worker()
br label %.exit
.mastercheck: ; preds = %entry
%nvptx_tid1 = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
%nvptx_num_threads2 = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
%nvptx_warp_size3 = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
%1 = sub nuw i32 %nvptx_warp_size3, 1
%2 = sub nuw i32 %nvptx_num_threads2, 1
%3 = xor i32 %1, -1
%master_tid = and i32 %2, %3
%4 = icmp eq i32 %nvptx_tid1, %master_tid
br i1 %4, label %.master, label %.exit
.master: ; preds = %.mastercheck
%nvptx_num_threads4 = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
%nvptx_warp_size5 = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
%thread_limit6 = sub nuw i32 %nvptx_num_threads4, %nvptx_warp_size5
call void @__kmpc_kernel_init(i32 %thread_limit6, i16 1)
call void @__kmpc_data_sharing_init_stack()
%5 = call i32 @__kmpc_global_thread_num(%struct.ident_t* @1)
store i32 %5, i32* %.threadid_temp., align 4
call void @__omp_outlined__(i32* %.threadid_temp., i32* %.zero.addr) call void @__omp_outlined__(i32* %.threadid_temp., i32* %.zero.addr)
br label %.termination.notifier call void @__kmpc_target_deinit(%struct.ident_t* @1, i1 false, i1 true)
ret void
.termination.notifier: ; preds = %.master
call void @__kmpc_kernel_deinit(i16 1)
call void @__kmpc_barrier_simple_spmd(%struct.ident_t* null, i32 0)
br label %.exit
.exit: ; preds = %.termination.notifier, %.mastercheck, %.worker worker.exit: ; preds = %entry
ret void ret void
} }
declare i32 @llvm.nvvm.read.ptx.sreg.tid.x() declare i32 @__kmpc_target_init(%struct.ident_t*, i1, i1, i1)
declare void @unknown()
declare i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
declare i32 @llvm.nvvm.read.ptx.sreg.warpsize()
declare void @__kmpc_kernel_init(i32, i16)
declare void @__kmpc_data_sharing_init_stack()
define internal void @__omp_outlined__(i32* noalias %.global_tid., i32* noalias %.bound_tid.) { define internal void @__omp_outlined__(i32* noalias %.global_tid., i32* noalias %.bound_tid.) {
entry: entry:
%.global_tid..addr = alloca i32*, align 8 %.global_tid..addr = alloca i32*, align 8
%.bound_tid..addr = alloca i32*, align 8 %.bound_tid..addr = alloca i32*, align 8
%captured_vars_addrs = alloca [0 x i8*], align 8 %captured_vars_addrs = alloca [0 x i8*], align 8
%captured_vars_addrs1 = alloca [0 x i8*], align 8 %captured_vars_addrs1 = alloca [0 x i8*], align 8
store i32* %.global_tid., i32** %.global_tid..addr, align 8 store i32* %.global_tid., i32** %.global_tid..addr, align 8
store i32* %.bound_tid., i32** %.bound_tid..addr, align 8 store i32* %.bound_tid., i32** %.bound_tid..addr, align 8
%0 = load i32*, i32** %.global_tid..addr, align 8 %0 = load i32*, i32** %.global_tid..addr, align 8
%1 = load i32, i32* %0, align 4 %1 = load i32, i32* %0, align 4
%2 = bitcast [0 x i8*]* %captured_vars_addrs to i8** %2 = bitcast [0 x i8*]* %captured_vars_addrs to i8**
call void @__kmpc_parallel_51(%struct.ident_t* @1, i32 %1, i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*)* @__omp_outlined__1 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__1_wrapper to i8*), i8** %2, i64 0) call void @__kmpc_parallel_51(%struct.ident_t* @1, i32 %1, i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*)* @__omp_outlined__1 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__1_wrapper to i8*), i8** %2, i64 0)
call void @bar() call void @bar()
call void @unknown()
%3 = bitcast [0 x i8*]* %captured_vars_addrs1 to i8** %3 = bitcast [0 x i8*]* %captured_vars_addrs1 to i8**
call void @__kmpc_parallel_51(%struct.ident_t* @1, i32 %1, i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*)* @__omp_outlined__3 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__3_wrapper to i8*), i8** %3, i64 0) call void @__kmpc_parallel_51(%struct.ident_t* @1, i32 %1, i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*)* @__omp_outlined__2 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__2_wrapper to i8*), i8** %3, i64 0)
ret void ret void
} }
define internal void @__omp_outlined__1(i32* noalias %.global_tid., i32* noalias %.bound_tid.) { define internal void @__omp_outlined__1(i32* noalias %.global_tid., i32* noalias %.bound_tid.) {
entry: entry:
%.global_tid..addr = alloca i32*, align 8 %.global_tid..addr = alloca i32*, align 8
%.bound_tid..addr = alloca i32*, align 8 %.bound_tid..addr = alloca i32*, align 8
store i32* %.global_tid., i32** %.global_tid..addr, align 8 store i32* %.global_tid., i32** %.global_tid..addr, align 8
store i32* %.bound_tid., i32** %.bound_tid..addr, align 8 store i32* %.bound_tid., i32** %.bound_tid..addr, align 8
ret void ret void
} }
define internal void @__omp_outlined__1_wrapper(i16 zeroext %0, i32 %1) { define internal void @__omp_outlined__1_wrapper(i16 zeroext %0, i32 %1) {
entry: entry:
%.addr = alloca i16, align 2 %.addr = alloca i16, align 2
%.addr1 = alloca i32, align 4 %.addr1 = alloca i32, align 4
%.zero.addr = alloca i32, align 4 %.zero.addr = alloca i32, align 4
%global_args = alloca i8**, align 8 %global_args = alloca i8**, align 8
store i32 0, i32* %.zero.addr, align 4 store i32 0, i32* %.zero.addr, align 4
store i16 %0, i16* %.addr, align 2 store i16 %0, i16* %.addr, align 2
store i32 %1, i32* %.addr1, align 4 store i32 %1, i32* %.addr1, align 4
call void @__kmpc_get_shared_variables(i8*** %global_args) call void @__kmpc_get_shared_variables(i8*** %global_args)
call void @__omp_outlined__1(i32* %.addr1, i32* %.zero.addr) call void @__omp_outlined__1(i32* %.addr1, i32* %.zero.addr)
ret void ret void
} }
declare void @__kmpc_get_shared_variables(i8***) declare void @__kmpc_get_shared_variables(i8***)
declare void @__kmpc_parallel_51(%struct.ident_t*, i32, i32, i32, i32, i8*, i8*, i8**, i64) declare void @__kmpc_parallel_51(%struct.ident_t*, i32, i32, i32, i32, i8*, i8*, i8**, i64)
define hidden void @bar() { define hidden void @bar() {
entry: entry:
%captured_vars_addrs = alloca [0 x i8*], align 8 %captured_vars_addrs = alloca [0 x i8*], align 8
%0 = call i32 @__kmpc_global_thread_num(%struct.ident_t* @1) %0 = call i32 @__kmpc_global_thread_num(%struct.ident_t* @2)
%1 = bitcast [0 x i8*]* %captured_vars_addrs to i8** %1 = bitcast [0 x i8*]* %captured_vars_addrs to i8**
call void @__kmpc_parallel_51(%struct.ident_t* @1, i32 %0, i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*)* @__omp_outlined__2 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__2_wrapper to i8*), i8** %1, i64 0) call void @__kmpc_parallel_51(%struct.ident_t* @2, i32 %0, i32 1, i32 -1, i32 -1, i8* bitcast (void (i32*, i32*)* @__omp_outlined__3 to i8*), i8* bitcast (void (i16, i32)* @__omp_outlined__3_wrapper to i8*), i8** %1, i64 0)
ret void ret void
} }
define internal void @__omp_outlined__2(i32* noalias %.global_tid., i32* noalias %.bound_tid.) { define internal void @__omp_outlined__2(i32* noalias %.global_tid., i32* noalias %.bound_tid.) {
entry: entry:
%.global_tid..addr = alloca i32*, align 8 %.global_tid..addr = alloca i32*, align 8
%.bound_tid..addr = alloca i32*, align 8 %.bound_tid..addr = alloca i32*, align 8
store i32* %.global_tid., i32** %.global_tid..addr, align 8 store i32* %.global_tid., i32** %.global_tid..addr, align 8
store i32* %.bound_tid., i32** %.bound_tid..addr, align 8 store i32* %.bound_tid., i32** %.bound_tid..addr, align 8
ret void ret void
} }
define internal void @__omp_outlined__2_wrapper(i16 zeroext %0, i32 %1) { define internal void @__omp_outlined__2_wrapper(i16 zeroext %0, i32 %1) {
entry: entry:
%.addr = alloca i16, align 2 %.addr = alloca i16, align 2
%.addr1 = alloca i32, align 4 %.addr1 = alloca i32, align 4
%.zero.addr = alloca i32, align 4 %.zero.addr = alloca i32, align 4
%global_args = alloca i8**, align 8 %global_args = alloca i8**, align 8
store i32 0, i32* %.zero.addr, align 4 store i32 0, i32* %.zero.addr, align 4
store i16 %0, i16* %.addr, align 2 store i16 %0, i16* %.addr, align 2
store i32 %1, i32* %.addr1, align 4 store i32 %1, i32* %.addr1, align 4
call void @__kmpc_get_shared_variables(i8*** %global_args) call void @__kmpc_get_shared_variables(i8*** %global_args)
call void @__omp_outlined__2(i32* %.addr1, i32* %.zero.addr) call void @__omp_outlined__2(i32* %.addr1, i32* %.zero.addr)
ret void ret void
} }
declare i32 @__kmpc_global_thread_num(%struct.ident_t*) declare i32 @__kmpc_global_thread_num(%struct.ident_t*)
declare void @__kmpc_target_deinit(%struct.ident_t*, i1, i1)
define internal void @__omp_outlined__3(i32* noalias %.global_tid., i32* noalias %.bound_tid.) { define internal void @__omp_outlined__3(i32* noalias %.global_tid., i32* noalias %.bound_tid.) {
entry: entry:
%.global_tid..addr = alloca i32*, align 8 %.global_tid..addr = alloca i32*, align 8
%.bound_tid..addr = alloca i32*, align 8 %.bound_tid..addr = alloca i32*, align 8
store i32* %.global_tid., i32** %.global_tid..addr, align 8 store i32* %.global_tid., i32** %.global_tid..addr, align 8
store i32* %.bound_tid., i32** %.bound_tid..addr, align 8 store i32* %.bound_tid., i32** %.bound_tid..addr, align 8
ret void ret void
} }
define internal void @__omp_outlined__3_wrapper(i16 zeroext %0, i32 %1) { define internal void @__omp_outlined__3_wrapper(i16 zeroext %0, i32 %1) {
entry: entry:
%.addr = alloca i16, align 2 %.addr = alloca i16, align 2
%.addr1 = alloca i32, align 4 %.addr1 = alloca i32, align 4
%.zero.addr = alloca i32, align 4 %.zero.addr = alloca i32, align 4
%global_args = alloca i8**, align 8 %global_args = alloca i8**, align 8
store i32 0, i32* %.zero.addr, align 4 store i32 0, i32* %.zero.addr, align 4
store i16 %0, i16* %.addr, align 2 store i16 %0, i16* %.addr, align 2
store i32 %1, i32* %.addr1, align 4 store i32 %1, i32* %.addr1, align 4
call void @__kmpc_get_shared_variables(i8*** %global_args) call void @__kmpc_get_shared_variables(i8*** %global_args)
call void @__omp_outlined__3(i32* %.addr1, i32* %.zero.addr) call void @__omp_outlined__3(i32* %.addr1, i32* %.zero.addr)
ret void ret void
} }
declare void @__kmpc_kernel_deinit(i16) !omp_offload.info = !{!0}
declare void @__kmpc_barrier_simple_spmd(%struct.ident_t*, i32)
declare i1 @__kmpc_kernel_parallel(i8**)
declare void @__kmpc_kernel_end_parallel()
!nvvm.annotations = !{!1} !nvvm.annotations = !{!1}
!llvm.module.flags = !{!2, !3} !llvm.module.flags = !{!2, !3}
!1 = !{void ()* @__omp_offloading_50_6dfa0f01_foo_l6, !"kernel", i32 1} !0 = !{i32 0, i32 66305, i32 555956, !"foo", i32 7, i32 0}
!1 = !{void ()* @__omp_offloading_10301_87b2c_foo_l7, !"kernel", i32 1}
!2 = !{i32 7, !"openmp", i32 50} !2 = !{i32 7, !"openmp", i32 50}
!3 = !{i32 7, !"openmp-device", i32 50} !3 = !{i32 7, !"openmp-device", i32 50}