DPC++ Runtime
Runtime libraries for oneAPI DPC++
online_compiler.cpp
Go to the documentation of this file.
1 //==----------- online_compiler.cpp ----------------------------------------==//
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 
10 #include <CL/sycl/detail/pi.hpp>
12 
13 #include <cstring>
14 
15 #include "ocloc_api.h"
16 
18 namespace sycl {
19 namespace ext {
20 namespace intel {
21 namespace experimental {
22 namespace detail {
23 
24 static std::vector<const char *>
26  bool Is64Bit, const std::string &DeviceStepping,
27  const std::string &UserArgs) {
28  std::vector<const char *> Args = {"ocloc", "-q", "-spv_only", "-device"};
29 
30  if (DeviceType == sycl::info::device_type::gpu) {
31  switch (DeviceArch) {
32  case device_arch::gpu_gen9_5:
33  Args.push_back("cfl");
34  break;
35 
36  case device_arch::gpu_gen11:
37  Args.push_back("icllp");
38  break;
39 
40  default:
41  Args.push_back("skl");
42  }
43  } else {
44  // TODO: change that to generic device when ocloc adds support for it.
45  // For now "skl" is used as the lowest arch with GEN9 arch.
46  Args.push_back("skl");
47  }
48 
49  if (DeviceStepping != "") {
50  Args.push_back("-revision_id");
51  Args.push_back(DeviceStepping.c_str());
52  }
53 
54  Args.push_back(Is64Bit ? "-64" : "-32");
55 
56  if (UserArgs != "") {
57  Args.push_back("-options");
58  Args.push_back(UserArgs.c_str());
59  }
60 
61  return Args;
62 }
63 
78 static std::vector<byte>
79 compileToSPIRV(const std::string &Source, sycl::info::device_type DeviceType,
80  device_arch DeviceArch, bool Is64Bit,
81  const std::string &DeviceStepping, void *&CompileToSPIRVHandle,
82  void *&FreeSPIRVOutputsHandle,
83  const std::vector<std::string> &UserArgs) {
84 
85  if (!CompileToSPIRVHandle) {
86 #ifdef __SYCL_RT_OS_WINDOWS
87  static const std::string OclocLibraryName = "ocloc64.dll";
88 #else
89  static const std::string OclocLibraryName = "libocloc.so";
90 #endif
91  void *OclocLibrary = sycl::detail::pi::loadOsLibrary(OclocLibraryName);
92  if (!OclocLibrary)
93  throw online_compile_error("Cannot load ocloc library: " +
94  OclocLibraryName);
95  void *OclocVersionHandle =
96  sycl::detail::pi::getOsLibraryFuncAddress(OclocLibrary, "oclocVersion");
97  // The initial versions of ocloc library did not have the oclocVersion()
98  // function. Those versions had the same API as the first version of ocloc
99  // library having that oclocVersion() function.
100  int LoadedVersion = ocloc_version_t::OCLOC_VERSION_1_0;
101  if (OclocVersionHandle) {
102  decltype(::oclocVersion) *OclocVersionFunc =
103  reinterpret_cast<decltype(::oclocVersion) *>(OclocVersionHandle);
104  LoadedVersion = OclocVersionFunc();
105  }
106  // The loaded library with version (A.B) is compatible with expected API/ABI
107  // version (X.Y) used here if A == B and B >= Y.
108  int LoadedVersionMajor = LoadedVersion >> 16;
109  int LoadedVersionMinor = LoadedVersion & 0xffff;
110  int CurrentVersionMajor = ocloc_version_t::OCLOC_VERSION_CURRENT >> 16;
111  int CurrentVersionMinor = ocloc_version_t::OCLOC_VERSION_CURRENT & 0xffff;
112  if (LoadedVersionMajor != CurrentVersionMajor ||
113  LoadedVersionMinor < CurrentVersionMinor)
114  throw online_compile_error(
115  std::string("Found incompatible version of ocloc library: (") +
116  std::to_string(LoadedVersionMajor) + "." +
117  std::to_string(LoadedVersionMinor) +
118  "). The supported versions are (" +
119  std::to_string(CurrentVersionMajor) +
120  ".N), where (N >= " + std::to_string(CurrentVersionMinor) + ").");
121 
122  CompileToSPIRVHandle =
123  sycl::detail::pi::getOsLibraryFuncAddress(OclocLibrary, "oclocInvoke");
124  if (!CompileToSPIRVHandle)
125  throw online_compile_error("Cannot load oclocInvoke() function");
126  FreeSPIRVOutputsHandle = sycl::detail::pi::getOsLibraryFuncAddress(
127  OclocLibrary, "oclocFreeOutput");
128  if (!FreeSPIRVOutputsHandle)
129  throw online_compile_error("Cannot load oclocFreeOutput() function");
130  }
131 
132  std::string CombinedUserArgs;
133  for (auto UserArg : UserArgs) {
134  if (UserArg == "")
135  continue;
136  if (CombinedUserArgs != "")
137  CombinedUserArgs = CombinedUserArgs + " " + UserArg;
138  else
139  CombinedUserArgs = UserArg;
140  }
141  std::vector<const char *> Args = detail::prepareOclocArgs(
142  DeviceType, DeviceArch, Is64Bit, DeviceStepping, CombinedUserArgs);
143 
144  uint32_t NumOutputs = 0;
145  byte **Outputs = nullptr;
146  size_t *OutputLengths = nullptr;
147  char **OutputNames = nullptr;
148 
149  const byte *Sources[] = {reinterpret_cast<const byte *>(Source.c_str())};
150  const char *SourceName = "main.cl";
151  const uint64_t SourceLengths[] = {Source.length() + 1};
152 
153  Args.push_back("-file");
154  Args.push_back(SourceName);
155 
156  decltype(::oclocInvoke) *OclocInvokeFunc =
157  reinterpret_cast<decltype(::oclocInvoke) *>(CompileToSPIRVHandle);
158  int CompileError =
159  OclocInvokeFunc(Args.size(), Args.data(), 1, Sources, SourceLengths,
160  &SourceName, 0, nullptr, nullptr, nullptr, &NumOutputs,
161  &Outputs, &OutputLengths, &OutputNames);
162 
163  std::vector<byte> SpirV;
164  std::string CompileLog;
165  for (uint32_t I = 0; I < NumOutputs; I++) {
166  size_t NameLen = strlen(OutputNames[I]);
167  if (NameLen >= 4 && strstr(OutputNames[I], ".spv") != nullptr &&
168  Outputs[I] != nullptr) {
169  assert(SpirV.size() == 0 && "More than one SPIR-V output found.");
170  SpirV = std::vector<byte>(Outputs[I], Outputs[I] + OutputLengths[I]);
171  } else if (!strcmp(OutputNames[I], "stdout.log")) {
172  CompileLog = std::string(reinterpret_cast<const char *>(Outputs[I]));
173  }
174  }
175 
176  // Try to free memory before reporting possible error.
177  decltype(::oclocFreeOutput) *OclocFreeOutputFunc =
178  reinterpret_cast<decltype(::oclocFreeOutput) *>(FreeSPIRVOutputsHandle);
179  int MemFreeError =
180  OclocFreeOutputFunc(&NumOutputs, &Outputs, &OutputLengths, &OutputNames);
181 
182  if (CompileError)
183  throw online_compile_error("ocloc reported compilation errors: {\n" +
184  CompileLog + "\n}");
185  if (SpirV.empty())
186  throw online_compile_error(
187  "Unexpected output: ocloc did not return SPIR-V");
188  if (MemFreeError)
189  throw online_compile_error("ocloc cannot safely free resources");
190 
191  return SpirV;
192 }
193 } // namespace detail
194 
195 template <>
196 template <>
197 __SYCL_EXPORT std::vector<byte>
199  const std::string &Source, const std::vector<std::string> &UserArgs) {
200 
201  if (OutputFormatVersion != std::pair<int, int>{0, 0}) {
202  std::string Version = std::to_string(OutputFormatVersion.first) + ", " +
203  std::to_string(OutputFormatVersion.second);
204  throw online_compile_error(std::string("The output format version (") +
205  Version + ") is not supported yet");
206  }
207 
208  return detail::compileToSPIRV(Source, DeviceType, DeviceArch, Is64Bit,
209  DeviceStepping, CompileToSPIRVHandle,
210  FreeSPIRVOutputsHandle, UserArgs);
211 }
212 
213 template <>
214 template <>
215 __SYCL_EXPORT std::vector<byte> online_compiler<source_language::cm>::compile(
216  const std::string &Source, const std::vector<std::string> &UserArgs) {
217 
218  if (OutputFormatVersion != std::pair<int, int>{0, 0}) {
219  std::string Version = std::to_string(OutputFormatVersion.first) + ", " +
220  std::to_string(OutputFormatVersion.second);
221  throw online_compile_error(std::string("The output format version (") +
222  Version + ") is not supported yet");
223  }
224 
225  std::vector<std::string> CMUserArgs = UserArgs;
226  CMUserArgs.push_back("-cmc");
227  return detail::compileToSPIRV(Source, DeviceType, DeviceArch, Is64Bit,
228  DeviceStepping, CompileToSPIRVHandle,
229  FreeSPIRVOutputsHandle, CMUserArgs);
230 }
231 } // namespace experimental
232 } // namespace intel
233 } // namespace ext
234 
235 namespace ext {
236 namespace __SYCL2020_DEPRECATED(
237  "use 'ext::intel::experimental' instead") intel {
238  using namespace ext::intel::experimental;
239 } // namespace intel
240 } // namespace ext
241 
243  "use 'ext::intel::experimental' instead") INTEL {
244  using namespace ext::intel::experimental;
245 } // namespace INTEL
246 } // namespace sycl
247 } // __SYCL_INLINE_NAMESPACE(cl)
oclocInvoke
SIGNATURE oclocInvoke(uint32_t NumArgs, const char *Argv[], uint32_t NumSources, const uint8_t **DataSources, const uint64_t *LenSources, const char **NameSources, uint32_t NumInputHeaders, const uint8_t **DataInputHeaders, const uint64_t *LenInputHeaders, const char **NameInputHeaders, uint32_t *NumOutputs, uint8_t ***DataOutputs, uint64_t **LenOutputs, char ***NameOutputs)
Invokes ocloc API using C interface.
__SYCL2020_DEPRECATED
#define __SYCL2020_DEPRECATED(message)
Definition: defines_elementary.hpp:56
OCLOC_VERSION_CURRENT
@ OCLOC_VERSION_CURRENT
latest known version
Definition: ocloc_api.h:22
cl::sycl::ext::intel::experimental::device_arch
Definition: online_compiler.hpp:30
cl::sycl::info::device_type
device_type
Definition: info_desc.hpp:180
os_util.hpp
online_compiler.hpp
cl::sycl::detail::pi::getOsLibraryFuncAddress
void * getOsLibraryFuncAddress(void *Library, const std::string &FunctionName)
Definition: posix_pi.cpp:34
sycl
Definition: invoke_simd.hpp:68
cl::sycl::ext::intel::experimental::detail::prepareOclocArgs
static std::vector< const char * > prepareOclocArgs(sycl::info::device_type DeviceType, device_arch DeviceArch, bool Is64Bit, const std::string &DeviceStepping, const std::string &UserArgs)
Definition: online_compiler.cpp:25
pi.hpp
cl::sycl::compile
kernel_bundle< bundle_state::object > compile(const kernel_bundle< bundle_state::input > &InputBundle, const std::vector< device > &Devs, const property_list &PropList={})
Definition: kernel_bundle.hpp:632
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
cl::sycl::ext::intel::experimental::online_compile_error
Represents an error happend during online compilation.
Definition: online_compiler.hpp:63
oclocVersion
SIGNATURE oclocVersion()
Returns the current version of ocloc.
oclocFreeOutput
SIGNATURE oclocFreeOutput(uint32_t *NumOutputs, uint8_t ***DataOutputs, uint64_t **LenOutputs, char ***NameOutputs)
Frees results of oclocInvoke.
cl::sycl::ext::intel::experimental::detail::compileToSPIRV
static std::vector< byte > compileToSPIRV(const std::string &Source, sycl::info::device_type DeviceType, device_arch DeviceArch, bool Is64Bit, const std::string &DeviceStepping, void *&CompileToSPIRVHandle, void *&FreeSPIRVOutputsHandle, const std::vector< std::string > &UserArgs)
Compiles the given source Source to SPIR-V IL and returns IL as a vector of bytes.
Definition: online_compiler.cpp:79
ocloc_api.h
cl::sycl::instead
std::uint8_t instead
Definition: aliases.hpp:69
OCLOC_VERSION_1_0
@ OCLOC_VERSION_1_0
version 1.0
Definition: ocloc_api.h:21
cl::sycl::detail::pi::loadOsLibrary
void * loadOsLibrary(const std::string &Library)
Definition: posix_pi.cpp:20
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12