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,51 @@ +//===-- 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. + + If serialization fails then the method should return `std::nullopt`. + + The `module` argument must be a GPU Module Op. 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 GPUTargetArrayAttr : TypedArrayAttrBase; + +def GPUNonEmptyTargetArrayAttr : + ConfinedAttr]>; + +#endif // GPU_COMPILATIONATTRINTERFACES diff --git a/mlir/include/mlir/Dialect/GPU/IR/CompilationInterfaces.h b/mlir/include/mlir/Dialect/GPU/IR/CompilationInterfaces.h new file mode 100644 --- /dev/null +++ b/mlir/include/mlir/Dialect/GPU/IR/CompilationInterfaces.h @@ -0,0 +1,90 @@ +//===-- CompilationInterfaces.h - GPU compilation interfaces ---*- 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 file defines interfaces for GPU compilation. +// +//===----------------------------------------------------------------------===// + +#ifndef MLIR_DIALECT_GPU_IR_COMPILATIONINTERFACES_H +#define MLIR_DIALECT_GPU_IR_COMPILATIONINTERFACES_H + +#include "mlir/IR/Attributes.h" + +namespace mlir { +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. + typedef enum { + offload, /// The process should produce an offloading representation. For + /// the NVVM & ROCDL targets this option produces LLVM IR. + assembly, /// The process should produce assembly code. + binary /// The process should produce a binary. + } CompilationTarget; + + /// Constructor initializing the toolkit path, the list of files to link to, + /// extra command line options & the compilation target. The default + /// compilation target is `binary`. + TargetOptions(StringRef toolkitPath = {}, + ArrayRef linkFiles = {}, StringRef cmdOptions = {}, + CompilationTarget compilationTarget = binary); + + /// Returns the typeID. + TypeID getTypeID() const; + + /// Returns the toolkit path. + StringRef getToolkitPath() const; + + /// Returns the files to link to. + ArrayRef getLinkFiles() const; + + /// Returns the command line options. + StringRef getCmdOptions() const; + + /// Returns a tokenization of the command line options. + std::pair> + tokenizeCmdOptions() const; + + /// Returns the compilation target. + CompilationTarget getCompilationTarget() 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 linkFiles = {}, StringRef cmdOptions = {}, + CompilationTarget compilationTarget = binary); + + /// Path to the target toolkit. + std::string toolkitPath; + + /// List of files to link with the LLVM module. + SmallVector linkFiles; + + /// An optional set of command line options to be used by the compilation + /// process. + std::string cmdOptions; + + /// Compilation process target representation. + CompilationTarget compilationTarget; + +private: + TypeID typeID; +}; +} // namespace gpu +} // namespace mlir + +MLIR_DECLARE_EXPLICIT_TYPE_ID(::mlir::gpu::TargetOptions) + +#include "mlir/Dialect/GPU/IR/CompilationAttrInterfaces.h.inc" + +#endif // MLIR_DIALECT_GPU_IR_COMPILATIONINTERFACES_H 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 @@ -16,6 +16,7 @@ #include "mlir/Bytecode/BytecodeOpInterface.h" #include "mlir/Dialect/DLTI/Traits.h" +#include "mlir/Dialect/GPU/IR/CompilationInterfaces.h" #include "mlir/IR/Builders.h" #include "mlir/IR/BuiltinTypes.h" #include "mlir/IR/Dialect.h" 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 @@ -28,7 +28,9 @@ #include "mlir/Interfaces/SideEffectInterfaces.h" #include "mlir/Transforms/InliningUtils.h" #include "llvm/ADT/TypeSwitch.h" +#include "llvm/Support/CommandLine.h" #include "llvm/Support/ErrorHandling.h" +#include "llvm/Support/StringSaver.h" using namespace mlir; using namespace mlir::gpu; @@ -1811,6 +1813,53 @@ results.add(context); } +//===----------------------------------------------------------------------===// +// GPU target options +//===----------------------------------------------------------------------===// + +TargetOptions::TargetOptions(StringRef toolkitPath, + ArrayRef linkFiles, + StringRef cmdOptions, + CompilationTarget compilationTarget) + : TargetOptions(TypeID::get(), toolkitPath, linkFiles, + cmdOptions, compilationTarget) {} + +TargetOptions::TargetOptions(TypeID typeID, StringRef toolkitPath, + ArrayRef linkFiles, + StringRef cmdOptions, + CompilationTarget compilationTarget) + : toolkitPath(toolkitPath.str()), linkFiles(linkFiles), + cmdOptions(cmdOptions.str()), compilationTarget(compilationTarget), + typeID(typeID) {} + +TypeID TargetOptions::getTypeID() const { return typeID; } + +StringRef TargetOptions::getToolkitPath() const { return toolkitPath; } + +ArrayRef TargetOptions::getLinkFiles() const { return linkFiles; } + +StringRef TargetOptions::getCmdOptions() const { return cmdOptions; } + +std::pair> +TargetOptions::tokenizeCmdOptions() const { + std::pair> options; + llvm::StringSaver stringSaver(options.first); +#ifdef _WIN32 + llvm::cl::TokenizeWindowsCommandLine(cmdOptions, stringSaver, options.second, + /*MarkEOLs=*/false); +#else + llvm::cl::TokenizeGNUCommandLine(cmdOptions, stringSaver, options.second, + /*MarkEOLs=*/false); +#endif // _WIN32 + return options; +} + +TargetOptions::CompilationTarget TargetOptions::getCompilationTarget() const { + return compilationTarget; +} + +MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::gpu::TargetOptions) + #include "mlir/Dialect/GPU/IR/GPUOpInterfaces.cpp.inc" #include "mlir/Dialect/GPU/IR/GPUOpsEnums.cpp.inc" @@ -1819,3 +1868,5 @@ #define GET_OP_CLASSES #include "mlir/Dialect/GPU/IR/GPUOps.cpp.inc" + +#include "mlir/Dialect/GPU/IR/CompilationAttrInterfaces.cpp.inc"