DPC++ Runtime
Runtime libraries for oneAPI DPC++
kernel_compiler_opencl.cpp
Go to the documentation of this file.
1 //==-- kernel_compiler_opencl.cpp OpenCL kernel compilation support -==//
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 
9 #include <sycl/detail/pi.hpp> // getOsLibraryFuncAddress
10 #include <sycl/exception.hpp> // make_error_code
11 
13 
14 #include "../online_compiler/ocloc_api.h"
15 #include "../split_string.hpp"
16 
17 #include <cstring> // strlen
18 #include <numeric> // for std::accumulate
19 #include <sstream>
20 
21 namespace sycl {
22 inline namespace _V1 {
23 namespace ext::oneapi::experimental {
24 namespace detail {
25 
26 // ensures the OclocLibrary has the right version, etc.
28  void *OclocVersionHandle =
30  // The initial versions of ocloc library did not have the oclocVersion()
31  // function. Those versions had the same API as the first version of ocloc
32  // library having that oclocVersion() function.
33  int LoadedVersion = ocloc_version_t::OCLOC_VERSION_1_0;
34  if (OclocVersionHandle) {
35  decltype(::oclocVersion) *OclocVersionFunc =
36  reinterpret_cast<decltype(::oclocVersion) *>(OclocVersionHandle);
37  LoadedVersion = OclocVersionFunc();
38  }
39  // The loaded library with version (A.B) is compatible with expected API/ABI
40  // version (X.Y) used here if A == B and B >= Y.
41  int LoadedVersionMajor = LoadedVersion >> 16;
42  int LoadedVersionMinor = LoadedVersion & 0xffff;
43  int CurrentVersionMajor = ocloc_version_t::OCLOC_VERSION_CURRENT >> 16;
44  int CurrentVersionMinor = ocloc_version_t::OCLOC_VERSION_CURRENT & 0xffff;
45  if (LoadedVersionMajor != CurrentVersionMajor ||
46  LoadedVersionMinor < CurrentVersionMinor) {
47  throw sycl::exception(
49  std::string("Found incompatible version of ocloc library: (") +
50  std::to_string(LoadedVersionMajor) + "." +
51  std::to_string(LoadedVersionMinor) +
52  "). The supported versions are (" +
53  std::to_string(CurrentVersionMajor) +
54  ".N), where (N >= " + std::to_string(CurrentVersionMinor) + ").");
55  }
56 }
57 
58 static void *OclocLibrary = nullptr;
59 
60 // load the ocloc shared library, check it.
62 #ifdef __SYCL_RT_OS_WINDOWS
63  static const std::string OclocLibraryName = "ocloc64.dll";
64 #else
65  static const std::string OclocLibraryName = "libocloc.so";
66 #endif
67  void *tempPtr = OclocLibrary;
68  if (tempPtr == nullptr) {
69  tempPtr = sycl::detail::pi::loadOsLibrary(OclocLibraryName);
70 
71  if (tempPtr == nullptr)
73  "Unable to load ocloc library " + OclocLibraryName);
74 
75  checkOclocLibrary(tempPtr);
76 
77  OclocLibrary = tempPtr;
78  }
79 
80  return OclocLibrary;
81 }
82 
84  // Already loaded?
85  if (OclocLibrary != nullptr)
86  return true;
87 
88  try {
89  // loads and checks version
91  return true;
92  } catch (...) {
93  return false;
94  }
95 }
96 
97 using voidPtr = void *;
98 
99 void SetupLibrary(voidPtr &oclocInvokeHandle, voidPtr &oclocFreeOutputHandle,
100  std::error_code the_errc) {
101  if (!oclocInvokeHandle) {
102  if (OclocLibrary == nullptr)
104 
105  oclocInvokeHandle =
107  if (!oclocInvokeHandle)
108  throw sycl::exception(the_errc, "Cannot load oclocInvoke() function");
109 
110  oclocFreeOutputHandle = sycl::detail::pi::getOsLibraryFuncAddress(
111  OclocLibrary, "oclocFreeOutput");
112  if (!oclocFreeOutputHandle)
113  throw sycl::exception(the_errc, "Cannot load oclocFreeOutput() function");
114  }
115 }
116 
117 std::string IPVersionsToString(const std::vector<uint32_t> IPVersionVec) {
118  std::stringstream ss;
119  bool amFirst = true;
120  for (uint32_t ipVersion : IPVersionVec) {
121  // if any device is not intelGPU, bail.
122  if (ipVersion < 0x02000000)
123  return "";
124 
125  if (!amFirst)
126  ss << ",";
127  amFirst = false;
128  ss << ipVersion;
129  }
130  return ss.str();
131 }
132 
133 spirv_vec_t OpenCLC_to_SPIRV(const std::string &Source,
134  const std::vector<uint32_t> &IPVersionVec,
135  const std::vector<std::string> &UserArgs,
136  std::string *LogPtr) {
137  std::vector<std::string> CMUserArgs = UserArgs;
138  CMUserArgs.push_back("-cmc");
139 
140  // handles into ocloc shared lib
141  static void *oclocInvokeHandle = nullptr;
142  static void *oclocFreeOutputHandle = nullptr;
144 
145  SetupLibrary(oclocInvokeHandle, oclocFreeOutputHandle, build_errc);
146 
147  // assemble ocloc args
148  std::string CombinedUserArgs =
149  std::accumulate(UserArgs.begin(), UserArgs.end(), std::string(""),
150  [](const std::string &acc, const std::string &s) {
151  return acc + s + " ";
152  });
153 
154  std::vector<const char *> Args = {"ocloc", "-q", "-spv_only", "-options",
155  CombinedUserArgs.c_str()};
156 
157  uint32_t NumOutputs = 0;
158  uint8_t **Outputs = nullptr;
159  uint64_t *OutputLengths = nullptr;
160  char **OutputNames = nullptr;
161 
162  const uint8_t *Sources[] = {
163  reinterpret_cast<const uint8_t *>(Source.c_str())};
164  const char *SourceName = "main.cl";
165  const uint64_t SourceLengths[] = {Source.length() + 1};
166 
167  Args.push_back("-file");
168  Args.push_back(SourceName);
169 
170  // device
171  std::string IPVersionsStr = IPVersionsToString(IPVersionVec);
172  if (!IPVersionsStr.empty()) {
173  Args.push_back("-device");
174  Args.push_back(IPVersionsStr.c_str());
175  }
176 
177  // invoke
178  decltype(::oclocInvoke) *OclocInvokeFunc =
179  reinterpret_cast<decltype(::oclocInvoke) *>(oclocInvokeHandle);
180  int CompileError =
181  OclocInvokeFunc(Args.size(), Args.data(), 1, Sources, SourceLengths,
182  &SourceName, 0, nullptr, nullptr, nullptr, &NumOutputs,
183  &Outputs, &OutputLengths, &OutputNames);
184 
185  // gather the results ( the SpirV and the Log)
186  spirv_vec_t SpirV;
187  std::string CompileLog;
188  for (uint32_t i = 0; i < NumOutputs; i++) {
189  size_t NameLen = strlen(OutputNames[i]);
190  if (NameLen >= 4 && strstr(OutputNames[i], ".spv") != nullptr &&
191  Outputs[i] != nullptr) {
192  assert(SpirV.size() == 0 && "More than one SPIR-V output found.");
193  SpirV = spirv_vec_t(Outputs[i], Outputs[i] + OutputLengths[i]);
194  } else if (!strcmp(OutputNames[i], "stdout.log")) {
195  if (OutputLengths[i] > 0) {
196  const char *LogText = reinterpret_cast<const char *>(Outputs[i]);
197  CompileLog.append(LogText, OutputLengths[i]);
198  if (LogPtr != nullptr)
199  LogPtr->append(LogText, OutputLengths[i]);
200  }
201  }
202  }
203 
204  // Try to free memory before reporting possible error.
205  decltype(::oclocFreeOutput) *OclocFreeOutputFunc =
206  reinterpret_cast<decltype(::oclocFreeOutput) *>(oclocFreeOutputHandle);
207  int MemFreeError =
208  OclocFreeOutputFunc(&NumOutputs, &Outputs, &OutputLengths, &OutputNames);
209 
210  if (CompileError)
211  throw sycl::exception(build_errc, "ocloc reported compilation errors: {\n" +
212  CompileLog + "\n}");
213 
214  if (SpirV.empty())
215  throw sycl::exception(build_errc,
216  "Unexpected output: ocloc did not return SPIR-V");
217 
218  if (MemFreeError)
219  throw sycl::exception(build_errc, "ocloc cannot safely free resources");
220 
221  return SpirV;
222 }
223 
224 std::string InvokeOclocQuery(uint32_t IPVersion, const char *identifier) {
225 
226  std::string QueryLog = "";
227 
228  // handles into ocloc shared lib
229  static void *oclocInvokeHandle = nullptr;
230  static void *oclocFreeOutputHandle = nullptr;
232 
233  SetupLibrary(oclocInvokeHandle, oclocFreeOutputHandle, the_errc);
234 
235  uint32_t NumOutputs = 0;
236  uint8_t **Outputs = nullptr;
237  uint64_t *OutputLengths = nullptr;
238  char **OutputNames = nullptr;
239 
240  std::vector<const char *> Args = {"ocloc", "query"};
241  std::vector<uint32_t> IPVersionVec{IPVersion};
242  std::string IPVersionsStr = IPVersionsToString(IPVersionVec);
243  if (!IPVersionsStr.empty()) {
244  Args.push_back("-device");
245  Args.push_back(IPVersionsStr.c_str());
246  }
247  Args.push_back(identifier);
248 
249  decltype(::oclocInvoke) *OclocInvokeFunc =
250  reinterpret_cast<decltype(::oclocInvoke) *>(oclocInvokeHandle);
251 
252  int InvokeError = OclocInvokeFunc(
253  Args.size(), Args.data(), 0, nullptr, 0, nullptr, 0, nullptr, nullptr,
254  nullptr, &NumOutputs, &Outputs, &OutputLengths, &OutputNames);
255 
256  // Gather the results.
257  for (uint32_t i = 0; i < NumOutputs; i++) {
258  if (!strcmp(OutputNames[i], "stdout.log")) {
259  if (OutputLengths[i] > 0) {
260  const char *LogText = reinterpret_cast<const char *>(Outputs[i]);
261  QueryLog.append(LogText, OutputLengths[i]);
262  }
263  }
264  }
265 
266  // Try to free memory before reporting possible error.
267  decltype(::oclocFreeOutput) *OclocFreeOutputFunc =
268  reinterpret_cast<decltype(::oclocFreeOutput) *>(oclocFreeOutputHandle);
269  int MemFreeError =
270  OclocFreeOutputFunc(&NumOutputs, &Outputs, &OutputLengths, &OutputNames);
271 
272  if (InvokeError)
273  throw sycl::exception(the_errc,
274  "ocloc reported errors: {\n" + QueryLog + "\n}");
275 
276  if (MemFreeError)
277  throw sycl::exception(the_errc, "ocloc cannot safely free resources");
278 
279  return QueryLog;
280 }
281 
282 bool OpenCLC_Feature_Available(const std::string &Feature, uint32_t IPVersion) {
283  static std::string FeatureLog = "";
284  if (FeatureLog.empty()) {
285  try {
286  FeatureLog = InvokeOclocQuery(IPVersion, "CL_DEVICE_OPENCL_C_FEATURES");
287  } catch (sycl::exception &) {
288  return false;
289  }
290  }
291 
292  // Allright, we have FeatureLog, so let's find that feature!
293  return (FeatureLog.find(Feature) != std::string::npos);
294 }
295 
297  const ext::oneapi::experimental::cl_version &Version, uint32_t IPVersion) {
298  static std::string VersionLog = "";
299  if (VersionLog.empty()) {
300  try {
301  VersionLog =
302  InvokeOclocQuery(IPVersion, "CL_DEVICE_OPENCL_C_ALL_VERSIONS");
303  } catch (sycl::exception &) {
304  return false;
305  }
306  }
307 
308  // Have VersionLog, will search.
309  // "OpenCL C":1.0.0 "OpenCL C":1.1.0 "OpenCL C":1.2.0 "OpenCL C":3.0.0
310  std::stringstream ss;
311  ss << Version.major << "." << Version.minor << "." << Version.patch;
312  return VersionLog.find(ss.str()) != std::string::npos;
313 }
314 
316  const std::string &Name, ext::oneapi::experimental::cl_version *VersionPtr,
317  uint32_t IPVersion) {
319  static std::string ExtensionByVersionLog = "";
320  if (ExtensionByVersionLog.empty()) {
321  try {
322  ExtensionByVersionLog =
323  InvokeOclocQuery(IPVersion, "CL_DEVICE_EXTENSIONS_WITH_VERSION");
324  } catch (sycl::exception &) {
325  return false;
326  }
327  }
328 
329  // ExtensionByVersionLog is ready. Time to find Name, and update VersionPtr.
330  // cl_khr_byte_addressable_store:1.0.0 cl_khr_device_uuid:1.0.0 ...
331  size_t where = ExtensionByVersionLog.find(Name);
332  if (where == std::string::npos) {
333  return false;
334  } // not there
335 
336  size_t colon = ExtensionByVersionLog.find(':', where);
337  if (colon == std::string::npos) {
338  throw sycl::exception(
339  rt_errc,
340  "trouble parsing query returned from CL_DEVICE_EXTENSIONS_WITH_VERSION "
341  "- extension not followed by colon (:)");
342  }
343  colon++; // move it forward
344 
345  size_t space = ExtensionByVersionLog.find(' ', colon); // could be npos
346 
347  size_t count = (space == std::string::npos) ? space : (space - colon);
348 
349  std::string versionStr = ExtensionByVersionLog.substr(colon, count);
350  std::vector<std::string> versionVec =
351  sycl::detail::split_string(versionStr, '.');
352  if (versionVec.size() != 3) {
353  throw sycl::exception(
354  rt_errc,
355  "trouble parsing query returned from "
356  "CL_DEVICE_EXTENSIONS_WITH_VERSION - version string unexpected: " +
357  versionStr);
358  }
359 
360  VersionPtr->major = std::stoi(versionVec[0]);
361  VersionPtr->minor = std::stoi(versionVec[1]);
362  VersionPtr->patch = std::stoi(versionVec[2]);
363 
364  return true;
365 }
366 
367 std::string OpenCLC_Profile(uint32_t IPVersion) {
368  try {
369  std::string result = InvokeOclocQuery(IPVersion, "CL_DEVICE_PROFILE");
370  // NOTE: result has \n\n amended. Clean it up.
371  // TODO: remove this once the ocloc query is fixed.
372  result.erase(std::remove_if(result.begin(), result.end(),
373  [](char c) {
374  return !std::isprint(c) || std::isspace(c);
375  }),
376  result.end());
377 
378  return result;
379  } catch (sycl::exception &) {
380  return "";
381  }
382 }
383 
384 } // namespace detail
385 } // namespace ext::oneapi::experimental
386 } // namespace _V1
387 } // namespace sycl
void * loadOsLibrary(const std::string &Library)
Definition: posix_pi.cpp:20
void * getOsLibraryFuncAddress(void *Library, const std::string &FunctionName)
Definition: posix_pi.cpp:47
std::vector< std::string > split_string(std::string_view str, char delimeter)
bool OpenCLC_Feature_Available(const std::string &Feature, uint32_t IPVersion)
std::string InvokeOclocQuery(uint32_t IPVersion, const char *identifier)
bool OpenCLC_Supports_Extension(const std::string &Name, ext::oneapi::experimental::cl_version *VersionPtr, uint32_t IPVersion)
spirv_vec_t OpenCLC_to_SPIRV(const std::string &Source, const std::vector< uint32_t > &IPVersionVec, const std::vector< std::string > &UserArgs, std::string *LogPtr)
std::string IPVersionsToString(const std::vector< uint32_t > IPVersionVec)
bool OpenCLC_Supports_Version(const ext::oneapi::experimental::cl_version &Version, uint32_t IPVersion)
void SetupLibrary(voidPtr &oclocInvokeHandle, voidPtr &oclocFreeOutputHandle, std::error_code the_errc)
std::error_code make_error_code(sycl::errc E) noexcept
Constructs an error code using e and sycl_category()
Definition: exception.cpp:64
Definition: access.hpp:18
error_code
Definition: defs.hpp:70
SIGNATURE oclocVersion()
Returns the current version of ocloc.
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.
@ OCLOC_VERSION_CURRENT
latest known version
Definition: ocloc_api.h:22
@ OCLOC_VERSION_1_0
version 1.0
Definition: ocloc_api.h:21
SIGNATURE oclocFreeOutput(uint32_t *NumOutputs, uint8_t ***DataOutputs, uint64_t **LenOutputs, char ***NameOutputs)
Frees results of oclocInvoke.
C++ wrapper of extern "C" PI interfaces.
where_expression< simd_mask< _Tp, _Abi >, simd< _Tp, _Abi > > where(const typename simd< _Tp, _Abi >::mask_type &, simd< _Tp, _Abi > &) noexcept