iree: Segmentation fault when compiling some models with `--iree-codegen-llvmgpu-use-mma-sync`
What happened?
When compiling some PyTorch models with --iree-codegen-llvmgpu-use-mma-sync
, seeing this crash:
Please report issues to https://github.com/openxla/iree/issues and include the crash backtrace.
#0 0x00007f4ed960bd9b llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) /work/third_party/llvm-project/llvm/lib/Support/Unix/Signals.inc:602:13
#1 0x00007f4ed9609fd0 llvm::sys::RunSignalHandlers() /work/third_party/llvm-project/llvm/lib/Support/Signals.cpp:105:18
#2 0x00007f4ed960c42f SignalHandler(int) /work/third_party/llvm-project/llvm/lib/Support/Unix/Signals.inc:413:1
#3 0x00007f4ed9adb980 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x12980)
#4 0x00007f4ed6d56417 mlir::detail::IROperandBase::getOwner() const /work/third_party/llvm-project/mlir/include/mlir/IR/UseDefLists.h:40:40
#5 0x00007f4ed6d56417 mlir::ValueUserIterator<mlir::ResultRange::UseIterator, mlir::OpOperand>::mapElement(mlir::OpOperand&) const /work/third_party/llvm-project/mlir/include/mlir/IR/UseDefLists.h:298:66
#6 0x00007f4ed6d56417 llvm::mapped_iterator_base<mlir::ValueUserIterator<mlir::ResultRange::UseIterator, mlir::OpOperand>, mlir::ResultRange::UseIterator, mlir::Operation*>::operator*() const /work/third_party/llvm-project/llvm/include/llvm/ADT/STLExtras.h:487:49
#7 0x00007f4ed6d56417 mlir::iree_compiler::getVectorContractOpOperandIdForVectorReadOp(mlir::Operation*) /work/compiler/src/iree/compiler/Codegen/Utils/GPUUtils.cpp:619:32
#8 0x00007f4ed6d56417 mlir::iree_compiler::getMmaNativeVectorSize(mlir::Operation*) /work/compiler/src/iree/compiler/Codegen/Utils/GPUUtils.cpp:669:9
#9 0x00007f4ed5cbbbec mlir::iree_compiler::populateVectorUnrollPatterns(mlir::RewritePatternSet&, bool)::$_0::operator()(mlir::Operation*) const /work/compiler/src/iree/compiler/Codegen/LLVMGPU/LLVMGPUTensorCoreVectorization.cpp:0:33
#10 0x00007f4ed5cbbbec std::_Function_handler<std::optional<llvm::SmallVector<long, 6u> > (mlir::Operation*), mlir::iree_compiler::populateVectorUnrollPatterns(mlir::RewritePatternSet&, bool)::$_0>::_M_invoke(std::_Any_data const&, mlir::Operation*&&) /usr/bin/../lib/gcc/x86_64-linux-gnu/9/../../../../include/c++/9/bits/std_function.h:285:9
#11 0x00007f4ed71c899a std::_Optional_base_impl<llvm::SmallVector<long, 6u>, std::_Optional_base<llvm::SmallVector<long, 6u>, false, false> >::_M_is_engaged() const /usr/bin/../lib/gcc/x86_64-linux-gnu/9/../../../../include/c++/9/optional:432:58
#12 0x00007f4ed71c899a std::optional<llvm::SmallVector<long, 6u> >::operator bool() const /usr/bin/../lib/gcc/x86_64-linux-gnu/9/../../../../include/c++/9/optional:913:22
#13 0x00007f4ed71c899a getTargetShape(mlir::vector::UnrollVectorOptions const&, mlir::Operation*) /work/third_party/llvm-project/mlir/lib/Dialect/Vector/Transforms/VectorUnroll.cpp:152:8
#14 0x00007f4ed71c7e9b std::_Optional_base_impl<llvm::SmallVector<long, 6u>, std::_Optional_base<llvm::SmallVector<long, 6u>, false, false> >::_M_is_engaged() const /usr/bin/../lib/gcc/x86_64-linux-gnu/9/../../../../include/c++/9/optional:432:58
#15 0x00007f4ed71c7e9b std::optional<llvm::SmallVector<long, 6u> >::operator bool() const /usr/bin/../lib/gcc/x86_64-linux-gnu/9/../../../../include/c++/9/optional:913:22
#16 0x00007f4ed71c7e9b (anonymous namespace)::UnrollTransferReadPattern::matchAndRewrite(mlir::vector::TransferReadOp, mlir::PatternRewriter&) const /work/third_party/llvm-project/mlir/lib/Dialect/Vector/Transforms/VectorUnroll.cpp:194:10
#17 0x00007f4ed8e9a113 mlir::PatternApplicator::matchAndRewrite(mlir::Operation*, mlir::PatternRewriter&, llvm::function_ref<bool (mlir::Pattern const&)>, llvm::function_ref<void (mlir::Pattern const&)>, llvm::function_ref<mlir::LogicalResult (mlir::Pattern const&)>) /work/third_party/llvm-project/mlir/lib/Rewrite/PatternApplicator.cpp:200:25
#18 0x00007f4ed8e66436 (anonymous namespace)::GreedyPatternRewriteDriver::processWorklist() /work/third_party/llvm-project/mlir/lib/Transforms/Utils/GreedyPatternRewriteDriver.cpp:229:19
#19 0x00007f4ed8e6315d (anonymous namespace)::RegionPatternRewriteDriver::simplify() && /work/third_party/llvm-project/mlir/lib/Transforms/Utils/GreedyPatternRewriteDriver.cpp:458:15
#20 0x00007f4ed8e6315d mlir::applyPatternsAndFoldGreedily(mlir::Region&, mlir::FrozenRewritePatternSet const&, mlir::GreedyRewriteConfig) /work/third_party/llvm-project/mlir/lib/Transforms/Utils/GreedyPatternRewriteDriver.cpp:487:47
#21 0x00007f4ed5cbb8d0 mlir::LogicalResult::failed() const /work/third_party/llvm-project/mlir/include/mlir/Support/LogicalResult.h:44:33
#22 0x00007f4ed5cbb8d0 mlir::applyPatternsAndFoldGreedily(mlir::Operation*, mlir::FrozenRewritePatternSet const&, mlir::GreedyRewriteConfig) /work/third_party/llvm-project/mlir/include/mlir/Transforms/GreedyPatternRewriteDriver.h:115:70
#23 0x00007f4ed5cbb8d0 mlir::iree_compiler::(anonymous namespace)::LLVMGPUTensorCoreVectorizationPass::runOnOperation() /work/compiler/src/iree/compiler/Codegen/LLVMGPU/LLVMGPUTensorCoreVectorization.cpp:113:18
#24 0x00007f4ed8ec7c96 mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int)::$_7::operator()() const /work/third_party/llvm-project/mlir/lib/Pass/Pass.cpp:482:17
#25 0x00007f4ed8ec7c96 void llvm::function_ref<void ()>::callback_fn<mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int)::$_7>(long) /work/third_party/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:45:12
#26 0x00007f4ed8ec7c96 llvm::function_ref<void ()>::operator()() const /work/third_party/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:68:12
#27 0x00007f4ed8ec7c96 void mlir::MLIRContext::executeAction<mlir::PassExecutionAction, mlir::Pass&>(llvm::function_ref<void ()>, llvm::ArrayRef<mlir::IRUnit>, mlir::Pass&) /work/third_party/llvm-project/mlir/include/mlir/IR/MLIRContext.h:275:7
#28 0x00007f4ed8ec7c96 mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) /work/third_party/llvm-project/mlir/lib/Pass/Pass.cpp:476:21
#29 0x00007f4ed8ec8558 mlir::LogicalResult::failed() const /work/third_party/llvm-project/mlir/include/mlir/Support/LogicalResult.h:44:33
#30 0x00007f4ed8ec8558 mlir::failed(mlir::LogicalResult) /work/third_party/llvm-project/mlir/include/mlir/Support/LogicalResult.h:72:58
#31 0x00007f4ed8ec8558 mlir::detail::OpToOpPassAdaptor::runPipeline(mlir::OpPassManager&, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::PipelineParentInfo const*) /work/third_party/llvm-project/mlir/lib/Pass/Pass.cpp:548:9
#32 0x00007f4ed8ecdece mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::$_15::operator()(mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo&) const /work/third_party/llvm-project/mlir/lib/Pass/Pass.cpp:773:5
#33 0x00007f4ed8ec9c0b mlir::LogicalResult::failed() const /work/third_party/llvm-project/mlir/include/mlir/Support/LogicalResult.h:44:33
#34 0x00007f4ed8ec9c0b mlir::failed(mlir::LogicalResult) /work/third_party/llvm-project/mlir/include/mlir/Support/LogicalResult.h:72:58
#35 0x00007f4ed8ec9c0b mlir::LogicalResult mlir::failableParallelForEach<__gnu_cxx::__normal_iterator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo*, std::vector<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo, std::allocator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo> > >, mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::$_15&>(mlir::MLIRContext*, __gnu_cxx::__normal_iterator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo*, std::vector<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo, std::allocator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo> > >, __gnu_cxx::__normal_iterator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo*, std::vector<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo, std::allocator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo> > >, mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::$_15&) /work/third_party/llvm-project/mlir/include/mlir/IR/Threading.h:46:11
#36 0x00007f4ed8ec9c0b mlir::LogicalResult mlir::failableParallelForEach<std::vector<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo, std::allocator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo> >&, mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::$_15&>(mlir::MLIRContext*, std::vector<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo, std::allocator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo> >&, mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::$_15&) /work/third_party/llvm-project/mlir/include/mlir/IR/Threading.h:92:10
#37 0x00007f4ed8ec9c0b mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool) /work/third_party/llvm-project/mlir/lib/Pass/Pass.cpp:778:14
#38 0x00007f4ed8ec7de8 mlir::detail::OpToOpPassAdaptor::runOnOperation(bool) /work/third_party/llvm-project/mlir/lib/Pass/Pass.cpp:669:5
#39 0x00007f4ed8ec7de8 mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int)::$_7::operator()() const /work/third_party/llvm-project/mlir/lib/Pass/Pass.cpp:480:20
#40 0x00007f4ed8ec7de8 void llvm::function_ref<void ()>::callback_fn<mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int)::$_7>(long) /work/third_party/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:45:12
#41 0x00007f4ed8ec7de8 llvm::function_ref<void ()>::operator()() const /work/third_party/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:68:12
#42 0x00007f4ed8ec7de8 void mlir::MLIRContext::executeAction<mlir::PassExecutionAction, mlir::Pass&>(llvm::function_ref<void ()>, llvm::ArrayRef<mlir::IRUnit>, mlir::Pass&) /work/third_party/llvm-project/mlir/include/mlir/IR/MLIRContext.h:275:7
#43 0x00007f4ed8ec7de8 mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) /work/third_party/llvm-project/mlir/lib/Pass/Pass.cpp:476:21
#44 0x00007f4ed8ec8558 mlir::LogicalResult::failed() const /work/third_party/llvm-project/mlir/include/mlir/Support/LogicalResult.h:44:33
#45 0x00007f4ed8ec8558 mlir::failed(mlir::LogicalResult) /work/third_party/llvm-project/mlir/include/mlir/Support/LogicalResult.h:72:58
#46 0x00007f4ed8ec8558 mlir::detail::OpToOpPassAdaptor::runPipeline(mlir::OpPassManager&, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::PipelineParentInfo const*) /work/third_party/llvm-project/mlir/lib/Pass/Pass.cpp:548:9
#47 0x00007f4ed8ecdece mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::$_15::operator()(mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo&) const /work/third_party/llvm-project/mlir/lib/Pass/Pass.cpp:773:5
#48 0x00007f4ed8ec9c0b mlir::LogicalResult::failed() const /work/third_party/llvm-project/mlir/include/mlir/Support/LogicalResult.h:44:33
#49 0x00007f4ed8ec9c0b mlir::failed(mlir::LogicalResult) /work/third_party/llvm-project/mlir/include/mlir/Support/LogicalResult.h:72:58
#50 0x00007f4ed8ec9c0b mlir::LogicalResult mlir::failableParallelForEach<__gnu_cxx::__normal_iterator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo*, std::vector<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo, std::allocator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo> > >, mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::$_15&>(mlir::MLIRContext*, __gnu_cxx::__normal_iterator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo*, std::vector<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo, std::allocator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo> > >, __gnu_cxx::__normal_iterator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo*, std::vector<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo, std::allocator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo> > >, mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::$_15&) /work/third_party/llvm-project/mlir/include/mlir/IR/Threading.h:46:11
#51 0x00007f4ed8ec9c0b mlir::LogicalResult mlir::failableParallelForEach<std::vector<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo, std::allocator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo> >&, mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::$_15&>(mlir::MLIRContext*, std::vector<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo, std::allocator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo> >&, mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::$_15&) /work/third_party/llvm-project/mlir/include/mlir/IR/Threading.h:92:10
#52 0x00007f4ed8ec9c0b mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool) /work/third_party/llvm-project/mlir/lib/Pass/Pass.cpp:778:14
#53 0x00007f4ed8ec7de8 mlir::detail::OpToOpPassAdaptor::runOnOperation(bool) /work/third_party/llvm-project/mlir/lib/Pass/Pass.cpp:669:5
#54 0x00007f4ed8ec7de8 mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int)::$_7::operator()() const /work/third_party/llvm-project/mlir/lib/Pass/Pass.cpp:480:20
#55 0x00007f4ed8ec7de8 void llvm::function_ref<void ()>::callback_fn<mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int)::$_7>(long) /work/third_party/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:45:12
#56 0x00007f4ed8ec7de8 llvm::function_ref<void ()>::operator()() const /work/third_party/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:68:12
#57 0x00007f4ed8ec7de8 void mlir::MLIRContext::executeAction<mlir::PassExecutionAction, mlir::Pass&>(llvm::function_ref<void ()>, llvm::ArrayRef<mlir::IRUnit>, mlir::Pass&) /work/third_party/llvm-project/mlir/include/mlir/IR/MLIRContext.h:275:7
#58 0x00007f4ed8ec7de8 mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) /work/third_party/llvm-project/mlir/lib/Pass/Pass.cpp:476:21
#59 0x00007f4ed8ec8558 mlir::LogicalResult::failed() const /work/third_party/llvm-project/mlir/include/mlir/Support/LogicalResult.h:44:33
#60 0x00007f4ed8ec8558 mlir::failed(mlir::LogicalResult) /work/third_party/llvm-project/mlir/include/mlir/Support/LogicalResult.h:72:58
#61 0x00007f4ed8ec8558 mlir::detail::OpToOpPassAdaptor::runPipeline(mlir::OpPassManager&, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::PipelineParentInfo const*) /work/third_party/llvm-project/mlir/lib/Pass/Pass.cpp:548:9
#62 0x00007f4ed8eccd21 mlir::LogicalResult llvm::function_ref<mlir::LogicalResult (mlir::OpPassManager&, mlir::Operation*)>::callback_fn<mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int)::$_6>(long, mlir::OpPassManager&, mlir::Operation*) /work/third_party/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:45:5
#63 0x00007f4ed5cb0ae8 mlir::LogicalResult::failed() const /work/third_party/llvm-project/mlir/include/mlir/Support/LogicalResult.h:44:33
#64 0x00007f4ed5cb0ae8 mlir::failed(mlir::LogicalResult) /work/third_party/llvm-project/mlir/include/mlir/Support/LogicalResult.h:72:58
#65 0x00007f4ed5cb0ae8 mlir::iree_compiler::(anonymous namespace)::LLVMGPULowerExecutableTargetPass::runOnOperation() /work/compiler/src/iree/compiler/Codegen/LLVMGPU/LLVMGPULowerExecutableTarget.cpp:191:7
#66 0x00007f4ed8ec7c96 mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int)::$_7::operator()() const /work/third_party/llvm-project/mlir/lib/Pass/Pass.cpp:482:17
#67 0x00007f4ed8ec7c96 void llvm::function_ref<void ()>::callback_fn<mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int)::$_7>(long) /work/third_party/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:45:12
#68 0x00007f4ed8ec7c96 llvm::function_ref<void ()>::operator()() const /work/third_party/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:68:12
#69 0x00007f4ed8ec7c96 void mlir::MLIRContext::executeAction<mlir::PassExecutionAction, mlir::Pass&>(llvm::function_ref<void ()>, llvm::ArrayRef<mlir::IRUnit>, mlir::Pass&) /work/third_party/llvm-project/mlir/include/mlir/IR/MLIRContext.h:275:7
#70 0x00007f4ed8ec7c96 mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) /work/third_party/llvm-project/mlir/lib/Pass/Pass.cpp:476:21
#71 0x00007f4ed8ec8558 mlir::LogicalResult::failed() const /work/third_party/llvm-project/mlir/include/mlir/Support/LogicalResult.h:44:33
#72 0x00007f4ed8ec8558 mlir::failed(mlir::LogicalResult) /work/third_party/llvm-project/mlir/include/mlir/Support/LogicalResult.h:72:58
#73 0x00007f4ed8ec8558 mlir::detail::OpToOpPassAdaptor::runPipeline(mlir::OpPassManager&, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::PipelineParentInfo const*) /work/third_party/llvm-project/mlir/lib/Pass/Pass.cpp:548:9
#74 0x00007f4ed8eccd21 mlir::LogicalResult llvm::function_ref<mlir::LogicalResult (mlir::OpPassManager&, mlir::Operation*)>::callback_fn<mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int)::$_6>(long, mlir::OpPassManager&, mlir::Operation*) /work/third_party/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:45:5
#75 0x00007f4ed537506b mlir::LogicalResult::failed() const /work/third_party/llvm-project/mlir/include/mlir/Support/LogicalResult.h:44:33
#76 0x00007f4ed537506b mlir::failed(mlir::LogicalResult) /work/third_party/llvm-project/mlir/include/mlir/Support/LogicalResult.h:72:58
#77 0x00007f4ed537506b mlir::iree_compiler::IREE::HAL::TranslateTargetExecutableVariantsPass::runOnOperation() /work/compiler/src/iree/compiler/Dialect/HAL/Transforms/TranslateExecutables.cpp:68:9
#78 0x00007f4ed8ec7c96 mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int)::$_7::operator()() const /work/third_party/llvm-project/mlir/lib/Pass/Pass.cpp:482:17
#79 0x00007f4ed8ec7c96 void llvm::function_ref<void ()>::callback_fn<mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int)::$_7>(long) /work/third_party/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:45:12
#80 0x00007f4ed8ec7c96 llvm::function_ref<void ()>::operator()() const /work/third_party/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:68:12
#81 0x00007f4ed8ec7c96 void mlir::MLIRContext::executeAction<mlir::PassExecutionAction, mlir::Pass&>(llvm::function_ref<void ()>, llvm::ArrayRef<mlir::IRUnit>, mlir::Pass&) /work/third_party/llvm-project/mlir/include/mlir/IR/MLIRContext.h:275:7
#82 0x00007f4ed8ec7c96 mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) /work/third_party/llvm-project/mlir/lib/Pass/Pass.cpp:476:21
#83 0x00007f4ed8ec8558 mlir::LogicalResult::failed() const /work/third_party/llvm-project/mlir/include/mlir/Support/LogicalResult.h:44:33
#84 0x00007f4ed8ec8558 mlir::failed(mlir::LogicalResult) /work/third_party/llvm-project/mlir/include/mlir/Support/LogicalResult.h:72:58
#85 0x00007f4ed8ec8558 mlir::detail::OpToOpPassAdaptor::runPipeline(mlir::OpPassManager&, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::PipelineParentInfo const*) /work/third_party/llvm-project/mlir/lib/Pass/Pass.cpp:548:9
#86 0x00007f4ed8ecdece mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::$_15::operator()(mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo&) const /work/third_party/llvm-project/mlir/lib/Pass/Pass.cpp:773:5
#87 0x00007f4ed8ec9c0b mlir::LogicalResult::failed() const /work/third_party/llvm-project/mlir/include/mlir/Support/LogicalResult.h:44:33
#88 0x00007f4ed8ec9c0b mlir::failed(mlir::LogicalResult) /work/third_party/llvm-project/mlir/include/mlir/Support/LogicalResult.h:72:58
#89 0x00007f4ed8ec9c0b mlir::LogicalResult mlir::failableParallelForEach<__gnu_cxx::__normal_iterator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo*, std::vector<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo, std::allocator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo> > >, mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::$_15&>(mlir::MLIRContext*, __gnu_cxx::__normal_iterator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo*, std::vector<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo, std::allocator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo> > >, __gnu_cxx::__normal_iterator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo*, std::vector<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo, std::allocator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo> > >, mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::$_15&) /work/third_party/llvm-project/mlir/include/mlir/IR/Threading.h:46:11
#90 0x00007f4ed8ec9c0b mlir::LogicalResult mlir::failableParallelForEach<std::vector<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo, std::allocator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo> >&, mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::$_15&>(mlir::MLIRContext*, std::vector<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo, std::allocator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo> >&, mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::$_15&) /work/third_party/llvm-project/mlir/include/mlir/IR/Threading.h:92:10
#91 0x00007f4ed8ec9c0b mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool) /work/third_party/llvm-project/mlir/lib/Pass/Pass.cpp:778:14
#92 0x00007f4ed8ec7de8 mlir::detail::OpToOpPassAdaptor::runOnOperation(bool) /work/third_party/llvm-project/mlir/lib/Pass/Pass.cpp:669:5
#93 0x00007f4ed8ec7de8 mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int)::$_7::operator()() const /work/third_party/llvm-project/mlir/lib/Pass/Pass.cpp:480:20
#94 0x00007f4ed8ec7de8 void llvm::function_ref<void ()>::callback_fn<mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int)::$_7>(long) /work/third_party/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:45:12
#95 0x00007f4ed8ec7de8 llvm::function_ref<void ()>::operator()() const /work/third_party/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:68:12
#96 0x00007f4ed8ec7de8 void mlir::MLIRContext::executeAction<mlir::PassExecutionAction, mlir::Pass&>(llvm::function_ref<void ()>, llvm::ArrayRef<mlir::IRUnit>, mlir::Pass&) /work/third_party/llvm-project/mlir/include/mlir/IR/MLIRContext.h:275:7
#97 0x00007f4ed8ec7de8 mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) /work/third_party/llvm-project/mlir/lib/Pass/Pass.cpp:476:21
#98 0x00007f4ed8ec8558 mlir::LogicalResult::failed() const /work/third_party/llvm-project/mlir/include/mlir/Support/LogicalResult.h:44:33
#99 0x00007f4ed8ec8558 mlir::failed(mlir::LogicalResult) /work/third_party/llvm-project/mlir/include/mlir/Support/LogicalResult.h:72:58
#100 0x00007f4ed8ec8558 mlir::detail::OpToOpPassAdaptor::runPipeline(mlir::OpPassManager&, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::PipelineParentInfo const*) /work/third_party/llvm-project/mlir/lib/Pass/Pass.cpp:548:9
#101 0x00007f4ed8eccd21 mlir::LogicalResult llvm::function_ref<mlir::LogicalResult (mlir::OpPassManager&, mlir::Operation*)>::callback_fn<mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int)::$_6>(long, mlir::OpPassManager&, mlir::Operation*) /work/third_party/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:45:5
#102 0x00007f4ed5375c49 mlir::LogicalResult::failed() const /work/third_party/llvm-project/mlir/include/mlir/Support/LogicalResult.h:44:33
#103 0x00007f4ed5375c49 mlir::failed(mlir::LogicalResult) /work/third_party/llvm-project/mlir/include/mlir/Support/LogicalResult.h:72:58
#104 0x00007f4ed5375c49 mlir::iree_compiler::IREE::HAL::TranslateExecutablesPass::runOnOperation() /work/compiler/src/iree/compiler/Dialect/HAL/Transforms/TranslateExecutables.cpp:125:9
#105 0x00007f4ed8ec7c96 mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int)::$_7::operator()() const /work/third_party/llvm-project/mlir/lib/Pass/Pass.cpp:482:17
#106 0x00007f4ed8ec7c96 void llvm::function_ref<void ()>::callback_fn<mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int)::$_7>(long) /work/third_party/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:45:12
#107 0x00007f4ed8ec7c96 llvm::function_ref<void ()>::operator()() const /work/third_party/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:68:12
#108 0x00007f4ed8ec7c96 void mlir::MLIRContext::executeAction<mlir::PassExecutionAction, mlir::Pass&>(llvm::function_ref<void ()>, llvm::ArrayRef<mlir::IRUnit>, mlir::Pass&) /work/third_party/llvm-project/mlir/include/mlir/IR/MLIRContext.h:275:7
#109 0x00007f4ed8ec7c96 mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) /work/third_party/llvm-project/mlir/lib/Pass/Pass.cpp:476:21
#110 0x00007f4ed8ec8558 mlir::LogicalResult::failed() const /work/third_party/llvm-project/mlir/include/mlir/Support/LogicalResult.h:44:33
#111 0x00007f4ed8ec8558 mlir::failed(mlir::LogicalResult) /work/third_party/llvm-project/mlir/include/mlir/Support/LogicalResult.h:72:58
#112 0x00007f4ed8ec8558 mlir::detail::OpToOpPassAdaptor::runPipeline(mlir::OpPassManager&, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::PipelineParentInfo const*) /work/third_party/llvm-project/mlir/lib/Pass/Pass.cpp:548:9
#113 0x00007f4ed8ecdece mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::$_15::operator()(mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo&) const /work/third_party/llvm-project/mlir/lib/Pass/Pass.cpp:773:5
#114 0x00007f4ed8ecdf7f mlir::LogicalResult::failed() const /work/third_party/llvm-project/mlir/include/mlir/Support/LogicalResult.h:44:33
#115 0x00007f4ed8ecdf7f mlir::failed(mlir::LogicalResult) /work/third_party/llvm-project/mlir/include/mlir/Support/LogicalResult.h:72:58
#116 0x00007f4ed8ecdf7f mlir::failableParallelForEach<__gnu_cxx::__normal_iterator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo*, std::vector<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo, std::allocator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo> > >, mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::$_15&>(mlir::MLIRContext*, __gnu_cxx::__normal_iterator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo*, std::vector<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo, std::allocator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo> > >, __gnu_cxx::__normal_iterator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo*, std::vector<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo, std::allocator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo> > >, mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::$_15&)::{lambda()#1}::operator()() const /work/third_party/llvm-project/mlir/include/mlir/IR/Threading.h:62:11
#117 0x00007f4ed8ecdf7f std::_Function_handler<void (), mlir::failableParallelForEach<__gnu_cxx::__normal_iterator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo*, std::vector<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo, std::allocator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo> > >, mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::$_15&>(mlir::MLIRContext*, __gnu_cxx::__normal_iterator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo*, std::vector<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo, std::allocator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo> > >, __gnu_cxx::__normal_iterator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo*, std::vector<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo, std::allocator<mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::OpPMInfo> > >, mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool)::$_15&)::{lambda()#1}>::_M_invoke(std::_Any_data const&) /usr/bin/../lib/gcc/x86_64-linux-gnu/9/../../../../include/c++/9/bits/std_function.h:300:2
#118 0x00007f4ed3c32656 std::__shared_ptr<std::promise<void>, (__gnu_cxx::_Lock_policy)2>::get() const /usr/bin/../lib/gcc/x86_64-linux-gnu/9/../../../../include/c++/9/bits/shared_ptr_base.h:1310:16
#119 0x00007f4ed3c32656 std::__shared_ptr_access<std::promise<void>, (__gnu_cxx::_Lock_policy)2, false, false>::_M_get() const /usr/bin/../lib/gcc/x86_64-linux-gnu/9/../../../../include/c++/9/bits/shared_ptr_base.h:1021:66
#120 0x00007f4ed3c32656 std::__shared_ptr_access<std::promise<void>, (__gnu_cxx::_Lock_policy)2, false, false>::operator->() const /usr/bin/../lib/gcc/x86_64-linux-gnu/9/../../../../include/c++/9/bits/shared_ptr_base.h:1015:9
#121 0x00007f4ed3c32656 llvm::ThreadPool::createTaskAndFuture(std::function<void ()>)::{lambda()#1}::operator()() const /work/third_party/llvm-project/llvm/include/llvm/Support/ThreadPool.h:136:15
#122 0x00007f4ed3c32656 std::_Function_handler<void (), llvm::ThreadPool::createTaskAndFuture(std::function<void ()>)::{lambda()#1}>::_M_invoke(std::_Any_data const&) /usr/bin/../lib/gcc/x86_64-linux-gnu/9/../../../../include/c++/9/bits/std_function.h:300:2
#123 0x00007f4ed95bb3e2 llvm::ThreadPool::processTasks(llvm::ThreadPoolTaskGroup*) /work/third_party/llvm-project/llvm/lib/Support/ThreadPool.cpp:104:5
#124 0x00007f4ed95bc1cc void llvm::thread::GenericThreadProxy<std::tuple<llvm::ThreadPool::grow(int)::$_0> >(void*) /usr/bin/../lib/gcc/x86_64-linux-gnu/9/../../../../include/c++/9/bits/unique_ptr.h:291:12
#125 0x00007f4ed95bc1cc void* llvm::thread::ThreadProxy<std::tuple<llvm::ThreadPool::grow(int)::$_0> >(void*) /work/third_party/llvm-project/llvm/include/llvm/Support/thread.h:55:5
#126 0x00007f4ed9ad06db start_thread (/lib/x86_64-linux-gnu/libpthread.so.0+0x76db)
#127 0x00007f4ecf5e561f clone (/lib/x86_64-linux-gnu/libc.so.6+0x12161f)
Segmentation fault (core dumped)
Steps to reproduce your issue
iree-compile --output-format=vm-bytecode \
--iree-hal-target-backends=cuda \
--iree-input-type=none \
--iree-hal-cuda-llvm-target-arch=sm_80 \
--iree-codegen-llvmgpu-use-mma-sync \
linalg.mlir -o linalg.vmfb
What component(s) does this issue relate to?
Compiler
Version information
Commit https://github.com/openxla/iree/commit/35bdd9d16f9d91a9531c7d913b0fb242010e953f
Additional context
No response
About this issue
- Original URL
- State: closed
- Created a year ago
- Comments: 22 (20 by maintainers)
Commits related to this issue
- Fix batch_matmul + transpose using nvgpu.mma.sync.1688.f32 (#13075) (#13170) This PR fixes the segmentation fault on #13075 for `linalg.batch_matmul` with transpose on the output (see comment [here... — committed to iree-org/iree by manishucsd a year ago
- Fix batch_matmul + transpose using nvgpu.mma.sync.1688.f32 (#13075) (#13170) This PR fixes the segmentation fault on #13075 for `linalg.batch_matmul` with transpose on the output (see comment [here... — committed to iree-org/iree by manishucsd a year ago
- Fix batch_matmul + transpose using nvgpu.mma.sync.1688.f32 (#13075) (#13170) This PR fixes the segmentation fault on #13075 for `linalg.batch_matmul` with transpose on the output (see comment [here... — committed to NatashaKnk/iree by manishucsd a year ago
Tested all benchmarks in https://github.com/openxla/iree/pull/13208. All are passing with some nice latency reductions! All clear to enable by default.
Paraphrasing Stella’s comment from discord:
It’s just a python script now. https://github.com/openxla/iree/blob/main/integrations/tensorflow/python_projects/iree_tf/iree/tools/tf/scripts/iree_import_tf/__main__.py You can follow https://github.com/openxla/iree/tree/main/integrations/tensorflow#quick-development-setup but will need to use tf-nightly. The one the CI uses is here: https://github.com/openxla/iree/blob/main/integrations/tensorflow/test/requirements.txt (install using
pip install -r integrations/tensorflow/test/requirements.txt
).