[mlir][gpu] Add the GPU offloading handler attribute interface.

[fmorac]

Summary

For an explanation of these patches see D154153.

Commit message:
This patch adds the OffloadingTranslationAttrTrait trait and the
GPUOffloadingLLVMTranslationAttrInterface attribute interface.

The interface indicates that an attribute is a GPU offloading translation
attribute. These kinds of attributes must implement an interface for handling
the translation of GPU offloading operations like gpu.binary & gpu.launch_func.

The interface is meant to be used as an opaque interface for embedding GPU
binaries into LLVM modules and launching GPU kernels.

Depends on D154129

[fmorac]

Rebasing.

[krzysz00]

I'm going to assume this interface is sufficient.

[fmorac]

Rebasing.

[mehdi_amini]

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

71

I think more context would be appreciated here about the role of the "manager" and the context where it is used (translation)

I'm not convinced that "ObjectManager" is expressive enough in terms of what this provides.

The name should likely include "Translation" somehow, GPUKernelOffloadingLLVMTranslationAttrInterface would the first thing I can come up with right now, thoughts?

87

The description say that this method launches a kernel: shall we say instead "Translate a gpu launch func operation to a LLVM IR target specific instructions sequence".

94

Shall we take a GPULaunchFuncOp instead of Operation* here? Same for the binaryOp being a GPUBinaryOp?

[fmorac]

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

71

I'll add it.

I settled in ObjectManager because it can be used to implement many concepts for managing GPU objects, like fat binaries & fat bin launch. However I agree the name needs to reflect it's for LLVM translation, what about GPUBinaryHandlerLLVMTranslationAttrInterface

87

I'll change it.

94

The reason these are like this is I couldn't break C++ dependencies and have nvptx_target & amdgpu_target in gpu, as

nvptx_target depends on TargetAttrInterface that depends on the definitions of those operations that depends on attributes (fwd declarations didn't work).

If we move, gpu.nvptx_target to nvvm.target, then there are no issues and the above methods can use GPULaunchFuncOp & GPUBinaryOp.

[mehdi_amini]

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

71

Works for me.

94

Is this the same discussion we're having in the parent about the location of the NVPTXTarget attribute?

[fmorac]

Switched name from GPUObjectManagerAttrInterface to GPUBinaryHandlerLLVMTranslationAttrInterface.

[fmorac]

Rebasing.

[fmorac]

Remove remnants of the ObjectManager name.

[mehdi_amini]

LG!

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

87

I don't think this one was done? It still says Launches a kernel

[fmorac]

Improved docs & removed unnecessary tablegen code. Also switched the name from Binary Handler to Offloading Handler.

[mehdi_amini]

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

68

Can we keep the TableGen name identical to the C++ name here? It got me confused on subsequence patches to read one and the other.

[fmorac]

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

68

Yes, I'll change I just used the GPU prefix because I saw it was customary with other Ops & Attributes.

[fmorac]

Add the OffloadingTranslationAttrTrait trait. This traits decouples gpu.binary op
from LLVM translation, making gpu.binary usable by other translation mechanisms like SPIRV.

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

 //===-- CompilationAttrInterfaces.td - GPU compilation interfaces ---------===//
 //
 // 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 file defines interfaces for GPU target attributes.
+// This file defines interfaces for GPU target attributes & GPU object manager
+// attributes.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef GPU_COMPILATIONATTRINTERFACES
 #define GPU_COMPILATIONATTRINTERFACES
 
 include "mlir/IR/AttrTypeBase.td"
 include "mlir/IR/OpBase.td"

 }
 
 def GPUTargetArrayAttr :
   TypedArrayAttrBase<GPUTargetAttr, "Array of GPU target attributes">;
 
 def GPUNonEmptyTargetArrayAttr :
   ConfinedAttr<GPUTargetArrayAttr, [ArrayMinCount<1>]>;
 
+//===----------------------------------------------------------------------===//
+// GPU object manager attribute interface.
+//===----------------------------------------------------------------------===//
+
+def GPUBinaryHandlerLLVMTranslationAttrInterface :
+    AttrInterface<"BinaryHandlerLLVMTranslationAttrInterface"> {
+  let description = [{
+    Interface for GPU object manager attributes. Attributes implementing this
+    interface manage the interaction between GPU objects and host IR.
+  }];
+  let cppNamespace = "::mlir::gpu";
+  let methods = [
+    InterfaceMethod<[{
+        Embeds a GPU object into a host LLVM module. The operation expected by
+        this method must be a GPU BinaryOp.
+
+        All attributes implementing this interface must implement this method.
+        If the method fails then it must return `failure`.
+      }],
+      "LogicalResult", "embedBinary",
+      (ins "Operation*":$binaryOp, "llvm::IRBuilderBase&":$hostBuilder,
+           "LLVM::ModuleTranslation&":$hostModuleTranslation)
+    >,
+    InterfaceMethod<[{
+        Launches a kernel inside a binary. The first argument must be a GPU
+        LaunchFuncOp, while the second one a GPU BinaryOp.
+
+        All attributes implementing this interface must implement this method.
+        If the method fails then it must return `failure`.
+      }],
+      "LogicalResult", "launchKernel",
+      (ins "Operation*":$launchFunc, "Operation*":$binaryOp,
+           "llvm::IRBuilderBase&":$hostBuilder,
+           "LLVM::ModuleTranslation&":$hostModuleTranslation)
+    >
+  ];
+}
+
+def ImplementsBinaryHandlerLLVMTranslationAttrInterface : AttrConstraint<
+  CPred<"isa<::mlir::gpu::BinaryHandlerLLVMTranslationAttrInterface>($_self)">,
+  "Attribute implementing the `BinaryHandlerLLVMTranslationAttrInterface` interface."
+>;
+
+def GPUBinaryHandlerLLVMTranslationAttr :
+    ConfinedAttr<AnyAttr, [ImplementsBinaryHandlerLLVMTranslationAttrInterface]> {
+  let description = [{
+    Generic compilation attribute implementing the `BinaryHandlerLLVMTranslationAttrInterface`
+    interface.
+  }];
+}
+
 #endif // GPU_COMPILATIONATTRINTERFACES

mlir/include/mlir/Dialect/GPU/IR/CompilationInterfaces.h

 //
 //===----------------------------------------------------------------------===//
 
 #ifndef MLIR_DIALECT_GPU_IR_COMPILATIONINTERFACES_H
 #define MLIR_DIALECT_GPU_IR_COMPILATIONINTERFACES_H
 
 #include "mlir/IR/Attributes.h"
 
+namespace llvm {
+class IRBuilderBase;
+}
+
 namespace mlir {
+namespace LLVM {
+class ModuleTranslation;
+}
 namespace gpu {
 /// This class serves as an opaque interface for passing options to the
 /// `TargetAttrInterface` methods. Users of this class must implement the
 /// `classof` method as well as using the macros `MLIR_*_EXPLICIT_TYPE_ID` to
 /// ensure type safeness. Targets are free to ignore these options.
 class TargetOptions {
 public:
   /// The target representation of the compilation process.