However, when I do change it to run the correct file i'm getting errors.
make gpu-mma-run
LLVM ERROR: Cannot select: intrinsic %llvm.nvvm.wmma.m16n16k16.store.d.row.stride.f32
PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace. Stack dump:
0. Program arguments: ../../llvm/build/bin/mlir-opt -pass-pipeline=builtin.module(gpu.module(strip-debuginfo,convert-gpu-to-nvvm,gpu-to-cubin))
1. Running pass 'Function Pass Manager' on module 'LLVMDialectModule'.
2. Running pass 'NVPTX DAG->DAG Pattern Instruction Selection' on function '@kernel1' #0 0x000055cd7a88aadf llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (../../llvm/build/bin/mlir-opt+0x214adf)
#1 0x000055cd7a888594 SignalHandler(int) Signals.cpp:0:0
#2 0x00007fa43ee66420 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x14420) #3 0x00007fa43d02a00b raise /build/glibc-SzIz7B/glibc-2.31/signal/../sysdeps/unix/sysv/linux/raise.c:51:1
#4 0x00007fa43d009859 abort /build/glibc-SzIz7B/glibc-2.31/stdlib/abort.c:81:7 #5 0x000055cd7a80b3aa llvm::ConvertUTF8toUTF32(unsigned char const**, unsigned char const*, unsigned int**, unsigned int*, llvm::ConversionFlags) (.cold) ConvertUTF.cpp:0:0
#6 0x000055cd7c51533d llvm::SelectionDAGISel::CannotYetSelect(llvm::SDNode*) (../../llvm/build/bin/mlir-opt+0x1e9f33d) #7 0x000055cd7c517ca2 llvm::SelectionDAGISel::SelectCodeCommon(llvm::SDNode*, unsigned char const*, unsigned int) (../../llvm/build/bin/mlir-opt+0x1ea1ca2) #8 0x000055cd7c512a90 llvm::SelectionDAGISel::DoInstructionSelection() (../../llvm/build/bin/mlir-opt+0x1e9ca90) #9 0x000055cd7c520651 llvm::SelectionDAGISel::CodeGenAndEmitDAG() (../../llvm/build/bin/mlir-opt+0x1eaa651)
#10 0x000055cd7c524a78 llvm::SelectionDAGISel::SelectAllBasicBlocks(llvm::Function const&) (../../llvm/build/bin/mlir-opt+0x1eaea78) #11 0x000055cd7c5269ff llvm::SelectionDAGISel::runOnMachineFunction(llvm::MachineFunction&) (.part.0) SelectionDAGISel.cpp:0:0
#12 0x000055cd7c73344e llvm::MachineFunctionPass::runOnFunction(llvm::Function&) (.part.0) MachineFunctionPass.cpp:0:0 #13 0x000055cd7d350985 llvm::FPPassManager::runOnFunction(llvm::Function&) (../../llvm/build/bin/mlir-opt+0x2cda985)
#14 0x000055cd7d350bc9 llvm::FPPassManager::runOnModule(llvm::Module&) (../../llvm/build/bin/mlir-opt+0x2cdabc9)
#15 0x000055cd7d351082 llvm::legacy::PassManagerImpl::run(llvm::Module&) (../../llvm/build/bin/mlir-opt+0x2cdb082)
#16 0x000055cd7ac89d33 mlir::gpu::SerializeToBlobPass::translateToISA[abi:cxx11](llvm::Module&, llvm::TargetMachine&) (../../llvm/build/bin/mlir-opt+0x613d33)
#17 0x000055cd7ac89f7a mlir::gpu::SerializeToBlobPass::runOnOperation() (../../llvm/build/bin/mlir-opt+0x613f7a)
#18 0x000055cd7c00b6ba mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) (../../llvm/build/bin/mlir-opt+0x19956ba)
#19 0x000055cd7c00bcb9 mlir::detail::OpToOpPassAdaptor::runPipeline(mlir::OpPassManager&, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::Pipe
lineParentInfo const*) (../../llvm/build/bin/mlir-opt+0x1995cb9)
#20 0x000055cd7c00a936 mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool) (../../llvm/build/bin/mlir-opt+0x1994936)
#21 0x000055cd7c00b59a mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) (../../llvm/build/bin/mlir-opt+0x199559a)
#22 0x000055cd7c00bcb9 mlir::detail::OpToOpPassAdaptor::runPipeline(mlir::OpPassManager&, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::Pipe
lineParentInfo const*) (../../llvm/build/bin/mlir-opt+0x1995cb9)
#23 0x000055cd7c00c6d9 mlir::PassManager::run(mlir::Operation*) (../../llvm/build/bin/mlir-opt+0x19966d9)
#24 0x000055cd7bffd05b performActions(llvm::raw_ostream&, bool, bool, llvm::SourceMgr&, mlir::MLIRContext*, llvm::function_ref<mlir::LogicalResult (mlir::PassManager&)>, bool, bool) (.constprop.0) MlirOptMai
n.cpp:0:0
#25 0x000055cd7bffd7e1 processBuffer(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, bool, bool, bool, bool, bool, bool, llvm::function_ref<mlir::LogicalResu
lt (mlir::PassManager&)>, mlir::DialectRegistry&, llvm::ThreadPool*) MlirOptMain.cpp:0:0
#26 0x000055cd7bffdbd3 mlir::LogicalResult llvm::function_ref<mlir::LogicalResult (std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::raw_ostream&)>::callback_fn<mlir::MlirOp
tMain(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::function_ref<mlir::LogicalResult (mlir::PassManager&)>, mlir::DialectRegistry&, bool, bool, bool,
bool, bool, bool, bool)::'lambda'(std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::raw_ostream&)>(long, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::Memor
yBuffer>>, llvm::raw_ostream&) MlirOptMain.cpp:0:0
#27 0x000055cd7c0b7ed4 mlir::splitAndProcessBuffer(std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::function_ref<mlir::LogicalResult (std::unique_ptr<llvm::MemoryBuffer, st
d::default_delete<llvm::MemoryBuffer>>, llvm::raw_ostream&)>, llvm::raw_ostream&, bool, bool) (../../llvm/build/bin/mlir-opt+0x1a41ed4)
#28 0x000055cd7bffc2a1 mlir::MlirOptMain(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, mlir::PassPipelineCLParser const&, mlir::DialectRegistry&, bool, boo
l, bool, bool, bool, bool, bool, bool) (../../llvm/build/bin/mlir-opt+0x19862a1)
#29 0x000055cd7bffdfd7 mlir::MlirOptMain(int, char**, llvm::StringRef, mlir::DialectRegistry&, bool) (../../llvm/build/bin/mlir-opt+0x1987fd7)
#30 0x000055cd7a812386 main (../../llvm/build/bin/mlir-opt+0x19c386)
#31 0x00007fa43d00b083 __libc_start_main /build/glibc-SzIz7B/glibc-2.31/csu/../csu/libc-start.c:342:3
#32 0x000055cd7a86642e _start (../../llvm/build/bin/mlir-opt+0x1f042e)
Aborted (core dumped) Error: entry point not found make: *** [makefile:73: gpu-mma-run] Error 1
abhi@abhi-MS-7C56:~/courses/ece527-SOC-Design/final_project/buddy-mlir/examples/MLIRGPU$ make gpu-mma-run LLVM ERROR: Cannot select: intrinsic %llvm.nvvm.wmma.m16n16k16.store.d.row.stride.f32 PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace. Stack dump:
0. Program arguments: ../../llvm/build/bin/mlir-opt -pass-pipeline=builtin.module(gpu.module(strip-debuginfo,convert-gpu-to-nvvm,gpu-to-cubin)) 1. Running pass 'Function Pass Manager' on module 'LLVMDialectModule'. 2. Running pass 'NVPTX DAG->DAG Pattern Instruction Selection' on function '@kernel1' #0 0x00005633cdd1aadf llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (../../llvm/build/bin/mlir-opt+0x214adf) #1 0x00005633cdd18594 SignalHandler(int) Signals.cpp:0:0
#2 0x00007f0246ca9420 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x14420) #3 0x00007f0244e6d00b raise /build/glibc-SzIz7B/glibc-2.31/signal/../sysdeps/unix/sysv/linux/raise.c:51:1 #4 0x00007f0244e4c859 abort /build/glibc-SzIz7B/glibc-2.31/stdlib/abort.c:81:7 #5 0x00005633cdc9b3aa llvm::ConvertUTF8toUTF32(unsigned char const**, unsigned char const*, unsigned int**, unsigned int*, llvm::ConversionFlags) (.cold) ConvertUTF.cpp:0:0
#6 0x00005633cf9a533d llvm::SelectionDAGISel::CannotYetSelect(llvm::SDNode*) (../../llvm/build/bin/mlir-opt+0x1e9f33d) #7 0x00005633cf9a7ca2 llvm::SelectionDAGISel::SelectCodeCommon(llvm::SDNode*, unsigned char const*, unsigned int) (../../llvm/build/bin/mlir-opt+0x1ea1ca2) #8 0x00005633cf9a2a90 llvm::SelectionDAGISel::DoInstructionSelection() (../../llvm/build/bin/mlir-opt+0x1e9ca90) #9 0x00005633cf9b0651 llvm::SelectionDAGISel::CodeGenAndEmitDAG() (../../llvm/build/bin/mlir-opt+0x1eaa651)
#10 0x00005633cf9b4a78 llvm::SelectionDAGISel::SelectAllBasicBlocks(llvm::Function const&) (../../llvm/build/bin/mlir-opt+0x1eaea78) #11 0x00005633cf9b69ff llvm::SelectionDAGISel::runOnMachineFunction(llvm::MachineFunction&) (.part.0) SelectionDAGISel.cpp:0:0 #12 0x00005633cfbc344e llvm::MachineFunctionPass::runOnFunction(llvm::Function&) (.part.0) MachineFunctionPass.cpp:0:0 #13 0x00005633d07e0985 llvm::FPPassManager::runOnFunction(llvm::Function&) (../../llvm/build/bin/mlir-opt+0x2cda985) #14 0x00005633d07e0bc9 llvm::FPPassManager::runOnModule(llvm::Module&) (../../llvm/build/bin/mlir-opt+0x2cdabc9) #15 0x00005633d07e1082 llvm::legacy::PassManagerImpl::run(llvm::Module&) (../../llvm/build/bin/mlir-opt+0x2cdb082) #16 0x00005633ce119d33 mlir::gpu::SerializeToBlobPass::translateToISA[abi:cxx11](llvm::Module&, llvm::TargetMachine&) (../../llvm/build/bin/mlir-opt+0x613d33)
#17 0x00005633ce119f7a mlir::gpu::SerializeToBlobPass::runOnOperation() (../../llvm/build/bin/mlir-opt+0x613f7a) #18 0x00005633cf49b6ba mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) (../../llvm/build/bin/mlir-opt+0x19956ba) #19 0x00005633cf49bcb9 mlir::detail::OpToOpPassAdaptor::runPipeline(mlir::OpPassManager&, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::PipelineParentInfo const*) (../../llvm/build/bin/mlir-opt+0x1995cb9)
#20 0x00005633cf49a936 mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool) (../../llvm/build/bin/mlir-opt+0x1994936) #21 0x00005633cf49b59a mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) (../../llvm/build/bin/mlir-opt+0x199559a) #22 0x00005633cf49bcb9 mlir::detail::OpToOpPassAdaptor::runPipeline(mlir::OpPassManager&, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::PipelineParentInfo const*) (../../llvm/build/bin/mlir-opt+0x1995cb9) #23 0x00005633cf49c6d9 mlir::PassManager::run(mlir::Operation*) (../../llvm/build/bin/mlir-opt+0x19966d9)
#24 0x00005633cf48d05b performActions(llvm::raw_ostream&, bool, bool, llvm::SourceMgr&, mlir::MLIRContext*, llvm::function_ref<mlir::LogicalResult (mlir::PassManager&)>, bool, bool) (.constprop.0) MlirOptMain.cpp:0:0#25 0x00005633cf48d7e1 processBuffer(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, bool, bool, bool, bool, bool, bool, llvm::function_ref<mlir::LogicalResult (mlir::PassManager&)>, mlir::DialectRegistry&, llvm::ThreadPool*) MlirOptMain.cpp:0:0
#26 0x00005633cf48dbd3 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>>, llvm::function_ref<mlir::LogicalResult (mlir::PassManager&)>, mlir::DialectRegistry&, bool, bool, bool, bool, bool, bool, bool)::'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 #27 0x00005633cf547ed4 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) (../../llvm/build/bin/mlir-opt+0x1a41ed4)
#28 0x00005633cf48c2a1 mlir::MlirOptMain(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, mlir::PassPipelineCLParser const&, mlir::DialectRegistry&, bool, bool, bool, bool, bool, bool, bool, bool) (../../llvm/build/bin/mlir-opt+0x19862a1) #29 0x00005633cf48dfd7 mlir::MlirOptMain(int, char**, llvm::StringRef, mlir::DialectRegistry&, bool) (../../llvm/build/bin/mlir-opt+0x1987fd7) #30 0x00005633cdca2386 main (../../llvm/build/bin/mlir-opt+0x19c386)
#31 0x00007f0244e4e083 __libc_start_main /build/glibc-SzIz7B/glibc-2.31/csu/../csu/libc-start.c:342:3 #32 0x00005633cdcf642e _start (../../llvm/build/bin/mlir-opt+0x1f042e) Aborted (core dumped) Error: entry point not found
make: *** [makefile:73: gpu-mma-run] Error 1
so it looks like this example is broken.