SparseTensor Dialect test sparse-matmul-2-4-prune.mlir crashed with Assertion 'this->_M_is_engaged()' failed.

[Ag-Cu]

system: Ubuntu 20.04.6 LTS (Focal Fossa)
current git head: #76697
build command:

cmake -G Ninja ../llvm \
   -DLLVM_ENABLE_PROJECTS=mlir \
   -DLLVM_BUILD_EXAMPLES=ON \
   -DLLVM_TARGETS_TO_BUILD="Native;NVPTX;AMDGPU" \
   -DCMAKE_BUILD_TYPE=Release \
   -DLLVM_ENABLE_ASSERTIONS=ON

cmake --build . --target check-mlir

reproduced with:
cd llvm-project/mlir/test/Integration/Dialect/SparseTensor/GPU/CUDA/sm80-lt

mlir-opt sparse-matmul-2-4-prune.mlir --sparsifier="enable-gpu-libgen gpu-triple=nvptx64-nvidia-cuda gpu-chip=sm_80 gpu-features=+ptx71 gpu-format=%gpu_compilation_format enable-runtime-library=true" 

will get:

/usr/include/c++/9/optional:438: constexpr _Tp& std::_Optional_base_impl<_Tp, _Dp>::_M_get() [with _Tp = mlir::gpu::CompilationTarget; _Dp = std::_Optional_base<mlir::gpu::CompilationTarget, true, true>]: Assertion 'this->_M_is_engaged()' failed.
PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace.
Stack dump:
0.      Program arguments: mlir-opt sparse-matmul-2-4-prune.mlir "--sparsifier=enable-gpu-libgen gpu-triple=nvptx64-nvidia-cuda gpu-chip=sm_80 gpu-features=+ptx71 gpu-format=%gpu_compilation_format enable-runtime-library=true"
 #0 0x00005581a18a2bcf llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (.localalias) (/home/yt/llvm-project/build/bin/mlir-opt+0x1685bcf)
 #1 0x00005581a189fc24 SignalHandler(int) (/home/yt/llvm-project/build/bin/mlir-opt+0x1682c24)
 #2 0x00007fe7d4206420 __restore_rt (/usr/lib/x86_64-linux-gnu/libpthread.so.0+0x14420)
 #3 0x00007fe7d3ca300b raise (/usr/lib/x86_64-linux-gnu/libc.so.6+0x4300b)
 #4 0x00007fe7d3c82859 abort (/usr/lib/x86_64-linux-gnu/libc.so.6+0x22859)
 #5 0x00005581a201ec2b (anonymous namespace)::GpuModuleToBinaryPass::runOnOperation() (/home/yt/llvm-project/build/bin/mlir-opt+0x1e01c2b)
 #6 0x00005581a459d3ae mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) (.localalias) (/home/yt/llvm-project/build/bin/mlir-opt+0x43803ae)
 #7 0x00005581a459d88a mlir::detail::OpToOpPassAdaptor::runPipeline(mlir::OpPassManager&, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::PipelineParentInfo const*) (.localalias) (/home/yt/llvm-project/build/bin/mlir-opt+0x438088a)
 #8 0x00005581a459e054 mlir::PassManager::run(mlir::Operation*) (/home/yt/llvm-project/build/bin/mlir-opt+0x4381054)
 #9 0x00005581a458ed43 performActions(llvm::raw_ostream&, std::shared_ptr<llvm::SourceMgr> const&, mlir::MLIRContext*, mlir::MlirOptMainConfig const&) (/home/yt/llvm-project/build/bin/mlir-opt+0x4371d43)
#10 0x00005581a45902c4 processBuffer(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer> >, mlir::MlirOptMainConfig const&, mlir::DialectRegistry&, llvm::ThreadPool*) (/home/yt/llvm-project/build/bin/mlir-opt+0x43732c4)
#11 0x00005581a4590474 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&) (/home/yt/llvm-project/build/bin/mlir-opt+0x4373474)
#12 0x00005581a468c964 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) (/home/yt/llvm-project/build/bin/mlir-opt+0x446f964)
#13 0x00005581a45888d7 mlir::MlirOptMain(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer> >, mlir::DialectRegistry&, mlir::MlirOptMainConfig const&) (.localalias) (/home/yt/llvm-project/build/bin/mlir-opt+0x436b8d7)
#14 0x00005581a45905c1 mlir::MlirOptMain(int, char**, llvm::StringRef, llvm::StringRef, mlir::DialectRegistry&) (.localalias) (/home/yt/llvm-project/build/bin/mlir-opt+0x43735c1)
#15 0x00005581a4590a46 mlir::MlirOptMain(int, char**, llvm::StringRef, mlir::DialectRegistry&) (/home/yt/llvm-project/build/bin/mlir-opt+0x4373a46)
#16 0x00005581a17d9293 main (/home/yt/llvm-project/build/bin/mlir-opt+0x15bc293)
#17 0x00007fe7d3c84083 __libc_start_main (/usr/lib/x86_64-linux-gnu/libc.so.6+0x24083)
#18 0x00005581a18725fe _start (/home/yt/llvm-project/build/bin/mlir-opt+0x16555fe)
Aborted 

