diff --git a/mlir/include/mlir/Dialect/GPU/IR/CMakeLists.txt b/mlir/include/mlir/Dialect/GPU/IR/CMakeLists.txt --- a/mlir/include/mlir/Dialect/GPU/IR/CMakeLists.txt +++ b/mlir/include/mlir/Dialect/GPU/IR/CMakeLists.txt @@ -16,6 +16,11 @@ mlir_tablegen(GPUOpsEnums.cpp.inc -gen-enum-defs) add_public_tablegen_target(MLIRGPUOpsEnumsGen) +set(LLVM_TARGET_DEFINITIONS CompilationAttrInterfaces.td) +mlir_tablegen(CompilationAttrInterfaces.h.inc -gen-attr-interface-decls) +mlir_tablegen(CompilationAttrInterfaces.cpp.inc -gen-attr-interface-defs) +add_public_tablegen_target(MLIRGPUCompilationAttrInterfacesIncGen) + set(LLVM_TARGET_DEFINITIONS GPUOps.td) mlir_tablegen(GPUOpsAttributes.h.inc -gen-attrdef-decls -attrdefs-dialect=gpu) mlir_tablegen(GPUOpsAttributes.cpp.inc -gen-attrdef-defs -attrdefs-dialect=gpu) diff --git a/mlir/include/mlir/Dialect/GPU/IR/CompilationAttrInterfaces.td b/mlir/include/mlir/Dialect/GPU/IR/CompilationAttrInterfaces.td new file mode 100644 --- /dev/null +++ b/mlir/include/mlir/Dialect/GPU/IR/CompilationAttrInterfaces.td @@ -0,0 +1,64 @@ +//===-- 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. +// +//===----------------------------------------------------------------------===// + +#ifndef GPU_COMPILATIONATTRINTERFACES +#define GPU_COMPILATIONATTRINTERFACES + +include "mlir/IR/AttrTypeBase.td" +include "mlir/IR/OpBase.td" + +//===----------------------------------------------------------------------===// +// GPU target attribute interface. +//===----------------------------------------------------------------------===// + +def GPUTargetAttrInterface : AttrInterface<"TargetAttrInterface"> { + let description = [{ + Interface for GPU target attributes. Attributes implementing this interface + compile GPU modules into binary objects, providing an opaque interface to + hide implementation details. + }]; + let cppNamespace = "::mlir::gpu"; + let methods = [ + InterfaceMethod<[{ + Serializes a GPU module to a string containing a representation of the + module. + + All attributes implementing this interface must implement this method. + If serialization fails then the method should return `std::nullopt`. + + The `options` argument is meant to be used for passing additional + options that are not in the attribute. + }], + "std::optional>", "serializeToObject", + (ins "Operation*":$module, "const gpu::TargetOptions&":$options) + > + ]; +} + +def ImplementsTargetAttrInterface : AttrConstraint< + CPred<"$_self.hasTrait<::mlir::gpu::TargetAttrInterface::Trait>()">, + "Attribute implementing the `TargetAttrInterface` interface." +>; + +def GPUTargetAttr : ConfinedAttr { + let description = [{ + Generic target attribute implementing the `TargetAttrInterface` interface. + }]; +} + +def GPUTargetArrayAttr : + TypedArrayAttrBase; + +def GPUNonEmptyTargetArrayAttr : + ConfinedAttr]>; + +#endif // GPU_COMPILATIONATTRINTERFACES diff --git a/mlir/include/mlir/Dialect/GPU/IR/GPUDialect.h b/mlir/include/mlir/Dialect/GPU/IR/GPUDialect.h --- a/mlir/include/mlir/Dialect/GPU/IR/GPUDialect.h +++ b/mlir/include/mlir/Dialect/GPU/IR/GPUDialect.h @@ -179,9 +179,45 @@ using SparseDnTensorHandleType = SparseHandleType; using SparseSpMatHandleType = SparseHandleType; +/// 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. +class TargetOptions { +public: + /// Constructor initializing the toolkit path and the list of bitcode files. + TargetOptions(StringRef toolkitPath = {}, + ArrayRef bitcodeFiles = {}); + + /// Returns the typeID. + TypeID getTypeID() const; + + /// Returns the toolkit path. + StringRef getToolkitPath() const; + + /// Returns the bitcode files to link to. + ArrayRef getBitcodeFiles() const; + +protected: + /// Derived classes must use this constructor to initialize `typeID` to the + /// appropiate value: ie. `TargetOptions(TypeID::get())`. + TargetOptions(TypeID typeID, StringRef toolkitPath = {}, + ArrayRef bitcodeFiles = {}); + + /// Path to the target toolkit. + std::string toolkitPath; + + /// List of files to link with the LLVM module. + SmallVector bitcodeFiles; + +private: + TypeID typeID; +}; } // namespace gpu } // namespace mlir +MLIR_DECLARE_EXPLICIT_TYPE_ID(::mlir::gpu::TargetOptions) + #include "mlir/Dialect/GPU/IR/GPUOpsEnums.h.inc" #include "mlir/Dialect/GPU/IR/GPUOpsDialect.h.inc" @@ -190,6 +226,8 @@ #include "mlir/Dialect/SCF/IR/DeviceMappingInterface.h" +#include "mlir/Dialect/GPU/IR/CompilationAttrInterfaces.h.inc" + #define GET_ATTRDEF_CLASSES #include "mlir/Dialect/GPU/IR/GPUOpsAttributes.h.inc" diff --git a/mlir/lib/Dialect/GPU/CMakeLists.txt b/mlir/lib/Dialect/GPU/CMakeLists.txt --- a/mlir/lib/Dialect/GPU/CMakeLists.txt +++ b/mlir/lib/Dialect/GPU/CMakeLists.txt @@ -32,6 +32,7 @@ MLIRGPUOpsAttributesIncGen MLIRGPUOpsEnumsGen MLIRGPUOpInterfacesIncGen + MLIRGPUCompilationAttrInterfacesIncGen LINK_LIBS PUBLIC MLIRArithDialect diff --git a/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp b/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp --- a/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp +++ b/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp @@ -1754,6 +1754,29 @@ results.add(context); } +//===----------------------------------------------------------------------===// +// GPU Compilation options +//===----------------------------------------------------------------------===// + +TargetOptions::TargetOptions(StringRef toolkitPath, + ArrayRef bitcodeFiles) + : TargetOptions(TypeID::get(), toolkitPath, bitcodeFiles) {} + +TargetOptions::TargetOptions(TypeID typeID, StringRef toolkitPath, + ArrayRef bitcodeFiles) + : toolkitPath(toolkitPath.str()), bitcodeFiles(bitcodeFiles), + typeID(typeID) {} + +TypeID TargetOptions::getTypeID() const { return typeID; } + +StringRef TargetOptions::getToolkitPath() const { return toolkitPath; } + +ArrayRef TargetOptions::getBitcodeFiles() const { + return bitcodeFiles; +} + +MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::gpu::TargetOptions) + #include "mlir/Dialect/GPU/IR/GPUOpInterfaces.cpp.inc" #include "mlir/Dialect/GPU/IR/GPUOpsEnums.cpp.inc" @@ -1762,3 +1785,5 @@ #define GET_OP_CLASSES #include "mlir/Dialect/GPU/IR/GPUOps.cpp.inc" + +#include "mlir/Dialect/GPU/IR/CompilationAttrInterfaces.cpp.inc"