DPC++ Runtime
Runtime libraries for oneAPI Data Parallel C++
program_impl.cpp
Go to the documentation of this file.
1 //==----- program_impl.cpp --- SYCL program implementation -----------------==//
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 
11 #include <CL/sycl/detail/pi.h>
12 #include <CL/sycl/kernel.hpp>
14 #include <detail/config.hpp>
15 #include <detail/kernel_impl.hpp>
16 #include <detail/program_impl.hpp>
18 
19 #include <algorithm>
20 #include <fstream>
21 #include <list>
22 #include <memory>
23 #include <mutex>
24 
26 namespace sycl {
27 namespace detail {
28 
29 program_impl::program_impl(ContextImplPtr Context,
30  const property_list &PropList)
31  : program_impl(Context, Context->get_info<info::context::devices>(),
32  PropList) {}
33 
35  std::vector<device> DeviceList,
36  const property_list &PropList)
37  : MContext(Context), MDevices(DeviceList), MPropList(PropList) {
38  if (Context->getDevices().size() > 1) {
40  "multiple devices within a context are not supported with "
41  "sycl::program and sycl::kernel",
43  }
44 }
45 
47  std::vector<std::shared_ptr<program_impl>> ProgramList,
48  std::string LinkOptions, const property_list &PropList)
49  : MState(program_state::linked), MPropList(PropList),
50  MLinkOptions(LinkOptions), MBuildOptions(LinkOptions) {
51  // Verify arguments
52  if (ProgramList.empty()) {
53  throw runtime_error("Non-empty vector of programs expected",
55  }
56 
57  // Sort the programs to avoid deadlocks due to locking multiple mutexes &
58  // verify that all programs are unique.
59  std::sort(ProgramList.begin(), ProgramList.end());
60  auto It = std::unique(ProgramList.begin(), ProgramList.end());
61  if (It != ProgramList.end()) {
62  throw runtime_error("Attempting to link a program with itself",
64  }
65 
66  MContext = ProgramList[0]->MContext;
67  if (MContext->getDevices().size() > 1) {
69  "multiple devices within a context are not supported with "
70  "sycl::program and sycl::kernel",
72  }
73  MDevices = ProgramList[0]->MDevices;
74  std::vector<device> DevicesSorted;
75  if (!is_host()) {
76  DevicesSorted = sort_devices_by_cl_device_id(MDevices);
77  }
78  check_device_feature_support<info::device::is_linker_available>(MDevices);
79  std::list<std::lock_guard<std::mutex>> Locks;
80  for (const auto &Prg : ProgramList) {
81  Locks.emplace_back(Prg->MMutex);
82  Prg->throw_if_state_is_not(program_state::compiled);
83  if (Prg->MContext != MContext) {
84  throw invalid_object_error(
85  "Not all programs are associated with the same context",
87  }
88  if (!is_host()) {
89  std::vector<device> PrgDevicesSorted =
90  sort_devices_by_cl_device_id(Prg->MDevices);
91  if (PrgDevicesSorted != DevicesSorted) {
92  throw invalid_object_error(
93  "Not all programs are associated with the same devices",
95  }
96  }
97  }
98 
99  if (!is_host()) {
100  std::vector<RT::PiDevice> Devices(get_pi_devices());
101  std::vector<RT::PiProgram> Programs;
102  bool NonInterOpToLink = false;
103  for (const auto &Prg : ProgramList) {
104  if (!Prg->MLinkable && NonInterOpToLink)
105  continue;
106  NonInterOpToLink |= !Prg->MLinkable;
107  Programs.push_back(Prg->MProgram);
108  }
109  const detail::plugin &Plugin = getPlugin();
111  MContext->getHandleRef(), Devices.size(), Devices.data(),
112  LinkOptions.c_str(), Programs.size(), Programs.data(), nullptr, nullptr,
113  &MProgram);
115  }
116 }
117 
119  pi_native_handle InteropProgram)
120  : program_impl(Context, InteropProgram, nullptr) {
121  MIsInterop = true;
122 }
123 
125  pi_native_handle InteropProgram,
126  RT::PiProgram Program)
127  : MProgram(Program), MContext(Context), MLinkable(true) {
128  const detail::plugin &Plugin = getPlugin();
129  if (MProgram == nullptr) {
130  assert(InteropProgram &&
131  "No InteropProgram/PiProgram defined with piextProgramFromNative");
132  // Translate the raw program handle into PI program.
134  InteropProgram, MContext->getHandleRef(), false, &MProgram);
135  } else
136  Plugin.call<PiApiKind::piProgramRetain>(Program);
137 
138  // TODO handle the case when cl_program build is in progress
139  pi_uint32 NumDevices;
141  MProgram, PI_PROGRAM_INFO_NUM_DEVICES, sizeof(pi_uint32), &NumDevices,
142  nullptr);
143  std::vector<RT::PiDevice> PiDevices(NumDevices);
145  sizeof(RT::PiDevice) * NumDevices,
146  PiDevices.data(), nullptr);
147  std::vector<device> SyclContextDevices =
148  MContext->get_info<info::context::devices>();
149 
150  // Keep only the subset of the devices (associated with context) that
151  // were actually used to create the program.
152  // This is possible when clCreateProgramWithBinary is used.
153  auto NewEnd = std::remove_if(
154  SyclContextDevices.begin(), SyclContextDevices.end(),
155  [&PiDevices](const sycl::device &Dev) {
156  return PiDevices.end() ==
157  std::find(PiDevices.begin(), PiDevices.end(),
158  detail::getSyclObjImpl(Dev)->getHandleRef());
159  });
160  SyclContextDevices.erase(NewEnd, SyclContextDevices.end());
161  MDevices = SyclContextDevices;
162  RT::PiDevice Device = getSyclObjImpl(MDevices[0])->getHandleRef();
163  assert(!MDevices.empty() && "No device found for this program");
164  // TODO check build for each device instead
165  cl_program_binary_type BinaryType;
167  MProgram, Device, CL_PROGRAM_BINARY_TYPE, sizeof(cl_program_binary_type),
168  &BinaryType, nullptr);
169  if (BinaryType == CL_PROGRAM_BINARY_TYPE_NONE) {
170  throw invalid_object_error(
171  "The native program passed to the program constructor has to be either "
172  "compiled or linked",
174  }
175  size_t Size = 0;
177  MProgram, Device, CL_PROGRAM_BUILD_OPTIONS, 0, nullptr, &Size);
178  std::vector<char> OptionsVector(Size);
179  Plugin.call<PiApiKind::piProgramGetBuildInfo>(MProgram, Device,
180  CL_PROGRAM_BUILD_OPTIONS, Size,
181  OptionsVector.data(), nullptr);
182  std::string Options(OptionsVector.begin(), OptionsVector.end());
183  switch (BinaryType) {
184  case CL_PROGRAM_BINARY_TYPE_NONE:
185  assert(false);
186  break;
187  case CL_PROGRAM_BINARY_TYPE_COMPILED_OBJECT:
188  MState = program_state::compiled;
189  MCompileOptions = Options;
190  MBuildOptions = Options;
191  break;
192  case CL_PROGRAM_BINARY_TYPE_LIBRARY:
193  case CL_PROGRAM_BINARY_TYPE_EXECUTABLE:
194  MState = program_state::linked;
195  MLinkOptions = "";
196  MBuildOptions = Options;
197  }
198 }
199 
201  : program_impl(Context, reinterpret_cast<pi_native_handle>(nullptr),
202  ProgramManager::getInstance().getPiProgramFromPiKernel(
203  Kernel, Context)) {
204  MIsInterop = true;
205 }
206 
208  // TODO catch an exception and put it to list of asynchronous exceptions
209  if (!is_host() && MProgram != nullptr) {
210  const detail::plugin &Plugin = getPlugin();
211  Plugin.call<PiApiKind::piProgramRelease>(MProgram);
212  }
213 }
214 
215 cl_program program_impl::get() const {
216  throw_if_state_is(program_state::none);
217  if (is_host()) {
218  throw invalid_object_error(
219  "This instance of program doesn't support OpenCL interoperability.",
221  }
223  return pi::cast<cl_program>(MProgram);
224 }
225 
226 void program_impl::compile_with_kernel_name(std::string KernelName,
227  std::string CompileOptions,
228  OSModuleHandle M) {
229  std::lock_guard<std::mutex> Lock(MMutex);
230  throw_if_state_is_not(program_state::none);
231  MProgramModuleHandle = M;
232  if (!is_host()) {
233  create_pi_program_with_kernel_name(
234  M, KernelName,
235  /*JITCompilationIsRequired=*/(!CompileOptions.empty()));
236  compile(CompileOptions);
237  }
238  MState = program_state::compiled;
239 }
240 
241 void program_impl::compile_with_source(std::string KernelSource,
242  std::string CompileOptions) {
243  std::lock_guard<std::mutex> Lock(MMutex);
244  throw_if_state_is_not(program_state::none);
245  // TODO should it throw if it's host?
246  if (!is_host()) {
247  create_cl_program_with_source(KernelSource);
248  compile(CompileOptions);
249  }
250  MState = program_state::compiled;
251  MIsInterop = true;
252 }
253 
254 void program_impl::build_with_kernel_name(std::string KernelName,
255  std::string BuildOptions,
256  OSModuleHandle Module) {
257  std::lock_guard<std::mutex> Lock(MMutex);
258  throw_if_state_is_not(program_state::none);
259  MProgramModuleHandle = Module;
260  if (!is_host()) {
261  MProgramAndKernelCachingAllowed = true;
262  MBuildOptions = BuildOptions;
265  detail::getSyclObjImpl(get_devices()[0]), KernelName, this,
266  /*JITCompilationIsRequired=*/(!BuildOptions.empty()));
267  const detail::plugin &Plugin = getPlugin();
268  Plugin.call<PiApiKind::piProgramRetain>(MProgram);
269  }
270  MState = program_state::linked;
271 }
272 
273 void program_impl::build_with_source(std::string KernelSource,
274  std::string BuildOptions) {
275  std::lock_guard<std::mutex> Lock(MMutex);
276  throw_if_state_is_not(program_state::none);
277  // TODO should it throw if it's host?
278  if (!is_host()) {
279  create_cl_program_with_source(KernelSource);
280  build(BuildOptions);
281  }
282  MState = program_state::linked;
283  MIsInterop = true;
284 }
285 
286 void program_impl::link(std::string LinkOptions) {
287  std::lock_guard<std::mutex> Lock(MMutex);
288  throw_if_state_is_not(program_state::compiled);
289  if (!is_host()) {
290  check_device_feature_support<info::device::is_linker_available>(MDevices);
291  std::vector<RT::PiDevice> Devices(get_pi_devices());
292  const detail::plugin &Plugin = getPlugin();
293  const char *LinkOpts = SYCLConfig<SYCL_PROGRAM_LINK_OPTIONS>::get();
294  if (!LinkOpts) {
295  LinkOpts = LinkOptions.c_str();
296  }
298  MContext->getHandleRef(), Devices.size(), Devices.data(), LinkOpts,
299  /*num_input_programs*/ 1, &MProgram, nullptr, nullptr, &MProgram);
301  MLinkOptions = LinkOptions;
302  MBuildOptions = LinkOptions;
303  }
304  MState = program_state::linked;
305 }
306 
307 bool program_impl::has_kernel(std::string KernelName,
308  bool IsCreatedFromSource) const {
309  throw_if_state_is(program_state::none);
310  if (is_host()) {
311  return !IsCreatedFromSource;
312  }
313 
314  std::vector<RT::PiDevice> Devices(get_pi_devices());
315  pi_uint64 function_ptr;
316  const detail::plugin &Plugin = getPlugin();
317 
318  RT::PiResult Err = PI_SUCCESS;
319  for (RT::PiDevice Device : Devices) {
321  Device, MProgram, KernelName.c_str(), &function_ptr);
322  if (Err != PI_SUCCESS && Err != PI_FUNCTION_ADDRESS_IS_NOT_AVAILABLE &&
323  Err != PI_INVALID_KERNEL_NAME)
324  throw runtime_error(
325  "Error from piextGetDeviceFunctionPointer when called by program",
326  Err);
328  return true;
329  }
330 
331  return false;
332 }
333 
334 kernel program_impl::get_kernel(std::string KernelName,
335  std::shared_ptr<program_impl> PtrToSelf,
336  bool IsCreatedFromSource) const {
337  throw_if_state_is(program_state::none);
338  if (is_host()) {
339  if (IsCreatedFromSource)
340  throw invalid_object_error("This instance of program is a host instance",
342 
343  return createSyclObjFromImpl<kernel>(
344  std::make_shared<kernel_impl>(MContext, PtrToSelf));
345  }
346  return createSyclObjFromImpl<kernel>(
347  std::make_shared<kernel_impl>(get_pi_kernel(KernelName), MContext,
348  PtrToSelf, IsCreatedFromSource, nullptr));
349 }
350 
351 std::vector<std::vector<char>> program_impl::get_binaries() const {
352  throw_if_state_is(program_state::none);
353  if (is_host())
354  return {};
355 
356  std::vector<std::vector<char>> Result;
357  const detail::plugin &Plugin = getPlugin();
358  std::vector<size_t> BinarySizes(MDevices.size());
361  sizeof(size_t) * BinarySizes.size(), BinarySizes.data(), nullptr);
362 
363  std::vector<char *> Pointers;
364  for (size_t I = 0; I < BinarySizes.size(); ++I) {
365  Result.emplace_back(BinarySizes[I]);
366  Pointers.push_back(Result[I].data());
367  }
369  sizeof(char *) * Pointers.size(),
370  Pointers.data(), nullptr);
371  return Result;
372 }
373 
374 void program_impl::create_cl_program_with_source(const std::string &Source) {
375  assert(!MProgram && "This program already has an encapsulated cl_program");
376  const char *Src = Source.c_str();
377  size_t Size = Source.size();
378  const detail::plugin &Plugin = getPlugin();
379  RT::PiResult Err =
381  MContext->getHandleRef(), 1, &Src, &Size, &MProgram);
382 
383  if (Err == PI_INVALID_OPERATION) {
384  throw feature_not_supported(
385  "program::compile_with_source is not supported by the selected backend",
387  }
388 
389  if (Err != PI_SUCCESS) {
390  Plugin.reportPiError(Err, "create_cl_program_with_source()");
391  }
392 }
393 
394 void program_impl::compile(const std::string &Options) {
395  check_device_feature_support<info::device::is_compiler_available>(MDevices);
396  std::vector<RT::PiDevice> Devices(get_pi_devices());
397  const detail::plugin &Plugin = getPlugin();
398  const char *CompileOpts = SYCLConfig<SYCL_PROGRAM_COMPILE_OPTIONS>::get();
399  if (!CompileOpts) {
400  CompileOpts = Options.c_str();
401  }
402  RT::PiResult Err = Plugin.call_nocheck<PiApiKind::piProgramCompile>(
403  MProgram, Devices.size(), Devices.data(), CompileOpts, 0, nullptr,
404  nullptr, nullptr, nullptr);
405 
406  if (Err != PI_SUCCESS) {
407  throw compile_program_error(
408  "Program compilation error:\n" +
409  ProgramManager::getProgramBuildLog(MProgram, MContext),
410  Err);
411  }
412  MCompileOptions = Options;
413  MBuildOptions = Options;
414 }
415 
416 void program_impl::build(const std::string &Options) {
417  check_device_feature_support<info::device::is_compiler_available>(MDevices);
418  std::vector<RT::PiDevice> Devices(get_pi_devices());
419  const detail::plugin &Plugin = getPlugin();
421  RT::PiResult Err = Plugin.call_nocheck<PiApiKind::piProgramBuild>(
422  MProgram, Devices.size(), Devices.data(), Options.c_str(), nullptr,
423  nullptr);
424 
425  if (Err != PI_SUCCESS) {
426  throw compile_program_error(
427  "Program build error:\n" +
428  ProgramManager::getProgramBuildLog(MProgram, MContext),
429  Err);
430  }
431  MBuildOptions = Options;
432 }
433 
434 std::vector<RT::PiDevice> program_impl::get_pi_devices() const {
435  std::vector<RT::PiDevice> PiDevices;
436  for (const auto &Device : MDevices) {
437  PiDevices.push_back(getSyclObjImpl(Device)->getHandleRef());
438  }
439  return PiDevices;
440 }
441 
442 RT::PiKernel program_impl::get_pi_kernel(const std::string &KernelName) const {
443  RT::PiKernel Kernel = nullptr;
444 
445  if (is_cacheable()) {
446  std::tie(Kernel, std::ignore, std::ignore) =
448  MProgramModuleHandle, detail::getSyclObjImpl(get_context()),
449  detail::getSyclObjImpl(get_devices()[0]), KernelName, this);
451  } else {
452  const detail::plugin &Plugin = getPlugin();
453  RT::PiResult Err = Plugin.call_nocheck<PiApiKind::piKernelCreate>(
454  MProgram, KernelName.c_str(), &Kernel);
455  if (Err == PI_INVALID_KERNEL_NAME) {
456  throw invalid_object_error(
457  "This instance of program does not contain the kernel requested",
458  Err);
459  }
460  Plugin.checkPiResult(Err);
461 
462  // Some PI Plugins (like OpenCL) require this call to enable USM
463  // For others, PI will turn this into a NOP.
465  sizeof(pi_bool), &PI_TRUE);
466  }
467 
468  return Kernel;
469 }
470 
471 std::vector<device>
472 program_impl::sort_devices_by_cl_device_id(std::vector<device> Devices) {
473  std::sort(Devices.begin(), Devices.end(),
474  [](const device &id1, const device &id2) {
475  return (detail::getSyclObjImpl(id1)->getHandleRef() <
476  detail::getSyclObjImpl(id2)->getHandleRef());
477  });
478  return Devices;
479 }
480 
481 void program_impl::throw_if_state_is(program_state State) const {
482  if (MState == State) {
483  throw invalid_object_error("Invalid program state", PI_INVALID_PROGRAM);
484  }
485 }
486 
487 void program_impl::throw_if_state_is_not(program_state State) const {
488  if (MState != State) {
489  throw invalid_object_error("Invalid program state", PI_INVALID_PROGRAM);
490  }
491 }
492 
493 void program_impl::create_pi_program_with_kernel_name(
494  OSModuleHandle Module, const std::string &KernelName,
495  bool JITCompilationIsRequired) {
496  assert(!MProgram && "This program already has an encapsulated PI program");
497  ProgramManager &PM = ProgramManager::getInstance();
498  const device FirstDevice = get_devices()[0];
499  RTDeviceBinaryImage &Img = PM.getDeviceImage(
500  Module, KernelName, get_context(), FirstDevice, JITCompilationIsRequired);
501  MProgram = PM.createPIProgram(Img, get_context(), {FirstDevice});
502 }
503 
504 template <>
505 cl_uint program_impl::get_info<info::program::reference_count>() const {
506  if (is_host()) {
507  throw invalid_object_error("This instance of program is a host instance",
509  }
510  pi_uint32 Result;
511  const detail::plugin &Plugin = getPlugin();
512  Plugin.call<PiApiKind::piProgramGetInfo>(MProgram,
514  sizeof(pi_uint32), &Result, nullptr);
515  return Result;
516 }
517 
518 template <> context program_impl::get_info<info::program::context>() const {
519  return get_context();
520 }
521 
522 template <>
523 std::vector<device> program_impl::get_info<info::program::devices>() const {
524  return get_devices();
525 }
526 
527 void program_impl::set_spec_constant_impl(const char *Name, const void *ValAddr,
528  size_t ValSize) {
529  if (MState != program_state::none)
531  "Invalid program state", PI_INVALID_PROGRAM);
532  // Reuse cached programs lock as opposed to introducing a new lock.
533  auto LockGuard = MContext->getKernelProgramCache().acquireCachedPrograms();
534  spec_constant_impl &SC = SpecConstRegistry[Name];
535  SC.set(ValSize, ValAddr);
536 }
537 
539  RT::PiProgram NativePrg) const {
540  // iterate via all specialization constants the program's image depends on,
541  // and set each to current runtime value (if any)
545 
546  auto LockGuard = Ctx->getKernelProgramCache().acquireCachedPrograms();
547  NativePrg = NativePrg ? NativePrg : getHandleRef();
548 
549  for (SCItTy SCIt : SCRange) {
550  auto SCEntry = SpecConstRegistry.find((*SCIt)->Name);
551  if (SCEntry == SpecConstRegistry.end())
552  // spec constant has not been set in user code - SPIR-V will use default
553  continue;
554  const spec_constant_impl &SC = SCEntry->second;
555  assert(SC.isSet() && "uninitialized spec constant");
556  pi::ByteArray Descriptors = pi::DeviceBinaryProperty(*SCIt).asByteArray();
557  // First 8 bytes are consumed by size of the property
558  assert(Descriptors.size() > 8 && "Unexpected property size");
559  // Expected layout is vector of 3-component tuples (flattened into a vector
560  // of scalars), where each tuple consists of: ID of a scalar spec constant,
561  // (which might be a member of the composite); offset, which is used to
562  // calculate location of scalar member within the composite or zero for
563  // scalar spec constants; size of a spec constant
564  assert(((Descriptors.size() - 8) / sizeof(std::uint32_t)) % 3 == 0 &&
565  "unexpected layout of composite spec const descriptors");
566  auto *It = reinterpret_cast<const std::uint32_t *>(&Descriptors[8]);
567  auto *End = reinterpret_cast<const std::uint32_t *>(&Descriptors[0] +
568  Descriptors.size());
569  while (It != End) {
571  NativePrg, /* ID */ It[0], /* Size */ It[2],
572  SC.getValuePtr() + /* Offset */ It[1]);
573  It += 3;
574  }
575  }
576 }
577 
579  const auto &Plugin = getPlugin();
580  if (Plugin.getBackend() == backend::opencl)
581  Plugin.call<PiApiKind::piProgramRetain>(MProgram);
582  pi_native_handle Handle;
583  Plugin.call<PiApiKind::piextProgramGetNativeHandle>(MProgram, &Handle);
584  return Handle;
585 }
586 
587 } // namespace detail
588 } // namespace sycl
589 } // __SYCL_INLINE_NAMESPACE(cl)
piKernelCreate
pi_result piKernelCreate(pi_program program, const char *kernel_name, pi_kernel *ret_kernel)
Definition: pi_esimd_emulator.cpp:1035
cl::sycl::detail::spec_constant_impl
Definition: spec_constant_impl.hpp:24
cl::sycl::detail::program_impl::get_context
context get_context() const
Definition: program_impl.hpp:259
property_list.hpp
PI_SUCCESS
@ PI_SUCCESS
Definition: pi.h:82
cl::sycl::detail::pi::DeviceBinaryProperty::asByteArray
ByteArray asByteArray() const
Definition: pi.cpp:640
pi.h
cl::sycl::detail::ContextImplPtr
std::shared_ptr< detail::context_impl > ContextImplPtr
Definition: memory_manager.hpp:31
cl::sycl::detail::program_impl::build_with_kernel_name
void build_with_kernel_name(std::string KernelName, std::string BuildOptions, OSModuleHandle M)
Builds the SYCL kernel function into encapsulated raw program.
Definition: program_impl.cpp:254
cl::sycl::detail::program_impl::getPlugin
const plugin & getPlugin() const
Definition: program_impl.hpp:266
cl::sycl::detail::SYCLConfig::get
static const char * get()
Definition: config.hpp:109
pi_bool
pi_uint32 pi_bool
Definition: pi.h:70
cl::sycl::detail::program_impl::has_kernel
bool has_kernel(std::string KernelName, bool IsCreatedFromSource) const
Checks if kernel is available for this program.
Definition: program_impl.cpp:307
cl::sycl::info::device
device
Definition: info_desc.hpp:50
PI_INVALID_OPERATION
@ PI_INVALID_OPERATION
Definition: pi.h:84
cl::sycl::detail::RTDeviceBinaryImage
Definition: device_binary_image.hpp:20
config.hpp
piProgramLink
pi_result piProgramLink(pi_context context, pi_uint32 num_devices, const pi_device *device_list, const char *options, pi_uint32 num_input_programs, const pi_program *input_programs, void(*pfn_notify)(pi_program program, void *user_data), void *user_data, pi_program *ret_program)
Definition: pi_opencl.cpp:743
PI_PROGRAM_INFO_BINARY_SIZES
@ PI_PROGRAM_INFO_BINARY_SIZES
Definition: pi.h:317
cl::sycl::backend::opencl
@ opencl
cl::sycl::detail::getPlugin
static const plugin & getPlugin(backend Backend)
Definition: backend.cpp:32
piextProgramGetNativeHandle
pi_result piextProgramGetNativeHandle(pi_program program, pi_native_handle *nativeHandle)
Gets the native handle of a PI program object.
Definition: pi_esimd_emulator.cpp:1026
piProgramRetain
pi_result piProgramRetain(pi_program program)
Definition: pi_esimd_emulator.cpp:1022
cl::sycl::detail::ProgramManager::getProgramBuildLog
static std::string getProgramBuildLog(const RT::PiProgram &Program, const ContextImplPtr Context)
Definition: program_manager.cpp:612
cl::sycl::detail::program_impl::program_impl
program_impl()=delete
cl::sycl::detail::pi::PiDevice
::pi_device PiDevice
Definition: pi.hpp:102
_pi_result
_pi_result
Definition: pi.h:81
cl::sycl::detail::pi::DeviceBinaryImage::getSpecConstants
const PropertyRange & getSpecConstants() const
Gets the iterator range over specialization constants in this binary image.
Definition: pi.hpp:351
cl::sycl::detail::program_impl::link
void link(std::string LinkOptions="")
Links encapsulated raw program.
Definition: program_impl.cpp:286
cl::sycl::errc::feature_not_supported
@ feature_not_supported
cl::sycl::detail::program_impl
Definition: program_impl.hpp:37
piProgramCompile
pi_result piProgramCompile(pi_program program, pi_uint32 num_devices, const pi_device *device_list, const char *options, pi_uint32 num_input_headers, const pi_program *input_headers, const char **header_include_names, void(*pfn_notify)(pi_program program, void *user_data), void *user_data)
cl::sycl::detail::LockGuard
std::lock_guard< SpinLock > LockGuard
Definition: global_handler.cpp:30
PI_USM_INDIRECT_ACCESS
@ PI_USM_INDIRECT_ACCESS
indicates that the kernel might access data through USM ptrs
Definition: pi.h:1253
cl::sycl::property_list
Objects of the property_list class are containers for the SYCL properties.
Definition: property_list.hpp:26
cl::sycl::detail::plugin::call_nocheck
RT::PiResult call_nocheck(ArgsT... Args) const
Calls the PiApi, traces the call, and returns the result.
Definition: plugin.hpp:145
cl::sycl::detail::spec_constant_impl::isSet
bool isSet() const
Definition: spec_constant_impl.hpp:34
_pi_kernel
Implementation of a PI Kernel for CUDA.
Definition: pi_cuda.hpp:578
piProgramGetInfo
pi_result piProgramGetInfo(pi_program program, pi_program_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret)
Definition: pi_esimd_emulator.cpp:994
cl::sycl::detail::program_impl::is_cacheable
bool is_cacheable() const
Definition: program_impl.hpp:344
compile_program_error
cl::sycl::detail::program_impl::get
cl_program get() const
Returns a valid cl_program instance.
Definition: program_impl.cpp:215
piextProgramCreateWithNativeHandle
pi_result piextProgramCreateWithNativeHandle(pi_native_handle nativeHandle, pi_context context, bool pluginOwnsNativeHandle, pi_program *program)
Creates PI program object from a native handle.
Definition: pi_esimd_emulator.cpp:1030
piProgramRelease
pi_result piProgramRelease(pi_program program)
Definition: pi_esimd_emulator.cpp:1024
pi_uint32
uint32_t pi_uint32
Definition: pi.h:68
kernel.hpp
program_impl.hpp
cl::sycl::device
The SYCL device class encapsulates a single SYCL device on which kernels may be executed.
Definition: device.hpp:35
cl::sycl::detail::program_impl::flush_spec_constants
void flush_spec_constants(const RTDeviceBinaryImage &Img, RT::PiProgram NativePrg=nullptr) const
Takes current values of specialization constants and "injects" them into the underlying native progra...
Definition: program_impl.cpp:538
cl::sycl::detail::program_impl::getNative
pi_native_handle getNative() const
Returns the native plugin handle.
Definition: program_impl.cpp:578
piProgramGetBuildInfo
pi_result piProgramGetBuildInfo(pi_program program, pi_device device, cl_program_build_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret)
Definition: pi_esimd_emulator.cpp:1017
piclProgramCreateWithSource
pi_result piclProgramCreateWithSource(pi_context context, pi_uint32 count, const char **strings, const size_t *lengths, pi_program *ret_program)
Definition: pi_esimd_emulator.cpp:989
cl::sycl::detail::ProgramManager::getBuiltPIProgram
RT::PiProgram getBuiltPIProgram(OSModuleHandle M, const ContextImplPtr &ContextImpl, const DeviceImplPtr &DeviceImpl, const std::string &KernelName, const program_impl *Prg=nullptr, bool JITCompilationIsRequired=false)
Builds or retrieves from cache a program defining the kernel with given name.
Definition: program_manager.cpp:439
piKernelSetExecInfo
pi_result piKernelSetExecInfo(pi_kernel kernel, pi_kernel_exec_info value_name, size_t param_value_size, const void *param_value)
API to set attributes controlling kernel execution.
Definition: pi_esimd_emulator.cpp:1495
cl::sycl::detail::program_impl::get_devices
std::vector< device > get_devices() const
Definition: program_impl.hpp:272
cl::sycl::detail::program_impl::~program_impl
~program_impl()
Definition: program_impl.cpp:207
cl::sycl::detail::plugin::call
void call(ArgsT... Args) const
Calls the API, traces the call, checks the result.
Definition: plugin.hpp:187
piKernelRetain
pi_result piKernelRetain(pi_kernel kernel)
Definition: pi_esimd_emulator.cpp:1067
cl::sycl::detail::pi::DeviceBinaryProperty
Definition: pi.hpp:227
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
cl::sycl::detail::pi::DeviceBinaryImage::PropertyRange
Definition: pi.hpp:249
PI_PROGRAM_INFO_REFERENCE_COUNT
@ PI_PROGRAM_INFO_REFERENCE_COUNT
Definition: pi.h:312
cl::sycl::detail::program_impl::compile_with_source
void compile_with_source(std::string KernelSource, std::string CompileOptions="")
Compiles the OpenCL C kernel function defined by source string.
Definition: program_impl.cpp:241
pi_uint64
uint64_t pi_uint64
Definition: pi.h:69
cl::sycl::detail::tie
auto tie(Ts &... Args)
Definition: tuple.hpp:40
cl::sycl::detail::program_impl::build_with_source
void build_with_source(std::string KernelSource, std::string BuildOptions="")
Builds the OpenCL C kernel function defined by source code.
Definition: program_impl.cpp:273
cl::sycl::detail::ProgramManager::getInstance
static ProgramManager & getInstance()
Definition: program_manager.cpp:63
_pi_program
Implementation of PI Program on CUDA Module object.
Definition: pi_cuda.hpp:523
PI_INVALID_KERNEL_NAME
@ PI_INVALID_KERNEL_NAME
Definition: pi.h:83
cl::sycl::detail::plugin
The plugin class provides a unified interface to the underlying low-level runtimes for the device-agn...
Definition: plugin.hpp:89
cl::sycl::detail::program_impl::compile_with_kernel_name
void compile_with_kernel_name(std::string KernelName, std::string CompileOptions, OSModuleHandle Module)
Compiles the SYCL kernel function into the encapsulated raw program.
Definition: program_impl.cpp:226
cl::sycl::detail::spec_constant_impl::set
void set(size_t Size, const void *Val)
Definition: spec_constant_impl.cpp:23
pi_native_handle
uintptr_t pi_native_handle
Definition: pi.h:72
cl::sycl::detail::pi::ByteArray::size
std::size_t size() const
Definition: pi.hpp:217
cl::sycl::detail::ProgramManager
Definition: program_manager.hpp:70
cl::sycl::detail::program_impl::get_binaries
std::vector< std::vector< char > > get_binaries() const
Returns built program binaries.
Definition: program_impl.cpp:351
PI_FUNCTION_ADDRESS_IS_NOT_AVAILABLE
@ PI_FUNCTION_ADDRESS_IS_NOT_AVAILABLE
PI_FUNCTION_ADDRESS_IS_NOT_AVAILABLE indicates a fallback method determines the function exists but i...
Definition: pi.h:115
cl::sycl::detail::OSModuleHandle
intptr_t OSModuleHandle
Uniquely identifies an operating system module (executable or a dynamic library)
Definition: os_util.hpp:48
PI_INVALID_VALUE
@ PI_INVALID_VALUE
Definition: pi.h:87
cl::sycl::detail::getSyclObjImpl
decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject)
Definition: common.hpp:198
cl::sycl::detail::plugin::checkPiResult
void checkPiResult(RT::PiResult pi_result) const
Checks return value from PI calls.
Definition: plugin.hpp:115
kernel_desc.hpp
cl::sycl::detail::program_impl::get_kernel
kernel get_kernel(std::string KernelName, std::shared_ptr< program_impl > PtrToSelf, bool IsCreatedFromSource) const
Returns a SYCL kernel for the SYCL kernel function defined by kernel name.
Definition: program_impl.cpp:334
piextGetDeviceFunctionPointer
pi_result piextGetDeviceFunctionPointer(pi_device device, pi_program program, const char *function_name, pi_uint64 *function_pointer_ret)
Retrieves a device function pointer to a user-defined function.
Definition: pi_esimd_emulator.cpp:1404
cl::sycl::context
The context class represents a SYCL context on which kernel functions may be executed.
Definition: context.hpp:35
cl::sycl::detail::pi::ByteArray
Definition: pi.hpp:211
cl::sycl::cl_uint
std::uint32_t cl_uint
Definition: aliases.hpp:83
cl::sycl::detail::spec_constant_impl::getValuePtr
const char * getValuePtr() const
Definition: spec_constant_impl.hpp:33
common.hpp
cl::sycl::ext::oneapi::experimental::spec_const_error
Definition: spec_constant.hpp:30
kernel_impl.hpp
cl::sycl::detail::pi::DeviceBinaryImage::PropertyRange::ConstIterator
Definition: pi.hpp:253
PI_PROGRAM_INFO_DEVICES
@ PI_PROGRAM_INFO_DEVICES
Definition: pi.h:315
cl::sycl::detail::program_impl::is_host
bool is_host() const
Definition: program_impl.hpp:139
spec_constant_impl.hpp
PI_PROGRAM_INFO_NUM_DEVICES
@ PI_PROGRAM_INFO_NUM_DEVICES
Definition: pi.h:314
piextProgramSetSpecializationConstant
pi_result piextProgramSetSpecializationConstant(pi_program prog, pi_uint32 spec_id, size_t spec_size, const void *spec_value)
Sets a specialization constant to a specific value.
Definition: pi_esimd_emulator.cpp:1500
cl::sycl::detail::program_impl::getHandleRef
RT::PiProgram & getHandleRef()
Definition: program_impl.hpp:133
cl::sycl::detail::program_impl::set_spec_constant_impl
void set_spec_constant_impl(const char *Name, const void *ValAddr, size_t ValSize)
Definition: program_impl.cpp:527
cl::sycl::detail::ProgramManager::flushSpecConstants
void flushSpecConstants(const program_impl &Prg, pi::PiProgram NativePrg=nullptr, const RTDeviceBinaryImage *Img=nullptr)
Resolves given program to a device binary image and requests the program to flush constants the image...
Definition: program_manager.cpp:1200
cl::sycl::kernel
Provides an abstraction of a SYCL kernel.
Definition: kernel.hpp:67
cl::sycl::detail::plugin::reportPiError
void reportPiError(RT::PiResult pi_result, const char *context) const
Definition: plugin.hpp:124
PI_TRUE
const pi_bool PI_TRUE
Definition: pi.h:486
piProgramBuild
pi_result piProgramBuild(pi_program program, pi_uint32 num_devices, const pi_device *device_list, const char *options, void(*pfn_notify)(pi_program program, void *user_data), void *user_data)
cl::sycl::detail::ProgramManager::getOrCreateKernel
std::tuple< RT::PiKernel, std::mutex *, RT::PiProgram > getOrCreateKernel(OSModuleHandle M, const ContextImplPtr &ContextImpl, const DeviceImplPtr &DeviceImpl, const std::string &KernelName, const program_impl *Prg)
Definition: program_manager.cpp:535
PI_INVALID_PROGRAM
@ PI_INVALID_PROGRAM
Definition: pi.h:94
PI_PROGRAM_INFO_BINARIES
@ PI_PROGRAM_INFO_BINARIES
Definition: pi.h:318
cl::sycl::info::context::devices
@ devices
_pi_device
PI device mapping to a CUdevice.
Definition: pi_cuda.hpp:71
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12