online_compiler.hpp

Go to the documentation of this file.

//===------- 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 <CL/sycl/detail/defines_elementary.hpp> // for __SYCL_INLINE_NAMESPACE
#include <CL/sycl/detail/export.hpp>             // for __SYCL_EXPORT
#include <CL/sycl/device.hpp>
 
#include <string>
#include <vector>
 
__SYCL_INLINE_NAMESPACE(cl) {
namespace sycl {
namespace ext {
namespace intel {
namespace experimental {
 
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) {}
 
  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
  };
 
  enum cpu {
    cpu_any = 1,
  };
 
  enum fpga {
    fpga_any = 1,
  };
 
  operator int() { return Val; }
 
private:
  int Val;
};
 
class online_compile_error : public sycl::exception {
public:
  online_compile_error() = default;
  online_compile_error(const std::string &Msg) : sycl::exception(Msg) {}
};
 
enum class source_language { opencl_c = 0, cm = 1 };
 
template <source_language Lang> class online_compiler {
public:
  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("") {}
 
  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("") {}
 
  // 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("") {}
 
  template <typename... Tys>
  std::vector<byte> compile(const std::string &src, const Tys &... args);
 
  online_compiler<Lang> &setOutputFormat(compiled_code_format fmt) {
    OutputFormat = fmt;
    return *this;
  }
 
  online_compiler<Lang> &setOutputFormatVersion(int major, int minor) {
    OutputFormatVersion = {major, minor};
    return *this;
  }
 
  online_compiler<Lang> &setTargetDeviceType(sycl::info::device_type type) {
    DeviceType = type;
    return *this;
  }
 
  online_compiler<Lang> &setTargetDeviceArch(device_arch arch) {
    DeviceArch = arch;
    return *this;
  }
 
  online_compiler<Lang> &set32bitTarget() {
    Is64Bit = false;
    return *this;
  };
 
  online_compiler<Lang> &set64bitTarget() {
    Is64Bit = true;
    return *this;
  };
 
  online_compiler<Lang> &setTargetDeviceStepping(const std::string &id) {
    DeviceStepping = id;
    return *this;
  }
 
private:
  compiled_code_format OutputFormat;
 
  std::pair<int, int> OutputFormatVersion;
 
  sycl::info::device_type DeviceType;
 
  device_arch DeviceArch;
 
  bool Is64Bit;
 
  std::string DeviceStepping;
 
  void *CompileToSPIRVHandle = nullptr;
  void *FreeSPIRVOutputsHandle = nullptr;
};
 
// Specializations of the online_compiler class and 'compile' function for
// particular languages and parameter types.
 
template <>
template <>
__SYCL_EXPORT std::vector<byte>
online_compiler<source_language::opencl_c>::compile(
    const std::string &src, const std::vector<std::string> &options);
 
template <>
template <>
std::vector<byte>
online_compiler<source_language::opencl_c>::compile(const std::string &src) {
  return compile(src, std::vector<std::string>{});
}
 
template <>
template <>
__SYCL_EXPORT std::vector<byte> online_compiler<source_language::cm>::compile(
    const std::string &src, const std::vector<std::string> &options);
 
template <>
template <>
std::vector<byte>
online_compiler<source_language::cm>::compile(const std::string &src) {
  return compile(src, std::vector<std::string>{});
}
 
} // namespace experimental
} // namespace intel
} // namespace ext
 
namespace ext {
namespace __SYCL2020_DEPRECATED(
    "use 'ext::intel::experimental' instead") intel {
  using namespace ext::intel::experimental;
} // namespace intel
} // namespace ext
 
} // namespace sycl
} // __SYCL_INLINE_NAMESPACE(cl)