search

microsoft / triton-shared

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

Only asserts on scalars are currently supported #17

Open manbearian opened 9 months ago

manbearian commented 9 months ago

created from #7.

Asserts on tensors are NYI. (Yes this is an assert while compiling tt.assert :)

repros.zip

triton-shared-opt -triton-to-linalg 14.mlir triton-shared-opt -triton-to-linalg 17.mlir triton-shared-opt -triton-to-linalg 23.mlir triton-shared-opt -triton-to-linalg 61.mlir triton-shared-opt -triton-to-linalg 65.mlir triton-shared-opt -triton-to-linalg 85.mlir triton-shared-opt -triton-to-linalg 94.mlir

Error output:

+++/home/ianb/test/ttirs_linalg_failed/14.mlir
/home/ianb/test/ttirs_linalg_failed/14.mlir: 40:5: error: other value used in masked load produced by unsupported instruction
    tt.assert %27, "index out of bounds: 0 <= tmp5 < 65", "<frozen importlib._bootstrap_external>", "_call_with_frames_removed", 883 : tensor<8x128xi1>
    ^
triton-shared-opt: /home/ianb/src/triton/third_party/triton_shared/lib/Conversion/TritonToLinalg/TritonToLinalg.cpp:452: virtual mlir::LogicalResult (anonymous namespace)::AssertConverter::matchAndRewrite(triton::AssertOp, mlir::OpConversionPattern<mlir::triton::AssertOp>::OpAdaptor, mlir::ConversionPatternRewriter &) const: Assertion `condVal && condVal.getType().isa<mlir::IntegerType>() && "Only asserts on scalars are currently supported"' 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/14.mlir
 #0 0x000055a0999ab37b 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 0x000055a0999a90b4 SignalHandler(int) Signals.cpp:0:0
 #2 0x00007f7ee27071f0 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x141f0)
 #3 0x00007f7ee21b1fbb raise ./signal/../sysdeps/unix/sysv/linux/raise.c:50:1
 #4 0x00007f7ee2197864 abort ./stdlib/abort.c:81:7
 #5 0x00007f7ee2197749 get_sysdep_segment_value ./intl/loadmsgcat.c:509:8
 #6 0x00007f7ee2197749 _nl_load_domain ./intl/loadmsgcat.c:970:34
 #7 0x00007f7ee21a93d6 (/lib/x86_64-linux-gnu/libc.so.6+0x383d6)
 #8 0x000055a096534daf (build/cmake.linux-x86_64-cpython-3.8/third_party/triton_shared/tools/triton-shared-opt/triton-shared-opt+0x13bfdaf)
 #9 0x000055a09641dbe9 mlir::OpConversionPattern<mlir::triton::AssertOp>::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:537:3
#10 0x000055a098afe5b1 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)
#11 0x000055a098b4a6b2 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)
#12 0x000055a098b0ad19 (anonymous namespace)::OperationLegalizer::legalize(mlir::Operation*, mlir::ConversionPatternRewriter&) DialectConversion.cpp:0:0
#13 0x000055a098b0b2f0 (anonymous namespace)::OperationConverter::convertOperations(llvm::ArrayRef<mlir::Operation*>, llvm::function_ref<void (mlir::Diagnostic&)>) DialectConversion.cpp:0:0
#14 0x000055a098b0d6b0 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)
#15 0x000055a0965271b2 (anonymous namespace)::TritonToLinalgPass::runOnOperation() /home/ianb/src/triton/third_party/triton_shared/lib/Conversion/TritonToLinalg/TritonToLinalgPass.cpp:194:16
#16 0x000055a09655d991 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)
#17 0x000055a09655e1e1 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)
#18 0x000055a09655ed4a 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)
#19 0x000055a09654ef7b performActions(llvm::raw_ostream&, std::shared_ptr<llvm::SourceMgr> const&, mlir::MLIRContext*, mlir::MlirOptMainConfig const&) MlirOptMain.cpp:0:0
#20 0x000055a09654fab5 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
#21 0x000055a09654fba0 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
#22 0x000055a0998964c5 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)
#23 0x000055a09654daa3 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)
#24 0x000055a09654fed3 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)
#25 0x000055a095550d5b main /home/ianb/src/triton/third_party/triton_shared/tools/triton-shared-opt/triton-shared-opt.cpp:16:33
#26 0x00007f7ee2199565 __libc_start_main ./csu/../csu/libc-start.c:332:16
#27 0x000055a095550c5e _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