I don't know if this is a bug or if I'm missing some dependencies.

[llvmbot]

@llvm/issue-subscribers-mlir-sparse

Author: Yin Tong (Ag-Cu)

system: Ubuntu 20.04.6 LTS (Focal Fossa) current git head: #76697 build command: ``` cmake -G Ninja ../llvm \ -DLLVM_ENABLE_PROJECTS=mlir \ -DLLVM_BUILD_EXAMPLES=ON \ -DLLVM_TARGETS_TO_BUILD="Native;NVPTX;AMDGPU" \ -DCMAKE_BUILD_TYPE=Release \ -DLLVM_ENABLE_ASSERTIONS=ON ``` `cmake --build . --target check-mlir`

reproduced with:
cd llvm-project/mlir/test/Integration/Dialect/SparseTensor/GPU/CUDA/sm80-lt

mlir-opt sparse-matmul-2-4-prune.mlir --sparsifier="enable-gpu-libgen gpu-triple=nvptx64-nvidia-cuda gpu-chip=sm_80 gpu-features=+ptx71 gpu-format=%gpu_compilation_format enable-runtime-library=true" 

will get:

/usr/include/c++/9/optional:438: constexpr _Tp& std::_Optional_base_impl<_Tp, _Dp>::_M_get() [with _Tp = mlir::gpu::CompilationTarget; _Dp = std::_Optional_base<mlir::gpu::CompilationTarget, true, true>]: Assertion 'this->_M_is_engaged()' failed.
PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace.
Stack dump:
0.      Program arguments: mlir-opt sparse-matmul-2-4-prune.mlir "--sparsifier=enable-gpu-libgen gpu-triple=nvptx64-nvidia-cuda gpu-chip=sm_80 gpu-features=+ptx71 gpu-format=%gpu_compilation_format enable-runtime-library=true"
 #<!-- -->0 0x00005581a18a2bcf llvm::sys::PrintStackTrace(llvm::raw_ostream&amp;, int) (.localalias) (/home/yt/llvm-project/build/bin/mlir-opt+0x1685bcf)
 #<!-- -->1 0x00005581a189fc24 SignalHandler(int) (/home/yt/llvm-project/build/bin/mlir-opt+0x1682c24)
 #<!-- -->2 0x00007fe7d4206420 __restore_rt (/usr/lib/x86_64-linux-gnu/libpthread.so.0+0x14420)
 #<!-- -->3 0x00007fe7d3ca300b raise (/usr/lib/x86_64-linux-gnu/libc.so.6+0x4300b)
 #<!-- -->4 0x00007fe7d3c82859 abort (/usr/lib/x86_64-linux-gnu/libc.so.6+0x22859)
 #<!-- -->5 0x00005581a201ec2b (anonymous namespace)::GpuModuleToBinaryPass::runOnOperation() (/home/yt/llvm-project/build/bin/mlir-opt+0x1e01c2b)
 #<!-- -->6 0x00005581a459d3ae mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) (.localalias) (/home/yt/llvm-project/build/bin/mlir-opt+0x43803ae)
 #<!-- -->7 0x00005581a459d88a mlir::detail::OpToOpPassAdaptor::runPipeline(mlir::OpPassManager&amp;, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::PipelineParentInfo const*) (.localalias) (/home/yt/llvm-project/build/bin/mlir-opt+0x438088a)
 #<!-- -->8 0x00005581a459e054 mlir::PassManager::run(mlir::Operation*) (/home/yt/llvm-project/build/bin/mlir-opt+0x4381054)
 #<!-- -->9 0x00005581a458ed43 performActions(llvm::raw_ostream&amp;, std::shared_ptr&lt;llvm::SourceMgr&gt; const&amp;, mlir::MLIRContext*, mlir::MlirOptMainConfig const&amp;) (/home/yt/llvm-project/build/bin/mlir-opt+0x4371d43)
