[llvm] [mlir] [mlir][nvvm] extract NVVM inliner interface from LLVM dialect (PR #206557)

[Jeremy Kun via llvm-commits]

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

>From 58c6225715ec85d29ba8f869ed68ee15c843216d Mon Sep 17 00:00:00 2001
From: Jeremy Kun <jkun at google.com>
Date: Mon, 29 Jun 2026 11:27:07 -0700
Subject: [PATCH] [NFC][mlir][nvvm] extract NVVM inliner interface from LLVM
 dialect

---
 .../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, 12 insertions(+), 16 deletions(-)

diff --git a/mlir/include/mlir/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.h b/mlir/include/mlir/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.h
index 69cc2e32285b6..472a049d567fd 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.h
+++ b/mlir/include/mlir/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.h
@@ -24,12 +24,6 @@ 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 13de1e7f2b2ad..14406ab55f138 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -28,6 +28,7 @@
 #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"
@@ -6330,6 +6331,15 @@ 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<
@@ -6344,6 +6354,7 @@ 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 63046b1ced914..fff166d39dad9 100644
--- a/mlir/lib/Dialect/LLVMIR/Transforms/CMakeLists.txt
+++ b/mlir/lib/Dialect/LLVMIR/Transforms/CMakeLists.txt
@@ -17,6 +17,5 @@ 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 680fafcc099f2..657835ae54149 100644
--- a/mlir/lib/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.cpp
+++ b/mlir/lib/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.cpp
@@ -14,7 +14,6 @@
 #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"
@@ -861,9 +860,3 @@ 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 2f55296f424cd..974b5f533860a 100644
--- a/mlir/lib/RegisterAllDialects.cpp
+++ b/mlir/lib/RegisterAllDialects.cpp
@@ -176,7 +176,6 @@ 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 12b8ee89ffb04..e9bb52e88ace7 100644
--- a/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel
+++ b/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel
@@ -5570,7 +5570,6 @@ cc_library(
         ":LLVMDialect",
         ":LLVMIRTransformsLegalizeForExport",
         ":LLVMPassIncGen",
-        ":NVVMDialect",
         ":Pass",
         ":TransformUtils",
         ":ViewLikeInterface",
@@ -6412,6 +6411,7 @@ cc_library(
         ":GPUDialect",
         ":IR",
         ":InferIntRangeInterface",
+        ":InliningUtils",
         ":LLVMDialect",
         ":NVVMOpsIncGen",
         ":NVVMRequiresSMTraitsIncGen",