Skip to content

[MLIR] convert-gpu-to-nvvm crashes with "index.mul created with unregistered dialect" #171303

@clearlove2077

Description

@clearlove2077

Hi, I found a crash bug when lowering programs with MLIR. Here is a code example that reproduces the issue:

module {
  gpu.module @vector_add_kernel {
    gpu.func @vector_add_local_tile_multi_elem_per_thread_half(%arg0: memref<?xf16>, %arg1: i32, %arg2: memref<?xf16>, %arg3: memref<?xf16>, %arg4: f16, %arg5: f16, %arg6: f16) {
      %c8_i32 = arith.constant 8 : i32
      %global_id_x = gpu.global_id  x
      %thread_id_x = gpu.thread_id  x
      %block_id_x = gpu.block_id  x
      %0 = arith.index_cast %global_id_x : index to i32
      %1 = arith.index_cast %thread_id_x : index to i32
      %2 = arith.index_cast %block_id_x : index to i32
      %3 = arith.muli %2, %0 : i32
      %4 = arith.addi %1, %3 : i32
      %5 = arith.divsi %arg1, %c8_i32 : i32
      %6 = arith.cmpi slt, %4, %5 : i32
      cf.cond_br %6, ^bb1, ^bb2
    ^bb1:  // pred: ^bb0
      %7 = arith.muli %4, %c8_i32 : i32
      %8 = arith.index_cast %7 : i32 to index
      %9 = vector.load %arg2[%8] : memref<?xf16>, vector<8xf16>
      %10 = vector.load %arg3[%8] : memref<?xf16>, vector<8xf16>
      %11 = arith.extf %arg4 : f16 to f32
      %12 = arith.extf %arg5 : f16 to f32
      %13 = arith.extf %arg6 : f16 to f32
      %14 = vector.broadcast %11 : f32 to vector<8xf32>
      %15 = vector.broadcast %12 : f32 to vector<8xf32>
      %16 = vector.broadcast %13 : f32 to vector<8xf32>
      %17 = arith.truncf %14 : vector<8xf32> to vector<8xf16>
      %18 = arith.truncf %15 : vector<8xf32> to vector<8xf16>
      %19 = arith.truncf %16 : vector<8xf32> to vector<8xf16>
      %20 = arith.mulf %9, %17 : vector<8xf16>
      %21 = arith.mulf %10, %18 : vector<8xf16>
      %22 = arith.addf %21, %19 : vector<8xf16>
      %23 = arith.addf %20, %22 : vector<8xf16>
      vector.store %23, %arg0[%8] : memref<?xf16>, vector<8xf16>
      cf.br ^bb2
    ^bb2:  // 2 preds: ^bb0, ^bb1
      gpu.return
    }
  }
}

Used Command

mlir-opt ./input.mlir  --convert-gpu-to-nvvm
LLVM ERROR: index.mul created with unregistered dialect. If this is intended, please call allowUnregisteredDialects() on the MLIRContext, or use -allow-unregistered-dialect with the MLIR tool used.
PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace and instructions to reproduce the bug.
Stack dump:
0.      Program arguments: bin/mlir-opt -gpu-kernel-outlining --convert-scf-to-cf --convert-gpu-to-nvvm -reconcile-unrealized-casts --nvvm-attach-target=chip=sm_75 --gpu-module-to-binary -gpu-to-llvm --convert-cf-to-llvm --convert-index-to-llvm --convert-arith-to-llvm --convert-math-to-llvm --convert-func-to-llvm --finalize-memref-to-llvm -convert-nvvm-to-llvm -reconcile-unrealized-casts /tmp/crash.mlir --mlir-print-ir-after-all
 #0 0x000055f8dadf1e6d llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) /llvm/lib/Support/Unix/Signals.inc:834:11
 #1 0x000055f8dadf239b PrintStackTraceSignalHandler(void*) /llvm/lib/Support/Unix/Signals.inc:916:1
 #2 0x000055f8dadf026f llvm::sys::RunSignalHandlers() /llvm/lib/Support/Signals.cpp:104:5
 #3 0x000055f8dadf2a39 SignalHandler(int, siginfo_t*, void*) /llvm/lib/Support/Unix/Signals.inc:426:38
 #4 0x00007ff6f5b13520 (/lib/x86_64-linux-gnu/libc.so.6+0x42520)
 #5 0x00007ff6f5b679fc __pthread_kill_implementation ./nptl/./nptl/pthread_kill.c:44:76
 #6 0x00007ff6f5b679fc __pthread_kill_internal ./nptl/./nptl/pthread_kill.c:78:10
 #7 0x00007ff6f5b679fc pthread_kill ./nptl/./nptl/pthread_kill.c:89:10
 #8 0x00007ff6f5b13476 gsignal ./signal/../sysdeps/posix/raise.c:27:6
 #9 0x00007ff6f5af97f3 abort ./stdlib/./stdlib/abort.c:81:7
