[mlir][interfaces] Add the `TargetInfo` attribute interface #78073

fabianmcg · 2024-01-13T23:52:13Z

This patch adds the TargetInfo attribute interface to the set of DLTI interfaces. Target information attributes provide essential information on the compilation target. This information includes the target triple identifier, the target chip identifier, and a string representation of the target features.

This patch also adds this new interface to the NVVM GPU target attributes.

llvmbot · 2024-01-14T00:01:35Z

@llvm/pr-subscribers-mlir-gpu
@llvm/pr-subscribers-mlir-dlti
@llvm/pr-subscribers-mlir-llvm

@llvm/pr-subscribers-mlir

Author: Fabian Mora (fabianmcg)

Changes

This patch adds the TargetInfo attribute interface to the set of DLTI interfaces. Target information attributes provide essential information on the compilation target. This information includes the target triple identifier, the target chip identifier, and a string representation of the target features.

This patch also adds this new interface to the NVVM and ROCDL GPU target attributes.

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

8 Files Affected:

(modified) mlir/include/mlir/Dialect/LLVMIR/NVVMDialect.h (+1)
(modified) mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td (+4-1)
(modified) mlir/include/mlir/Dialect/LLVMIR/ROCDLDialect.h (+1)
(modified) mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td (+4-2)
(modified) mlir/include/mlir/Interfaces/DataLayoutInterfaces.td (+33)
(modified) mlir/lib/Dialect/LLVMIR/CMakeLists.txt (+2)
(modified) mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp (+8)
(modified) mlir/lib/Dialect/LLVMIR/IR/ROCDLDialect.cpp (+8)

diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMDialect.h b/mlir/include/mlir/Dialect/LLVMIR/NVVMDialect.h
index 08019e77ae6af8..1a55d08be9edc2 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMDialect.h
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMDialect.h
@@ -19,6 +19,7 @@
 #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
 #include "mlir/IR/Dialect.h"
 #include "mlir/IR/OpDefinition.h"
+#include "mlir/Interfaces/DataLayoutInterfaces.h"
 #include "mlir/Interfaces/SideEffectInterfaces.h"
 #include "llvm/IR/IntrinsicsNVPTX.h"
 
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index c5f68a2ebe3952..0bbbde6270cd69 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -17,6 +17,7 @@ include "mlir/IR/EnumAttr.td"
 include "mlir/Dialect/GPU/IR/CompilationAttrInterfaces.td"
 include "mlir/Dialect/LLVMIR/LLVMOpBase.td"
 include "mlir/Interfaces/SideEffectInterfaces.td"
+include "mlir/Interfaces/DataLayoutInterfaces.td"
 include "mlir/Dialect/LLVMIR/BasicPtxBuilderInterface.td"
 
 def LLVM_PointerGlobal : LLVM_PointerInAddressSpace<1>;
@@ -1894,7 +1895,9 @@ def NVVM_WgmmaMmaAsyncOp : NVVM_Op<"wgmma.mma_async",
 // NVVM target attribute.
 //===----------------------------------------------------------------------===//
 