#<!-- -->10 0x00005581a45902c4 processBuffer(llvm::raw_ostream&amp;, std::unique_ptr&lt;llvm::MemoryBuffer, std::default_delete&lt;llvm::MemoryBuffer&gt; &gt;, mlir::MlirOptMainConfig const&amp;, mlir::DialectRegistry&amp;, llvm::ThreadPool*) (/home/yt/llvm-project/build/bin/mlir-opt+0x43732c4)
#<!-- -->11 0x00005581a4590474 mlir::LogicalResult llvm::function_ref&lt;mlir::LogicalResult (std::unique_ptr&lt;llvm::MemoryBuffer, std::default_delete&lt;llvm::MemoryBuffer&gt; &gt;, llvm::raw_ostream&amp;)&gt;::callback_fn&lt;mlir::MlirOptMain(llvm::raw_ostream&amp;, std::unique_ptr&lt;llvm::MemoryBuffer, std::default_delete&lt;llvm::MemoryBuffer&gt; &gt;, mlir::DialectRegistry&amp;, mlir::MlirOptMainConfig const&amp;)::'lambda'(std::unique_ptr&lt;llvm::MemoryBuffer, std::default_delete&lt;llvm::MemoryBuffer&gt; &gt;, llvm::raw_ostream&amp;)&gt;(long, std::unique_ptr&lt;llvm::MemoryBuffer, std::default_delete&lt;llvm::MemoryBuffer&gt; &gt;, llvm::raw_ostream&amp;) (/home/yt/llvm-project/build/bin/mlir-opt+0x4373474)
#<!-- -->12 0x00005581a468c964 mlir::splitAndProcessBuffer(std::unique_ptr&lt;llvm::MemoryBuffer, std::default_delete&lt;llvm::MemoryBuffer&gt; &gt;, llvm::function_ref&lt;mlir::LogicalResult (std::unique_ptr&lt;llvm::MemoryBuffer, std::default_delete&lt;llvm::MemoryBuffer&gt; &gt;, llvm::raw_ostream&amp;)&gt;, llvm::raw_ostream&amp;, bool, bool) (/home/yt/llvm-project/build/bin/mlir-opt+0x446f964)
#<!-- -->13 0x00005581a45888d7 mlir::MlirOptMain(llvm::raw_ostream&amp;, std::unique_ptr&lt;llvm::MemoryBuffer, std::default_delete&lt;llvm::MemoryBuffer&gt; &gt;, mlir::DialectRegistry&amp;, mlir::MlirOptMainConfig const&amp;) (.localalias) (/home/yt/llvm-project/build/bin/mlir-opt+0x436b8d7)
#<!-- -->14 0x00005581a45905c1 mlir::MlirOptMain(int, char**, llvm::StringRef, llvm::StringRef, mlir::DialectRegistry&amp;) (.localalias) (/home/yt/llvm-project/build/bin/mlir-opt+0x43735c1)
#<!-- -->15 0x00005581a4590a46 mlir::MlirOptMain(int, char**, llvm::StringRef, mlir::DialectRegistry&amp;) (/home/yt/llvm-project/build/bin/mlir-opt+0x4373a46)
#<!-- -->16 0x00005581a17d9293 main (/home/yt/llvm-project/build/bin/mlir-opt+0x15bc293)
#<!-- -->17 0x00007fe7d3c84083 __libc_start_main (/usr/lib/x86_64-linux-gnu/libc.so.6+0x24083)
#<!-- -->18 0x00005581a18725fe _start (/home/yt/llvm-project/build/bin/mlir-opt+0x16555fe)
Aborted 

I don't know if this is a bug or if I'm missing some dependencies.

[sott0n]

The issue stems from an incorrect usage of gpu-format=%gpu_compilation_format when encountering a null value, as %gpu_compilation_format is a variable access.

The gpu_compilation_format variable is set at mlir/test/Integration/Dialect/SparseTensor/GPU/CUDA/lit.local.cfg :

config.substitutions.append(("%gpu_compilation_format", config.gpu_compilation_format))

and the config.gpu_compilation_format variable is set at mlir/test/lit.site.cfg.py.in :

config.gpu_compilation_format = "@MLIR_GPU_COMPILATION_TEST_FORMAT@"

and the MLIR_GPU_COMPILATION_TEST_FORMAT is set at mlir/test/CMakeLists.txt:

 set(MLIR_GPU_COMPILATION_TEST_FORMAT "fatbin" CACHE STRING
        "The GPU compilation format used by the tests.")