[Mlir-commits] [llvm] [mlir] [mlir][nvvm] Revert #206557 (PR #206770)

[Jeremy Kun]

https://github.com/j2kun updated https://github.com/llvm/llvm-project/pull/206770

>From c8bc4508741a8d9874631bfee7c069fca4474c43 Mon Sep 17 00:00:00 2001
From: Jeremy Kun <jkun at google.com>
Date: Tue, 30 Jun 2026 09:51:22 -0700
Subject: [PATCH] Revert "[NFC][mlir][nvvm] extract NVVM inliner interface from
 LLVM dialect"

This reverts commit 05eca6fc82f152aaacfd265ffa17cde7f7cee1ed.
---
 .../Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.h  |  6 ++++++
 mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp            | 11 -----------
 mlir/lib/Dialect/LLVMIR/Transforms/CMakeLists.txt     |  1 +
 .../LLVMIR/Transforms/InlinerInterfaceImpl.cpp        |  7 +++++++
 mlir/lib/RegisterAllDialects.cpp                      |  1 +
 utils/bazel/llvm-project-overlay/mlir/BUILD.bazel     |  2 +-
 6 files changed, 16 insertions(+), 12 deletions(-)

diff --git a/mlir/include/mlir/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.h b/mlir/include/mlir/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.h
index 472a049d567fd..69cc2e32285b6 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.h
+++ b/mlir/include/mlir/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.h
@@ -24,6 +24,12 @@ void registerInlinerInterface(DialectRegistry &registry);
 
 } // namespace LLVM
 
+namespace NVVM {
+/// Register the `NVVMInlinerInterface` implementation of
+/// `DialectInlinerInterface` with the NVVM dialect.
+void registerInlinerInterface(DialectRegistry &registry);
+} // namespace NVVM
+
 } // namespace mlir
 
 #endif // MLIR_DIALECT_LLVMIR_TRANSFORMS_INLINERINTERFACEIMPL_H
diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
index 14406ab55f138..13de1e7f2b2ad 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -28,7 +28,6 @@
 #include "mlir/IR/Operation.h"
 #include "mlir/IR/OperationSupport.h"
 #include "mlir/IR/Types.h"
-#include "mlir/Transforms/InliningUtils.h"
 #include "llvm/ADT/STLExtras.h"
 #include "llvm/ADT/TypeSwitch.h"
 #include "llvm/IR/IRBuilder.h"
@@ -6331,15 +6330,6 @@ LogicalResult Tcgen05LdRedOp::verify() {
 // NVVMDialect initialization, type parsing, and registration.
 //===----------------------------------------------------------------------===//
 
-namespace {
-struct NVVMInlinerInterface final : DialectInlinerInterface {
-  using DialectInlinerInterface::DialectInlinerInterface;
-  bool isLegalToInline(Operation *, Region *, bool, IRMapping &) const final {
-    return true;
-  }
-};
-} // namespace
-
 // TODO: This should be the llvm.nvvm dialect once this is supported.
 void NVVMDialect::initialize() {
   addOperations<
@@ -6354,7 +6344,6 @@ void NVVMDialect::initialize() {
   // Support unknown operations because not all NVVM operations are
   // registered.
   allowUnknownOperations();
-  addInterfaces<NVVMInlinerInterface>();
   declarePromisedInterface<ConvertToLLVMPatternInterface, NVVMDialect>();
   declarePromisedInterface<gpu::TargetAttrInterface, NVVMTargetAttr>();
 }
diff --git a/mlir/lib/Dialect/LLVMIR/Transforms/CMakeLists.txt b/mlir/lib/Dialect/LLVMIR/Transforms/CMakeLists.txt
index fff166d39dad9..63046b1ced914 100644
--- a/mlir/lib/Dialect/LLVMIR/Transforms/CMakeLists.txt
+++ b/mlir/lib/Dialect/LLVMIR/Transforms/CMakeLists.txt
@@ -17,5 +17,6 @@ add_mlir_dialect_library(MLIRLLVMIRTransforms
   MLIRLLVMDialect
   MLIRPass
   MLIRTransforms
+  MLIRNVVMDialect
   MLIROpenMPDialect
   )
diff --git a/mlir/lib/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.cpp b/mlir/lib/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.cpp
index 657835ae54149..680fafcc099f2 100644
--- a/mlir/lib/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.cpp
+++ b/mlir/lib/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.cpp
@@ -14,6 +14,7 @@
 #include "mlir/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.h"
 #include "mlir/Analysis/SliceWalk.h"
 #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
+#include "mlir/Dialect/LLVMIR/NVVMDialect.h"
 #include "mlir/IR/Matchers.h"
 #include "mlir/Interfaces/DataLayoutInterfaces.h"
 #include "mlir/Interfaces/ViewLikeInterface.h"
@@ -860,3 +861,9 @@ void mlir::LLVM::registerInlinerInterface(DialectRegistry &registry) {
     dialect->addInterfaces<LLVMInlinerInterface>();
   });
 }
+
+void mlir::NVVM::registerInlinerInterface(DialectRegistry &registry) {
+  registry.addExtension(+[](MLIRContext *ctx, NVVM::NVVMDialect *dialect) {
+    dialect->addInterfaces<LLVMInlinerInterface>();
+  });
+}
diff --git a/mlir/lib/RegisterAllDialects.cpp b/mlir/lib/RegisterAllDialects.cpp
index 974b5f533860a..2f55296f424cd 100644
--- a/mlir/lib/RegisterAllDialects.cpp
+++ b/mlir/lib/RegisterAllDialects.cpp
@@ -176,6 +176,7 @@ void mlir::registerAllDialects(DialectRegistry &registry) {
   gpu::registerIndexedAccessOpInterfaceExternalModels(registry);
   gpu::registerValueBoundsOpInterfaceExternalModels(registry);
   LLVM::registerInlinerInterface(registry);
+  NVVM::registerInlinerInterface(registry);
   linalg::registerAllDialectInterfaceImplementations(registry);
   linalg::registerRuntimeVerifiableOpInterfaceExternalModels(registry);
   memref::registerAllocationOpInterfaceExternalModels(registry);
diff --git a/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel b/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel
index e9bb52e88ace7..12b8ee89ffb04 100644
--- a/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel
+++ b/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel
@@ -5570,6 +5570,7 @@ cc_library(
         ":LLVMDialect",
         ":LLVMIRTransformsLegalizeForExport",
         ":LLVMPassIncGen",
+        ":NVVMDialect",
         ":Pass",
         ":TransformUtils",
         ":ViewLikeInterface",
@@ -6411,7 +6412,6 @@ cc_library(
         ":GPUDialect",
         ":IR",
         ":InferIntRangeInterface",
-        ":InliningUtils",
         ":LLVMDialect",
         ":NVVMOpsIncGen",
         ":NVVMRequiresSMTraitsIncGen",