-def NVVM_TargettAttr : NVVM_Attr<"NVVMTarget", "target"> {
+def NVVM_TargettAttr : NVVM_Attr<"NVVMTarget", "target", [
+    DeclareAttrInterfaceMethods<TargetInfoAttrInterface>
+  ]> {
   let description = [{
     GPU target attribute for controlling compilation of NVIDIA targets. All
     parameters decay into default values if not present.
diff --git a/mlir/include/mlir/Dialect/LLVMIR/ROCDLDialect.h b/mlir/include/mlir/Dialect/LLVMIR/ROCDLDialect.h
index c2a82ffc1c43cf..fa1131a463e1ab 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/ROCDLDialect.h
+++ b/mlir/include/mlir/Dialect/LLVMIR/ROCDLDialect.h
@@ -26,6 +26,7 @@
 #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
 #include "mlir/IR/Dialect.h"
 #include "mlir/IR/OpDefinition.h"
+#include "mlir/Interfaces/DataLayoutInterfaces.h"
 #include "mlir/Interfaces/SideEffectInterfaces.h"
 
 ///// Ops /////
diff --git a/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td b/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td
index 48b830ae34f292..a492709c299544 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td
@@ -15,6 +15,7 @@
 
 include "mlir/Dialect/GPU/IR/CompilationAttrInterfaces.td"
 include "mlir/Dialect/LLVMIR/LLVMOpBase.td"
+include "mlir/Interfaces/DataLayoutInterfaces.td"
 include "mlir/Interfaces/SideEffectInterfaces.td"
 
 //===----------------------------------------------------------------------===//
@@ -608,8 +609,9 @@ def ROCDL_CvtSrFp8F32Op :
 // ROCDL target attribute.
 //===----------------------------------------------------------------------===//
 
-def ROCDL_TargettAttr :
-    ROCDL_Attr<"ROCDLTarget", "target"> {
+def ROCDL_TargettAttr : ROCDL_Attr<"ROCDLTarget", "target", [
+    DeclareAttrInterfaceMethods<TargetInfoAttrInterface>
+  ]> {
   let description = [{
     ROCDL target attribute for controlling compilation of AMDGPU targets. All
     parameters decay into default values if not present.
diff --git a/mlir/include/mlir/Interfaces/DataLayoutInterfaces.td b/mlir/include/mlir/Interfaces/DataLayoutInterfaces.td
index a8def967fffcfa..eac9521aadc11e 100644
--- a/mlir/include/mlir/Interfaces/DataLayoutInterfaces.td
+++ b/mlir/include/mlir/Interfaces/DataLayoutInterfaces.td
@@ -188,6 +188,39 @@ def DataLayoutSpecInterface : AttrInterface<"DataLayoutSpecInterface"> {
   }];
 }
 
+def TargetInfoAttrInterface : AttrInterface<"TargetInfoAttrInterface"> {
+  let cppNamespace = "::mlir";
+
+  let description = [{
+    Attribute interface describing target information.
+
+    Target information attributes provide essential information on the
+    compilation target. This information includes the target triple identifier,
+    the target chip identifier, and a string representation of the target features.
+  }];
+
+  let methods = [
+    InterfaceMethod<
+      /*description=*/"Returns the target triple identifier.",
+      /*retTy=*/"::mlir::StringRef",
+      /*methodName=*/"getTargetTriple",
+      /*args=*/(ins)
+    >,
+    InterfaceMethod<
+      /*description=*/"Returns the target chip identifier.",
+      /*retTy=*/"::mlir::StringRef",
+      /*methodName=*/"getTargetChip",
+      /*args=*/(ins)
+    >,
+    InterfaceMethod<
+      /*description=*/"Returns the target features as a string.",
+      /*retTy=*/"std::string",
+      /*methodName=*/"getTargetFeatures",
+      /*args=*/(ins)
+    >
+  ];
+}
+
 //===----------------------------------------------------------------------===//
 // Operation interface
 //===----------------------------------------------------------------------===//
diff --git a/mlir/lib/Dialect/LLVMIR/CMakeLists.txt b/mlir/lib/Dialect/LLVMIR/CMakeLists.txt
index b00259677697a5..00b78e30ee8b09 100644
--- a/mlir/lib/Dialect/LLVMIR/CMakeLists.txt
+++ b/mlir/lib/Dialect/LLVMIR/CMakeLists.txt
@@ -61,6 +61,7 @@ add_mlir_dialect_library(MLIRNVVMDialect
   LINK_LIBS PUBLIC
   MLIRIR
   MLIRLLVMDialect
+  MLIRDataLayoutInterfaces
   MLIRSideEffectInterfaces
   )
 
@@ -83,5 +84,6 @@ add_mlir_dialect_library(MLIRROCDLDialect
   LINK_LIBS PUBLIC
   MLIRIR
   MLIRLLVMDialect
+  MLIRDataLayoutInterfaces
   MLIRSideEffectInterfaces
   )
diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
index aa49c4dc31fbc0..b73504ac4969af 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -1106,6 +1106,14 @@ NVVMTargetAttr::verify(function_ref<InFlightDiagnostic()> emitError,
   return success();
 }
 
+StringRef NVVMTargetAttr::getTargetTriple() const { return getTriple(); }
+
+StringRef NVVMTargetAttr::getTargetChip() const { return getChip(); }
+
+std::string NVVMTargetAttr::getTargetFeatures() const {
+  return getFeatures().str();
+}
+
 #define GET_OP_CLASSES
 #include "mlir/Dialect/LLVMIR/NVVMOps.cpp.inc"
 
diff --git a/mlir/lib/Dialect/LLVMIR/IR/ROCDLDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/ROCDLDialect.cpp
index 26e46b31ddc018..8b10c48718a3f8 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/ROCDLDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/ROCDLDialect.cpp
@@ -295,6 +295,14 @@ ROCDLTargetAttr::verify(function_ref<InFlightDiagnostic()> emitError,
   return success();
 }
 
+StringRef ROCDLTargetAttr::getTargetTriple() const { return getTriple(); }
+
+StringRef ROCDLTargetAttr::getTargetChip() const { return getChip(); }
+
+std::string ROCDLTargetAttr::getTargetFeatures() const {
+  return getFeatures().str();
+}
+
 #define GET_OP_CLASSES
 #include "mlir/Dialect/LLVMIR/ROCDLOps.cpp.inc"

This patch adds the offloading translation attribute. This attribute uses LLVM offloading infrastructure to embed GPU binaries in the IR. At the program start, the LLVM offloading mechanism registers kernels and variables with the runtime library: CUDA RT, HIP RT, or LibOMPTarget. The offloading mechanism relies on the runtime library to dispatch the correct kernel based on the registered symbols. This patch is 3/4 on introducing the OffloadEmbeddingAttr GPU translation attribute. Note: Ignore the base commits; those are being reviewed in PRs llvm#78057, llvm#78098, and llvm#78073.

This patch adds the TargetInfo attribute interface to the set of DLTI interfaces. Target information attributes provide essential information on the compilation target. This information includes the target triple identifier, the target chip identifier, and a string representation of the target features. This patch also adds this new interface to the NVVM and ROCDL GPU target attributes.

joker-eph · 2025-04-20T12:36:04Z

mlir/include/mlir/Dialect/DLTI/Transforms/Passes.h

+/// Sets the target specs using the target attached to the module.
+LogicalResult setTargetSpecsFromTarget(Operation *op);


Suggested change

/// Sets the target specs using the target attached to the module.

LogicalResult setTargetSpecsFromTarget(Operation *op);

/// Sets the target specs using the target attached to the operation.

LogicalResult setTargetSpecsFromTarget(DataLayoutOpInterface op);

joker-eph · 2025-04-20T12:40:03Z

mlir/include/mlir/Interfaces/DataLayoutInterfaces.td

+    InterfaceMethod<
+      /*description=*/"Returns the target triple identifier.",
+      /*retTy=*/"::llvm::StringRef",
+      /*methodName=*/"getTargetTriple",
+      /*args=*/(ins)
+    >,
+    InterfaceMethod<
+      /*description=*/"Returns the target chip identifier.",
+      /*retTy=*/"::llvm::StringRef",
+      /*methodName=*/"getTargetChip",
+      /*args=*/(ins)
+    >,
+    InterfaceMethod<
+      /*description=*/"Returns the target features as a string.",
+      /*retTy=*/"std::string",
+      /*methodName=*/"getTargetFeatures",
+      /*args=*/(ins)
+    >,


Is this entirely "LLVM centric"?

joker-eph · 2025-04-20T12:42:12Z

mlir/lib/Target/LLVM/NVVM/Target.cpp

 } // namespace

 // Register the NVVM dialect, the NVVM translation & the target interface.
 void mlir::NVVM::registerNVVMTargetInterfaceExternalModels(
    DialectRegistry &registry) {
  registry.addExtension(+[](MLIRContext *ctx, NVVM::NVVMDialect *dialect) {
    NVVMTargetAttr::attachInterface<NVVMTargetAttrImpl>(*ctx);
+    NVVMTargetAttr::attachInterface<NVVMTargetInfoAttrImpl>(*ctx);


There needs a promise on the attribute I believe.

joker-eph · 2025-04-20T13:01:56Z

mlir/lib/Target/LLVM/NVVM/Target.cpp

+  if (failed(info))
+    return failure();
+  spec.dataLayout =
+      translateDataLayout(info->getDataLayout(), attribute.getContext());


Despite the method name, this does only set the data layout?

joker-eph · 2025-04-20T13:02:52Z

mlir/include/mlir/Interfaces/DataLayoutInterfaces.td

+      /*retTy=*/"::llvm::LogicalResult",
+      /*methodName=*/"setTargetSpec",
+      /*args=*/(ins "::mlir::TargetSpec&":$spec)


Suggested change

/*retTy=*/"::llvm::LogicalResult",

/*methodName=*/"setTargetSpec",

/*args=*/(ins "::mlir::TargetSpec&":$spec)

/*retTy=*/"::mlir::FailureOr<::mlir::TargetSpec>",

/*methodName=*/"getTargetSpec",

/*args=*/(ins)

fabianmcg requested review from ftynse, zero9178 and qcolombet January 14, 2024 00:00

fabianmcg marked this pull request as ready for review January 14, 2024 00:01

llvmbot added mlir:llvm mlir labels Jan 14, 2024

fabianmcg mentioned this pull request Jan 15, 2024

[mlir][gpu] Add the OffloadEmbeddingAttr offloading translation attr #78117

Draft

fabianmcg added a commit to fabianmcg/llvm-project that referenced this pull request Jan 16, 2024

Base commit, PR llvm#78073

24aa668

fabianmcg and others added 2 commits April 20, 2025 10:28

update pr

ae4b14b

fabianmcg force-pushed the target-info branch from 61c8809 to ae4b14b Compare April 20, 2025 10:43

fabianmcg requested a review from grypp as a code owner April 20, 2025 10:43

llvmbot added mlir:gpu mlir:dlti labels Apr 20, 2025

fabianmcg marked this pull request as draft April 20, 2025 10:44

joker-eph reviewed Apr 20, 2025

View reviewed changes

rengolin requested a review from rolfmorel April 20, 2025 15:58

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

[mlir][interfaces] Add the `TargetInfo` attribute interface #78073

[mlir][interfaces] Add the `TargetInfo` attribute interface #78073

Uh oh!

fabianmcg commented Jan 13, 2024 •

edited

Loading

Uh oh!

llvmbot commented Jan 14, 2024 •

edited

Loading

Uh oh!

joker-eph Apr 20, 2025

Uh oh!

joker-eph Apr 20, 2025

Uh oh!

joker-eph Apr 20, 2025

Uh oh!

joker-eph Apr 20, 2025

Uh oh!

joker-eph Apr 20, 2025

Uh oh!

Uh oh!

		/// Sets the target specs using the target attached to the module.
		LogicalResult setTargetSpecsFromTarget(Operation *op);

[mlir][interfaces] Add the TargetInfo attribute interface #78073

Are you sure you want to change the base?

[mlir][interfaces] Add the TargetInfo attribute interface #78073

Uh oh!

Conversation

fabianmcg commented Jan 13, 2024 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

llvmbot commented Jan 14, 2024 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

joker-eph Apr 20, 2025

Choose a reason for hiding this comment

Uh oh!

joker-eph Apr 20, 2025

Choose a reason for hiding this comment

Uh oh!

joker-eph Apr 20, 2025

Choose a reason for hiding this comment

Uh oh!

joker-eph Apr 20, 2025

Choose a reason for hiding this comment

Uh oh!

joker-eph Apr 20, 2025

Choose a reason for hiding this comment

Uh oh!

Uh oh!

[mlir][interfaces] Add the `TargetInfo` attribute interface #78073

[mlir][interfaces] Add the `TargetInfo` attribute interface #78073

fabianmcg commented Jan 13, 2024 •

edited

Loading

llvmbot commented Jan 14, 2024 •

edited

Loading