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