search

microsoft / triton-shared

Shared Middle-Layer for Triton Compilation
MIT License
142 stars 27 forks source link

Assert: No support for multiple modulos within an expression #14

Closed manbearian closed 8 months ago

manbearian commented 9 months ago

created from #7.

Limitation in supporting modulo of pointers.

repro.zip

triton-shared-opt -triton-to-linalg 11.mlir triton-shared-opt -triton-to-linalg 12.mlir triton-shared-opt -triton-to-linalg 16.mlir triton-shared-opt -triton-to-linalg 21.mlir triton-shared-opt -triton-to-linalg 36.mlir triton-shared-opt -triton-to-linalg 38.mlir triton-shared-opt -triton-to-linalg 44.mlir triton-shared-opt -triton-to-linalg 66.mlir triton-shared-opt -triton-to-linalg 82.mlir triton-shared-opt -triton-to-linalg 92.mlir

Error output:

+++/home/ianb/test/ttirs_linalg_failed/11.mlir
triton-shared-opt: /home/ianb/src/triton/third_party/triton_shared/lib/Analysis/PtrAnalysis.cpp:369: static void mlir::triton::PtrAnalysis::visitOperandRem(arith::RemSIOp, mlir::triton::PtrState &, const mlir::Location, mlir::ConversionPatternRewriter &, const llvm::SmallDenseMap<Value, PtrState> &): Assertion `state.getRank() == 1 && !state.modulos.back().has_value() && "No support for multiple modulos within an expression"' failed.
PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace.
Stack dump:
0.  Program arguments: build/cmake.linux-x86_64-cpython-3.8/third_party/triton_shared/tools/triton-shared-opt/triton-shared-opt -triton-to-linalg /home/ianb/test/ttirs_linalg_failed/11.mlir
 #0 0x0000563d41ad837b llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (build/cmake.linux-x86_64-cpython-3.8/third_party/triton_shared/tools/triton-shared-opt/triton-shared-opt+0x483637b)
 #1 0x0000563d41ad60b4 SignalHandler(int) Signals.cpp:0:0
 #2 0x00007f84a2ad41f0 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x141f0)
 #3 0x00007f84a257efbb raise ./signal/../sysdeps/unix/sysv/linux/raise.c:50:1
 #4 0x00007f84a2564864 abort ./stdlib/abort.c:81:7
 #5 0x00007f84a2564749 get_sysdep_segment_value ./intl/loadmsgcat.c:509:8
 #6 0x00007f84a2564749 _nl_load_domain ./intl/loadmsgcat.c:970:34
 #7 0x00007f84a25763d6 (/lib/x86_64-linux-gnu/libc.so.6+0x383d6)
 #8 0x0000563d3e72e543 mlir::triton::PtrAnalysis::visitOperandRem(mlir::arith::RemSIOp, mlir::triton::PtrState&, mlir::Location, mlir::ConversionPatternRewriter&, llvm::SmallDenseMap<mlir::Value, mlir::triton::PtrState, 4u, llvm::DenseMapInfo<mlir::Value, void>, llvm::detail::DenseMapPair<mlir::Value, mlir::triton::PtrState>> const&) /home/ianb/src/triton/third_party/triton_shared/lib/Analysis/PtrAnalysis.cpp:366:3
 #9 0x0000563d3e72dff2 mlir::triton::PtrAnalysis::visitOperand(mlir::Value, mlir::triton::PtrState&, mlir::Location, mlir::ConversionPatternRewriter&, llvm::SmallDenseMap<mlir::Value, mlir::triton::PtrState, 4u, llvm::DenseMapInfo<mlir::Value, void>, llvm::detail::DenseMapPair<mlir::Value, mlir::triton::PtrState>> const&) /home/ianb/src/triton/third_party/triton_shared/lib/Analysis/PtrAnalysis.cpp:0:5
#10 0x0000563d3e72e164 llvm::SmallVectorTemplateCommon<mlir::OpFoldResult, void>::getFirstEl() const /home/ianb/.triton/llvm/llvm+mlir-17.0.0-x86_64-linux-gnu-ubuntu-18.04-release/include/llvm/ADT/SmallVector.h:133:46
#11 0x0000563d3e72e164 llvm::SmallVectorTemplateCommon<mlir::OpFoldResult, void>::SmallVectorTemplateCommon(unsigned long) /home/ianb/.triton/llvm/llvm+mlir-17.0.0-x86_64-linux-gnu-ubuntu-18.04-release/include/llvm/ADT/SmallVector.h:138:49
#12 0x0000563d3e72e164 llvm::SmallVectorTemplateBase<mlir::OpFoldResult, true>::SmallVectorTemplateBase(unsigned long) /home/ianb/.triton/llvm/llvm+mlir-17.0.0-x86_64-linux-gnu-ubuntu-18.04-release/include/llvm/ADT/SmallVector.h:491:42
#13 0x0000563d3e72e164 llvm::SmallVectorImpl<mlir::OpFoldResult>::SmallVectorImpl(unsigned int) /home/ianb/.triton/llvm/llvm+mlir-17.0.0-x86_64-linux-gnu-ubuntu-18.04-release/include/llvm/ADT/SmallVector.h:592:9
#14 0x0000563d3e72e164 llvm::SmallVector<mlir::OpFoldResult, 6u>::SmallVector() /home/ianb/.triton/llvm/llvm+mlir-17.0.0-x86_64-linux-gnu-ubuntu-18.04-release/include/llvm/ADT/SmallVector.h:1202:19
#15 0x0000563d3e72e164 mlir::triton::PtrState::PtrState() /home/ianb/src/triton/third_party/triton_shared/include/triton-shared/Analysis/PtrAnalysis.h:41:7
#16 0x0000563d3e72e164 mlir::triton::PtrAnalysis::visitOperandMul(mlir::arith::MulIOp, mlir::triton::PtrState&, mlir::Location, mlir::ConversionPatternRewriter&, llvm::SmallDenseMap<mlir::Value, mlir::triton::PtrState, 4u, llvm::DenseMapInfo<mlir::Value, void>, llvm::detail::DenseMapPair<mlir::Value, mlir::triton::PtrState>> const&) /home/ianb/src/triton/third_party/triton_shared/lib/Analysis/PtrAnalysis.cpp:356:12
#17 0x0000563d3e72de62 mlir::triton::PtrAnalysis::visitOperand(mlir::Value, mlir::triton::PtrState&, mlir::Location, mlir::ConversionPatternRewriter&, llvm::SmallDenseMap<mlir::Value, mlir::triton::PtrState, 4u, llvm::DenseMapInfo<mlir::Value, void>, llvm::detail::DenseMapPair<mlir::Value, mlir::triton::PtrState>> const&) /home/ianb/src/triton/third_party/triton_shared/lib/Analysis/PtrAnalysis.cpp:0:5
#18 0x0000563d3e72ed17 mlir::triton::PtrAnalysis::visitOperandBroadcast(mlir::triton::BroadcastOp, mlir::triton::PtrState&, mlir::Location, mlir::ConversionPatternRewriter&, llvm::SmallDenseMap<mlir::Value, mlir::triton::PtrState, 4u, llvm::DenseMapInfo<mlir::Value, void>, llvm::detail::DenseMapPair<mlir::Value, mlir::triton::PtrState>> const&) /home/ianb/src/triton/third_party/triton_shared/lib/Analysis/PtrAnalysis.cpp:437:24
#19 0x0000563d3e72df12 mlir::triton::PtrAnalysis::visitOperand(mlir::Value, mlir::triton::PtrState&, mlir::Location, mlir::ConversionPatternRewriter&, llvm::SmallDenseMap<mlir::Value, mlir::triton::PtrState, 4u, llvm::DenseMapInfo<mlir::Value, void>, llvm::detail::DenseMapPair<mlir::Value, mlir::triton::PtrState>> const&) /home/ianb/src/triton/third_party/triton_shared/lib/Analysis/PtrAnalysis.cpp:0:5
#20 0x0000563d3e72da70 llvm::SmallVectorBase<unsigned int>::size() const /home/ianb/.triton/llvm/llvm+mlir-17.0.0-x86_64-linux-gnu-ubuntu-18.04-release/include/llvm/ADT/SmallVector.h:91:32
#21 0x0000563d3e72da70 mlir::triton::PtrState::getRank() const /home/ianb/src/triton/third_party/triton_shared/lib/Analysis/PtrAnalysis.cpp:48:3
#22 0x0000563d3e72da70 mlir::triton::PtrAnalysis::visitOperandAdd(mlir::arith::AddIOp, mlir::triton::PtrState&, mlir::Location, mlir::ConversionPatternRewriter&, llvm::SmallDenseMap<mlir::Value, mlir::triton::PtrState, 4u, llvm::DenseMapInfo<mlir::Value, void>, llvm::detail::DenseMapPair<mlir::Value, mlir::triton::PtrState>> const&) /home/ianb/src/triton/third_party/triton_shared/lib/Analysis/PtrAnalysis.cpp:341:17
#23 0x0000563d3e72dde1 mlir::triton::PtrAnalysis::visitOperand(mlir::Value, mlir::triton::PtrState&, mlir::Location, mlir::ConversionPatternRewriter&, llvm::SmallDenseMap<mlir::Value, mlir::triton::PtrState, 4u, llvm::DenseMapInfo<mlir::Value, void>, llvm::detail::DenseMapPair<mlir::Value, mlir::triton::PtrState>> const&) /home/ianb/src/triton/third_party/triton_shared/lib/Analysis/PtrAnalysis.cpp:0:5
#24 0x0000563d3e72fc16 mlir::Value::operator bool() const /home/ianb/.triton/llvm/llvm+mlir-17.0.0-x86_64-linux-gnu-ubuntu-18.04-release/include/mlir/IR/Value.h:117:43
#25 0x0000563d3e72fc16 mlir::triton::PtrAnalysis::visitOperandAddptr(mlir::triton::AddPtrOp, mlir::triton::PtrState&, mlir::Location, mlir::ConversionPatternRewriter&, llvm::SmallDenseMap<mlir::Value, mlir::triton::PtrState, 4u, llvm::DenseMapInfo<mlir::Value, void>, llvm::detail::DenseMapPair<mlir::Value, mlir::triton::PtrState>> const&) /home/ianb/src/triton/third_party/triton_shared/lib/Analysis/PtrAnalysis.cpp:528:3
#26 0x0000563d3e730714 mlir::triton::PtrAnalysis::rewriteAddptrOp(mlir::triton::AddPtrOp, mlir::ConversionPatternRewriter&, llvm::SmallDenseMap<mlir::Value, mlir::triton::PtrState, 4u, llvm::DenseMapInfo<mlir::Value, void>, llvm::detail::DenseMapPair<mlir::Value, mlir::triton::PtrState>>&) /home/ianb/src/triton/third_party/triton_shared/lib/Analysis/PtrAnalysis.cpp:689:26
#27 0x0000563d3e65c3e1 llvm::SmallDenseMap<mlir::Value, mlir::triton::PtrState, 4u, llvm::DenseMapInfo<mlir::Value, void>, llvm::detail::DenseMapPair<mlir::Value, mlir::triton::PtrState>>::~SmallDenseMap() /home/ianb/.triton/llvm/llvm+mlir-17.0.0-x86_64-linux-gnu-ubuntu-18.04-release/include/llvm/ADT/DenseMap.h:960:11
#28 0x0000563d3e65c3e1 (anonymous namespace)::AddPtrConverter::matchAndRewrite(mlir::triton::AddPtrOp, mlir::triton::AddPtrOpAdaptor, mlir::ConversionPatternRewriter&) const /home/ianb/src/triton/third_party/triton_shared/lib/Conversion/TritonToLinalg/TritonToLinalg.cpp:436:3
#29 0x0000563d3e54e2a7 mlir::OpConversionPattern<mlir::triton::AddPtrOp>::matchAndRewrite(mlir::Operation*, llvm::ArrayRef<mlir::Value>, mlir::ConversionPatternRewriter&) const /home/ianb/.triton/llvm/llvm+mlir-17.0.0-x86_64-linux-gnu-ubuntu-18.04-release/include/mlir/Transforms/DialectConversion.h:536:73
#30 0x0000563d40c2b5b1 mlir::ConversionPattern::matchAndRewrite(mlir::Operation*, mlir::PatternRewriter&) const (build/cmake.linux-x86_64-cpython-3.8/third_party/triton_shared/tools/triton-shared-opt/triton-shared-opt+0x39895b1)
#31 0x0000563d40c776b2 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&)>) (build/cmake.linux-x86_64-cpython-3.8/third_party/triton_shared/tools/triton-shared-opt/triton-shared-opt+0x39d56b2)
#32 0x0000563d40c37d19 (anonymous namespace)::OperationLegalizer::legalize(mlir::Operation*, mlir::ConversionPatternRewriter&) DialectConversion.cpp:0:0
#33 0x0000563d40c382f0 (anonymous namespace)::OperationConverter::convertOperations(llvm::ArrayRef<mlir::Operation*>, llvm::function_ref<void (mlir::Diagnostic&)>) DialectConversion.cpp:0:0
#34 0x0000563d40c3a6b0 mlir::applyFullConversion(mlir::Operation*, mlir::ConversionTarget&, mlir::FrozenRewritePatternSet const&) (build/cmake.linux-x86_64-cpython-3.8/third_party/triton_shared/tools/triton-shared-opt/triton-shared-opt+0x39986b0)
#35 0x0000563d3e6541b2 (anonymous namespace)::TritonToLinalgPass::runOnOperation() /home/ianb/src/triton/third_party/triton_shared/lib/Conversion/TritonToLinalg/TritonToLinalgPass.cpp:194:16
#36 0x0000563d3e68a991 mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) (build/cmake.linux-x86_64-cpython-3.8/third_party/triton_shared/tools/triton-shared-opt/triton-shared-opt+0x13e8991)
#37 0x0000563d3e68b1e1 mlir::detail::OpToOpPassAdaptor::runPipeline(mlir::OpPassManager&, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::PipelineParentInfo const*) (build/cmake.linux-x86_64-cpython-3.8/third_party/triton_shared/tools/triton-shared-opt/triton-shared-opt+0x13e91e1)
#38 0x0000563d3e68bd4a mlir::PassManager::run(mlir::Operation*) (build/cmake.linux-x86_64-cpython-3.8/third_party/triton_shared/tools/triton-shared-opt/triton-shared-opt+0x13e9d4a)
#39 0x0000563d3e67bf7b performActions(llvm::raw_ostream&, std::shared_ptr<llvm::SourceMgr> const&, mlir::MLIRContext*, mlir::MlirOptMainConfig const&) MlirOptMain.cpp:0:0
#40 0x0000563d3e67cab5 processBuffer(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, mlir::MlirOptMainConfig const&, mlir::DialectRegistry&, llvm::ThreadPool*) MlirOptMain.cpp:0:0
#41 0x0000563d3e67cba0 mlir::LogicalResult llvm::function_ref<mlir::LogicalResult (std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::raw_ostream&)>::callback_fn<mlir::MlirOptMain(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, mlir::DialectRegistry&, mlir::MlirOptMainConfig const&)::'lambda'(std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::raw_ostream&)>(long, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::raw_ostream&) MlirOptMain.cpp:0:0
#42 0x0000563d419c34c5 mlir::splitAndProcessBuffer(std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::function_ref<mlir::LogicalResult (std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::raw_ostream&)>, llvm::raw_ostream&, bool, bool) (build/cmake.linux-x86_64-cpython-3.8/third_party/triton_shared/tools/triton-shared-opt/triton-shared-opt+0x47214c5)
#43 0x0000563d3e67aaa3 mlir::MlirOptMain(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, mlir::DialectRegistry&, mlir::MlirOptMainConfig const&) (build/cmake.linux-x86_64-cpython-3.8/third_party/triton_shared/tools/triton-shared-opt/triton-shared-opt+0x13d8aa3)
#44 0x0000563d3e67ced3 mlir::MlirOptMain(int, char**, llvm::StringRef, mlir::DialectRegistry&) (build/cmake.linux-x86_64-cpython-3.8/third_party/triton_shared/tools/triton-shared-opt/triton-shared-opt+0x13daed3)
#45 0x0000563d3d67dd5b main /home/ianb/src/triton/third_party/triton_shared/tools/triton-shared-opt/triton-shared-opt.cpp:16:33
#46 0x00007f84a2566565 __libc_start_main ./csu/../csu/libc-start.c:332:16
#47 0x0000563d3d67dc5e _start (build/cmake.linux-x86_64-cpython-3.8/third_party/triton_shared/tools/triton-shared-opt/triton-shared-opt+0x3dbc5e)
find: ‘build/cmake.linux-x86_64-cpython-3.8/third_party/triton_shared/tools/triton-shared-opt/triton-shared-opt’ terminated by signal 6
manbearian commented 9 months ago

@nhat-nguyen FYI