online_compiler.cpp

//==----------- online_compiler.cpp ----------------------------------------==//
//
// 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
//
//===----------------------------------------------------------------------===//

#define __SYCL_ONLINE_COMPILER_CPP

#include <sycl/detail/os_util.hpp>
#include <sycl/detail/ur.hpp>
#include <sycl/ext/intel/experimental/online_compiler.hpp>

#include <cstring>

#include "ocloc_api.h"

namespace sycl {
inline namespace _V1 {
namespace ext::intel::experimental {
namespace detail {

using namespace sycl::detail;

static std::vector<const char *>
prepareOclocArgs(sycl::info::device_type DeviceType, device_arch DeviceArch,
                 bool Is64Bit, string_view DeviceStepping,
                 const std::string &UserArgs) {
  std::vector<const char *> Args = {"ocloc", "-q", "-spv_only", "-device"};

  if (DeviceType == sycl::info::device_type::gpu) {
    switch (DeviceArch) {
    case device_arch::gpu_gen9:
      Args.push_back("skl");
      break;

    case device_arch::gpu_gen9_5:
      Args.push_back("cfl");
      break;

    case device_arch::gpu_gen11:
      Args.push_back("icllp");
      break;

    case device_arch::gpu_gen12:
      Args.push_back("tgllp");
      break;

    default:
      Args.push_back("tgllp");
    }
  } else {
    // TODO: change that to generic device when ocloc adds support for it.
    // For now "tgllp" is used as the option supported on all known GPU RT.
    Args.push_back("tgllp");
  }

  if (DeviceStepping != "") {
    Args.push_back("-revision_id");
    Args.push_back(DeviceStepping.data());
  }

  Args.push_back(Is64Bit ? "-64" : "-32");

  if (UserArgs != "") {
    Args.push_back("-options");
    Args.push_back(UserArgs.c_str());
  }

  return Args;
}

/// Compiles the given source \p Source to SPIR-V IL and returns IL as a vector
/// of bytes.
/// @param Source - Either OpenCL or CM source code.
/// @param DeviceType - SYCL device type, e.g. cpu, gpu, accelerator, etc.
/// @param DeviceArch - More detailed info on the target device architecture.
/// @param Is64Bit - If set to true, specifies the 64-bit architecture.
///                  Otherwise, 32-bit is assumed.
/// @param DeviceStepping - implementation specific target device stepping.
/// @param CompileToSPIRVHandle - Output parameter. It is set to the address
///                               of the library function doing the compilation.
/// @param FreeSPIRVOutputsHandle - Output parameter. It is set to the address
///                                 of the library function freeing memory
///                                 allocated during the compilation.
/// @param UserArgs - User's options to ocloc compiler.
static std::vector<byte>
compileToSPIRV(string_view Src, sycl::info::device_type DeviceType,
               device_arch DeviceArch, bool Is64Bit, string_view DeviceStepping,
               void *&CompileToSPIRVHandle, void *&FreeSPIRVOutputsHandle,
               const std::vector<std::string> &UserArgs) {
  std::string Source{Src.data()};

  if (!CompileToSPIRVHandle) {
#ifdef __SYCL_RT_OS_WINDOWS
    static const std::string OclocLibraryName = "ocloc64.dll";
#else
    static const std::string OclocLibraryName = "libocloc.so";
#endif
    void *OclocLibrary = sycl::detail::ur::loadOsLibrary(OclocLibraryName);
    if (!OclocLibrary)
      throw online_compile_error("Cannot load ocloc library: " +
                                 OclocLibraryName);
    void *OclocVersionHandle =
        sycl::detail::ur::getOsLibraryFuncAddress(OclocLibrary, "oclocVersion");
    // The initial versions of ocloc library did not have the oclocVersion()
    // function. Those versions had the same API as the first version of ocloc
    // library having that oclocVersion() function.
    int LoadedVersion = ocloc_version_t::OCLOC_VERSION_1_0;
    if (OclocVersionHandle) {
      decltype(::oclocVersion) *OclocVersionFunc =
          reinterpret_cast<decltype(::oclocVersion) *>(OclocVersionHandle);
      LoadedVersion = OclocVersionFunc();
    }
    // The loaded library with version (A.B) is compatible with expected API/ABI
    // version (X.Y) used here if A == B and B >= Y.
    int LoadedVersionMajor = LoadedVersion >> 16;
    int LoadedVersionMinor = LoadedVersion & 0xffff;
    int CurrentVersionMajor = ocloc_version_t::OCLOC_VERSION_CURRENT >> 16;
    int CurrentVersionMinor = ocloc_version_t::OCLOC_VERSION_CURRENT & 0xffff;
    if (LoadedVersionMajor != CurrentVersionMajor ||
        LoadedVersionMinor < CurrentVersionMinor)
      throw online_compile_error(
          std::string("Found incompatible version of ocloc library: (") +
          std::to_string(LoadedVersionMajor) + "." +
          std::to_string(LoadedVersionMinor) +
          "). The supported versions are (" +
          std::to_string(CurrentVersionMajor) +
          ".N), where (N >= " + std::to_string(CurrentVersionMinor) + ").");

    CompileToSPIRVHandle =
        sycl::detail::ur::getOsLibraryFuncAddress(OclocLibrary, "oclocInvoke");
    if (!CompileToSPIRVHandle)
      throw online_compile_error("Cannot load oclocInvoke() function");
    FreeSPIRVOutputsHandle = sycl::detail::ur::getOsLibraryFuncAddress(
        OclocLibrary, "oclocFreeOutput");
    if (!FreeSPIRVOutputsHandle)
      throw online_compile_error("Cannot load oclocFreeOutput() function");
  }

  std::string CombinedUserArgs;
  for (auto UserArg : UserArgs) {
    if (UserArg == "")
      continue;
    if (CombinedUserArgs != "")
      CombinedUserArgs = CombinedUserArgs + " " + UserArg;
    else
      CombinedUserArgs = UserArg;
  }
  std::vector<const char *> Args = detail::prepareOclocArgs(
      DeviceType, DeviceArch, Is64Bit, DeviceStepping, CombinedUserArgs);

  uint32_t NumOutputs = 0;
  byte **Outputs = nullptr;
  uint64_t *OutputLengths = nullptr;
  char **OutputNames = nullptr;

  const byte *Sources[] = {reinterpret_cast<const byte *>(Source.c_str())};
  const char *SourceName = "main.cl";
  const uint64_t SourceLengths[] = {Source.length() + 1};

  Args.push_back("-file");
  Args.push_back(SourceName);

  decltype(::oclocInvoke) *OclocInvokeFunc =
      reinterpret_cast<decltype(::oclocInvoke) *>(CompileToSPIRVHandle);
  int CompileError =
      OclocInvokeFunc(Args.size(), Args.data(), 1, Sources, SourceLengths,
                      &SourceName, 0, nullptr, nullptr, nullptr, &NumOutputs,
                      &Outputs, &OutputLengths, &OutputNames);

  std::vector<byte> SpirV;
  std::string CompileLog;
  for (uint32_t I = 0; I < NumOutputs; I++) {
    size_t NameLen = strlen(OutputNames[I]);
    if (NameLen >= 4 && strstr(OutputNames[I], ".spv") != nullptr &&
        Outputs[I] != nullptr) {
      assert(SpirV.size() == 0 && "More than one SPIR-V output found.");
      SpirV = std::vector<byte>(Outputs[I], Outputs[I] + OutputLengths[I]);
    } else if (!strcmp(OutputNames[I], "stdout.log")) {
      CompileLog = std::string(reinterpret_cast<const char *>(Outputs[I]));
    }
  }

  // Try to free memory before reporting possible error.
  decltype(::oclocFreeOutput) *OclocFreeOutputFunc =
      reinterpret_cast<decltype(::oclocFreeOutput) *>(FreeSPIRVOutputsHandle);
  int MemFreeError =
      OclocFreeOutputFunc(&NumOutputs, &Outputs, &OutputLengths, &OutputNames);

  if (CompileError)
    throw online_compile_error("ocloc reported compilation errors: {\n" +
                               CompileLog + "\n}");
  if (SpirV.empty())
    throw online_compile_error(
        "Unexpected output: ocloc did not return SPIR-V");
  if (MemFreeError)
    throw online_compile_error("ocloc cannot safely free resources");

  return SpirV;
}
} // namespace detail

template <source_language Lang>
__SYCL_EXPORT std::vector<byte> online_compiler<Lang>::compile_impl(
    detail::string_view Src, const std::vector<detail::string_view> &Options,
    std::pair<int, int> OutputFormatVersion, sycl::info::device_type DeviceType,
    device_arch DeviceArch, bool Is64Bit, detail::string_view DeviceStepping,
    void *&CompileToSPIRVHandle, void *&FreeSPIRVOutputsHandle) {

  if (OutputFormatVersion != std::pair<int, int>{0, 0}) {
    std::string Version = std::to_string(OutputFormatVersion.first) + ", " +
                          std::to_string(OutputFormatVersion.second);
    throw online_compile_error(std::string("The output format version (") +
                               Version + ") is not supported yet");
  }

  std::vector<std::string> UserArgs;
  for (auto &&Opt : Options)
    UserArgs.emplace_back(Opt.data());

  if constexpr (Lang == source_language::cm)
    UserArgs.push_back("-cmc");

  return detail::compileToSPIRV(Src, DeviceType, DeviceArch, Is64Bit,
                                DeviceStepping, CompileToSPIRVHandle,
                                FreeSPIRVOutputsHandle, UserArgs);
}

template __SYCL_EXPORT std::vector<byte>
online_compiler<source_language::opencl_c>::compile_impl(
    detail::string_view Src, const std::vector<detail::string_view> &Options,
    std::pair<int, int> OutputFormatVersion, sycl::info::device_type DeviceType,
    device_arch DeviceArch, bool Is64Bit, detail::string_view DeviceStepping,
    void *&CompileToSPIRVHandle, void *&FreeSPIRVOutputsHandle);

template __SYCL_EXPORT std::vector<byte>
online_compiler<source_language::cm>::compile_impl(
    detail::string_view Src, const std::vector<detail::string_view> &Options,
    std::pair<int, int> OutputFormatVersion, sycl::info::device_type DeviceType,
    device_arch DeviceArch, bool Is64Bit, detail::string_view DeviceStepping,
    void *&CompileToSPIRVHandle, void *&FreeSPIRVOutputsHandle);
} // namespace ext::intel::experimental

namespace ext {
namespace __SYCL2020_DEPRECATED(
    "use 'ext::intel::experimental' instead") intel {
using namespace ext::intel::experimental;
}
} // namespace ext

namespace __SYCL2020_DEPRECATED(
    "use 'ext::intel::experimental' instead") INTEL {
using namespace ext::intel::experimental;
}
} // namespace _V1
} // namespace sycl