[llvm] InferAddressSpaces: Make getPredicatedAddrSpace less confusing (PR #104052)

Thu Aug 15 03:46:11 PDT 2024

llvm-ci wrote:

LLVM Buildbot has detected a new failure on builder `mlir-rocm-mi200` running on `mi200-buildbot` while building `llvm` at step 6 "test-build-check-mlir-build-only-check-mlir".

Full details are available at: https://lab.llvm.org/buildbot/#/builders/177/builds/3221

Here is the relevant piece of the build log for the reference:
```
Step 6 (test-build-check-mlir-build-only-check-mlir) failure: test (failure)
******************** TEST 'MLIR :: Integration/GPU/ROCM/printf.mlir' FAILED ********************
Exit Code: 2

Command Output (stdout):
--
# RUN: at line 1
/vol/worker/mi200-buildbot/mlir-rocm-mi200/build/bin/mlir-opt /vol/worker/mi200-buildbot/mlir-rocm-mi200/llvm-project/mlir/test/Integration/GPU/ROCM/printf.mlir  | /vol/worker/mi200-buildbot/mlir-rocm-mi200/build/bin/mlir-opt -pass-pipeline='builtin.module(gpu.module(strip-debuginfo,convert-gpu-to-rocdl{index-bitwidth=32 runtime=HIP}),rocdl-attach-target{chip=gfx906})'  | /vol/worker/mi200-buildbot/mlir-rocm-mi200/build/bin/mlir-opt -gpu-to-llvm -reconcile-unrealized-casts -gpu-module-to-binary  | /vol/worker/mi200-buildbot/mlir-rocm-mi200/build/bin/mlir-cpu-runner    --shared-libs=/vol/worker/mi200-buildbot/mlir-rocm-mi200/build/lib/libmlir_rocm_runtime.so    --shared-libs=/vol/worker/mi200-buildbot/mlir-rocm-mi200/build/lib/libmlir_runner_utils.so    --entry-point-result=void  | /vol/worker/mi200-buildbot/mlir-rocm-mi200/build/bin/FileCheck /vol/worker/mi200-buildbot/mlir-rocm-mi200/llvm-project/mlir/test/Integration/GPU/ROCM/printf.mlir
# executed command: /vol/worker/mi200-buildbot/mlir-rocm-mi200/build/bin/mlir-opt /vol/worker/mi200-buildbot/mlir-rocm-mi200/llvm-project/mlir/test/Integration/GPU/ROCM/printf.mlir
# executed command: /vol/worker/mi200-buildbot/mlir-rocm-mi200/build/bin/mlir-opt '-pass-pipeline=builtin.module(gpu.module(strip-debuginfo,convert-gpu-to-rocdl{index-bitwidth=32 runtime=HIP}),rocdl-attach-target{chip=gfx906})'
# executed command: /vol/worker/mi200-buildbot/mlir-rocm-mi200/build/bin/mlir-opt -gpu-to-llvm -reconcile-unrealized-casts -gpu-module-to-binary
# .---command stderr------------
# | mlir-opt: /vol/worker/mi200-buildbot/mlir-rocm-mi200/llvm-project/llvm/include/llvm/Support/Casting.h:566: decltype(auto) llvm::cast(const From &) [To = llvm::Instruction, From = llvm::Value]: Assertion `isa<To>(Val) && "cast<Ty>() argument of incompatible type!"' failed.
# | PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace.
# | Stack dump:
# | 0.	Program arguments: /vol/worker/mi200-buildbot/mlir-rocm-mi200/build/bin/mlir-opt -gpu-to-llvm -reconcile-unrealized-casts -gpu-module-to-binary
# | 1.	Running pass "require<GlobalsAA>,function(invalidate<AAManager>),require<ProfileSummaryAnalysis>,cgscc(devirt<4>(InlinerPass,PostOrderFunctionAttrsPass<skip-non-recursive-function-attrs>,OpenMPOptCGSCCPass,function(AMDGPUPromoteKernelArgumentsPass,InferAddressSpacesPass,AMDGPULowerKernelAttributesPass,AMDGPUPromoteAllocaToVectorPass),function<eager-inv;no-rerun>(SROAPass<modify-cfg>,EarlyCSEPass<memssa>,SpeculativeExecutionPass<only-if-divergent-target>,JumpThreadingPass,CorrelatedValuePropagationPass,SimplifyCFGPass<bonus-inst-threshold=1;no-forward-switch-cond;switch-range-to-icmp;no-switch-to-lookup;keep-loops;no-hoist-common-insts;no-sink-common-insts;speculate-blocks;simplify-cond-branch;no-speculate-unpredictables>,InstCombinePass<max-iterations=1;no-use-loop-info;no-verify-fixpoint>,AggressiveInstCombinePass,LibCallsShrinkWrapPass,AMDGPUUseNativeCallsPass,AMDGPUSimplifyLibCallsPass,TailCallElimPass,SimplifyCFGPass<bonus-inst-threshold=1;no-forward-switch-cond;switch-range-to-icmp;no-switch-to-lookup;keep-loops;no-hoist-common-insts;no-sink-common-insts;speculate-blocks;simplify-cond-branch;no-speculate-unpredictables>,ReassociatePass,ConstraintEliminationPass,loop-mssa(LoopInstSimplifyPass,LoopSimplifyCFGPass,LICMPass<no-allowspeculation>,LoopRotatePass<header-duplication;no-prepare-for-lto>,LICMPass<allowspeculation>,SimpleLoopUnswitchPass<no-nontrivial;trivial>),SimplifyCFGPass<bonus-inst-threshold=1;no-forward-switch-cond;switch-range-to-icmp;no-switch-to-lookup;keep-loops;no-hoist-common-insts;no-sink-common-insts;speculate-blocks;simplify-cond-branch;no-speculate-unpredictables>,InstCombinePass<max-iterations=1;no-use-loop-info;no-verify-fixpoint>,loop(LoopIdiomRecognizePass,IndVarSimplifyPass,SimpleLoopUnswitchPass<no-nontrivial;trivial>,LoopDeletionPass,LoopFullUnrollPass),SROAPass<modify-cfg>,VectorCombinePass,MergedLoadStoreMotionPass<no-split-footer-bb>,GVNPass<>,SCCPPass,BDCEPass,InstCombinePass<max-iterations=1;no-use-loop-info;no-verify-fixpoint>,AMDGPUUseNativeCallsPass,AMDGPUSimplifyLibCallsPass,JumpThreadingPass,CorrelatedValuePropagationPass,ADCEPass,MemCpyOptPass,DSEPass,MoveAutoInitPass,loop-mssa(LICMPass<allowspeculation>),CoroElidePass,SimplifyCFGPass<bonus-inst-threshold=1;no-forward-switch-cond;switch-range-to-icmp;no-switch-to-lookup;keep-loops;hoist-common-insts;sink-common-insts;speculate-blocks;simplify-cond-branch;no-speculate-unpredictables>,InstCombinePass<max-iterations=1;no-use-loop-info;no-verify-fixpoint>,AMDGPUUseNativeCallsPass,AMDGPUSimplifyLibCallsPass),PostOrderFunctionAttrsPass,function(require<ShouldNotRunFunctionPassesAnalysis>),CoroSplitPass)),function(invalidate<ShouldNotRunFunctionPassesAnalysis>),cgscc(devirt<4>())" on module "LLVMDialectModule"
# | 2.	Running pass "cgscc(devirt<4>(InlinerPass,PostOrderFunctionAttrsPass<skip-non-recursive-function-attrs>,OpenMPOptCGSCCPass,function(AMDGPUPromoteKernelArgumentsPass,InferAddressSpacesPass,AMDGPULowerKernelAttributesPass,AMDGPUPromoteAllocaToVectorPass),function<eager-inv;no-rerun>(SROAPass<modify-cfg>,EarlyCSEPass<memssa>,SpeculativeExecutionPass<only-if-divergent-target>,JumpThreadingPass,CorrelatedValuePropagationPass,SimplifyCFGPass<bonus-inst-threshold=1;no-forward-switch-cond;switch-range-to-icmp;no-switch-to-lookup;keep-loops;no-hoist-common-insts;no-sink-common-insts;speculate-blocks;simplify-cond-branch;no-speculate-unpredictables>,InstCombinePass<max-iterations=1;no-use-loop-info;no-verify-fixpoint>,AggressiveInstCombinePass,LibCallsShrinkWrapPass,AMDGPUUseNativeCallsPass,AMDGPUSimplifyLibCallsPass,TailCallElimPass,SimplifyCFGPass<bonus-inst-threshold=1;no-forward-switch-cond;switch-range-to-icmp;no-switch-to-lookup;keep-loops;no-hoist-common-insts;no-sink-common-insts;speculate-blocks;simplify-cond-branch;no-speculate-unpredictables>,ReassociatePass,ConstraintEliminationPass,loop-mssa(LoopInstSimplifyPass,LoopSimplifyCFGPass,LICMPass<no-allowspeculation>,LoopRotatePass<header-duplication;no-prepare-for-lto>,LICMPass<allowspeculation>,SimpleLoopUnswitchPass<no-nontrivial;trivial>),SimplifyCFGPass<bonus-inst-threshold=1;no-forward-switch-cond;switch-range-to-icmp;no-switch-to-lookup;keep-loops;no-hoist-common-insts;no-sink-common-insts;speculate-blocks;simplify-cond-branch;no-speculate-unpredictables>,InstCombinePass<max-iterations=1;no-use-loop-info;no-verify-fixpoint>,loop(LoopIdiomRecognizePass,IndVarSimplifyPass,SimpleLoopUnswitchPass<no-nontrivial;trivial>,LoopDeletionPass,LoopFullUnrollPass),SROAPass<modify-cfg>,VectorCombinePass,MergedLoadStoreMotionPass<no-split-footer-bb>,GVNPass<>,SCCPPass,BDCEPass,InstCombinePass<max-iterations=1;no-use-loop-info;no-verify-fixpoint>,AMDGPUUseNativeCallsPass,AMDGPUSimplifyLibCallsPass,JumpThreadingPass,CorrelatedValuePropagationPass,ADCEPass,MemCpyOptPass,DSEPass,MoveAutoInitPass,loop-mssa(LICMPass<allowspeculation>),CoroElidePass,SimplifyCFGPass<bonus-inst-threshold=1;no-forward-switch-cond;switch-range-to-icmp;no-switch-to-lookup;keep-loops;hoist-common-insts;sink-common-insts;speculate-blocks;simplify-cond-branch;no-speculate-unpredictables>,InstCombinePass<max-iterations=1;no-use-loop-info;no-verify-fixpoint>,AMDGPUUseNativeCallsPass,AMDGPUSimplifyLibCallsPass),PostOrderFunctionAttrsPass,function(require<ShouldNotRunFunctionPassesAnalysis>),CoroSplitPass))" on module "LLVMDialectModule"
# | 3.	Running pass "InferAddressSpacesPass" on function "__ockl_printf_append_string_n"
# |  #0 0x0000555ef1551d78 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/vol/worker/mi200-buildbot/mlir-rocm-mi200/build/bin/mlir-opt+0x3f58d78)
# |  #1 0x0000555ef154f86e llvm::sys::RunSignalHandlers() (/vol/worker/mi200-buildbot/mlir-rocm-mi200/build/bin/mlir-opt+0x3f5686e)
# |  #2 0x0000555ef155273d SignalHandler(int) Signals.cpp:0:0
# |  #3 0x00007fcd3b77d420 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x14420)
# |  #4 0x00007fcd3b24000b raise (/lib/x86_64-linux-gnu/libc.so.6+0x4300b)
# |  #5 0x00007fcd3b21f859 abort (/lib/x86_64-linux-gnu/libc.so.6+0x22859)
# |  #6 0x00007fcd3b21f729 (/lib/x86_64-linux-gnu/libc.so.6+0x22729)
# |  #7 0x00007fcd3b230fd6 (/lib/x86_64-linux-gnu/libc.so.6+0x33fd6)
# |  #8 0x0000555ef628c770 (anonymous namespace)::InferAddressSpacesImpl::run(llvm::Function&) InferAddressSpaces.cpp:0:0
# |  #9 0x0000555ef628acc3 llvm::InferAddressSpacesPass::run(llvm::Function&, llvm::AnalysisManager<llvm::Function>&) (/vol/worker/mi200-buildbot/mlir-rocm-mi200/build/bin/mlir-opt+0x8c91cc3)
# | #10 0x0000555ef4b76d3d llvm::detail::PassModel<llvm::Function, llvm::InferAddressSpacesPass, llvm::AnalysisManager<llvm::Function>>::run(llvm::Function&, llvm::AnalysisManager<llvm::Function>&) AMDGPUTargetMachine.cpp:0:0
# | #11 0x0000555ef2a432aa llvm::PassManager<llvm::Function, llvm::AnalysisManager<llvm::Function>>::run(llvm::Function&, llvm::AnalysisManager<llvm::Function>&) (/vol/worker/mi200-buildbot/mlir-rocm-mi200/build/bin/mlir-opt+0x544a2aa)
# | #12 0x0000555ef49f8ead llvm::detail::PassModel<llvm::Function, llvm::PassManager<llvm::Function, llvm::AnalysisManager<llvm::Function>>, llvm::AnalysisManager<llvm::Function>>::run(llvm::Function&, llvm::AnalysisManager<llvm::Function>&) NVPTXTargetMachine.cpp:0:0
# | #13 0x0000555ef6822ac9 llvm::CGSCCToFunctionPassAdaptor::run(llvm::LazyCallGraph::SCC&, llvm::AnalysisManager<llvm::LazyCallGraph::SCC, llvm::LazyCallGraph&>&, llvm::LazyCallGraph&, llvm::CGSCCUpdateResult&) (/vol/worker/mi200-buildbot/mlir-rocm-mi200/build/bin/mlir-opt+0x9229ac9)
# | #14 0x0000555ef4b770dd llvm::detail::PassModel<llvm::LazyCallGraph::SCC, llvm::CGSCCToFunctionPassAdaptor, llvm::AnalysisManager<llvm::LazyCallGraph::SCC, llvm::LazyCallGraph&>, llvm::LazyCallGraph&, llvm::CGSCCUpdateResult&>::run(llvm::LazyCallGraph::SCC&, llvm::AnalysisManager<llvm::LazyCallGraph::SCC, llvm::LazyCallGraph&>&, llvm::LazyCallGraph&, llvm::CGSCCUpdateResult&) AMDGPUTargetMachine.cpp:0:0
# | #15 0x0000555ef681e200 llvm::PassManager<llvm::LazyCallGraph::SCC, llvm::AnalysisManager<llvm::LazyCallGraph::SCC, llvm::LazyCallGraph&>, llvm::LazyCallGraph&, llvm::CGSCCUpdateResult&>::run(llvm::LazyCallGraph::SCC&, llvm::AnalysisManager<llvm::LazyCallGraph::SCC, llvm::LazyCallGraph&>&, llvm::LazyCallGraph&, llvm::CGSCCUpdateResult&) (/vol/worker/mi200-buildbot/mlir-rocm-mi200/build/bin/mlir-opt+0x9225200)
# | #16 0x0000555ef55a251d llvm::detail::PassModel<llvm::LazyCallGraph::SCC, llvm::PassManager<llvm::LazyCallGraph::SCC, llvm::AnalysisManager<llvm::LazyCallGraph::SCC, llvm::LazyCallGraph&>, llvm::LazyCallGraph&, llvm::CGSCCUpdateResult&>, llvm::AnalysisManager<llvm::LazyCallGraph::SCC, llvm::LazyCallGraph&>, llvm::LazyCallGraph&, llvm::CGSCCUpdateResult&>::run(llvm::LazyCallGraph::SCC&, llvm::AnalysisManager<llvm::LazyCallGraph::SCC, llvm::LazyCallGraph&>&, llvm::LazyCallGraph&, llvm::CGSCCUpdateResult&) PassBuilder.cpp:0:0
# | #17 0x0000555ef6820dd8 llvm::DevirtSCCRepeatedPass::run(llvm::LazyCallGraph::SCC&, llvm::AnalysisManager<llvm::LazyCallGraph::SCC, llvm::LazyCallGraph&>&, llvm::LazyCallGraph&, llvm::CGSCCUpdateResult&) (/vol/worker/mi200-buildbot/mlir-rocm-mi200/build/bin/mlir-opt+0x9227dd8)
# | #18 0x0000555ef55c3f6d llvm::detail::PassModel<llvm::LazyCallGraph::SCC, llvm::DevirtSCCRepeatedPass, llvm::AnalysisManager<llvm::LazyCallGraph::SCC, llvm::LazyCallGraph&>, llvm::LazyCallGraph&, llvm::CGSCCUpdateResult&>::run(llvm::LazyCallGraph::SCC&, llvm::AnalysisManager<llvm::LazyCallGraph::SCC, llvm::LazyCallGraph&>&, llvm::LazyCallGraph&, llvm::CGSCCUpdateResult&) PassBuilder.cpp:0:0
# | #19 0x0000555ef681fbeb llvm::ModuleToPostOrderCGSCCPassAdaptor::run(llvm::Module&, llvm::AnalysisManager<llvm::Module>&) (/vol/worker/mi200-buildbot/mlir-rocm-mi200/build/bin/mlir-opt+0x9226beb)
# | #20 0x0000555ef55a27bd llvm::detail::PassModel<llvm::Module, llvm::ModuleToPostOrderCGSCCPassAdaptor, llvm::AnalysisManager<llvm::Module>>::run(llvm::Module&, llvm::AnalysisManager<llvm::Module>&) PassBuilder.cpp:0:0
# | #21 0x0000555ef2a41f7a llvm::PassManager<llvm::Module, llvm::AnalysisManager<llvm::Module>>::run(llvm::Module&, llvm::AnalysisManager<llvm::Module>&) (/vol/worker/mi200-buildbot/mlir-rocm-mi200/build/bin/mlir-opt+0x5448f7a)
# | #22 0x0000555ef5782ec0 llvm::ModuleInlinerWrapperPass::run(llvm::Module&, llvm::AnalysisManager<llvm::Module>&) (/vol/worker/mi200-buildbot/mlir-rocm-mi200/build/bin/mlir-opt+0x8189ec0)
# | #23 0x0000555ef55a98dd llvm::detail::PassModel<llvm::Module, llvm::ModuleInlinerWrapperPass, llvm::AnalysisManager<llvm::Module>>::run(llvm::Module&, llvm::AnalysisManager<llvm::Module>&) PassBuilder.cpp:0:0
# | #24 0x0000555ef2a41f7a llvm::PassManager<llvm::Module, llvm::AnalysisManager<llvm::Module>>::run(llvm::Module&, llvm::AnalysisManager<llvm::Module>&) (/vol/worker/mi200-buildbot/mlir-rocm-mi200/build/bin/mlir-opt+0x5448f7a)
# | #25 0x0000555ef55286d1 std::_Function_handler<llvm::Error (llvm::Module*), mlir::makeOptimizingTransformer(unsigned int, unsigned int, llvm::TargetMachine*)::$_0>::_M_invoke(std::_Any_data const&, llvm::Module*&&) OptUtils.cpp:0:0
# | #26 0x0000555ef4a3cc20 mlir::LLVM::ModuleToObject::optimizeModule(llvm::Module&, int) (/vol/worker/mi200-buildbot/mlir-rocm-mi200/build/bin/mlir-opt+0x7443c20)
# | #27 0x0000555ef4a3d9a8 mlir::LLVM::ModuleToObject::run() (/vol/worker/mi200-buildbot/mlir-rocm-mi200/build/bin/mlir-opt+0x74449a8)
# | #28 0x0000555ef3d3944e mlir::gpu::detail::TargetAttrInterfaceInterfaceTraits::FallbackModel<(anonymous namespace)::ROCDLTargetAttrImpl>::serializeToObject(mlir::gpu::detail::TargetAttrInterfaceInterfaceTraits::Concept const*, mlir::Attribute, mlir::Operation*, mlir::gpu::TargetOptions const&) Target.cpp:0:0
# | #29 0x0000555ef1bfdfe3 mlir::gpu::TargetAttrInterface::serializeToObject(mlir::Operation*, mlir::gpu::TargetOptions const&) const (/vol/worker/mi200-buildbot/mlir-rocm-mi200/build/bin/mlir-opt+0x4604fe3)
# | #30 0x0000555ef1c86d13 mlir::gpu::transformGpuModulesToBinaries(mlir::Operation*, mlir::gpu::OffloadingLLVMTranslationAttrInterface, mlir::gpu::TargetOptions const&) (/vol/worker/mi200-buildbot/mlir-rocm-mi200/build/bin/mlir-opt+0x468dd13)
# | #31 0x0000555ef1c87ebf (anonymous namespace)::GpuModuleToBinaryPass::runOnOperation() ModuleToBinary.cpp:0:0
...

```

https://github.com/llvm/llvm-project/pull/104052