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>
16 #include <sycl/kernel.hpp>
17 #include <sycl/property_list.hpp>
18 
19 #include <algorithm>
20 #include <fstream>
21 #include <list>
22 #include <memory>
23 #include <mutex>
24 
25 namespace sycl {
26 inline namespace _V1 {
27 namespace detail {
28 
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) {
39  throw sycl::exception(
40  sycl::errc::feature_not_supported,
41  "multiple devices within a context are not supported with "
42  "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",
54  PI_ERROR_INVALID_VALUE);
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",
63  PI_ERROR_INVALID_PROGRAM);
64  }
65 
66  MContext = ProgramList[0]->MContext;
67  if (MContext->getDevices().size() > 1) {
68  throw sycl::exception(
69  sycl::errc::feature_not_supported,
70  "multiple devices within a context are not supported with "
71  "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",
86  PI_ERROR_INVALID_PROGRAM);
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",
94  PI_ERROR_INVALID_PROGRAM);
95  }
96  }
97  }
98 
99  if (!is_host()) {
100  std::vector<sycl::detail::pi::PiDevice> Devices(get_pi_devices());
101  std::vector<sycl::detail::pi::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 PluginPtr &Plugin = getPlugin();
111  Plugin->call_nocheck<PiApiKind::piProgramLink>(
112  MContext->getHandleRef(), Devices.size(), Devices.data(),
113  LinkOptions.c_str(), Programs.size(), Programs.data(), nullptr,
114  nullptr, &MProgram);
115  Plugin->checkPiResult<compile_program_error>(Err);
116  }
117 }
118 
120  pi_native_handle InteropProgram)
121  : program_impl(Context, InteropProgram, nullptr) {
122  MIsInterop = true;
123 }
124 
126  pi_native_handle InteropProgram,
128  : MProgram(Program), MContext(Context), MLinkable(true) {
129  const PluginPtr &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;
141  Plugin->call<PiApiKind::piProgramGetInfo>(
142  MProgram, PI_PROGRAM_INFO_NUM_DEVICES, sizeof(pi_uint32), &NumDevices,
143  nullptr);
144  std::vector<sycl::detail::pi::PiDevice> PiDevices(NumDevices);
147  NumDevices,
148  PiDevices.data(), nullptr);
149 
150  std::vector<device> PlatformDevices =
151  MContext->getPlatformImpl()->get_devices();
152  // Keep only the subset of the devices (associated with context) that
153  // were actually used to create the program.
154  // This is possible when clCreateProgramWithBinary is used.
155  auto NewEnd = std::remove_if(
156  PlatformDevices.begin(), PlatformDevices.end(),
157  [&PiDevices](const sycl::device &Dev) {
158  return PiDevices.end() ==
159  std::find(PiDevices.begin(), PiDevices.end(),
160  detail::getSyclObjImpl(Dev)->getHandleRef());
161  });
162  PlatformDevices.erase(NewEnd, PlatformDevices.end());
163  MDevices = PlatformDevices;
164  assert(!MDevices.empty() && "No device found for this program");
165  sycl::detail::pi::PiDevice Device = PiDevices[0];
166  // TODO check build for each device instead
167  cl_program_binary_type BinaryType = PI_PROGRAM_BINARY_TYPE_NONE;
168  Plugin->call<PiApiKind::piProgramGetBuildInfo>(
169  MProgram, Device, PI_PROGRAM_BUILD_INFO_BINARY_TYPE,
170  sizeof(cl_program_binary_type), &BinaryType, nullptr);
171  if (BinaryType == PI_PROGRAM_BINARY_TYPE_NONE) {
172  throw invalid_object_error(
173  "The native program passed to the program constructor has to be either "
174  "compiled or linked",
175  PI_ERROR_INVALID_PROGRAM);
176  }
177  size_t Size = 0;
178  Plugin->call<PiApiKind::piProgramGetBuildInfo>(
179  MProgram, Device, PI_PROGRAM_BUILD_INFO_OPTIONS, 0, nullptr, &Size);
180  std::vector<char> OptionsVector(Size);
181  Plugin->call<PiApiKind::piProgramGetBuildInfo>(
182  MProgram, Device, PI_PROGRAM_BUILD_INFO_OPTIONS, Size,
183  OptionsVector.data(), nullptr);
184  std::string Options(OptionsVector.begin(), OptionsVector.end());
185  switch (BinaryType) {
187  MState = program_state::compiled;
188  MCompileOptions = Options;
189  MBuildOptions = Options;
190  return;
193  MState = program_state::linked;
194  MLinkOptions = "";
195  MBuildOptions = Options;
196  return;
197  }
198  assert(false && "BinaryType is invalid.");
199 }
200 
203  : program_impl(Context, reinterpret_cast<pi_native_handle>(nullptr),
204  ProgramManager::getInstance().getPiProgramFromPiKernel(
205  Kernel, Context)) {
206  MIsInterop = true;
207 }
208 
210  // TODO catch an exception and put it to list of asynchronous exceptions
211  if (!is_host() && MProgram != nullptr) {
212  const PluginPtr &Plugin = getPlugin();
213  Plugin->call<PiApiKind::piProgramRelease>(MProgram);
214  }
215 }
216 
217 cl_program program_impl::get() const {
218  throw_if_state_is(program_state::none);
219  if (is_host()) {
220  throw invalid_object_error(
221  "This instance of program doesn't support OpenCL interoperability.",
222  PI_ERROR_INVALID_PROGRAM);
223  }
224  getPlugin()->call<PiApiKind::piProgramRetain>(MProgram);
225  return pi::cast<cl_program>(MProgram);
226 }
227 
228 void program_impl::compile_with_kernel_name(std::string KernelName,
229  std::string CompileOptions) {
230  std::lock_guard<std::mutex> Lock(MMutex);
231  throw_if_state_is_not(program_state::none);
232  if (!is_host()) {
233  create_pi_program_with_kernel_name(
234  KernelName,
235  /*JITCompilationIsRequired=*/(!CompileOptions.empty()));
236  compile(CompileOptions);
237  }
238  MState = program_state::compiled;
239 }
240 
241 void program_impl::link(std::string LinkOptions) {
242  std::lock_guard<std::mutex> Lock(MMutex);
243  throw_if_state_is_not(program_state::compiled);
244  if (!is_host()) {
245  check_device_feature_support<info::device::is_linker_available>(MDevices);
246  std::vector<sycl::detail::pi::PiDevice> Devices(get_pi_devices());
247  const PluginPtr &Plugin = getPlugin();
248  const char *LinkOpts = SYCLConfig<SYCL_PROGRAM_LINK_OPTIONS>::get();
249  if (!LinkOpts) {
250  LinkOpts = LinkOptions.c_str();
251  }
252 
253  // Plugin resets MProgram with a new pi_program as a result of the call to
254  // "piProgramLink". Thus, we need to release MProgram before the call to
255  // piProgramLink.
256  if (MProgram != nullptr)
257  Plugin->call<PiApiKind::piProgramRelease>(MProgram);
258 
260  Plugin->call_nocheck<PiApiKind::piProgramLink>(
261  MContext->getHandleRef(), Devices.size(), Devices.data(), LinkOpts,
262  /*num_input_programs*/ 1, &MProgram, nullptr, nullptr, &MProgram);
263  Plugin->checkPiResult<compile_program_error>(Err);
264  MLinkOptions = LinkOptions;
265  MBuildOptions = LinkOptions;
266  }
267  MState = program_state::linked;
268 }
269 
270 bool program_impl::has_kernel(std::string KernelName,
271  bool IsCreatedFromSource) const {
272  throw_if_state_is(program_state::none);
273  if (is_host()) {
274  return !IsCreatedFromSource;
275  }
276 
277  std::vector<sycl::detail::pi::PiDevice> Devices(get_pi_devices());
278  pi_uint64 function_ptr;
279  const PluginPtr &Plugin = getPlugin();
280 
281  sycl::detail::pi::PiResult Err = PI_SUCCESS;
282  for (sycl::detail::pi::PiDevice Device : Devices) {
283  Err = Plugin->call_nocheck<PiApiKind::piextGetDeviceFunctionPointer>(
284  Device, MProgram, KernelName.c_str(), &function_ptr);
285  if (Err != PI_SUCCESS &&
286  Err != PI_ERROR_FUNCTION_ADDRESS_IS_NOT_AVAILABLE &&
287  Err != PI_ERROR_INVALID_KERNEL_NAME)
288  throw runtime_error(
289  "Error from piextGetDeviceFunctionPointer when called by program",
290  Err);
291  if (Err == PI_SUCCESS || Err == PI_ERROR_FUNCTION_ADDRESS_IS_NOT_AVAILABLE)
292  return true;
293  }
294 
295  return false;
296 }
297 
298 kernel program_impl::get_kernel(std::string KernelName,
299  std::shared_ptr<program_impl> PtrToSelf,
300  bool IsCreatedFromSource) const {
301  throw_if_state_is(program_state::none);
302  if (is_host()) {
303  if (IsCreatedFromSource)
304  throw invalid_object_error("This instance of program is a host instance",
305  PI_ERROR_INVALID_PROGRAM);
306 
307  return createSyclObjFromImpl<kernel>(
308  std::make_shared<kernel_impl>(MContext, PtrToSelf));
309  }
310  auto [Kernel, ArgMask] = get_pi_kernel_arg_mask_pair(KernelName);
311  return createSyclObjFromImpl<kernel>(std::make_shared<kernel_impl>(
312  Kernel, MContext, PtrToSelf, IsCreatedFromSource, nullptr, ArgMask));
313 }
314 
315 std::vector<std::vector<char>> program_impl::get_binaries() const {
316  throw_if_state_is(program_state::none);
317  if (is_host())
318  return {};
319 
320  std::vector<std::vector<char>> Result;
321  const PluginPtr &Plugin = getPlugin();
322  std::vector<size_t> BinarySizes(MDevices.size());
323  Plugin->call<PiApiKind::piProgramGetInfo>(
325  sizeof(size_t) * BinarySizes.size(), BinarySizes.data(), nullptr);
326 
327  std::vector<char *> Pointers;
328  for (size_t I = 0; I < BinarySizes.size(); ++I) {
329  Result.emplace_back(BinarySizes[I]);
330  Pointers.push_back(Result[I].data());
331  }
333  sizeof(char *) * Pointers.size(),
334  Pointers.data(), nullptr);
335  return Result;
336 }
337 
338 void program_impl::compile(const std::string &Options) {
339  check_device_feature_support<info::device::is_compiler_available>(MDevices);
340  std::vector<sycl::detail::pi::PiDevice> Devices(get_pi_devices());
341  const PluginPtr &Plugin = getPlugin();
342  const char *CompileOpts = SYCLConfig<SYCL_PROGRAM_COMPILE_OPTIONS>::get();
343  if (!CompileOpts) {
344  CompileOpts = Options.c_str();
345  }
347  Plugin->call_nocheck<PiApiKind::piProgramCompile>(
348  MProgram, Devices.size(), Devices.data(), CompileOpts, 0, nullptr,
349  nullptr, nullptr, nullptr);
350 
351  if (Err != PI_SUCCESS) {
352  throw compile_program_error(
353  "Program compilation error:\n" +
354  ProgramManager::getProgramBuildLog(MProgram, MContext),
355  Err);
356  }
357  MCompileOptions = Options;
358  MBuildOptions = Options;
359 }
360 
361 void program_impl::build(const std::string &Options) {
362  check_device_feature_support<info::device::is_compiler_available>(MDevices);
363  std::vector<sycl::detail::pi::PiDevice> Devices(get_pi_devices());
364  const PluginPtr &Plugin = getPlugin();
367  Plugin->call_nocheck<PiApiKind::piProgramBuild>(
368  MProgram, Devices.size(), Devices.data(), Options.c_str(), nullptr,
369  nullptr);
370 
371  if (Err != PI_SUCCESS) {
372  throw compile_program_error(
373  "Program build error:\n" +
374  ProgramManager::getProgramBuildLog(MProgram, MContext),
375  Err);
376  }
377  MBuildOptions = Options;
378 }
379 
380 std::vector<sycl::detail::pi::PiDevice> program_impl::get_pi_devices() const {
381  std::vector<sycl::detail::pi::PiDevice> PiDevices;
382  for (const auto &Device : MDevices) {
383  PiDevices.push_back(getSyclObjImpl(Device)->getHandleRef());
384  }
385  return PiDevices;
386 }
387 
388 std::pair<sycl::detail::pi::PiKernel, const KernelArgMask *>
389 program_impl::get_pi_kernel_arg_mask_pair(const std::string &KernelName) const {
390  std::pair<sycl::detail::pi::PiKernel, const KernelArgMask *> Result;
391 
392  const PluginPtr &Plugin = getPlugin();
394  Plugin->call_nocheck<PiApiKind::piKernelCreate>(
395  MProgram, KernelName.c_str(), &Result.first);
396  if (Err == PI_ERROR_INVALID_KERNEL_NAME) {
397  throw invalid_object_error(
398  "This instance of program does not contain the kernel requested",
399  Err);
400  }
401  Plugin->checkPiResult(Err);
402 
403  // Some PI Plugins (like OpenCL) require this call to enable USM
404  // For others, PI will turn this into a NOP.
405  Plugin->call<PiApiKind::piKernelSetExecInfo>(
406  Result.first, PI_USM_INDIRECT_ACCESS, sizeof(pi_bool), &PI_TRUE);
407 
408  return Result;
409 }
410 
411 std::vector<device>
412 program_impl::sort_devices_by_cl_device_id(std::vector<device> Devices) {
413  std::sort(Devices.begin(), Devices.end(),
414  [](const device &id1, const device &id2) {
415  return (detail::getSyclObjImpl(id1)->getHandleRef() <
416  detail::getSyclObjImpl(id2)->getHandleRef());
417  });
418  return Devices;
419 }
420 
421 void program_impl::throw_if_state_is(program_state State) const {
422  if (MState == State) {
423  throw invalid_object_error("Invalid program state",
424  PI_ERROR_INVALID_PROGRAM);
425  }
426 }
427 
428 void program_impl::throw_if_state_is_not(program_state State) const {
429  if (MState != State) {
430  throw invalid_object_error("Invalid program state",
431  PI_ERROR_INVALID_PROGRAM);
432  }
433 }
434 
435 void program_impl::create_pi_program_with_kernel_name(
436  const std::string &KernelName, bool JITCompilationIsRequired) {
437  assert(!MProgram && "This program already has an encapsulated PI program");
438  ProgramManager &PM = ProgramManager::getInstance();
439  const device FirstDevice = get_devices()[0];
440  RTDeviceBinaryImage &Img = PM.getDeviceImage(
441  KernelName, get_context(), FirstDevice, JITCompilationIsRequired);
442  MProgram = PM.createPIProgram(Img, get_context(), {FirstDevice});
443 }
444 
446  const RTDeviceBinaryImage &Img,
447  sycl::detail::pi::PiProgram NativePrg) const {
448  // iterate via all specialization constants the program's image depends on,
449  // and set each to current runtime value (if any)
453 
454  auto LockGuard = Ctx->getKernelProgramCache().acquireCachedPrograms();
455  NativePrg = NativePrg ? NativePrg : getHandleRef();
456 
457  for (SCItTy SCIt : SCRange) {
458  auto SCEntry = SpecConstRegistry.find((*SCIt)->Name);
459  if (SCEntry == SpecConstRegistry.end())
460  // spec constant has not been set in user code - SPIR-V will use default
461  continue;
462  const spec_constant_impl &SC = SCEntry->second;
463  assert(SC.isSet() && "uninitialized spec constant");
464  ByteArray Descriptors = DeviceBinaryProperty(*SCIt).asByteArray();
465 
466  // First 8 bytes are consumed by the size of the property.
467  Descriptors.dropBytes(8);
468 
469  // Expected layout is vector of 3-component tuples (flattened into a
470  // vector of scalars), where each tuple consists of: ID of a scalar spec
471  // constant, (which might be a member of the composite); offset, which
472  // is used to calculate location of scalar member within the composite
473  // or zero for scalar spec constants; size of a spec constant.
474  while (!Descriptors.empty()) {
475  auto [Id, Offset, Size] =
476  Descriptors.consume<uint32_t, uint32_t, uint32_t>();
477 
478  Ctx->getPlugin()->call<PiApiKind::piextProgramSetSpecializationConstant>(
479  NativePrg, Id, Size, SC.getValuePtr() + Offset);
480  }
481  }
482 }
483 
485  const auto &Plugin = getPlugin();
486  if (getContextImplPtr()->getBackend() == backend::opencl)
487  Plugin->call<PiApiKind::piProgramRetain>(MProgram);
488  pi_native_handle Handle;
489  Plugin->call<PiApiKind::piextProgramGetNativeHandle>(MProgram, &Handle);
490  return Handle;
491 }
492 
493 } // namespace detail
494 } // namespace _V1
495 } // namespace sycl
The context class represents a SYCL context on which kernel functions may be executed.
Definition: context.hpp:51
void dropBytes(std::size_t Bytes)
static ProgramManager & getInstance()
static std::string getProgramBuildLog(const sycl::detail::pi::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...
const PropertyRange & getSpecConstants() const
Gets the iterator range over specialization constants in this binary image.
static const char * get()
Definition: config.hpp:115
void link(std::string LinkOptions="")
Links encapsulated raw program.
void compile_with_kernel_name(std::string KernelName, std::string CompileOptions)
Compiles the SYCL kernel function into the encapsulated raw program.
ContextImplPtr getContextImplPtr() const
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 flush_spec_constants(const RTDeviceBinaryImage &Img, sycl::detail::pi::PiProgram NativePrg=nullptr) const
Takes current values of specialization constants and "injects" them into the underlying native progra...
const PluginPtr & 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.
sycl::detail::pi::PiProgram & getHandleRef()
The SYCL device class encapsulates a single SYCL device on which kernels may be executed.
Definition: device.hpp:59
Provides an abstraction of a SYCL kernel.
Definition: kernel.hpp:74
Objects of the property_list class are containers for the SYCL properties.
::pi_device PiDevice
Definition: pi.hpp:131
std::shared_ptr< sycl::detail::context_impl > ContextImplPtr
Definition: event_impl.hpp:33
decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject)
Definition: impl_utils.hpp:30
std::shared_ptr< plugin > PluginPtr
Definition: pi.hpp:48
std::lock_guard< SpinLock > LockGuard
Definition: access.hpp:18
pi_result piKernelCreate(pi_program program, const char *kernel_name, pi_kernel *ret_kernel)
Definition: pi_cuda.cpp:341
uintptr_t pi_native_handle
Definition: pi.h:206
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)
Definition: pi_cuda.cpp:310
_pi_result
Definition: pi.h:213
pi_uint32 pi_bool
Definition: pi.h:204
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_cuda.cpp:333
pi_result piProgramRetain(pi_program program)
Definition: pi_cuda.cpp:320
const pi_bool PI_TRUE
Definition: pi.h:664
@ PI_USM_INDIRECT_ACCESS
indicates that the kernel might access data through USM ptrs
Definition: pi.h:1572
pi_result piProgramRelease(pi_program program)
Definition: pi_cuda.cpp:324
uint64_t pi_uint64
Definition: pi.h:203
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_cuda.cpp:1051
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_cuda.cpp:1044
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_cuda.cpp:272
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:202
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)
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 piextProgramGetNativeHandle(pi_program program, pi_native_handle *nativeHandle)
Gets the native handle of a PI program object.
Definition: pi_cuda.cpp:328
@ PI_PROGRAM_INFO_NUM_DEVICES
Definition: pi.h:451
@ PI_PROGRAM_INFO_BINARY_SIZES
Definition: pi.h:454
@ PI_PROGRAM_INFO_BINARIES
Definition: pi.h:455
@ PI_PROGRAM_INFO_DEVICES
Definition: pi.h:452
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_cuda.cpp:844
@ PI_PROGRAM_BINARY_TYPE_LIBRARY
Definition: pi.h:254
@ PI_PROGRAM_BINARY_TYPE_EXECUTABLE
Definition: pi.h:255
@ PI_PROGRAM_BINARY_TYPE_COMPILED_OBJECT
Definition: pi.h:253
@ PI_PROGRAM_BINARY_TYPE_NONE
Definition: pi.h:252
@ PI_PROGRAM_BUILD_INFO_OPTIONS
Definition: pi.h:239
@ PI_PROGRAM_BUILD_INFO_BINARY_TYPE
Definition: pi.h:241