//===------- online_compiler.hpp - Online source compilation service ------===// // // 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 // //===----------------------------------------------------------------------===// #pragma once #include #include #include namespace sycl { namespace ext::libceed { using byte = unsigned char; enum class compiled_code_format { spir_v = 0 // the only format supported for now }; class device_arch { public: static constexpr int any = 0; device_arch(int Val) : Val(Val) {} // TODO1: the list must be extended with a bunch of new GPUs available. // TODO2: the list of supported GPUs grows rapidly. // The API must allow user to define the target GPU option even if it is // not listed in this enumerator below. enum gpu { gpu_any = 1, gpu_gen9 = 2, gpu_skl = gpu_gen9, gpu_gen9_5 = 3, gpu_kbl = gpu_gen9_5, gpu_cfl = gpu_gen9_5, gpu_gen11 = 4, gpu_icl = gpu_gen11, gpu_gen12 = 5, gpu_tgl = gpu_gen12, gpu_tgllp = gpu_gen12 }; enum cpu { cpu_any = 1, }; enum fpga { fpga_any = 1, }; operator int() { return Val; } private: int Val; }; /// Represents an error happend during online compilation. class online_compile_error : public sycl::exception { public: online_compile_error() = default; online_compile_error(const std::string &Msg) : sycl::exception(Msg) {} }; /// Designates a source language for the online compiler. enum class source_language { opencl_c = 0, cm = 1 }; /// Represents an online compiler for the language given as template /// parameter. template class online_compiler { public: /// Constructs online compiler which can target any device and produces /// given compiled code format. Produces 64-bit device code. /// The created compiler is "optimistic" - it assumes all applicable SYCL /// device capabilities are supported by the target device(s). online_compiler(compiled_code_format fmt = compiled_code_format::spir_v) : OutputFormat(fmt), OutputFormatVersion({0, 0}), DeviceType(sycl::info::device_type::all), DeviceArch(device_arch::any), Is64Bit(true), DeviceStepping("") {} /// Constructs online compiler which targets given architecture and produces /// given compiled code format. Produces 64-bit device code. /// Throws online_compile_error if values of constructor arguments are /// contradictory or not supported - e.g. if the source language is not /// supported for given device type. online_compiler(sycl::info::device_type dev_type, device_arch arch, compiled_code_format fmt = compiled_code_format::spir_v) : OutputFormat(fmt), OutputFormatVersion({0, 0}), DeviceType(dev_type), DeviceArch(arch), Is64Bit(true), DeviceStepping("") {} /// Constructs online compiler for the target specified by given SYCL device. // TODO: the initial version generates the generic code (SKL now), need // to do additional device::info calls to determine the device by it's // features. online_compiler(const sycl::device &) : OutputFormat(compiled_code_format::spir_v), OutputFormatVersion({0, 0}), DeviceType(sycl::info::device_type::all), DeviceArch(device_arch::any), Is64Bit(true), DeviceStepping("") {} /// Compiles given in-memory \c Lang source to a binary blob. Blob format, /// other parameters are set in the constructor by the compilation target /// specification parameters. /// Specialization for each language will provide exact signatures, which /// can be different for different languages. /// Throws online_compile_error if compilation is not successful. template std::vector compile(const std::string &src, const Tys &...args); /// Sets the compiled code format of the compilation target and returns *this. online_compiler &setOutputFormat(compiled_code_format fmt) { OutputFormat = fmt; return *this; } /// Sets the compiled code format version of the compilation target and /// returns *this. online_compiler &setOutputFormatVersion(int major, int minor) { OutputFormatVersion = {major, minor}; return *this; } /// Sets the device type of the compilation target and returns *this. online_compiler &setTargetDeviceType(sycl::info::device_type type) { DeviceType = type; return *this; } /// Sets the device architecture of the compilation target and returns *this. online_compiler &setTargetDeviceArch(device_arch arch) { DeviceArch = arch; return *this; } /// Makes the compilation target 32-bit and returns *this. online_compiler &set32bitTarget() { Is64Bit = false; return *this; }; /// Makes the compilation target 64-bit and returns *this. online_compiler &set64bitTarget() { Is64Bit = true; return *this; }; /// Sets implementation-defined target device stepping of the compilation /// target and returns *this. online_compiler &setTargetDeviceStepping(const std::string &id) { DeviceStepping = id; return *this; } private: /// Compiled code format. compiled_code_format OutputFormat; /// Compiled code format version - a pair of "major" and "minor" components std::pair OutputFormatVersion; /// Target device type sycl::info::device_type DeviceType; /// Target device architecture device_arch DeviceArch; /// Whether the target device architecture is 64-bit bool Is64Bit; /// Target device stepping (implementation defined) std::string DeviceStepping; /// Handles to helper functions used by the implementation. void *CompileToSPIRVHandle = nullptr; void *FreeSPIRVOutputsHandle = nullptr; }; // Specializations of the online_compiler class and 'compile' function for // particular languages and parameter types. /// Compiles the given OpenCL source. May throw \c online_compile_error. /// @param src - contents of the source. /// @param options - compilation options (implementation defined); standard /// OpenCL JIT compiler options must be supported. template <> template <> std::vector online_compiler::compile(const std::string &src, const std::vector &options); // /// Compiles the given OpenCL source. May throw \c online_compile_error. // /// @param src - contents of the source. // template <> // template <> // std::vector // online_compiler::compile(const std::string &src) { // return compile(src, std::vector{}); // } /// Compiles the given CM source \p src. /// @param src - contents of the source. /// @param options - compilation options (implementation defined). template <> template <> std::vector online_compiler::compile(const std::string &src, const std::vector &options); /// Compiles the given CM source \p src. // template <> // template <> // std::vector online_compiler::compile(const std::string &src) { // return compile(src, std::vector{}); // } } // namespace ext::libceed } // namespace sycl