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 | |
18 | namespace llvm { |
19 | class IRBuilderBase; |
20 | } |
21 | |
22 | namespace mlir { |
23 | class SymbolTable; |
24 | namespace LLVM { |
25 | class ModuleTranslation; |
26 | } |
27 | namespace gpu { |
28 | enum 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`. |
34 | template <typename ConcreteType> |
35 | class 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. |
45 | class TargetOptions { |
46 | public: |
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 | |
86 | protected: |
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 | |
112 | private: |
113 | TypeID typeID; |
114 | }; |
115 | } // namespace gpu |
116 | } // namespace mlir |
117 | |
118 | MLIR_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 | |