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

via llvm-commits llvm-commits at lists.llvm.org
Mon Jun 29 11:56:43 PDT 2026 

llvmorg-github-actions[bot] wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-mlir-llvm

Author: Jeremy Kun (j2kun)

<details>
<summary>Changes</summary>

marking as draft until #<!-- -->206520 is submitted

---
Full diff: https://github.com/llvm/llvm-project/pull/206557.diff


20 Files Affected:

- (modified) mlir/docs/Passes.md (+4) 
- (modified) mlir/include/mlir/Dialect/CMakeLists.txt (+1) 
- (modified) mlir/include/mlir/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.h (-6) 
- (modified) mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.h (-1) 
- (modified) mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.td (-5) 
- (added) mlir/include/mlir/Dialect/NVVM/CMakeLists.txt (+1) 
- (added) mlir/include/mlir/Dialect/NVVM/Transforms/CMakeLists.txt (+5) 
- (renamed) mlir/include/mlir/Dialect/NVVM/Transforms/OptimizeForNVVM.h (+7-7) 
- (added) mlir/include/mlir/Dialect/NVVM/Transforms/Passes.h (+26) 
- (added) mlir/include/mlir/Dialect/NVVM/Transforms/Passes.td (+18) 
- (modified) mlir/lib/Dialect/CMakeLists.txt (+1) 
- (modified) mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp (+11) 
- (modified) mlir/lib/Dialect/LLVMIR/Transforms/CMakeLists.txt (-2) 
- (modified) mlir/lib/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.cpp (-7) 
- (added) mlir/lib/Dialect/NVVM/CMakeLists.txt (+1) 
- (added) mlir/lib/Dialect/NVVM/Transforms/CMakeLists.txt (+16) 
- (renamed) mlir/lib/Dialect/NVVM/Transforms/OptimizeForNVVM.cpp (+6-6) 
- (modified) mlir/lib/RegisterAllDialects.cpp (-1) 
- (modified) mlir/lib/RegisterAllPasses.cpp (+2) 
- (modified) utils/bazel/llvm-project-overlay/mlir/BUILD.bazel (+30-1) 


``````````diff
diff --git a/mlir/docs/Passes.md b/mlir/docs/Passes.md
index 0226add6d0838..9e70da8a3ec97 100644
--- a/mlir/docs/Passes.md
+++ b/mlir/docs/Passes.md
@@ -92,6 +92,10 @@ This document describes the available MLIR passes and their contracts.
 
 [include "NVGPUPasses.md"]
 
+## 'nvvm' Dialect Passes
+
+[include "NVVMPasses.md"]
+
 ## 'quant' Dialect Passes
 
 [include "QuantPasses.md"]
diff --git a/mlir/include/mlir/Dialect/CMakeLists.txt b/mlir/include/mlir/Dialect/CMakeLists.txt
index d2505877e2dd0..d7ea0cc3df267 100644
--- a/mlir/include/mlir/Dialect/CMakeLists.txt
+++ b/mlir/include/mlir/Dialect/CMakeLists.txt
@@ -22,6 +22,7 @@ add_subdirectory(Shard)
 add_subdirectory(MLProgram)
 add_subdirectory(MPI)
 add_subdirectory(NVGPU)
+add_subdirectory(NVVM)
 add_subdirectory(OpenACC)
 add_subdirectory(OpenACCMPCommon)
 add_subdirectory(OpenMP)
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/include/mlir/Dialect/LLVMIR/Transforms/Passes.h b/mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.h
index c598ffe3602dc..16076fe314818 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.h
+++ b/mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.h
@@ -12,7 +12,6 @@
 #include "mlir/Dialect/LLVMIR/LLVMAttrs.h"
 #include "mlir/Dialect/LLVMIR/Transforms/AddComdats.h"
 #include "mlir/Dialect/LLVMIR/Transforms/LegalizeForExport.h"
-#include "mlir/Dialect/LLVMIR/Transforms/OptimizeForNVVM.h"
 #include "mlir/Dialect/LLVMIR/Transforms/RequestCWrappers.h"
 #include "mlir/Pass/Pass.h"
 
diff --git a/mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.td b/mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.td
index eaffd2978047b..73fc9cae23ac5 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.td
@@ -63,11 +63,6 @@ def LLVMUseDefaultVisibilityPass
 	   )}]>,
   ];
 }
