[Mlir-commits] [mlir] 86c4dfa - [mlir][gpu] Add GPU target attribute interface.
Fabian Mora
llvmlistbot at llvm.org
Tue Aug 8 06:11:02 PDT 2023
Author: Fabian Mora
Date: 2023-08-08T13:10:54Z
New Revision: 86c4dfa209b5afe7a5bf8e50efea191446f9e2ae
URL: https://github.com/llvm/llvm-project/commit/86c4dfa209b5afe7a5bf8e50efea191446f9e2ae
DIFF: https://github.com/llvm/llvm-project/commit/86c4dfa209b5afe7a5bf8e50efea191446f9e2ae.diff
LOG: [mlir][gpu] Add GPU target attribute interface.
**For an explanation of these patches see D154153.**
Commit message:
This patch adds the `GPUTargetAttrInterface` attribute interface, this interface
is meant to be used as an opaque interface for serializing GPU modules into
binary strings.
Reviewed By: mehdi_amini, krzysz00
Differential Revision: https://reviews.llvm.org/D154104
Added:
mlir/include/mlir/Dialect/GPU/IR/CompilationAttrInterfaces.td
mlir/include/mlir/Dialect/GPU/IR/CompilationInterfaces.h
Modified:
mlir/include/mlir/Dialect/GPU/IR/CMakeLists.txt
mlir/include/mlir/Dialect/GPU/IR/GPUDialect.h
mlir/lib/Dialect/GPU/CMakeLists.txt
mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
Removed:
################################################################################
diff --git a/mlir/include/mlir/Dialect/GPU/IR/CMakeLists.txt b/mlir/include/mlir/Dialect/GPU/IR/CMakeLists.txt
index feef123eb9c991..a16ae297400499 100644
--- 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.h.inc -gen-enum-decls)
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
index 00000000000000..a3c7e5a5284f06
--- /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<SmallVector<char, 0>>", "serializeToObject",
+ (ins "Operation*":$module, "const gpu::TargetOptions&":$options)>
+ ];
+}
+
+def GPUTargetArrayAttr : TypedArrayAttrBase<GPUTargetAttrInterface,
+ "array of GPU target attributes">;
+
+def GPUNonEmptyTargetArrayAttr :
+ ConfinedAttr<GPUTargetArrayAttr, [ArrayMinCount<1>]>;
+
+#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
index 00000000000000..090e272cb17e7c
--- /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<std::string> 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<std::string> getLinkFiles() const;
+
+ /// Returns the command line options.
+ StringRef getCmdOptions() const;
+
+ /// Returns a tokenization of the command line options.
+ std::pair<llvm::BumpPtrAllocator, SmallVector<const char *>>
+ 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<DerivedClass>())`.
+ TargetOptions(TypeID typeID, StringRef toolkitPath = {},
+ ArrayRef<std::string> linkFiles = {}, StringRef cmdOptions = {},
+ CompilationTarget compilationTarget = binary);
+
+ /// Path to the target toolkit.
+ std::string toolkitPath;
+
+ /// List of files to link with the LLVM module.
+ SmallVector<std::string> 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
index c27306cb775b13..adeb7b87ce57ce 100644
--- 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
index 4250e40eac491f..f3c518fd93066e 100644
--- a/mlir/lib/Dialect/GPU/CMakeLists.txt
+++ b/mlir/lib/Dialect/GPU/CMakeLists.txt
@@ -32,6 +32,7 @@ add_mlir_dialect_library(MLIRGPUDialect
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
index 77a2e01b5e0758..c89973c3290449 100644
--- 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 @@ void AllocOp::getCanonicalizationPatterns(RewritePatternSet &results,
results.add<SimplifyDimOfAllocOp>(context);
}
+//===----------------------------------------------------------------------===//
+// GPU target options
+//===----------------------------------------------------------------------===//
+
+TargetOptions::TargetOptions(StringRef toolkitPath,
+ ArrayRef<std::string> linkFiles,
+ StringRef cmdOptions,
+ CompilationTarget compilationTarget)
+ : TargetOptions(TypeID::get<TargetOptions>(), toolkitPath, linkFiles,
+ cmdOptions, compilationTarget) {}
+
+TargetOptions::TargetOptions(TypeID typeID, StringRef toolkitPath,
+ ArrayRef<std::string> 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<std::string> TargetOptions::getLinkFiles() const { return linkFiles; }
+
+StringRef TargetOptions::getCmdOptions() const { return cmdOptions; }
+
+std::pair<llvm::BumpPtrAllocator, SmallVector<const char *>>
+TargetOptions::tokenizeCmdOptions() const {
+ std::pair<llvm::BumpPtrAllocator, SmallVector<const char *>> 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 @@ void AllocOp::getCanonicalizationPatterns(RewritePatternSet &results,
#define GET_OP_CLASSES
#include "mlir/Dialect/GPU/IR/GPUOps.cpp.inc"
+
+#include "mlir/Dialect/GPU/IR/CompilationAttrInterfaces.cpp.inc"
More information about the Mlir-commits
mailing list