update pr

Author: fabianmcg

mlir/include/mlir/Dialect/DLTI/CMakeLists.txt

@@ -1,4 +1,5 @@
 add_subdirectory(TransformOps)
+add_subdirectory(Transforms)
 
 add_mlir_dialect(DLTI dlti)
 add_mlir_doc(DLTIAttrs DLTIDialect Dialects/ -gen-dialect-doc)

mlir/include/mlir/Dialect/DLTI/DLTIBase.td

@@ -39,6 +39,10 @@ def DLTI_Dialect : Dialect {
     constexpr const static ::llvm::StringLiteral
     kTargetDeviceDescAttrName = "dlti.target_device_spec";
 
+    // Top-level attribute name for target information.
+    constexpr const static ::llvm::StringLiteral
+    kTargetDescAttrName = "dlti.target";
+
     // Constants used in entries.
     constexpr const static ::llvm::StringLiteral
     kDataLayoutEndiannessKey = "dlti.endianness";

mlir/include/mlir/Dialect/DLTI/Traits.h

@@ -19,6 +19,10 @@ namespace impl {
 LogicalResult verifyHasDefaultDLTIDataLayoutTrait(Operation *op);
 DataLayoutSpecInterface getDataLayoutSpec(Operation *op);
 TargetSystemSpecInterface getTargetSystemSpec(Operation *op);
+TargetAttrInterface getTargetAttr(Operation *op);
+void setDataLayoutSpec(Operation *op, DataLayoutSpecInterface spec);
+void setTargetSystemSpec(Operation *op, TargetSystemSpecInterface spec);
+void setTargetAttr(Operation *op, TargetAttrInterface target);
 } // namespace impl
 
 /// Trait to be used by operations willing to use the implementation of the
@@ -39,11 +43,27 @@ class HasDefaultDLTIDataLayout
     return impl::getDataLayoutSpec(this->getOperation());
   }
 
+  /// Sets the data layout specification.
+  void setDataLayoutSpec(DataLayoutSpecInterface spec) {
+    impl::setDataLayoutSpec(this->getOperation(), spec);
+  }
   /// Returns the target system description specification as provided by DLTI
   /// dialect
   TargetSystemSpecInterface getTargetSystemSpec() {
     return impl::getTargetSystemSpec(this->getOperation());
   }
+  /// Sets the target system description specification.
+  void setTargetSystemSpec(TargetSystemSpecInterface spec) {
+    impl::setTargetSystemSpec(this->getOperation(), spec);
+  }
+  /// Returns the target information as provided by DLTI dialect.
+  TargetAttrInterface getTargetAttr() {
+    return impl::getTargetAttr(this->getOperation());
+  }
+  /// Sets the target information.
+  void setTargetAttr(TargetAttrInterface target) {
+    impl::setTargetAttr(this->getOperation(), target);
+  }
 };
 } // namespace mlir
 

mlir/include/mlir/Dialect/DLTI/Transforms/CMakeLists.txt

@@ -0,0 +1,7 @@
+set(LLVM_TARGET_DEFINITIONS Passes.td)
+mlir_tablegen(Passes.h.inc -gen-pass-decls -name DLTI)
+mlir_tablegen(Passes.capi.h.inc -gen-pass-capi-header --prefix DLTI)
+mlir_tablegen(Passes.capi.cpp.inc -gen-pass-capi-impl --prefix DLTI)
+add_public_tablegen_target(MLIRDLTIPassIncGen)
+
+add_mlir_doc(Passes DLTIPasses ./ -gen-pass-doc)

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

@@ -0,0 +1,30 @@
+//===- Passes.h - Pass Entrypoints ------------------------------*- 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
+//
+//===----------------------------------------------------------------------===//
+//
+// This header file defines prototypes that expose pass constructors.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef MLIR_DIALECT_DLTI_TRANSFORMS_PASSES_H
+#define MLIR_DIALECT_DLTI_TRANSFORMS_PASSES_H
+
+#include "mlir/Pass/Pass.h"
+
+namespace mlir {
+#define GEN_PASS_DECL
+#include "mlir/Dialect/DLTI/Transforms/Passes.h.inc"
+
+/// Generate the code for registering passes.
+#define GEN_PASS_REGISTRATION
+#include "mlir/Dialect/DLTI/Transforms/Passes.h.inc"
+
+/// Sets the target specs using the target attached to the module.
+LogicalResult setTargetSpecsFromTarget(Operation *op);
+} // namespace mlir
+
+#endif // MLIR_DIALECT_DLTI_TRANSFORMS_PASSES_H

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