#10 0x000055f8dacd8885 llvm::report_fatal_error(llvm::Twine const&, bool) /llvm/lib/Support/ErrorHandling.cpp:137:5
#11 0x000055f8e74f1e81 mlir::Operation::Operation(mlir::Location, mlir::OperationName, unsigned int, unsigned int, unsigned int, unsigned int, int, mlir::DictionaryAttr, mlir::OpaqueProperties, bool) /mlir/lib/IR/Operation.cpp:186:7
#12 0x000055f8e74f1682 mlir::Operation::create(mlir::Location, mlir::OperationName, mlir::TypeRange, mlir::ValueRange, mlir::DictionaryAttr, mlir::OpaqueProperties, mlir::BlockRange, unsigned int, unsigned int) /mlir/lib/IR/Operation.cpp:126:7
#13 0x000055f8e74f13ea mlir::Operation::create(mlir::Location, mlir::OperationName, mlir::TypeRange, mlir::ValueRange, mlir::NamedAttrList&&, mlir::OpaqueProperties, mlir::BlockRange, unsigned int, unsigned int) /mlir/lib/IR/Operation.cpp:78:3
#14 0x000055f8e74f114f mlir::Operation::create(mlir::Location, mlir::OperationName, mlir::TypeRange, mlir::ValueRange, mlir::NamedAttrList&&, mlir::OpaqueProperties, mlir::BlockRange, mlir::RegionRange, unsigned int) /mlir/lib/IR/Operation.cpp:59:14
#15 0x000055f8e74f0f0e mlir::Operation::create(mlir::OperationState const&) /mlir/lib/IR/Operation.cpp:36:19
#16 0x000055f8e737c3b1 mlir::OpBuilder::create(mlir::OperationState const&) /mlir/lib/IR/Builders.cpp:458:17
#17 0x000055f8e436729a mlir::index::MulOp::create(mlir::OpBuilder&, mlir::Location, mlir::Value, mlir::Value) /build-debug/tools/mlir/include/mlir/Dialect/Index/IR/IndexOps.cpp.inc:5542:50
#18 0x000055f8de0144be (anonymous namespace)::GpuGlobalIdRewriter::matchAndRewrite(mlir::gpu::GlobalIdOp, mlir::PatternRewriter&) const /mlir/lib/Dialect/GPU/Transforms/GlobalIdRewriter.cpp:32:16
#19 0x000055f8de014d2b mlir::detail::OpOrInterfaceRewritePatternBase<mlir::gpu::GlobalIdOp>::matchAndRewrite(mlir::Operation*, mlir::PatternRewriter&) const /mlir/include/mlir/IR/PatternMatch.h:299:12
#20 0x000055f8e6f57918 mlir::PatternApplicator::matchAndRewrite(mlir::Operation*, mlir::PatternRewriter&, llvm::function_ref<bool (mlir::Pattern const&)>, llvm::function_ref<void (mlir::Pattern const&)>, llvm::function_ref<llvm::LogicalResult (mlir::Pattern const&)>)::$_0::operator()() const /mlir/lib/Rewrite/PatternApplicator.cpp:223:31
#21 0x000055f8e6f575c5 void llvm::function_ref<void ()>::callback_fn<mlir::PatternApplicator::matchAndRewrite(mlir::Operation*, mlir::PatternRewriter&, llvm::function_ref<bool (mlir::Pattern const&)>, llvm::function_ref<void (mlir::Pattern const&)>, llvm::function_ref<llvm::LogicalResult (mlir::Pattern const&)>)::$_0>(long) /llvm/include/llvm/ADT/STLFunctionalExtras.h:46:5
#22 0x000055f8dacfc929 llvm::function_ref<void ()>::operator()() const /llvm/include/llvm/ADT/STLFunctionalExtras.h:69:5
#23 0x000055f8e6f5914b void mlir::MLIRContext::executeAction<mlir::ApplyPatternAction, mlir::Pattern const&>(llvm::function_ref<void ()>, llvm::ArrayRef<mlir::IRUnit>, mlir::Pattern const&) /mlir/include/mlir/IR/MLIRContext.h:291:3
#24 0x000055f8e6f559aa mlir::PatternApplicator::matchAndRewrite(mlir::Operation*, mlir::PatternRewriter&, llvm::function_ref<bool (mlir::Pattern const&)>, llvm::function_ref<void (mlir::Pattern const&)>, llvm::function_ref<llvm::LogicalResult (mlir::Pattern const&)>) /mlir/lib/Rewrite/PatternApplicator.cpp:242:9
#25 0x000055f8e6f104f5 (anonymous namespace)::GreedyPatternRewriteDriver::processWorklist() /mlir/lib/Transforms/Utils/GreedyPatternRewriteDriver.cpp:620:17
#26 0x000055f8e6f0f807 (anonymous namespace)::RegionPatternRewriteDriver::simplify(bool*) &&::$_2::operator()() const /mlir/lib/Transforms/Utils/GreedyPatternRewriteDriver.cpp:890:31
#27 0x000055f8e6f0f775 void llvm::function_ref<void ()>::callback_fn<(anonymous namespace)::RegionPatternRewriteDriver::simplify(bool*) &&::$_2>(long) /llvm/include/llvm/ADT/STLFunctionalExtras.h:46:5
#28 0x000055f8dacfc929 llvm::function_ref<void ()>::operator()() const /llvm/include/llvm/ADT/STLFunctionalExtras.h:69:5
#29 0x000055f8e6f0ee6b void mlir::MLIRContext::executeAction<(anonymous namespace)::GreedyPatternRewriteIteration, long&>(llvm::function_ref<void ()>, llvm::ArrayRef<mlir::IRUnit>, long&) /mlir/include/mlir/IR/MLIRContext.h:291:3
#30 0x000055f8e6f0d03e (anonymous namespace)::RegionPatternRewriteDriver::simplify(bool*) && /mlir/lib/Transforms/Utils/GreedyPatternRewriteDriver.cpp:903:3
#31 0x000055f8e6f0cc34 mlir::applyPatternsGreedily(mlir::Region&, mlir::FrozenRewritePatternSet const&, mlir::GreedyRewriteConfig, bool*) /mlir/lib/Transforms/Utils/GreedyPatternRewriteDriver.cpp:935:47
#32 0x000055f8db248500 mlir::applyPatternsGreedily(mlir::Operation*, mlir::FrozenRewritePatternSet const&, mlir::GreedyRewriteConfig, bool*) /mlir/include/mlir/Transforms/GreedyPatternRewriteDriver.h:224:15
#33 0x000055f8dd720cd1 (anonymous namespace)::LowerGpuOpsToNVVMOpsPass::runOnOperation() /mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp:376:18
#34 0x000055f8e7018664 mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int)::$_3::operator()() const /mlir/lib/Pass/Pass.cpp:0:17
#35 0x000055f8e7018605 void llvm::function_ref<void ()>::callback_fn<mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int)::$_3>(long) /llvm/include/llvm/ADT/STLFunctionalExtras.h:46:5

My MLIR version:

mlir-opt --version
LLVM (http://llvm.org/):
  LLVM version 22.0.0git
  Optimized build with assertions.

Metadata

Metadata

Assignees

Labels

crashPrefer [crash-on-valid] or [crash-on-invalid]mlir

Type

Projects

No projects

Milestone

No milestone

Relationships

None yet

Development

No branches or pull requests

Issue actions