|
| 1 | += SYCL Intel extension: Online Compilation |
| 2 | +Konstantin Bobrovskii < konstantin[email protected]>, John Pennycook < john[email protected]> |
| 3 | +v0.1 |
| 4 | +:source-highlighter: pygments |
| 5 | +:icons: font |
| 6 | +:dpcpp: pass:[DPC++] |
| 7 | + |
| 8 | +== Introduction |
| 9 | +This document describes an interface for online compilation from high-level languages, such as |
| 10 | +OpenCL, to a binary format, such as SPIR-V, loadable by SYCL backends. Unlike SYCL 2020 provisional's |
| 11 | +OpenCL backend online compilation interface, this interface is not bound to any particular backend and does |
| 12 | +not require available SYCL context for online compilation. |
| 13 | + |
| 14 | +This gives a flexibility to "cross-compile" to SPIR-V or other supported formats without any SYCL |
| 15 | +device or context available. The online compilation API uses the `online_compiler` class to access |
| 16 | +compilation services. Instances of the class are constructed based on a specification of the desired |
| 17 | +compilation target passed to the constructors - such as compiled code format, target architecture, |
| 18 | +etc. All the settings are optional, and by default the target is generic SPIR-V. |
| 19 | + |
| 20 | +This API is an Intel SYCL extension. |
| 21 | + |
| 22 | +== Status |
| 23 | + |
| 24 | +This extension has been deprecated. Although it is still supported in {dpcpp}, |
| 25 | +we expect that the interfaces defined in this specification will be removed in |
| 26 | +an upcoming {dpcpp} release. *Shipping software products should stop using |
| 27 | +APIs defined in this specification and use an alternative instead.* |
| 28 | + |
| 29 | +== Online compilation API |
| 30 | + |
| 31 | +All online compilation API elements reside in the `sycl::INTEL` namespace. |
| 32 | + |
| 33 | +=== Source language specification |
| 34 | + |
| 35 | +Elements of the enum designate the source language: |
| 36 | +[source,c++] |
| 37 | +----------------- |
| 38 | +enum class source_language { |
| 39 | + opencl_c, // OpenCL C language |
| 40 | + cm // Intel's C-for-Media language |
| 41 | +}; |
| 42 | +----------------- |
| 43 | + |
| 44 | +=== APIs to express compilation target characteristics |
| 45 | + |
| 46 | +The desired format of the compiled code: |
| 47 | +[source,c++] |
| 48 | +----------------- |
| 49 | +enum class compiled_code_format { |
| 50 | + spir_v |
| 51 | +}; |
| 52 | +----------------- |
| 53 | + |
| 54 | +Target device architecture: |
| 55 | +[source,c++] |
| 56 | +----------------- |
| 57 | +class device_arch { |
| 58 | +public: |
| 59 | + static constexpr int any = 0; // designates an unspecified architecture |
| 60 | + device_arch(int Val); |
| 61 | +
|
| 62 | + // GPU architecture IDs |
| 63 | + enum gpu { gpu_any = 1, ... }; |
| 64 | + // CPU architecture IDs |
| 65 | + enum cpu { cpu_any = 1, ... }; |
| 66 | + // FPGA architecture IDs |
| 67 | + enum fpga { fpga_any = 1, ... }; |
| 68 | +
|
| 69 | + // Converts this architecture representation to an integer. |
| 70 | + operator int(); |
| 71 | +}; |
| 72 | +----------------- |
| 73 | + |
| 74 | +=== Compiler API |
| 75 | + |
| 76 | +To compile a source, a user program must first construct an instance of the `sycl::ext::intel::online_compiler` class. Then pass the source as a `std::string` object to online compiler's `compile` function along with other relevant parameters. The `online_compiler` is templated by the source language, and the `compile` function is a variadic template function. Instantiations of the `online_compiler::compile` for different languages may have different sets of formal parameters. The `compile` function returns a binary blob - a `std::vector<unsigned char>` - with the device code compiled according to the compilation target specification provided at online compiler construction time. |
| 77 | + |
| 78 | +==== Online compiler |
| 79 | +[source,c++] |
| 80 | +----------------- |
| 81 | +template <source_language Lang> class online_compiler; |
| 82 | +----------------- |
| 83 | + |
| 84 | +==== Compilation target specification elements. |
| 85 | +[cols="40,60",options="header"] |
| 86 | +|=== |
| 87 | +|Element name and type |Description |
| 88 | + |
| 89 | +|`compiled_code_format` OutputFormat |
| 90 | +|Compiled code format. |
| 91 | + |
| 92 | +|`std::pair<int, int>` OutputFormatVersion |
| 93 | +|Compiled code format version - a pair of "major" and "minor" components. |
| 94 | + |
| 95 | +|`sycl::info::device_type` DeviceType |
| 96 | +|Target device type. |
| 97 | + |
| 98 | +|`device_arch` DeviceArch |
| 99 | +|Target device architecture. |
| 100 | + |
| 101 | +|`bool` Is64Bit |
| 102 | +|Whether the target device architecture is 64-bit. |
| 103 | + |
| 104 | +|`std::string` DeviceStepping |
| 105 | +|Target device stepping (implementation defined). |
| 106 | +|=== |
| 107 | + |
| 108 | +Online compiler construction or source compilation may be unsuccessful, in which case an instance |
| 109 | +of `sycl::ext::intel::online_compile_error` is thrown. For example, when some of the compilation |
| 110 | +target specification elements are not supported by the implementation, or there is a syntax error |
| 111 | +in the source program. |
| 112 | + |
| 113 | + |
| 114 | +==== `sycl::ext::intel::online_compiler` constructors. |
| 115 | +[cols="40,60",options="header"] |
| 116 | +|=== |
| 117 | +|Constructor |Description |
| 118 | + |
| 119 | +|`online_compiler(compiled_code_format fmt = compiled_code_format::spir_v)` |
| 120 | +| Constructs online compiler which can target any device and produces |
| 121 | + given compiled code format. Produced device code is 64-bit. OutputFormatVersion is |
| 122 | + implementation defined. The created compiler is "optimistic" - it assumes all applicable SYCL |
| 123 | + device capabilities are supported by the target device(s). |
| 124 | + |
| 125 | +|`online_compiler( |
| 126 | + sycl::info::device_type dev_type, |
| 127 | + device_arch arch, |
| 128 | + compiled_code_format fmt = compiled_code_format::spir_v)` |
| 129 | +| Constructor version which allows to specify target device type and architecture. |
| 130 | + |
| 131 | +|`online_compiler(const sycl::device &dev)` |
| 132 | +|Constructs online compiler for the target specified by given SYCL device. |
| 133 | +|=== |
| 134 | + |
| 135 | +==== The compilation function - `online_compiler::compile` |
| 136 | +It compiles given in-memory source to a binary blob. Blob format, |
| 137 | +other parameters are set in the constructor. Specialization for each language will provide exact |
| 138 | +signatures, which can be different for different languages.Throws `online_compile_error` if |
| 139 | +compilation is not successful. |
| 140 | +[source,c++] |
| 141 | +----------------- |
| 142 | +template <typename... Tys> |
| 143 | + std::vector<byte> compile(const std::string &src, const Tys&... args); |
| 144 | +----------------- |
| 145 | + |
| 146 | +Instantiations of the compilation function: |
| 147 | +[source,c++] |
| 148 | +----------------- |
| 149 | +/// Compiles given OpenCL source. May throw \c online_compile_error. |
| 150 | +/// @param src - contents of the source |
| 151 | +/// @param options - compilation options (implementation defined); standard |
| 152 | +/// OpenCL JIT compiler options must be supported |
| 153 | +template <> |
| 154 | +template <> |
| 155 | +std::vector<byte> online_compiler<source_language::opencl_c>::compile( |
| 156 | + const std::string &src, const std::vector<std::string> &options); |
| 157 | +
|
| 158 | +/// Compiles given CM source. |
| 159 | +template <> |
| 160 | +template <> |
| 161 | +std::vector<byte> online_compiler<source_language::cm>::compile( |
| 162 | + const std::string &src); |
| 163 | +
|
| 164 | +/// Compiles given CM source. |
| 165 | +/// @param options - compilation options (implementation defined) |
| 166 | +template <> |
| 167 | +template <> |
| 168 | +std::vector<byte> online_compiler<source_language::cm>::compile( |
| 169 | + const std::string &src, const std::vector<std::string> &options); |
| 170 | +----------------- |
| 171 | + |
| 172 | +== API usage example |
| 173 | +This example compiles an OpenCL source to a generic SPIR-V. |
| 174 | +[source,c++] |
| 175 | +----------------- |
| 176 | +#include "sycl/ext/intel/online_compiler.hpp" |
| 177 | +
|
| 178 | +#include <iostream> |
| 179 | +#include <vector> |
| 180 | +
|
| 181 | +static const char *kernelSource = R"===( |
| 182 | +__kernel void my_kernel(__global int *in, __global int *out) { |
| 183 | + size_t i = get_global_id(0); |
| 184 | + out[i] = in[i] + 1; |
| 185 | +} |
| 186 | +)==="; |
| 187 | +
|
| 188 | +using namespace sycl::INTEL; |
| 189 | +
|
| 190 | +int main(int argc, char **argv) { |
| 191 | + online_compiler<source_language::opencl_c> compiler; |
| 192 | + std::vector<byte> blob; |
| 193 | +
|
| 194 | + try { |
| 195 | + blob = compiler.compile( |
| 196 | + std::string(kernelSource), |
| 197 | + std::vector<std::string> { |
| 198 | + std::string("-cl-fast-relaxed-math") |
| 199 | + } |
| 200 | + ); |
| 201 | + } |
| 202 | + catch (online_compile_error &e) { |
| 203 | + std::cout << "compilation failed\n"; |
| 204 | + return 1; |
| 205 | + } |
| 206 | + return 0; |
| 207 | +} |
| 208 | +----------------- |
0 commit comments