@@ -0,0 +1,40 @@
+//===-- Passes.td - DLTI 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_DLTI_PASSES
+#define MLIR_DIALECT_DLTI_PASSES
+
+include "mlir/Pass/PassBase.td"
+
+def DltiSetTargetSpecsFromTarget: Pass<"dlti-set-target-specs", ""> {
+  let summary = "Sets DLTI target specs using a target.";
+  let description = [{
+    This pass potentially sets the following DLTI target specs in the current
+    operation:
+    - The data layout.
+    - The target system spec.
+
+    Example:
+
+    ```mlir
+    // Given the following input:
+    builtin.module @module_1 attributes {dlti.target = #my.target} {...}
+    // After applying the pass:
+    builtin.module @module_1 attributes {
+      dlti.target = #my.target,
+      dlti.target_system_spec = #my.system_spec,
+      dlti.dl_spec = #my.dl_spec
+    } {...}
+    ```
+  }];
+  let dependentDialects = [
+    "::mlir::DLTIDialect"
+  ];
+}
+
+#endif // MLIR_DIALECT_DLTI_PASSES

mlir/include/mlir/Dialect/GPU/IR/GPUOps.td

@@ -1388,7 +1388,8 @@ def GPU_BarrierOp : GPU_Op<"barrier"> {
 }
 
 def GPU_GPUModuleOp : GPU_Op<"module", [
-      IsolatedFromAbove, DataLayoutOpInterface, HasDefaultDLTIDataLayout,
+      IsolatedFromAbove,
+      DeclareOpInterfaceMethods<DataLayoutOpInterface, ["getTargetAttr", "setTargetAttr"]>,
       NoRegionArguments, SymbolTable, Symbol] # GraphRegionNoTerminator.traits> {
   let summary = "A top level compilation unit containing code to be run on a GPU.";
   let description = [{

mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td

@@ -3247,9 +3247,7 @@ def NVVM_Tcgen05StOp : NVVM_Op<"tcgen05.st"> {
 // NVVM target attribute.
 //===----------------------------------------------------------------------===//
 
-def NVVM_TargettAttr : NVVM_Attr<"NVVMTarget", "target", [
-    DeclareAttrInterfaceMethods<TargetInfoAttrInterface>
-  ]> {
+def NVVM_TargettAttr : NVVM_Attr<"NVVMTarget", "target"> {
   let description = [{
     GPU target attribute for controlling compilation of NVIDIA targets. All
     parameters decay into default values if not present.

mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td

@@ -1117,9 +1117,7 @@ def ROCDL_CvtSrFp8F32Op :
 // ROCDL target attribute.
 //===----------------------------------------------------------------------===//
 
-def ROCDL_TargettAttr : ROCDL_Attr<"ROCDLTarget", "target", [
-    DeclareAttrInterfaceMethods<TargetInfoAttrInterface>
-  ]> {
+def ROCDL_TargettAttr : ROCDL_Attr<"ROCDLTarget", "target"> {
   let description = [{
     ROCDL target attribute for controlling compilation of AMDGPU targets. All
     parameters decay into default values if not present.

mlir/include/mlir/InitAllPasses.h

@@ -23,6 +23,7 @@
 #include "mlir/Dialect/Async/Passes.h"
 #include "mlir/Dialect/Bufferization/Pipelines/Passes.h"
 #include "mlir/Dialect/Bufferization/Transforms/Passes.h"
+#include "mlir/Dialect/DLTI/Transforms/Passes.h"
 #include "mlir/Dialect/EmitC/Transforms/Passes.h"
 #include "mlir/Dialect/Func/Transforms/Passes.h"
 #include "mlir/Dialect/GPU/Pipelines/Passes.h"
@@ -75,6 +76,7 @@ inline void registerAllPasses() {
   bufferization::registerBufferizationPasses();
   func::registerFuncPasses();
   registerGPUPasses();
+  registerDLTIPasses();
   registerLinalgPasses();
   registerNVGPUPasses();
   registerSparseTensorPasses();