DPC++ Runtime
Runtime libraries for oneAPI DPC++
error_handling.cpp
Go to the documentation of this file.
1 //===------------------- enqueue_kernel.cpp ---------------------*- C++ -*-===//
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 // SYCL error handling of enqueue kernel operations
10 //
11 //===----------------------------------------------------------------------===//
12 
13 #include "error_handling.hpp"
14 
15 #include <detail/plugin.hpp>
16 #include <sycl/backend_types.hpp>
17 #include <sycl/detail/pi.hpp>
18 
19 namespace sycl {
20 inline namespace _V1 {
21 namespace detail::enqueue_kernel_launch {
22 
23 void handleInvalidWorkGroupSize(const device_impl &DeviceImpl, pi_kernel Kernel,
24  const NDRDescT &NDRDesc) {
25  sycl::platform Platform = DeviceImpl.get_platform();
26 
27  // Some of the error handling below is special for particular OpenCL
28  // versions. If this is an OpenCL backend, get the version.
29  bool IsOpenCL = false; // Backend is any OpenCL version
30  bool IsOpenCLV1x = false; // Backend is OpenCL 1.x
31  bool IsOpenCLVGE20 = false; // Backend is Greater or Equal to OpenCL 2.0
32  bool IsLevelZero = false; // Backend is any OneAPI Level 0 version
33  bool IsCuda = false; // Backend is CUDA
34  auto Backend = Platform.get_backend();
35  if (Backend == sycl::backend::opencl) {
36  std::string VersionString =
37  DeviceImpl.get_info<info::device::version>().substr(7, 3);
38  IsOpenCL = true;
39  IsOpenCLV1x = (VersionString.find("1.") == 0);
40  IsOpenCLVGE20 =
41  (VersionString.find("2.") == 0) || (VersionString.find("3.") == 0);
42  } else if (Backend == sycl::backend::ext_oneapi_level_zero) {
43  IsLevelZero = true;
44  } else if (Backend == sycl::backend::ext_oneapi_cuda) {
45  IsCuda = true;
46  }
47 
48  const PluginPtr &Plugin = DeviceImpl.getPlugin();
49  sycl::detail::pi::PiDevice Device = DeviceImpl.getHandleRef();
50 
51  size_t CompileWGSize[3] = {0};
52  Plugin->call<PiApiKind::piKernelGetGroupInfo>(
54  sizeof(size_t) * 3, CompileWGSize, nullptr);
55 
56  size_t MaxWGSize = 0;
57  Plugin->call<PiApiKind::piDeviceGetInfo>(Device,
59  sizeof(size_t), &MaxWGSize, nullptr);
60 
61  const bool HasLocalSize = (NDRDesc.LocalSize[0] != 0);
62 
63  if (CompileWGSize[0] != 0) {
64  if (CompileWGSize[0] > MaxWGSize || CompileWGSize[1] > MaxWGSize ||
65  CompileWGSize[2] > MaxWGSize)
66  throw sycl::exception(
68  "Submitting a kernel decorated with reqd_work_group_size attribute "
69  "to a device that does not support this work group size is invalid.");
70 
71  // OpenCL 1.x && 2.0:
72  // PI_ERROR_INVALID_WORK_GROUP_SIZE if local_work_size is NULL and the
73  // reqd_work_group_size attribute is used to declare the work-group size
74  // for kernel in the program source.
75  if (!HasLocalSize && (IsOpenCLV1x || IsOpenCLVGE20)) {
76  throw sycl::nd_range_error(
77  "OpenCL 1.x and 2.0 requires to pass local size argument even if "
78  "required work-group size was specified in the program source",
79  PI_ERROR_INVALID_WORK_GROUP_SIZE);
80  }
81  // PI_ERROR_INVALID_WORK_GROUP_SIZE if local_work_size is specified and does
82  // not match the required work-group size for kernel in the program source.
83  if (NDRDesc.LocalSize[0] != CompileWGSize[0] ||
84  NDRDesc.LocalSize[1] != CompileWGSize[1] ||
85  NDRDesc.LocalSize[2] != CompileWGSize[2])
86  throw sycl::nd_range_error(
87  "The specified local size {" + std::to_string(NDRDesc.LocalSize[2]) +
88  ", " + std::to_string(NDRDesc.LocalSize[1]) + ", " +
89  std::to_string(NDRDesc.LocalSize[0]) +
90  "} doesn't match the required work-group size specified "
91  "in the program source {" +
92  std::to_string(CompileWGSize[2]) + ", " +
93  std::to_string(CompileWGSize[1]) + ", " +
94  std::to_string(CompileWGSize[0]) + "}",
95  PI_ERROR_INVALID_WORK_GROUP_SIZE);
96  }
97 
98  if (HasLocalSize) {
99  size_t MaxThreadsPerBlock[3] = {};
100  Plugin->call<PiApiKind::piDeviceGetInfo>(
101  Device, PI_DEVICE_INFO_MAX_WORK_ITEM_SIZES, sizeof(MaxThreadsPerBlock),
102  MaxThreadsPerBlock, nullptr);
103 
104  for (size_t I = 0; I < 3; ++I) {
105  if (MaxThreadsPerBlock[I] < NDRDesc.LocalSize[I]) {
106  throw sycl::nd_range_error(
107  "The number of work-items in each dimension of a work-group cannot "
108  "exceed {" +
109  std::to_string(MaxThreadsPerBlock[0]) + ", " +
110  std::to_string(MaxThreadsPerBlock[1]) + ", " +
111  std::to_string(MaxThreadsPerBlock[2]) + "} for this device",
112  PI_ERROR_INVALID_WORK_GROUP_SIZE);
113  }
114  }
115  }
116 
117  if (IsOpenCLV1x) {
118  // OpenCL 1.x:
119  // PI_ERROR_INVALID_WORK_GROUP_SIZE if local_work_size is specified and
120  // the total number of work-items in the work-group computed as
121  // local_work_size[0] * ... * local_work_size[work_dim - 1] is greater
122  // than the value specified by PI_DEVICE_MAX_WORK_GROUP_SIZE in
123  // table 4.3
124  const size_t TotalNumberOfWIs =
125  NDRDesc.LocalSize[0] * NDRDesc.LocalSize[1] * NDRDesc.LocalSize[2];
126  if (TotalNumberOfWIs > MaxWGSize)
127  throw sycl::nd_range_error(
128  "Total number of work-items in a work-group cannot exceed " +
129  std::to_string(MaxWGSize),
130  PI_ERROR_INVALID_WORK_GROUP_SIZE);
131  } else if (IsOpenCLVGE20 || IsLevelZero) {
132  // OpenCL 2.x or OneAPI Level Zero:
133  // PI_ERROR_INVALID_WORK_GROUP_SIZE if local_work_size is specified and
134  // the total number of work-items in the work-group computed as
135  // local_work_size[0] * ... * local_work_size[work_dim - 1] is greater
136  // than the value specified by PI_KERNEL_GROUP_INFO_WORK_GROUP_SIZE in
137  // table 5.21.
138  size_t KernelWGSize = 0;
139  Plugin->call<PiApiKind::piKernelGetGroupInfo>(
140  Kernel, Device, PI_KERNEL_GROUP_INFO_WORK_GROUP_SIZE, sizeof(size_t),
141  &KernelWGSize, nullptr);
142  const size_t TotalNumberOfWIs =
143  NDRDesc.LocalSize[0] * NDRDesc.LocalSize[1] * NDRDesc.LocalSize[2];
144  if (TotalNumberOfWIs > KernelWGSize)
145  throw sycl::nd_range_error(
146  "Total number of work-items in a work-group cannot exceed " +
147  std::to_string(KernelWGSize) + " for this kernel",
148  PI_ERROR_INVALID_WORK_GROUP_SIZE);
149  } else {
150  // TODO: Should probably have something similar for the other backends
151  }
152 
153  if (HasLocalSize) {
154  // Is the global range size evenly divisible by the local workgroup size?
155  const bool NonUniformWGs =
156  (NDRDesc.LocalSize[0] != 0 &&
157  NDRDesc.GlobalSize[0] % NDRDesc.LocalSize[0] != 0) ||
158  (NDRDesc.LocalSize[1] != 0 &&
159  NDRDesc.GlobalSize[1] % NDRDesc.LocalSize[1] != 0) ||
160  (NDRDesc.LocalSize[2] != 0 &&
161  NDRDesc.GlobalSize[2] % NDRDesc.LocalSize[2] != 0);
162  // Is the local size of the workgroup greater than the global range size in
163  // any dimension?
164  if (IsOpenCL) {
165  const bool LocalExceedsGlobal =
166  NonUniformWGs && (NDRDesc.LocalSize[0] > NDRDesc.GlobalSize[0] ||
167  NDRDesc.LocalSize[1] > NDRDesc.GlobalSize[1] ||
168  NDRDesc.LocalSize[2] > NDRDesc.GlobalSize[2]);
169 
170  if (NonUniformWGs) {
171  if (IsOpenCLV1x) {
172  // OpenCL 1.x:
173  // PI_ERROR_INVALID_WORK_GROUP_SIZE if local_work_size is specified
174  // and number of workitems specified by global_work_size is not evenly
175  // divisible by size of work-group given by local_work_size
176  if (LocalExceedsGlobal)
177  throw sycl::nd_range_error("Local workgroup size cannot be greater "
178  "than global range in any dimension",
179  PI_ERROR_INVALID_WORK_GROUP_SIZE);
180  else
181  throw sycl::nd_range_error(
182  "Global_work_size must be evenly divisible by local_work_size. "
183  "Non-uniform work-groups are not supported by the target "
184  "device",
185  PI_ERROR_INVALID_WORK_GROUP_SIZE);
186  } else {
187  // OpenCL 2.x:
188  // PI_ERROR_INVALID_WORK_GROUP_SIZE if the program was compiled with
189  // –cl-uniform-work-group-size and the number of work-items specified
190  // by global_work_size is not evenly divisible by size of work-group
191  // given by local_work_size
192 
193  pi_program Program = nullptr;
194  Plugin->call<PiApiKind::piKernelGetInfo>(
195  Kernel, PI_KERNEL_INFO_PROGRAM, sizeof(pi_program), &Program,
196  nullptr);
197  size_t OptsSize = 0;
198  Plugin->call<PiApiKind::piProgramGetBuildInfo>(
199  Program, Device, PI_PROGRAM_BUILD_INFO_OPTIONS, 0, nullptr,
200  &OptsSize);
201  std::string Opts(OptsSize, '\0');
202  Plugin->call<PiApiKind::piProgramGetBuildInfo>(
203  Program, Device, PI_PROGRAM_BUILD_INFO_OPTIONS, OptsSize,
204  &Opts.front(), nullptr);
205  const bool HasStd20 = Opts.find("-cl-std=CL2.0") != std::string::npos;
206  const bool RequiresUniformWGSize =
207  Opts.find("-cl-uniform-work-group-size") != std::string::npos;
208  std::string LocalWGSize = std::to_string(NDRDesc.LocalSize[0]) +
209  ", " +
210  std::to_string(NDRDesc.LocalSize[1]) +
211  ", " + std::to_string(NDRDesc.LocalSize[2]);
212  std::string GlobalWGSize =
213  std::to_string(NDRDesc.GlobalSize[0]) + ", " +
214  std::to_string(NDRDesc.GlobalSize[1]) + ", " +
215  std::to_string(NDRDesc.GlobalSize[2]);
216  std::string message =
217  LocalExceedsGlobal
218  ? "Local work-group size {" + LocalWGSize +
219  "} is greater than global range size {" + GlobalWGSize +
220  "}. "
221  : "Global work size {" + GlobalWGSize +
222  "} is not evenly divisible by local work-group size {" +
223  LocalWGSize + "}. ";
224  if (!HasStd20)
225  throw sycl::nd_range_error(
226  message.append(
227  "Non-uniform work-groups are not allowed by "
228  "default. Underlying "
229  "OpenCL 2.x implementation supports this feature "
230  "and to enable "
231  "it, build device program with -cl-std=CL2.0"),
232  PI_ERROR_INVALID_WORK_GROUP_SIZE);
233  else if (RequiresUniformWGSize)
234  throw sycl::nd_range_error(
235  message.append(
236  "Non-uniform work-groups are not allowed by when "
237  "-cl-uniform-work-group-size flag is used. Underlying "
238  "OpenCL 2.x implementation supports this feature, but it "
239  "is "
240  "being "
241  "disabled by -cl-uniform-work-group-size build flag"),
242  PI_ERROR_INVALID_WORK_GROUP_SIZE);
243  // else unknown. fallback (below)
244  }
245  }
246  } else if (IsCuda) {
247  // CUDA:
248  // PI_ERROR_INVALID_WORK_GROUP_SIZE is returned when the kernel registers
249  // required for the launch config exceeds the maximum number of registers
250  // per block (PI_EXT_CODEPLAY_DEVICE_INFO_MAX_REGISTERS_PER_WORK_GROUP).
251  // This is if local_work_size[0] * ... * local_work_size[work_dim - 1]
252  // multiplied by PI_KERNEL_GROUP_INFO_NUM_REGS is greater than the value
253  // of PI_KERNEL_MAX_NUM_REGISTERS_PER_BLOCK. See Table 15: Technical
254  // Specifications per Compute Capability, for limitations.
255  const size_t TotalNumberOfWIs =
256  NDRDesc.LocalSize[0] * NDRDesc.LocalSize[1] * NDRDesc.LocalSize[2];
257 
258  uint32_t NumRegisters = 0;
259  Plugin->call<PiApiKind::piKernelGetGroupInfo>(
260  Kernel, Device, PI_KERNEL_GROUP_INFO_NUM_REGS, sizeof(NumRegisters),
261  &NumRegisters, nullptr);
262 
263  uint32_t MaxRegistersPerBlock =
264  DeviceImpl.get_info<ext::codeplay::experimental::info::device::
265  max_registers_per_work_group>();
266 
267  const bool HasExceededAvailableRegisters =
268  TotalNumberOfWIs * NumRegisters > MaxRegistersPerBlock;
269 
270  if (HasExceededAvailableRegisters) {
271  std::string message(
272  "Exceeded the number of registers available on the hardware.\n");
273  throw sycl::nd_range_error(
274  // Additional information which can be helpful to the user.
275  message.append(
276  "\tThe number registers per work-group cannot exceed " +
277  std::to_string(MaxRegistersPerBlock) +
278  " for this kernel on this device.\n"
279  "\tThe kernel uses " +
280  std::to_string(NumRegisters) +
281  " registers per work-item for a total of " +
282  std::to_string(TotalNumberOfWIs) +
283  " work-items per work-group.\n"),
284  PI_ERROR_INVALID_WORK_GROUP_SIZE);
285  }
286  } else {
287  // TODO: Decide what checks (if any) we need for the other backends
288  }
289  throw sycl::nd_range_error(
290  "Non-uniform work-groups are not supported by the target device",
291  PI_ERROR_INVALID_WORK_GROUP_SIZE);
292  }
293  // TODO: required number of sub-groups, OpenCL 2.1:
294  // PI_ERROR_INVALID_WORK_GROUP_SIZE if local_work_size is specified and is not
295  // consistent with the required number of sub-groups for kernel in the
296  // program source.
297 
298  // Fallback
299  constexpr pi_result Error = PI_ERROR_INVALID_WORK_GROUP_SIZE;
300  throw runtime_error(
301  "PI backend failed. PI backend returns: " + codeToString(Error), Error);
302 }
303 
304 void handleInvalidWorkItemSize(const device_impl &DeviceImpl,
305  const NDRDescT &NDRDesc) {
306 
307  const PluginPtr &Plugin = DeviceImpl.getPlugin();
308  sycl::detail::pi::PiDevice Device = DeviceImpl.getHandleRef();
309 
310  size_t MaxWISize[] = {0, 0, 0};
311 
312  Plugin->call<PiApiKind::piDeviceGetInfo>(
313  Device, PI_DEVICE_INFO_MAX_WORK_ITEM_SIZES, sizeof(MaxWISize), &MaxWISize,
314  nullptr);
315  for (unsigned I = 0; I < NDRDesc.Dims; I++) {
316  if (NDRDesc.LocalSize[I] > MaxWISize[I])
317  throw sycl::nd_range_error(
318  "Number of work-items in a work-group exceed limit for dimension " +
319  std::to_string(I) + " : " + std::to_string(NDRDesc.LocalSize[I]) +
320  " > " + std::to_string(MaxWISize[I]),
321  PI_ERROR_INVALID_WORK_ITEM_SIZE);
322  }
323 }
324 
325 void handleInvalidValue(const device_impl &DeviceImpl,
326  const NDRDescT &NDRDesc) {
327  const PluginPtr &Plugin = DeviceImpl.getPlugin();
328  sycl::detail::pi::PiDevice Device = DeviceImpl.getHandleRef();
329 
330  size_t MaxNWGs[] = {0, 0, 0};
331  Plugin->call<PiApiKind::piDeviceGetInfo>(
332  Device, PI_EXT_ONEAPI_DEVICE_INFO_MAX_WORK_GROUPS_3D, sizeof(MaxNWGs),
333  &MaxNWGs, nullptr);
334  for (unsigned int I = 0; I < NDRDesc.Dims; I++) {
335  size_t NWgs = NDRDesc.GlobalSize[I] / NDRDesc.LocalSize[I];
336  if (NWgs > MaxNWGs[I])
337  throw sycl::nd_range_error(
338  "Number of work-groups exceed limit for dimension " +
339  std::to_string(I) + " : " + std::to_string(NWgs) + " > " +
340  std::to_string(MaxNWGs[I]),
341  PI_ERROR_INVALID_VALUE);
342  }
343 
344  // fallback
345  constexpr pi_result Error = PI_ERROR_INVALID_VALUE;
346  throw runtime_error(
347  "Native API failed. Native API returns: " + codeToString(Error), Error);
348 }
349 
350 void handleErrorOrWarning(pi_result Error, const device_impl &DeviceImpl,
351  pi_kernel Kernel, const NDRDescT &NDRDesc) {
352  assert(Error != PI_SUCCESS &&
353  "Success is expected to be handled on caller side");
354  switch (Error) {
355  case PI_ERROR_INVALID_WORK_GROUP_SIZE:
356  return handleInvalidWorkGroupSize(DeviceImpl, Kernel, NDRDesc);
357 
358  case PI_ERROR_INVALID_KERNEL_ARGS:
359  throw sycl::nd_range_error(
360  "The kernel argument values have not been specified "
361  " OR "
362  "a kernel argument declared to be a pointer to a type.",
363  PI_ERROR_INVALID_KERNEL_ARGS);
364 
365  case PI_ERROR_INVALID_WORK_ITEM_SIZE:
366  return handleInvalidWorkItemSize(DeviceImpl, NDRDesc);
367 
368  case PI_ERROR_IMAGE_FORMAT_NOT_SUPPORTED:
369  throw sycl::nd_range_error(
370  "image object is specified as an argument value"
371  " and the image format is not supported by device associated"
372  " with queue",
373  PI_ERROR_IMAGE_FORMAT_NOT_SUPPORTED);
374 
375  case PI_ERROR_MISALIGNED_SUB_BUFFER_OFFSET:
376  throw sycl::nd_range_error(
377  "a sub-buffer object is specified as the value for an argument "
378  " that is a buffer object and the offset specified "
379  "when the sub-buffer object is created is not aligned "
380  "to CL_DEVICE_MEM_BASE_ADDR_ALIGN value for device associated"
381  " with queue",
382  PI_ERROR_MISALIGNED_SUB_BUFFER_OFFSET);
383 
384  case PI_ERROR_MEM_OBJECT_ALLOCATION_FAILURE:
385  throw sycl::nd_range_error(
386  "failure to allocate memory for data store associated with image"
387  " or buffer objects specified as arguments to kernel",
388  PI_ERROR_MEM_OBJECT_ALLOCATION_FAILURE);
389 
390  case PI_ERROR_INVALID_IMAGE_SIZE:
391  throw sycl::nd_range_error(
392  "image object is specified as an argument value and the image "
393  "dimensions (image width, height, specified or compute row and/or "
394  "slice pitch) are not supported by device associated with queue",
395  PI_ERROR_INVALID_IMAGE_SIZE);
396 
397  case PI_ERROR_INVALID_VALUE:
398  return handleInvalidValue(DeviceImpl, NDRDesc);
399 
400  case PI_ERROR_PLUGIN_SPECIFIC_ERROR:
401  // checkPiResult does all the necessary handling for
402  // PI_ERROR_PLUGIN_SPECIFIC_ERROR, making sure an error is thrown or not,
403  // depending on whether PI_ERROR_PLUGIN_SPECIFIC_ERROR contains an error or
404  // a warning. It also ensures that the contents of the error message buffer
405  // (used only by PI_ERROR_PLUGIN_SPECIFIC_ERROR) get handled correctly.
406  return DeviceImpl.getPlugin()->checkPiResult(Error);
407 
408  // TODO: Handle other error codes
409 
410  default:
411  throw runtime_error(
412  "Native API failed. Native API returns: " + codeToString(Error), Error);
413  }
414 }
415 
416 } // namespace detail::enqueue_kernel_launch
417 
418 namespace detail::kernel_get_group_info {
420  const PluginPtr &Plugin) {
421  assert(Error != PI_SUCCESS &&
422  "Success is expected to be handled on caller side");
423  switch (Error) {
424  case PI_ERROR_INVALID_VALUE:
425  if (Descriptor == CL_KERNEL_GLOBAL_WORK_SIZE)
426  throw sycl::exception(
428  "info::kernel_device_specific::global_work_size descriptor may only "
429  "be used if the device type is device_type::custom or if the kernel "
430  "is a built-in kernel.");
431  break;
432  // TODO: Handle other error codes
433  default:
434  Plugin->checkPiResult(Error);
435  break;
436  }
437 }
438 } // namespace detail::kernel_get_group_info
439 
440 } // namespace _V1
441 } // namespace sycl
sycl::range< 3 > GlobalSize
Definition: cg_types.hpp:129
sycl::range< 3 > LocalSize
Definition: cg_types.hpp:130
Param::return_type get_info() const
Queries this SYCL device for information requested by the template parameter param.
platform get_platform() const
Get associated SYCL platform.
const PluginPtr & getPlugin() const
sycl::detail::pi::PiDevice & getHandleRef()
Get reference to PI device.
Definition: device_impl.hpp:66
Encapsulates a SYCL platform on which kernels may be executed.
Definition: platform.hpp:109
backend get_backend() const noexcept
Returns the backend associated with this platform.
Definition: platform.cpp:57
void handleInvalidValue(const device_impl &DeviceImpl, const NDRDescT &NDRDesc)
void handleInvalidWorkGroupSize(const device_impl &DeviceImpl, pi_kernel Kernel, const NDRDescT &NDRDesc)
void handleInvalidWorkItemSize(const device_impl &DeviceImpl, const NDRDescT &NDRDesc)
void handleErrorOrWarning(pi_result Error, const device_impl &DeviceImpl, pi_kernel Kernel, const NDRDescT &NDRDesc)
Analyzes error code and arguments of piEnqueueKernelLaunch to emit user-friendly exception describing...
void handleErrorOrWarning(pi_result Error, pi_kernel_group_info Descriptor, const PluginPtr &Plugin)
Analyzes error code of piKernelGetGroupInfo.
std::string codeToString(pi_int32 code)
Definition: common.hpp:153
std::shared_ptr< plugin > PluginPtr
Definition: pi.hpp:48
std::string string
Definition: handler.hpp:426
std::error_code make_error_code(sycl::errc E) noexcept
Constructs an error code using e and sycl_category()
Definition: exception.cpp:107
Definition: access.hpp:18
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:216
@ PI_DEVICE_INFO_MAX_WORK_GROUP_SIZE
Definition: pi.h:302
@ PI_DEVICE_INFO_MAX_WORK_ITEM_SIZES
Definition: pi.h:301
@ PI_EXT_ONEAPI_DEVICE_INFO_MAX_WORK_GROUPS_3D
Definition: pi.h:416
_pi_program * pi_program
Definition: pi.h:1133
_pi_kernel_group_info
Definition: pi.h:505
@ PI_KERNEL_GROUP_INFO_WORK_GROUP_SIZE
Definition: pi.h:507
@ PI_KERNEL_GROUP_INFO_COMPILE_WORK_GROUP_SIZE
Definition: pi.h:508
@ PI_KERNEL_GROUP_INFO_NUM_REGS
Definition: pi.h:513
pi_result piDeviceGetInfo(pi_device device, pi_device_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret)
Returns requested info for provided native device Return PI_DEVICE_INFO_EXTENSION_DEVICELIB_ASSERT fo...
Definition: pi_cuda.cpp:78
@ PI_KERNEL_INFO_PROGRAM
Definition: pi.h:501
pi_result piKernelGetInfo(pi_kernel kernel, pi_kernel_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret)
Definition: pi_cuda.cpp:366
pi_result piKernelGetGroupInfo(pi_kernel kernel, pi_device device, pi_kernel_group_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret)
Definition: pi_cuda.cpp:508
@ PI_PROGRAM_BUILD_INFO_OPTIONS
Definition: pi.h:242
C++ wrapper of extern "C" PI interfaces.