-
-def NVVMOptimizeForTargetPass : Pass<"llvm-optimize-for-nvvm-target"> {
-  let summary = "Optimize NVVM IR";
-}
-
 def DIScopeForLLVMFuncOpPass : Pass<"ensure-debug-info-scope-on-llvm-func", "::mlir::ModuleOp"> {
   let summary = "Materialize LLVM debug info subprogram attribute on every LLVMFuncOp";
   let description = [{
diff --git a/mlir/include/mlir/Dialect/NVVM/CMakeLists.txt b/mlir/include/mlir/Dialect/NVVM/CMakeLists.txt
new file mode 100644
index 0000000000000..e31af32661164
--- /dev/null
+++ b/mlir/include/mlir/Dialect/NVVM/CMakeLists.txt
@@ -0,0 +1 @@
+add_subdirectory(Transforms)
diff --git a/mlir/include/mlir/Dialect/NVVM/Transforms/CMakeLists.txt b/mlir/include/mlir/Dialect/NVVM/Transforms/CMakeLists.txt
new file mode 100644
index 0000000000000..4bf064bdbe231
--- /dev/null
+++ b/mlir/include/mlir/Dialect/NVVM/Transforms/CMakeLists.txt
@@ -0,0 +1,5 @@
+set(LLVM_TARGET_DEFINITIONS Passes.td)
+mlir_tablegen(Passes.h.inc -gen-pass-decls -name NVVM)
+add_mlir_dialect_tablegen_target(MLIRNVVMPassIncGen)
+
+add_mlir_doc(Passes NVVMPasses ./ -gen-pass-doc)
diff --git a/mlir/include/mlir/Dialect/LLVMIR/Transforms/OptimizeForNVVM.h b/mlir/include/mlir/Dialect/NVVM/Transforms/OptimizeForNVVM.h
similarity index 54%
rename from mlir/include/mlir/Dialect/LLVMIR/Transforms/OptimizeForNVVM.h
rename to mlir/include/mlir/Dialect/NVVM/Transforms/OptimizeForNVVM.h
index 3f625a4e484b3..f02d7dbe9778c 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/Transforms/OptimizeForNVVM.h
+++ b/mlir/include/mlir/Dialect/NVVM/Transforms/OptimizeForNVVM.h
@@ -1,4 +1,4 @@
-//===- OptimizeForNVVM.h - Optimize LLVM IR for NVVM -*- C++ -*-===//
+//===- OptimizeForNVVM.h - Optimize LLVM IR for NVVM ------------*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,17 +6,17 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef MLIR_DIALECT_LLVMIR_TRANSFORMS_OPTIMIZENVVM_H
-#define MLIR_DIALECT_LLVMIR_TRANSFORMS_OPTIMIZENVVM_H
+#ifndef MLIR_DIALECT_NVVM_TRANSFORMS_OPTIMIZEFORNVVM_H
+#define MLIR_DIALECT_NVVM_TRANSFORMS_OPTIMIZEFORNVVM_H
 
 #include <memory>
 
 namespace mlir {
 class Pass;
-namespace LLVM {
+namespace NVVM {
 #define GEN_PASS_DECL_NVVMOPTIMIZEFORTARGETPASS
-#include "mlir/Dialect/LLVMIR/Transforms/Passes.h.inc"
-} // namespace LLVM
+#include "mlir/Dialect/NVVM/Transforms/Passes.h.inc"
+} // namespace NVVM
 } // namespace mlir
 
-#endif // MLIR_DIALECT_LLVMIR_TRANSFORMS_OPTIMIZENVVM_H
+#endif // MLIR_DIALECT_NVVM_TRANSFORMS_OPTIMIZEFORNVVM_H
diff --git a/mlir/include/mlir/Dialect/NVVM/Transforms/Passes.h b/mlir/include/mlir/Dialect/NVVM/Transforms/Passes.h
new file mode 100644
index 0000000000000..b81e5ae991016
--- /dev/null
+++ b/mlir/include/mlir/Dialect/NVVM/Transforms/Passes.h
@@ -0,0 +1,26 @@
+//===- Passes.h - NVVM Pass Construction and Registration -------*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef MLIR_DIALECT_NVVM_TRANSFORMS_PASSES_H
+#define MLIR_DIALECT_NVVM_TRANSFORMS_PASSES_H
+
+#include "mlir/Dialect/NVVM/Transforms/OptimizeForNVVM.h"
+#include "mlir/Pass/Pass.h"
+
+namespace mlir {
+namespace NVVM {
+
+/// Generate the code for registering passes.
+#define GEN_PASS_DECL
+#define GEN_PASS_REGISTRATION
+#include "mlir/Dialect/NVVM/Transforms/Passes.h.inc"
+
+} // namespace NVVM
+} // namespace mlir
+
+#endif // MLIR_DIALECT_NVVM_TRANSFORMS_PASSES_H
diff --git a/mlir/include/mlir/Dialect/NVVM/Transforms/Passes.td b/mlir/include/mlir/Dialect/NVVM/Transforms/Passes.td
new file mode 100644
index 0000000000000..b69b368e93e4e
--- /dev/null
+++ b/mlir/include/mlir/Dialect/NVVM/Transforms/Passes.td
@@ -0,0 +1,18 @@
+//===-- Passes.td - NVVM pass definition file --------------*- tablegen -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef MLIR_DIALECT_NVVM_TRANSFORMS_PASSES
+#define MLIR_DIALECT_NVVM_TRANSFORMS_PASSES
+
+include "mlir/Pass/PassBase.td"
+
+def NVVMOptimizeForTargetPass : Pass<"llvm-optimize-for-nvvm-target"> {
+  let summary = "Optimize NVVM IR";
+}
+
+#endif // MLIR_DIALECT_NVVM_TRANSFORMS_PASSES
diff --git a/mlir/lib/Dialect/CMakeLists.txt b/mlir/lib/Dialect/CMakeLists.txt
index 66f68c369f81f..55e75f3f0b9ce 100644
--- a/mlir/lib/Dialect/CMakeLists.txt
+++ b/mlir/lib/Dialect/CMakeLists.txt
@@ -22,6 +22,7 @@ add_subdirectory(Shard)
 add_subdirectory(MLProgram)
 add_subdirectory(MPI)
 add_subdirectory(NVGPU)
+add_subdirectory(NVVM)
 add_subdirectory(OpenACC)
 add_subdirectory(OpenACCMPCommon)
 add_subdirectory(OpenMP)
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 6ffa2d9a8a9e8..fff166d39dad9 100644
--- a/mlir/lib/Dialect/LLVMIR/Transforms/CMakeLists.txt
+++ b/mlir/lib/Dialect/LLVMIR/Transforms/CMakeLists.txt
@@ -5,7 +5,6 @@ add_mlir_dialect_library(MLIRLLVMIRTransforms
   DIScopeForLLVMFuncOp.cpp
   InlinerInterfaceImpl.cpp
   LegalizeForExport.cpp
-  OptimizeForNVVM.cpp
   RequestCWrappers.cpp
   UseDefaultVisibilityPass.cpp
 
@@ -18,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/Dialect/NVVM/CMakeLists.txt b/mlir/lib/Dialect/NVVM/CMakeLists.txt
new file mode 100644
index 0000000000000..e31af32661164
--- /dev/null
+++ b/mlir/lib/Dialect/NVVM/CMakeLists.txt
@@ -0,0 +1 @@
+add_subdirectory(Transforms)
diff --git a/mlir/lib/Dialect/NVVM/Transforms/CMakeLists.txt b/mlir/lib/Dialect/NVVM/Transforms/CMakeLists.txt
new file mode 100644
index 0000000000000..935fa5700c7e5
--- /dev/null
+++ b/mlir/lib/Dialect/NVVM/Transforms/CMakeLists.txt
@@ -0,0 +1,16 @@
+add_mlir_dialect_library(MLIRNVVMTransforms
+  OptimizeForNVVM.cpp
+
+  ADDITIONAL_HEADER_DIRS
+  ${MLIR_MAIN_INCLUDE_DIR}/mlir/Dialect/NVVM
+
+  DEPENDS
+  MLIRNVVMPassIncGen
+
+  LINK_LIBS PUBLIC
+  MLIRIR
+  MLIRLLVMDialect
+  MLIRNVVMDialect
+  MLIRPass
+  MLIRTransforms
+  )
diff --git a/mlir/lib/Dialect/LLVMIR/Transforms/OptimizeForNVVM.cpp b/mlir/lib/Dialect/NVVM/Transforms/OptimizeForNVVM.cpp
similarity index 94%
rename from mlir/lib/Dialect/LLVMIR/Transforms/OptimizeForNVVM.cpp
rename to mlir/lib/Dialect/NVVM/Transforms/OptimizeForNVVM.cpp
index 7f34f7d142d73..0994946d09686 100644
--- a/mlir/lib/Dialect/LLVMIR/Transforms/OptimizeForNVVM.cpp
+++ b/mlir/lib/Dialect/NVVM/Transforms/OptimizeForNVVM.cpp
@@ -1,4 +1,4 @@
-//===- OptimizeForNVVM.cpp - Optimize LLVM IR for NVVM ---------===//
+//===- OptimizeForNVVM.cpp - Optimize LLVM IR for NVVM --------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,7 +6,7 @@
 //
 //===----------------------------------------------------------------------===//
 
-#include "mlir/Dialect/LLVMIR/Transforms/OptimizeForNVVM.h"
+#include "mlir/Dialect/NVVM/Transforms/OptimizeForNVVM.h"
 
 #include "mlir/Dialect/LLVMIR/NVVMDialect.h"
 #include "mlir/IR/Builders.h"
@@ -15,10 +15,10 @@
 #include "mlir/Transforms/GreedyPatternRewriteDriver.h"
 
 namespace mlir {
-namespace LLVM {
+namespace NVVM {
 #define GEN_PASS_DEF_NVVMOPTIMIZEFORTARGETPASS
-#include "mlir/Dialect/LLVMIR/Transforms/Passes.h.inc"
-} // namespace LLVM
+#include "mlir/Dialect/NVVM/Transforms/Passes.h.inc"
+} // namespace NVVM
 } // namespace mlir
 
 using namespace mlir;
@@ -40,7 +40,7 @@ struct ExpandDivF16 : public OpRewritePattern<LLVM::FDivOp> {
 };
 
 struct NVVMOptimizeForTarget
-    : public LLVM::impl::NVVMOptimizeForTargetPassBase<NVVMOptimizeForTarget> {
+    : public NVVM::impl::NVVMOptimizeForTargetPassBase<NVVMOptimizeForTarget> {
   void runOnOperation() override;
 
   void getDependentDialects(DialectRegistry &registry) const override {
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/mlir/lib/RegisterAllPasses.cpp b/mlir/lib/RegisterAllPasses.cpp
index c645d737cb766..76aac7ae1f182 100644
--- a/mlir/lib/RegisterAllPasses.cpp
+++ b/mlir/lib/RegisterAllPasses.cpp
@@ -32,6 +32,7 @@
 #include "mlir/Dialect/Math/Transforms/Passes.h"
 #include "mlir/Dialect/MemRef/Transforms/Passes.h"
 #include "mlir/Dialect/NVGPU/Transforms/Passes.h"
+#include "mlir/Dialect/NVVM/Transforms/Passes.h"
 #include "mlir/Dialect/OpenACC/Transforms/Passes.h"
 #include "mlir/Dialect/OpenMP/Transforms/Passes.h"
 #include "mlir/Dialect/Quant/Transforms/Passes.h"
@@ -74,6 +75,7 @@ void mlir::registerAllPasses() {
   registerGPUPasses();
   registerLinalgPasses();
   registerNVGPUPasses();
+  NVVM::registerNVVMPasses();
   registerSparseTensorPasses();
   LLVM::registerLLVMPasses();
   LLVM::registerTargetLLVMIRTransformsPasses();
diff --git a/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel b/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel
index 442b3a6cd421c..0add5598e21f3 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",
@@ -6425,6 +6425,34 @@ cc_library(
     ],
 )
 
+gentbl_cc_library(
+    name = "NVVMPassIncGen",
+    strip_include_prefix = "include",
+    tbl_outs = {
+        "include/mlir/Dialect/NVVM/Transforms/Passes.h.inc": ["-gen-pass-decls -name NVVM"],
+    },
+    tblgen = ":mlir-tblgen",
+    td_file = "include/mlir/Dialect/NVVM/Transforms/Passes.td",
+    deps = [
+        ":PassBaseTdFiles",
+    ],
+)
+
+cc_library(
+    name = "NVVMTransforms",
+    srcs = glob(["lib/Dialect/NVVM/Transforms/*.cpp"]),
+    hdrs = glob(["include/mlir/Dialect/NVVM/Transforms/*.h"]),
+    includes = ["include"],
+    deps = [
+        ":IR",
+        ":LLVMDialect",
+        ":NVVMDialect",
+        ":NVVMPassIncGen",
+        ":Pass",
+        ":TransformUtils",
+    ],
+)
+
 cc_library(
     name = "TargetLLVM",
     srcs = ["lib/Target/LLVM/ModuleToObject.cpp"],
@@ -9806,6 +9834,7 @@ cc_library(
         ":MathTransforms",
         ":MemRefTransforms",
         ":NVGPUTransforms",
+        ":NVVMTransforms",
         ":OpenACCTransforms",
         ":OpenMPTransforms",
         ":QuantTransforms",

``````````

</details>


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

More information about the llvm-commits mailing list