[Mlir-commits] [mlir] [MLIR][NVGPU] Use NVVM enums in NVGPU dialect (PR #195812)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Tue May 12 06:52:39 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-mlir-nvgpu
@llvm/pr-subscribers-mlir-llvm
Author: Srinivasa Ravi (Wolfram70)
<details>
<summary>Changes</summary>
Updates the `nvgpu.rcp` Op to use the NVVM `FPRoundingModeAttr`
attribute instead of redefining the attribute in the NVGPU dialect.
Follows up https://github.com/llvm/llvm-project/pull/195811.
---
Patch is 23.27 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/195812.diff
10 Files Affected:
- (added) mlir/include/mlir/Dialect/LLVMIR/NVVMDialect.td (+94)
- (added) mlir/include/mlir/Dialect/LLVMIR/NVVMEnums.td (+72)
- (modified) mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td (+7-136)
- (modified) mlir/include/mlir/Dialect/NVGPU/IR/NVGPU.td (+2)
- (modified) mlir/include/mlir/Dialect/NVGPU/IR/NVGPUDialect.h (+1)
- (modified) mlir/include/mlir/Dialect/NVGPU/IR/NVGPUOps.td (+4-3)
- (modified) mlir/lib/Dialect/NVGPU/IR/CMakeLists.txt (+1)
- (modified) mlir/lib/Dialect/NVGPU/IR/NVGPUDialect.cpp (+9-4)
- (modified) mlir/test/Conversion/NVGPUToNVVM/nvgpu-to-nvvm.mlir (+1-1)
- (modified) mlir/test/Dialect/NVGPU/invalid.mlir (+6-6)
``````````diff
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMDialect.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMDialect.td
new file mode 100644
index 0000000000000..025e093ebd8b6
--- /dev/null
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMDialect.td
@@ -0,0 +1,94 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// This file contains the declaration of the NVVM IR dialect.
+///
+//===----------------------------------------------------------------------===//
+
+#ifndef NVVMIR_DIALECT
+#define NVVMIR_DIALECT
+
+include "mlir/IR/DialectBase.td"
+
+def NVVM_Dialect : Dialect {
+ let name = "nvvm";
+ let cppNamespace = "::mlir::NVVM";
+ let dependentDialects = ["LLVM::LLVMDialect"];
+ let hasOperationAttrVerify = 1;
+
+ let extraClassDeclaration = [{
+ /// Get the name of the attribute used to annotate external kernel
+ /// functions.
+ static StringRef getKernelFuncAttrName() { return "nvvm.kernel"; }
+ /// Get the name of the attribute used to annotate max threads required
+ /// per CTA for kernel functions.
+ static StringRef getMaxntidAttrName() { return "nvvm.maxntid"; }
+ /// Get the name of the metadata names for each dimension
+ static StringRef getMaxntidXName() { return "maxntidx"; }
+ static StringRef getMaxntidYName() { return "maxntidy"; }
+ static StringRef getMaxntidZName() { return "maxntidz"; }
+
+ /// Get the name of the attribute used to annotate exact threads required
+ /// per CTA for kernel functions.
+ static StringRef getReqntidAttrName() { return "nvvm.reqntid"; }
+ /// Get the name of the metadata names for each dimension
+ static StringRef getReqntidXName() { return "reqntidx"; }
+ static StringRef getReqntidYName() { return "reqntidy"; }
+ static StringRef getReqntidZName() { return "reqntidz"; }
+
+ /// Get the name of the attribute used to annotate exact CTAs required
+ /// per cluster for kernel functions.
+ static StringRef getClusterDimAttrName() { return "nvvm.cluster_dim"; }
+ /// Get the name of the metadata names for each dimension
+ static StringRef getClusterDimXName() { return "cluster_dim_x"; }
+ static StringRef getClusterDimYName() { return "cluster_dim_y"; }
+ static StringRef getClusterDimZName() { return "cluster_dim_z"; }
+
+ /// Get the name of the attribute used to annotate maximum number of
+ /// CTAs per cluster for kernel functions.
+ static StringRef getClusterMaxBlocksAttrName() { return "nvvm.cluster_max_blocks"; }
+
+ /// Get the name of the attribute used to annotate min CTA required
+ /// per SM for kernel functions.
+ static StringRef getMinctasmAttrName() { return "nvvm.minctasm"; }
+
+ /// Get the name of the attribute used to annotate max number of
+ /// registers that can be allocated per thread.
+ static StringRef getMaxnregAttrName() { return "nvvm.maxnreg"; }
+
+ /// Get the name of the attribute used to annotate kernel arguments that
+ /// are grid constants.
+ static StringRef getGridConstantAttrName() { return "nvvm.grid_constant"; }
+
+ /// Get the name of the attribute used to annotate the `.blocksareclusters`
+ /// PTX directive for kernel functions.
+ /// This attribute implies that the grid launch configuration for the
+ /// corresponding kernel function is specifying the number of clusters
+ /// instead of the number of thread blocks. This attribute is only
+ /// allowed for kernel functions and requires nvvm.reqntid and
+ /// nvvm.cluster_dim attributes.
+ static StringRef getBlocksAreClustersAttrName() { return "nvvm.blocksareclusters"; }
+
+ /// Get the name of the attribute used to annotate managed global variables.
+ static StringRef getManagedAttrName() { return "nvvm.managed"; }
+
+ /// Verify an attribute from this dialect on the argument at 'argIndex' for
+ /// the region at 'regionIndex' on the given operation. Returns failure if
+ /// the verification failed, success otherwise. This hook may optionally be
+ /// invoked from any operation containing a region.
+ LogicalResult verifyRegionArgAttribute(Operation *op,
+ unsigned regionIndex,
+ unsigned argIndex,
+ NamedAttribute argAttr) override;
+ }];
+
+ let useDefaultAttributePrinterParser = 1;
+}
+
+#endif // NVVMIR_DIALECT
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMEnums.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMEnums.td
new file mode 100644
index 0000000000000..42d196c5662d1
--- /dev/null
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMEnums.td
@@ -0,0 +1,72 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// This file contains the declaration of the NVVM IR enum attributes.
+///
+//===----------------------------------------------------------------------===//
+
+#ifndef NVVMIR_ENUMS
+#define NVVMIR_ENUMS
+
+include "mlir/Dialect/LLVMIR/NVVMDialect.td"
+include "mlir/IR/EnumAttr.td"
+
+// Attributes for the floating point rounding modes supported by PTX
+def FPRoundingModeNone : I32EnumAttrCase<"NONE", 0, "none">;
+def FPRoundingModeRN : I32EnumAttrCase<"RN", 1, "rn">;
+def FPRoundingModeRM : I32EnumAttrCase<"RM", 2, "rm">;
+def FPRoundingModeRP : I32EnumAttrCase<"RP", 3, "rp">;
+def FPRoundingModeRZ : I32EnumAttrCase<"RZ", 4, "rz">;
+def FPRoundingModeRNA : I32EnumAttrCase<"RNA", 5, "rna">;
+def FPRoundingModeRS : I32EnumAttrCase<"RS", 6, "rs">;
+
+def FPRoundingMode : I32EnumAttr<"FPRoundingMode", "NVVM FPRoundingMode kind",
+ [FPRoundingModeNone, FPRoundingModeRN, FPRoundingModeRM,
+ FPRoundingModeRP, FPRoundingModeRZ, FPRoundingModeRNA, FPRoundingModeRS]> {
+ let genSpecializedAttr = 0;
+ let cppNamespace = "::mlir::NVVM";
+}
+def FPRoundingModeAttr : EnumAttr<NVVM_Dialect, FPRoundingMode, "fp_rnd_mode"> {
+ let assemblyFormat = "`<` $value `>`";
+}
+
+def SaturationModeNone : I32EnumAttrCase<"NONE", 0, "none">;
+def SaturationModeFinite : I32EnumAttrCase<"SATFINITE", 1, "satfinite">;
+def SaturationModeSat : I32EnumAttrCase<"SAT", 2, "sat">;
+
+def SaturationMode : I32EnumAttr<"SaturationMode", "NVVM SaturationMode kind",
+ [SaturationModeNone, SaturationModeFinite, SaturationModeSat]> {
+ let genSpecializedAttr = 0;
+ let cppNamespace = "::mlir::NVVM";
+}
+def SaturationModeAttr : EnumAttr<NVVM_Dialect, SaturationMode, "sat_mode"> {
+ let summary = "Describes the saturation mode";
+ let description = [{
+ A `nvvm.sat_mode` attribute specifies the saturation mode for instructions
+ involving floating points or integers. It can be one of the following
+ values:
+ - `none`: No saturation is applied.
+ - `satfinite`: If the absolute value of input (ignoring sign) is greater
+ than the `MAX_NORM` of the specified destination format, then the result
+ is the sign-preserved `MAX_NORM` of the destination format and a positive
+ `MAX_NORM` in unsigned datatypes for which the destination sign is not
+ supported. If the input is `NaN`, then the result can be `NaN` or the
+ `MAX_NORM` of the destination format, depending on the format.
+ - `sat`: For integer destination types, this limits the value to `MININT..
+ MAXINT` and applies to both signed and unsigned integer datatypes. For
+ floating point destination types (applies to only `F16`, `F32`, and `F64`
+ types), this limits the value to the range `[0.0, 1.0]` and flushes NaN
+ results to positive zero.
+
+ [For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-cvt)
+}];
+ let assemblyFormat = "`<` $value `>`";
+}
+
+#endif // NVVMIR_ENUMS
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 51ff22dfdc65c..0d271acd862ba 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -13,17 +13,19 @@
#ifndef NVVMIR_OPS
#define NVVMIR_OPS
-include "mlir/IR/EnumAttr.td"
include "mlir/Dialect/GPU/IR/CompilationAttrInterfaces.td"
+include "mlir/Dialect/LLVMIR/BasicPtxBuilderInterface.td"
include "mlir/Dialect/LLVMIR/LLVMOpBase.td"
+include "mlir/Dialect/LLVMIR/LLVMTypes.td"
+include "mlir/Dialect/LLVMIR/NVVMDialect.td"
+include "mlir/Dialect/LLVMIR/NVVMEnums.td"
include "mlir/Dialect/LLVMIR/NVVMRequiresSMTraits.td"
include "mlir/Dialect/Ptr/IR/MemorySpaceInterfaces.td"
-include "mlir/Interfaces/SideEffectInterfaces.td"
-include "mlir/Dialect/LLVMIR/BasicPtxBuilderInterface.td"
+include "mlir/IR/CommonAttrConstraints.td"
+include "mlir/IR/EnumAttr.td"
include "mlir/Interfaces/InferIntRangeInterface.td"
include "mlir/Interfaces/InferTypeOpInterface.td"
-include "mlir/Dialect/LLVMIR/LLVMTypes.td"
-include "mlir/IR/CommonAttrConstraints.td"
+include "mlir/Interfaces/SideEffectInterfaces.td"
def LLVM_PointerGeneric : LLVM_PointerInAddressSpace<0>;
def LLVM_PointerGlobal : LLVM_PointerInAddressSpace<1>;
@@ -33,85 +35,6 @@ def LLVM_PointerLocal : LLVM_PointerInAddressSpace<5>;
def LLVM_PointerTensor : LLVM_PointerInAddressSpace<6>;
def LLVM_PointerSharedCluster : LLVM_PointerInAddressSpace<7>;
-//===----------------------------------------------------------------------===//
-// NVVM dialect definitions
-//===----------------------------------------------------------------------===//
-
-def NVVM_Dialect : Dialect {
- let name = "nvvm";
- let cppNamespace = "::mlir::NVVM";
- let dependentDialects = ["LLVM::LLVMDialect"];
- let hasOperationAttrVerify = 1;
-
- let extraClassDeclaration = [{
- /// Get the name of the attribute used to annotate external kernel
- /// functions.
- static StringRef getKernelFuncAttrName() { return "nvvm.kernel"; }
- /// Get the name of the attribute used to annotate max threads required
- /// per CTA for kernel functions.
- static StringRef getMaxntidAttrName() { return "nvvm.maxntid"; }
- /// Get the name of the metadata names for each dimension
- static StringRef getMaxntidXName() { return "maxntidx"; }
- static StringRef getMaxntidYName() { return "maxntidy"; }
- static StringRef getMaxntidZName() { return "maxntidz"; }
-
- /// Get the name of the attribute used to annotate exact threads required
- /// per CTA for kernel functions.
- static StringRef getReqntidAttrName() { return "nvvm.reqntid"; }
- /// Get the name of the metadata names for each dimension
- static StringRef getReqntidXName() { return "reqntidx"; }
- static StringRef getReqntidYName() { return "reqntidy"; }
- static StringRef getReqntidZName() { return "reqntidz"; }
-
- /// Get the name of the attribute used to annotate exact CTAs required
- /// per cluster for kernel functions.
- static StringRef getClusterDimAttrName() { return "nvvm.cluster_dim"; }
- /// Get the name of the metadata names for each dimension
- static StringRef getClusterDimXName() { return "cluster_dim_x"; }
- static StringRef getClusterDimYName() { return "cluster_dim_y"; }
- static StringRef getClusterDimZName() { return "cluster_dim_z"; }
-
- /// Get the name of the attribute used to annotate maximum number of
- /// CTAs per cluster for kernel functions.
- static StringRef getClusterMaxBlocksAttrName() { return "nvvm.cluster_max_blocks"; }
-
- /// Get the name of the attribute used to annotate min CTA required
- /// per SM for kernel functions.
- static StringRef getMinctasmAttrName() { return "nvvm.minctasm"; }
-
- /// Get the name of the attribute used to annotate max number of
- /// registers that can be allocated per thread.
- static StringRef getMaxnregAttrName() { return "nvvm.maxnreg"; }
-
- /// Get the name of the attribute used to annotate kernel arguments that
- /// are grid constants.
- static StringRef getGridConstantAttrName() { return "nvvm.grid_constant"; }
-
- /// Get the name of the attribute used to annotate the `.blocksareclusters`
- /// PTX directive for kernel functions.
- /// This attribute implies that the grid launch configuration for the
- /// corresponding kernel function is specifying the number of clusters
- /// instead of the number of thread blocks. This attribute is only
- /// allowed for kernel functions and requires nvvm.reqntid and
- /// nvvm.cluster_dim attributes.
- static StringRef getBlocksAreClustersAttrName() { return "nvvm.blocksareclusters"; }
-
- /// Get the name of the attribute used to annotate managed global variables.
- static StringRef getManagedAttrName() { return "nvvm.managed"; }
-
- /// Verify an attribute from this dialect on the argument at 'argIndex' for
- /// the region at 'regionIndex' on the given operation. Returns failure if
- /// the verification failed, success otherwise. This hook may optionally be
- /// invoked from any operation containing a region.
- LogicalResult verifyRegionArgAttribute(Operation *op,
- unsigned regionIndex,
- unsigned argIndex,
- NamedAttribute argAttr) override;
- }];
-
- let useDefaultAttributePrinterParser = 1;
-}
-
//===----------------------------------------------------------------------===//
// NVVM op definitions
//===----------------------------------------------------------------------===//
@@ -1917,58 +1840,6 @@ def NVVM_CpAsyncMBarrierArriveOp : NVVM_VoidIntrinsicOp<"cp.async.mbarrier.arriv
// NVVM Conversion Ops (for "cvt.*" family of PTX instructions)
//===----------------------------------------------------------------------===//
-// Attributes for the floating point rounding modes supported by PTX
-def FPRoundingModeNone : I32EnumAttrCase<"NONE", 0, "none">;
-def FPRoundingModeRN : I32EnumAttrCase<"RN", 1, "rn">;
-def FPRoundingModeRM : I32EnumAttrCase<"RM", 2, "rm">;
-def FPRoundingModeRP : I32EnumAttrCase<"RP", 3, "rp">;
-def FPRoundingModeRZ : I32EnumAttrCase<"RZ", 4, "rz">;
-def FPRoundingModeRNA : I32EnumAttrCase<"RNA", 5, "rna">;
-def FPRoundingModeRS : I32EnumAttrCase<"RS", 6, "rs">;
-
-def FPRoundingMode : I32EnumAttr<"FPRoundingMode", "NVVM FPRoundingMode kind",
- [FPRoundingModeNone, FPRoundingModeRN, FPRoundingModeRM,
- FPRoundingModeRP, FPRoundingModeRZ, FPRoundingModeRNA, FPRoundingModeRS]> {
- let genSpecializedAttr = 0;
- let cppNamespace = "::mlir::NVVM";
-}
-def FPRoundingModeAttr : EnumAttr<NVVM_Dialect, FPRoundingMode, "fp_rnd_mode"> {
- let assemblyFormat = "`<` $value `>`";
-}
-
-def SaturationModeNone : I32EnumAttrCase<"NONE", 0, "none">;
-def SaturationModeFinite : I32EnumAttrCase<"SATFINITE", 1, "satfinite">;
-def SaturationModeSat : I32EnumAttrCase<"SAT", 2, "sat">;
-
-def SaturationMode : I32EnumAttr<"SaturationMode", "NVVM SaturationMode kind",
- [SaturationModeNone, SaturationModeFinite, SaturationModeSat]> {
- let genSpecializedAttr = 0;
- let cppNamespace = "::mlir::NVVM";
-}
-def SaturationModeAttr : EnumAttr<NVVM_Dialect, SaturationMode, "sat_mode"> {
- let summary = "Describes the saturation mode";
- let description = [{
- A `nvvm.sat_mode` attribute specifies the saturation mode for instructions
- involving floating points or integers. It can be one of the following
- values:
- - `none`: No saturation is applied.
- - `satfinite`: If the absolute value of input (ignoring sign) is greater
- than the `MAX_NORM` of the specified destination format, then the result
- is the sign-preserved `MAX_NORM` of the destination format and a positive
- `MAX_NORM` in unsigned datatypes for which the destination sign is not
- supported. If the input is `NaN`, then the result can be `NaN` or th
- `MAX_NORM` of the destination format, depending on the format.
- - `sat`: For integer destination types, this limits the value to `MININT..
- MAXINT` and applies to both signed and unsigned integer datatypes. For
- floating point destination types (applies to only `F16`, `F32`, and `F64`
- types), this limits the value to the range `[0.0, 1.0]` and flushes NaN
- results to positive zero.
-
- [For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-cvt)
-}];
- let assemblyFormat = "`<` $value `>`";
-}
-
def NVVM_ConvertFloatToTF32Op : NVVM_Op<"convert.float.to.tf32"> {
let summary = "Convert the given float input to TF32";
let description = [{
diff --git a/mlir/include/mlir/Dialect/NVGPU/IR/NVGPU.td b/mlir/include/mlir/Dialect/NVGPU/IR/NVGPU.td
index 7f7a54cb0c57e..1c0d7bd1113ea 100644
--- a/mlir/include/mlir/Dialect/NVGPU/IR/NVGPU.td
+++ b/mlir/include/mlir/Dialect/NVGPU/IR/NVGPU.td
@@ -9,6 +9,7 @@
#ifndef MLIR_DIALECT_NVGPU_IR_NVGPU_TD
#define MLIR_DIALECT_NVGPU_IR_NVGPU_TD
+include "mlir/Dialect/LLVMIR/NVVMDialect.td"
include "mlir/Interfaces/InferTypeOpInterface.td"
include "mlir/Interfaces/SideEffectInterfaces.td"
include "mlir/IR/AttrTypeBase.td"
@@ -18,6 +19,7 @@ include "mlir/IR/EnumAttr.td"
def NVGPU_Dialect : Dialect {
let name = "nvgpu";
let cppNamespace = "::mlir::nvgpu";
+ let dependentDialects = ["NVVM::NVVMDialect"];
let description = [{
The `NVGPU` dialect provides a bridge between higher-level target-agnostic
dialects (GPU and Vector) and the lower-level target-specific dialect
diff --git a/mlir/include/mlir/Dialect/NVGPU/IR/NVGPUDialect.h b/mlir/include/mlir/Dialect/NVGPU/IR/NVGPUDialect.h
index ccee0de65caa5..41f134fd12925 100644
--- a/mlir/include/mlir/Dialect/NVGPU/IR/NVGPUDialect.h
+++ b/mlir/include/mlir/Dialect/NVGPU/IR/NVGPUDialect.h
@@ -14,6 +14,7 @@
#define MLIR_DIALECT_NVGPU_NVGPUDIALECT_H_
#include "mlir/Bytecode/BytecodeOpInterface.h"
+#include "mlir/Dialect/LLVMIR/NVVMDialect.h"
#include "mlir/IR/BuiltinTypes.h"
#include "mlir/IR/Dialect.h"
#include "mlir/IR/OpDefinition.h"
diff --git a/mlir/include/mlir/Dialect/NVGPU/IR/NVGPUOps.td b/mlir/include/mlir/Dialect/NVGPU/IR/NVGPUOps.td
index 4c11725405ea5..6ebb9577d1337 100644
--- a/mlir/include/mlir/Dialect/NVGPU/IR/NVGPUOps.td
+++ b/mlir/include/mlir/Dialect/NVGPU/IR/NVGPUOps.td
@@ -20,6 +20,7 @@
#ifndef MLIR_DIALECT_NVGPU_IR_NVGPUOPS_TD
#define MLIR_DIALECT_NVGPU_IR_NVGPUOPS_TD
+include "mlir/Dialect/LLVMIR/NVVMEnums.td"
include "mlir/Dialect/NVGPU/IR/NVGPU.td"
include "mlir/Dialect/NVGPU/IR/NVGPUTypes.td"
@@ -660,12 +661,12 @@ def NVGPU_RcpOp : NVGPU_Op<"rcp", [Pure,
The input and output must be of the same vector type and shape.
}];
let arguments = (ins VectorOfNonZeroRankOf<[F32]>:$in,
- DefaultValuedAttr<RcpRoundingModeAttr, "RcpRoundingMode::APPROX">:$rounding,
+ DefaultValuedAttr<FPRoundingModeAttr, "::mlir::NVVM::FPRoundingMode::NONE">:$rounding,
+ UnitAttr:$approx,
UnitAttr:$ftz);
let results = (outs VectorOfNonZeroRankOf<[F32]>:$out);
let assemblyFormat = [{
- $in `{` `rounding` `=` $rounding (`,` `ftz` $ftz^)? `}`
- attr-dict `:` type($out)
+ $in attr-dict `:` type($out)
}];
let hasVerifier = 1;
}
diff --git a/mlir/lib/Dialect/NVGPU/IR/CMakeLists.txt b/mlir/lib/Dialect/NVGPU/IR/CMakeLists.txt
index 10aa502ee67f8..94776845c5321 100644
--- a/mlir/lib/Dialect/NVGPU/IR/CMakeLists.txt
+++ b/mlir/lib/Dialect/NVGPU/IR/CMakeLists.txt
@@ -12,6 +12,7 @@ add_mlir_dialect_library(MLIRNVGPUDialect
LINK_LIBS PUBLIC
MLIRGPUDialect
+ MLIRNVVMDialect
MLIRIR
MLIRSideEffectInterfaces
)
diff --git a/mlir/lib/Dialect/NVGPU/IR/NVGPUDialect.cpp b/mlir/lib/Dialect/NVGPU/IR/NVGPUDialect.cpp
index b4d8270177544..6e1ed05c9d0d0 100644
--- a/mlir/lib/Dialect/NVGPU/IR/NVG...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/195812
More information about the Mlir-commits
mailing list