1//===-- CompilationInterfaces.h - GPU compilation interfaces ---*- C++ -*-===//
2//
3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4// See https://llvm.org/LICENSE.txt for license information.
5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6//
7//===----------------------------------------------------------------------===//
8//
9// This file defines interfaces for GPU compilation.
10//
11//===----------------------------------------------------------------------===//
12
13#ifndef MLIR_DIALECT_GPU_IR_COMPILATIONINTERFACES_H
14#define MLIR_DIALECT_GPU_IR_COMPILATIONINTERFACES_H
15
16#include "mlir/IR/Attributes.h"
17
18namespace llvm {
19class IRBuilderBase;
20}
21
22namespace mlir {
23class SymbolTable;
24namespace LLVM {
25class ModuleTranslation;
26}
27namespace gpu {
28enum class CompilationTarget : uint32_t;
29
30/// This class indicates that the attribute associated with this trait is a GPU
31/// offloading translation attribute. These kinds of attributes must implement
32/// an interface for handling the translation of GPU offloading operations like
33/// `gpu.binary` & `gpu.launch_func`.
34template <typename ConcreteType>
35class OffloadingTranslationAttrTrait
36 : public AttributeTrait::TraitBase<ConcreteType,
37 OffloadingTranslationAttrTrait> {
38 // TODO: Verify the attribute promises or implements the interface.
39};
40
41/// This class serves as an opaque interface for passing options to the
42/// `TargetAttrInterface` methods. Users of this class must implement the
43/// `classof` method as well as using the macros `MLIR_*_EXPLICIT_TYPE_ID` to
44/// ensure type safeness. Targets are free to ignore these options.
45class TargetOptions {
46public:
47 /// Constructor initializing the toolkit path, the list of files to link to,
48 /// extra command line options, the compilation target and a callback for
49 /// obtaining the parent symbol table. The default compilation target is
50 /// `Fatbin`.
51 TargetOptions(
52 StringRef toolkitPath = {}, ArrayRef<std::string> linkFiles = {},
53 StringRef cmdOptions = {},
54 CompilationTarget compilationTarget = getDefaultCompilationTarget(),
55 function_ref<SymbolTable *()> getSymbolTableCallback = {});
56
57 /// Returns the typeID.
58 TypeID getTypeID() const;
59
60 /// Returns the toolkit path.
61 StringRef getToolkitPath() const;
62
63 /// Returns the files to link to.
64 ArrayRef<std::string> getLinkFiles() const;
65
66 /// Returns the command line options.
67 StringRef getCmdOptions() const;
68
69 /// Returns a tokenization of the command line options.
70 std::pair<llvm::BumpPtrAllocator, SmallVector<const char *>>
71 tokenizeCmdOptions() const;
72
73 /// Returns the compilation target.
74 CompilationTarget getCompilationTarget() const;
75
76 /// Returns the result of the `getSymbolTableCallback` callback or a nullptr
77 /// if no callback was provided.
78 /// Note: The callback itself can return nullptr. It is up to the target how
79 /// to react to getting a nullptr, e.g., emitting an error or constructing the
80 /// table.
81 SymbolTable *getSymbolTable() const;
82
83 /// Returns the default compilation target: `CompilationTarget::Fatbin`.
84 static CompilationTarget getDefaultCompilationTarget();
85
86protected:
87 /// Derived classes must use this constructor to initialize `typeID` to the
88 /// appropiate value: ie. `TargetOptions(TypeID::get<DerivedClass>())`.
89 TargetOptions(
90 TypeID typeID, StringRef toolkitPath = {},
91 ArrayRef<std::string> linkFiles = {}, StringRef cmdOptions = {},
92 CompilationTarget compilationTarget = getDefaultCompilationTarget(),
93 function_ref<SymbolTable *()> getSymbolTableCallback = {});
94
95 /// Path to the target toolkit.
96 std::string toolkitPath;
97
98 /// List of files to link with the LLVM module.
99 SmallVector<std::string> linkFiles;
100
101 /// An optional set of command line options to be used by the compilation
102 /// process.
103 std::string cmdOptions;
104
105 /// Compilation process target format.
106 CompilationTarget compilationTarget;
107
108 /// Callback for obtaining the parent symbol table of all the GPU modules
109 /// being serialized.
110 function_ref<SymbolTable *()> getSymbolTableCallback;
111
112private:
113 TypeID typeID;
114};
115} // namespace gpu
116} // namespace mlir
117
118MLIR_DECLARE_EXPLICIT_TYPE_ID(::mlir::gpu::TargetOptions)
119
120#include "mlir/Dialect/GPU/IR/CompilationAttrInterfaces.h.inc"
121
122#endif // MLIR_DIALECT_GPU_IR_COMPILATIONINTERFACES_H
123

source code of mlir/include/mlir/Dialect/GPU/IR/CompilationInterfaces.h