DPC++ Runtime
Runtime libraries for oneAPI DPC++
pi_esimd_emulator.cpp
Go to the documentation of this file.
1 //===---------- pi_esimd_emulator.cpp - CM Emulation Plugin ---------------===//
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 
14 
15 #include <stdint.h>
16 
25 #include <CL/sycl/group.hpp>
26 #include <CL/sycl/id.hpp>
27 #include <CL/sycl/kernel.hpp>
28 #include <CL/sycl/nd_item.hpp>
29 #include <CL/sycl/range.hpp>
30 #include <sycl/ext/intel/esimd/common.hpp> // SLM_BTI
31 
32 #include <esimdemu_support.h>
33 
34 #include <cstdarg>
35 #include <cstdio>
36 #include <cstring>
37 #include <functional>
38 #include <map>
39 #include <memory>
40 #include <string>
41 #include <thread>
42 #include <utility>
43 
44 #include "pi_esimd_emulator.hpp"
45 
46 #define ARG_UNUSED(x) (void)x
47 
48 namespace {
49 
50 // Helper functions for unified 'Return' type declaration - imported
51 // from pi_level_zero.cpp
52 template <typename T, typename Assign>
53 pi_result getInfoImpl(size_t ParamValueSize, void *ParamValue,
54  size_t *ParamValueSizeRet, T Value, size_t ValueSize,
55  Assign &&AssignFunc) {
56  if (ParamValue != nullptr) {
57  if (ParamValueSize < ValueSize) {
58  return PI_INVALID_VALUE;
59  }
60  AssignFunc(ParamValue, Value, ValueSize);
61  }
62  if (ParamValueSizeRet != nullptr) {
63  *ParamValueSizeRet = ValueSize;
64  }
65  return PI_SUCCESS;
66 }
67 
68 template <typename T>
69 pi_result getInfo(size_t ParamValueSize, void *ParamValue,
70  size_t *ParamValueSizeRet, T Value) {
71  auto assignment = [](void *ParamValue, T Value, size_t ValueSize) {
72  ARG_UNUSED(ValueSize);
73  *static_cast<T *>(ParamValue) = Value;
74  };
75  return getInfoImpl(ParamValueSize, ParamValue, ParamValueSizeRet, Value,
76  sizeof(T), assignment);
77 }
78 
79 template <typename T>
80 pi_result getInfoArray(size_t ArrayLength, size_t ParamValueSize,
81  void *ParamValue, size_t *ParamValueSizeRet, T *Value) {
82  return getInfoImpl(ParamValueSize, ParamValue, ParamValueSizeRet, Value,
83  ArrayLength * sizeof(T), memcpy);
84 }
85 
86 template <>
87 pi_result getInfo<const char *>(size_t ParamValueSize, void *ParamValue,
88  size_t *ParamValueSizeRet, const char *Value) {
89  return getInfoArray(strlen(Value) + 1, ParamValueSize, ParamValue,
90  ParamValueSizeRet, Value);
91 }
92 
93 class ReturnHelper {
94 public:
95  ReturnHelper(size_t ArgParamValueSize, void *ArgParamValue,
96  size_t *ArgParamValueSizeRet)
97  : ParamValueSize(ArgParamValueSize), ParamValue(ArgParamValue),
98  ParamValueSizeRet(ArgParamValueSizeRet) {}
99 
100  template <class T> pi_result operator()(const T &t) {
101  return getInfo(ParamValueSize, ParamValue, ParamValueSizeRet, t);
102  }
103 
104 private:
105  size_t ParamValueSize;
106  void *ParamValue;
107  size_t *ParamValueSizeRet;
108 };
109 
110 } // anonymous namespace
111 
112 // Controls PI level tracing prints.
113 static bool PrintPiTrace = false;
114 
115 // Global variables used in PI_esimd_emulator
116 // Note we only create a simple pointer variables such that C++ RT won't
117 // deallocate them automatically at the end of the main program.
118 // The heap memory allocated for this global variable reclaimed only when
119 // Sycl RT calls piTearDown().
120 static sycl::detail::ESIMDEmuPluginOpaqueData *PiESimdDeviceAccess;
121 
122 // Single-entry cache for piPlatformsGet call.
124 // TODO/FIXME : Memory leak. Handle with 'piTearDown'.
125 static std::mutex *PiPlatformCacheLock = new std::mutex;
126 
127 // Mapping between surface index and CM-managed surface
128 static std::unordered_map<unsigned int, _pi_mem *> *PiESimdSurfaceMap =
129  new std::unordered_map<unsigned int, _pi_mem *>;
130 // TODO/FIXME : Memory leak. Handle with 'piTearDown'.
131 static std::mutex *PiESimdSurfaceMapLock = new std::mutex;
132 
133 // To be compared with ESIMD_EMULATOR_PLUGIN_OPAQUE_DATA_VERSION in device
134 // interface header file
135 #define ESIMDEmuPluginDataVersion 0
136 
137 // To be compared with ESIMD_DEVICE_INTERFACE_VERSION in device
138 // interface header file
139 #define ESIMDEmuPluginInterfaceVersion 1
140 
141 // For PI_DEVICE_INFO_DRIVER_VERSION info
142 static char ESimdEmuVersionString[32];
143 
144 // Global variables for PI_PLUGIN_SPECIFIC_ERROR
145 constexpr size_t MaxMessageSize = 256;
147 thread_local char ErrorMessage[MaxMessageSize];
148 
149 // Utility function for setting a message and warning
150 [[maybe_unused]] static void setErrorMessage(const char *message,
151  pi_result error_code) {
152  assert(strlen(message) <= MaxMessageSize);
153  strcpy(ErrorMessage, message);
154  ErrorMessageCode = error_code;
155 }
156 
157 // Returns plugin specific error and warning messages
159  *message = &ErrorMessage[0];
160  return ErrorMessageCode;
161 }
162 
163 using IDBuilder = sycl::detail::Builder;
164 
165 template <int NDims>
166 using KernelFunc = std::function<void(const sycl::nd_item<NDims> &)>;
167 
168 // Struct to wrap dimension info and lambda function to be invoked by
169 // CM Kernel launcher that only accepts raw function pointer for
170 // kernel execution. Function instances of 'InvokeKernel' un-wrap
171 // this struct instance and invoke lambda function ('Func')
172 template <int NDims> struct KernelInvocationContext {
174  const sycl::range<NDims> &LocalSize;
175  const sycl::range<NDims> &GlobalSize;
176  const sycl::id<NDims> &GlobalOffset;
177 };
178 
179 // A helper structure to create multi-dimensional range when
180 // dimensionality is given as a template parameter. `create` function
181 // in specializations accepts a template `Gen` function which
182 // generates range extent for a dimension given as an argument.
183 template <int NDims> struct RangeBuilder;
184 
185 template <> struct RangeBuilder<1> {
186  template <typename Gen> static sycl::range<1> create(Gen G) {
187  return sycl::range<1>{G(0)};
188  }
189 };
190 template <> struct RangeBuilder<2> {
191  template <typename Gen> static sycl::range<2> create(Gen G) {
192  return sycl::range<2>{G(0), G(1)};
193  }
194 };
195 template <> struct RangeBuilder<3> {
196  template <typename Gen> static sycl::range<3> create(Gen G) {
197  return sycl::range<3>{G(0), G(1), G(2)};
198  }
199 };
200 
201 // Function template to generate entry point of kernel execution as
202 // raw function pointer. CM kernel launcher executes one instance of
203 // this function per 'NDims'
204 template <int NDims> void InvokeKernel(KernelInvocationContext<NDims> *ctx) {
205 
206  sycl::range<NDims> GroupSize{
207  sycl::detail::InitializedVal<NDims, sycl::range>::template get<0>()};
208 
209  for (int i = 0; i < NDims; ++i) {
210  GroupSize[i] = ctx->GlobalSize[i] / ctx->LocalSize[i];
211  }
212 
213  const sycl::id<NDims> LocalID = RangeBuilder<NDims>::create(
214  [](int i) { return cm_support::get_thread_idx(i); });
215 
216  const sycl::id<NDims> GroupID = RangeBuilder<NDims>::create(
217  [](int i) { return cm_support::get_group_idx(i); });
218 
219  const sycl::group<NDims> Group = IDBuilder::createGroup<NDims>(
220  ctx->GlobalSize, ctx->LocalSize, GroupSize, GroupID);
221 
222  const sycl::id<NDims> GlobalID =
223  GroupID * ctx->LocalSize + LocalID + ctx->GlobalOffset;
224 
225  const sycl::item<NDims, /*Offset=*/true> GlobalItem =
226  IDBuilder::createItem<NDims, true>(ctx->GlobalSize, GlobalID,
227  ctx->GlobalOffset);
228 
229  const sycl::item<NDims, /*Offset=*/false> LocalItem =
230  IDBuilder::createItem<NDims, false>(ctx->LocalSize, LocalID);
231 
232  const sycl::nd_item<NDims> NDItem =
233  IDBuilder::createNDItem<NDims>(GlobalItem, LocalItem, Group);
234 
235  ctx->Func(NDItem);
236 }
237 
238 // Interface for lauching kernels using libcm from CM EMU project.
239 template <int DIMS> class libCMBatch {
240 private:
241  const KernelFunc<DIMS> &MKernel;
242  std::vector<uint32_t> GroupDim, SpaceDim;
243 
244 public:
246  : MKernel(Kernel), GroupDim{1, 1, 1}, SpaceDim{1, 1, 1} {}
247 
248  void runIterationSpace(const sycl::range<DIMS> &LocalSize,
249  const sycl::range<DIMS> &GlobalSize,
250  const sycl::id<DIMS> &GlobalOffset) {
251 
252  for (int I = 0; I < DIMS; I++) {
253  SpaceDim[I] = (uint32_t)LocalSize[I];
254  GroupDim[I] = (uint32_t)(GlobalSize[I] / LocalSize[I]);
255  }
256 
257  const auto InvokeKernelArg = KernelInvocationContext<DIMS>{
258  MKernel, LocalSize, GlobalSize, GlobalOffset};
259 
260  EsimdemuKernel{reinterpret_cast<fptrVoid>(InvokeKernel<DIMS>), GroupDim,
261  SpaceDim}
262  .launchMT(sizeof(InvokeKernelArg), &InvokeKernelArg);
263  }
264 };
265 
266 unsigned int sycl_get_cm_surface_index(void *PtrInput) {
267  _pi_mem *Surface = static_cast<_pi_mem *>(PtrInput);
268 
269  return Surface->SurfaceIndex;
270 }
271 
272 // Function to provide image info for kernel compilation using surface
273 // index without dependency on '_pi_image' definition
274 void sycl_get_cm_buffer_params(unsigned int IndexInput, char **BaseAddr,
275  uint32_t *Width, std::mutex **BufMtxLock) {
276  std::lock_guard<std::mutex> Lock{*PiESimdSurfaceMapLock};
277  auto MemIter = PiESimdSurfaceMap->find(IndexInput);
278 
279  assert(MemIter != PiESimdSurfaceMap->end() && "Invalid Surface Index");
280 
281  _pi_buffer *Buf = static_cast<_pi_buffer *>(MemIter->second);
282 
283  *BaseAddr = Buf->MapHostPtr;
284  *Width = static_cast<uint32_t>(Buf->Size);
285 
286  *BufMtxLock = &(Buf->SurfaceLock);
287 }
288 
289 // Function to provide image info for kernel compilation using surface
290 // index without dependency on '_pi_image' definition
291 void sycl_get_cm_image_params(unsigned int IndexInput, char **BaseAddr,
292  uint32_t *Width, uint32_t *Height, uint32_t *Bpp,
293  std::mutex **ImgMtxLock) {
294  std::lock_guard<std::mutex> Lock{*PiESimdSurfaceMapLock};
295  auto MemIter = PiESimdSurfaceMap->find(IndexInput);
296  assert(MemIter != PiESimdSurfaceMap->end() && "Invalid Surface Index");
297 
298  _pi_image *Img = static_cast<_pi_image *>(MemIter->second);
299 
300  *BaseAddr = Img->MapHostPtr;
301 
302  *Bpp = static_cast<uint32_t>(Img->BytesPerPixel);
303  *Width = static_cast<uint32_t>(Img->Width) * (*Bpp);
304  *Height = static_cast<uint32_t>(Img->Height);
305 
306  *ImgMtxLock = &(Img->SurfaceLock);
307 }
308 
311 sycl::detail::ESIMDDeviceInterface::ESIMDDeviceInterface() {
313  reserved = nullptr;
314 
315  /* From 'esimd_emulator_functions_v1.h' : Start */
316  cm_barrier_ptr = cm_support::barrier;
317  cm_sbarrier_ptr = cm_support::split_barrier;
318  cm_fence_ptr = cm_support::fence;
319 
320  sycl_get_surface_base_addr_ptr = cm_support::get_surface_base_addr;
321  __cm_emu_get_slm_ptr = cm_support::get_slm_base;
322  cm_slm_init_ptr = cm_support::init_slm;
323 
324  sycl_get_cm_surface_index_ptr = sycl_get_cm_surface_index;
325  sycl_get_cm_buffer_params_ptr = sycl_get_cm_buffer_params;
326  sycl_get_cm_image_params_ptr = sycl_get_cm_image_params;
327 
328  /* From 'esimd_emulator_functions_v1.h' : End */
329 }
330 
333 
334 static bool isNull(int NDims, const size_t *R) {
335  return ((0 == R[0]) && (NDims < 2 || 0 == R[1]) && (NDims < 3 || 0 == R[2]));
336 }
337 
338 // NDims is the number of dimensions in the ND-range. Kernels are
339 // normalized in the handler so that all kernels take an sycl::nd_item
340 // as argument (see StoreLambda in CL/sycl/handler.hpp). For kernels
341 // whose workgroup size (LocalWorkSize) is unspecified, InvokeImpl
342 // sets LocalWorkSize to {1, 1, 1}, i.e. each workgroup contains just
343 // one work item. CM emulator will run several workgroups in parallel
344 // depending on environment settings.
345 
346 template <int NDims> struct InvokeImpl {
347 
348  static sycl::range<NDims> get_range(const size_t *Array) {
349  if constexpr (NDims == 1)
350  return sycl::range<NDims>{Array[0]};
351  else if constexpr (NDims == 2)
352  return sycl::range<NDims>{Array[0], Array[1]};
353  else if constexpr (NDims == 3)
354  return sycl::range<NDims>{Array[0], Array[1], Array[2]};
355  }
356 
357  static void invoke(pi_kernel Kernel, const size_t *GlobalWorkOffset,
358  const size_t *GlobalWorkSize,
359  const size_t *LocalWorkSize) {
360  libCMBatch<NDims>{*reinterpret_cast<KernelFunc<NDims> *>(Kernel)}
361  .runIterationSpace(get_range(LocalWorkSize), get_range(GlobalWorkSize),
362  sycl::id<NDims>{get_range(GlobalWorkOffset)});
363  }
364 };
365 
366 extern "C" {
367 
368 #define DIE_NO_IMPLEMENTATION \
369  if (PrintPiTrace) { \
370  std::cerr << "Not Implemented : " << __FUNCTION__ \
371  << " - File : " << __FILE__; \
372  std::cerr << " / Line : " << __LINE__ << std::endl; \
373  } \
374  return PI_INVALID_OPERATION;
375 
376 #define CONTINUE_NO_IMPLEMENTATION \
377  if (PrintPiTrace) { \
378  std::cerr << "Warning : Not Implemented : " << __FUNCTION__ \
379  << " - File : " << __FILE__; \
380  std::cerr << " / Line : " << __LINE__ << std::endl; \
381  } \
382  return PI_SUCCESS;
383 
384 #define CASE_PI_UNSUPPORTED(not_supported) \
385  case not_supported: \
386  if (PrintPiTrace) { \
387  std::cerr << std::endl \
388  << "Unsupported PI case : " << #not_supported << " in " \
389  << __FUNCTION__ << ":" << __LINE__ << "(" << __FILE__ << ")" \
390  << std::endl; \
391  } \
392  return PI_INVALID_OPERATION;
393 
395  pi_uint32 *NumPlatforms) {
396  static bool PiPlatformCachePopulated = false;
397  static const char *PiTrace = std::getenv("SYCL_PI_TRACE");
398  static const int PiTraceValue = PiTrace ? std::stoi(PiTrace) : 0;
399 
400  if (PiTraceValue == -1) { // Means print all PI traces
401  PrintPiTrace = true;
402  }
403 
404  if (NumPlatforms) {
405  *NumPlatforms = 1;
406  }
407 
408  if (NumEntries == 0) {
410  if (Platforms != nullptr) {
411  if (PrintPiTrace) {
412  std::cerr << "Invalid Arguments for piPlatformsGet of esimd_emultor "
413  "(Platforms!=nullptr) while querying number of platforms"
414  << std::endl;
415  }
416  return PI_INVALID_VALUE;
417  }
418  return PI_SUCCESS;
419  }
420 
421  if (Platforms == nullptr && NumPlatforms == nullptr) {
422  return PI_INVALID_VALUE;
423  }
424 
425  std::lock_guard<std::mutex> Lock{*PiPlatformCacheLock};
428  PiPlatformCache->CmEmuVersion = std::string("0.0.1");
430  }
431 
432  if (Platforms && NumEntries > 0) {
433  *Platforms = PiPlatformCache;
434  }
435 
436  return PI_SUCCESS;
437 }
438 
440  size_t ParamValueSize, void *ParamValue,
441  size_t *ParamValueSizeRet) {
442  if (Platform == nullptr) {
443  return PI_INVALID_PLATFORM;
444  }
445  ReturnHelper ReturnValue(ParamValueSize, ParamValue, ParamValueSizeRet);
446 
447  switch (ParamName) {
449  return ReturnValue("Intel(R) ESIMD_EMULATOR/GPU");
450 
452  return ReturnValue("Intel(R) Corporation");
453 
455  return ReturnValue(Platform->CmEmuVersion.c_str());
456 
458  return ReturnValue("FULL_PROFILE");
459 
461  return ReturnValue("");
462 
463  default:
464  // TODO: implement other parameters
465  die("Unsupported ParamName in piPlatformGetInfo");
466  }
467 
468  return PI_SUCCESS;
469 }
470 
473 }
474 
477 }
478 
480  pi_uint32 NumEntries, pi_device *Devices,
481  pi_uint32 *NumDevices) {
482  if (Platform == nullptr) {
483  return PI_INVALID_PLATFORM;
484  }
485 
486  pi_result Res = Platform->populateDeviceCacheIfNeeded();
487  if (Res != PI_SUCCESS) {
488  return Res;
489  }
490 
491  // CM has single-root-GPU-device without sub-device support.
492  pi_uint32 DeviceCount = (DeviceType & PI_DEVICE_TYPE_GPU) ? 1 : 0;
493 
494  if (NumDevices) {
495  *NumDevices = DeviceCount;
496  }
497 
498  if (NumEntries == 0) {
500  if (Devices != nullptr) {
501  if (PrintPiTrace) {
502  std::cerr << "Invalid Arguments for piDevicesGet of esimd_emultor "
503  "(Devices!=nullptr) while querying number of platforms"
504  << std::endl;
505  }
506  return PI_INVALID_VALUE;
507  }
508  return PI_SUCCESS;
509  }
510 
511  if (DeviceCount == 0) {
513  return PI_SUCCESS;
514  }
515 
516  if (Devices) {
517  *Devices = Platform->PiDeviceCache.get();
518  }
519  return PI_SUCCESS;
520 }
521 
522 // Check the device cache and load it if necessary.
524  std::lock_guard<std::mutex> Lock(PiDeviceCacheMutex);
525 
526  if (DeviceCachePopulated) {
527  return PI_SUCCESS;
528  }
529  cm_support::CmDevice *CmDevice = nullptr;
530  // TODO FIXME Implement proper version checking and reporting:
531  // - version passed to cm_support::CreateCmDevice
532  // - CmEmuVersion
533  // - PluginVersion
534  // - ESIMDEmuPluginOpaqueData::version
535  //
536  // PI_DEVICE_INFO_DRIVER_VERSION could report the ESIMDDeviceInterface
537  // version, PI_PLATFORM_INFO_VERSION - the underlying libCM library version.
538  unsigned int Version = 0;
539 
540  int Result = cm_support::CreateCmDevice(CmDevice, Version);
541 
542  if (Result != cm_support::CM_SUCCESS) {
543  return PI_INVALID_DEVICE;
544  }
545 
546  // CM Device version info consists of two decimal numbers - major
547  // and minor. Minor is single-digit. Version info is encoded into a
548  // unsigned integer value = 100 * major + minor. Second from right
549  // digit in decimal must be zero as it is used as 'dot'
550  // REF - $CM_EMU/common/cm_version_defs.h - 'CURRENT_CM_VERSION'
551  // e.g. CM version 7.3 => Device version = 703
552 
553  if (((Version / 10) % 10) != 0) {
554  if (PrintPiTrace) {
555  std::cerr << "CM_EMU Device version info is incorrect : " << Version
556  << std::endl;
557  }
558  return PI_INVALID_DEVICE;
559  }
560 
561  std::ostringstream StrFormat;
562  StrFormat << (int)(Version / 100) << "." << (int)(Version % 10);
563 
564  std::unique_ptr<_pi_device> Device(
565  new _pi_device(this, CmDevice, StrFormat.str()));
566  PiDeviceCache = std::move(Device);
567  DeviceCachePopulated = true;
568  return PI_SUCCESS;
569 }
570 
572  if (Device == nullptr) {
573  return PI_INVALID_DEVICE;
574  }
575 
576  // CM supports only single device, which is root-device. 'Retain' is
577  // No-op.
578  return PI_SUCCESS;
579 }
580 
582  if (Device == nullptr) {
583  return PI_INVALID_DEVICE;
584  }
585 
586  // CM supports only single device, which is root-device. 'Release'
587  // is No-op.
588  return PI_SUCCESS;
589 }
590 
592  size_t ParamValueSize, void *ParamValue,
593  size_t *ParamValueSizeRet) {
594  ReturnHelper ReturnValue(ParamValueSize, ParamValue, ParamValueSizeRet);
595 
596  switch (ParamName) {
597  case PI_DEVICE_INFO_TYPE:
598  return ReturnValue(PI_DEVICE_TYPE_GPU);
600  return ReturnValue(pi_device{0});
602  return ReturnValue(Device->Platform);
603  case PI_DEVICE_INFO_NAME:
604  return ReturnValue("ESIMD_EMULATOR");
606  return ReturnValue(pi_bool{true});
614  return ReturnValue(ESimdEmuVersionString);
616  return ReturnValue("Intel(R) Corporation");
618  return ReturnValue(size_t{8192});
620  return ReturnValue(size_t{8192});
622  return ReturnValue(pi_bool{1});
624  // TODO : Populate return string accordingly - e.g. cl_khr_fp16,
625  // cl_khr_fp64, cl_khr_int64_base_atomics,
626  // cl_khr_int64_extended_atomics
627  return ReturnValue("");
629  return ReturnValue(Device->VersionStr.c_str());
630  case PI_DEVICE_INFO_BUILD_ON_SUBDEVICE: // emulator doesn't support partition
631  return ReturnValue(pi_bool{true});
633  return ReturnValue(pi_bool{false});
635  return ReturnValue(pi_bool{false});
637  return ReturnValue(pi_uint32{256});
639  return ReturnValue(pi_uint32{0});
641  return ReturnValue(pi_device_partition_property{0});
643  // '0x8086' : 'Intel HD graphics vendor ID'
644  return ReturnValue(pi_uint32{0x8086});
646  // Default SLM_MAX_SIZE from CM_EMU
647  return ReturnValue(pi_uint32{65536});
649  return ReturnValue(size_t{256});
651  // Imported from level_zero
652  return ReturnValue(pi_uint32{8});
656  // Default minimum values required by the SYCL specification.
657  return ReturnValue(size_t{2048});
659  return ReturnValue(pi_uint32{3});
661  return ReturnValue(pi_device_partition_property{0});
663  return ReturnValue("");
665  return ReturnValue(pi_queue_properties{PI_QUEUE_ON_DEVICE});
667  struct {
668  size_t Arr[3];
669  } MaxGroupSize = {{256, 256, 1}};
670  return ReturnValue(MaxGroupSize);
671  }
686  return ReturnValue(pi_uint32{1});
687 
688  // Imported from level_zero
694  pi_uint64 Supported = 0;
695  // TODO[1.0]: how to query for USM support now?
696  if (true) {
697  // TODO: Use ze_memory_access_capabilities_t
698  Supported = PI_USM_ACCESS | PI_USM_ATOMIC_ACCESS |
700  }
701  return ReturnValue(Supported);
702  }
704  return ReturnValue(
705  pi_uint32{sizeof(void *) * std::numeric_limits<unsigned char>::digits});
707  return ReturnValue(pi_uint32{1000});
709  return ReturnValue(pi_bool{true});
711  return ReturnValue(pi_bool{true});
715  return ReturnValue(pi_uint32{0});
718  return ReturnValue(size_t{0x80000000});
721  return ReturnValue(size_t{0});
724  return ReturnValue(pi_uint32{16});
727  return ReturnValue(size_t{32});
732  uint64_t FPValue = PI_FP_CORRECTLY_ROUNDED_DIVIDE_SQRT |
735  PI_FP_FMA;
736  return ReturnValue(pi_uint64{FPValue});
737  }
739  return ReturnValue(PI_DEVICE_MEM_CACHE_TYPE_READ_WRITE_CACHE);
741  // TODO : CHECK
742  return ReturnValue(pi_uint32{64});
744  // TODO : CHECK
745  return ReturnValue(pi_uint64{0});
747  // TODO : CHECK
748  return ReturnValue(pi_uint64{0});
750  // TODO : CHECK
751  return ReturnValue(pi_uint64{0});
753  // TODO : CHECK
754  return ReturnValue(pi_uint32{64});
756  // TODO : CHECK
757  return ReturnValue(PI_DEVICE_LOCAL_MEM_TYPE_LOCAL);
759  return ReturnValue(pi_bool{false});
761  // TODO : CHECK
762  return ReturnValue(size_t{0});
764  // TODO : CHECK
765  return ReturnValue("");
767  // TODO : CHECK
768  return ReturnValue(size_t{1024});
770  return ReturnValue(pi_bool{false});
772  return ReturnValue(pi_device_affinity_domain{0});
774  // TODO : CHECK
775  return ReturnValue(pi_uint64{0});
777  // TODO : CHECK
778  return ReturnValue(
781  return ReturnValue("FULL_PROFILE");
783  // TODO : CHECK
784  return ReturnValue(pi_uint32{0});
785 
790 
791  // Intel-specific extensions
806 
807  default:
809  }
810  return PI_SUCCESS;
811 }
812 
814  pi_uint32, pi_device *, pi_uint32 *) {
816 }
817 
820 }
821 
823  pi_device *) {
825 }
826 
828  pi_uint32 NumDevices, const pi_device *Devices,
829  void (*PFnNotify)(const char *ErrInfo,
830  const void *PrivateInfo, size_t CB,
831  void *UserData),
832  void *UserData, pi_context *RetContext) {
833  ARG_UNUSED(Properties);
834  ARG_UNUSED(PFnNotify);
835  ARG_UNUSED(UserData);
836 
837  if (NumDevices != 1) {
838  return PI_INVALID_VALUE;
839  }
840  if (Devices == nullptr) {
841  return PI_INVALID_DEVICE;
842  }
843  if (RetContext == nullptr) {
844  return PI_INVALID_VALUE;
845  }
846 
847  try {
849  *RetContext = new _pi_context(Devices[0]);
850  } catch (const std::bad_alloc &) {
851  return PI_OUT_OF_HOST_MEMORY;
852  } catch (...) {
853  return PI_ERROR_UNKNOWN;
854  }
855  return PI_SUCCESS;
856 }
857 
859  size_t *) {
861 }
862 
864  pi_context_extended_deleter, void *) {
866 }
867 
870 }
871 
873  const pi_device *, bool,
874  pi_context *) {
876 }
877 
879  if (Context == nullptr) {
880  return PI_INVALID_CONTEXT;
881  }
882 
883  ++(Context->RefCount);
884 
885  return PI_SUCCESS;
886 }
887 
889  if (Context == nullptr || (Context->RefCount <= 0)) {
890  return PI_INVALID_CONTEXT;
891  }
892 
893  if (--(Context->RefCount) == 0) {
896  std::lock_guard<std::mutex> Lock(Context->Addr2CmBufferSVMLock);
897  for (auto &Entry : Context->Addr2CmBufferSVM) {
898  Context->Device->CmDevicePtr->DestroyBufferSVM(Entry.second);
899  }
900  delete Context;
901  }
902 
903  return PI_SUCCESS;
904 }
905 
908  if (HostPtr == nullptr) {
909  if (PrintPiTrace) {
910  std::cerr << "HostPtr argument is required for "
911  "PI_MEM_FLAGS_HOST_PTR_USE/COPY"
912  << std::endl;
913  }
914  return false;
915  }
916  // COPY and USE are mutually exclusive
919  if (PrintPiTrace) {
920  std::cerr
921  << "PI_MEM_FLAGS_HOST_PTR_USE and _COPY cannot be used together"
922  << std::endl;
923  }
924  return false;
925  }
926  }
927  return true;
928 }
929 
931  pi_queue_properties Properties, pi_queue *Queue) {
932  ARG_UNUSED(Device);
933 
934  if (Properties & PI_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE) {
935  // TODO : Support Out-of-order Queue
936  *Queue = nullptr;
938  }
939 
940  cm_support::CmQueue *CmQueue;
941 
942  int Result = Context->Device->CmDevicePtr->CreateQueue(CmQueue);
943  if (Result != cm_support::CM_SUCCESS) {
944  return PI_INVALID_CONTEXT;
945  }
946 
947  try {
948  *Queue = new _pi_queue(Context, CmQueue);
949  } catch (const std::bad_alloc &) {
950  return PI_OUT_OF_HOST_MEMORY;
951  } catch (...) {
952  return PI_ERROR_UNKNOWN;
953  }
954 
955  return PI_SUCCESS;
956 }
957 
958 pi_result piQueueGetInfo(pi_queue, pi_queue_info, size_t, void *, size_t *) {
960 }
961 
963  if (Queue == nullptr) {
964  return PI_INVALID_QUEUE;
965  }
966  ++(Queue->RefCount);
967  return PI_SUCCESS;
968 }
969 
971  if ((Queue == nullptr) || (Queue->CmQueuePtr == nullptr)) {
972  return PI_INVALID_QUEUE;
973  }
974 
975  if (--(Queue->RefCount) == 0) {
976  // CM's 'DestoryQueue' is no-op
977  // Queue->Context->Device->CmDevicePTr->DestroyQueue(Queue->CmQueuePtr);
978  delete Queue;
979  }
980 
981  return PI_SUCCESS;
982 }
983 
985  // No-op as enqueued commands with ESIMD_EMULATOR plugin are blocking
986  // ones that do not return until their completion - kernel execution
987  // and memory read.
989 }
990 
992  // No-op as enqueued commands with ESIMD_EMULATOR plugin are blocking
993  // ones that do not return until their completion - kernel execution
994  // and memory read.
996 }
997 
1000 }
1001 
1003  pi_device, bool, pi_queue *) {
1005 }
1006 
1008  void *HostPtr, pi_mem *RetMem,
1009  const pi_mem_properties *properties) {
1010  ARG_UNUSED(properties);
1011 
1012  if ((Flags & PI_MEM_FLAGS_ACCESS_RW) == 0) {
1013  if (PrintPiTrace) {
1014  std::cerr << "Invalid memory attribute for piMemBufferCreate"
1015  << std::endl;
1016  }
1017  return PI_INVALID_OPERATION;
1018  }
1019 
1020  if (Context == nullptr) {
1021  return PI_INVALID_CONTEXT;
1022  }
1023  if (RetMem == nullptr) {
1024  return PI_INVALID_VALUE;
1025  }
1026 
1027  // Flag & HostPtr argument sanity check
1028  if (!Context->checkSurfaceArgument(Flags, HostPtr)) {
1029  return PI_INVALID_OPERATION;
1030  }
1031 
1032  char *MapBasePtr = nullptr;
1033  cm_surface_ptr_t CmBuf;
1034  cm_support::SurfaceIndex *CmIndex;
1035  int Status = cm_support::CM_FAILURE;
1036 
1037  if (Flags & PI_MEM_FLAGS_HOST_PTR_USE) {
1039  Status = Context->Device->CmDevicePtr->CreateBufferUP(
1040  static_cast<unsigned int>(Size), HostPtr, CmBuf.UPBufPtr);
1041  CmBuf.UPBufPtr->GetIndex(CmIndex);
1042  } else {
1044  Status = Context->Device->CmDevicePtr->CreateBuffer(
1045  static_cast<unsigned int>(Size), CmBuf.RegularBufPtr);
1046  CmBuf.RegularBufPtr->GetIndex(CmIndex);
1047 
1048  if (Flags & PI_MEM_FLAGS_HOST_PTR_COPY) {
1049  CmBuf.RegularBufPtr->WriteSurface(
1050  reinterpret_cast<const unsigned char *>(HostPtr), nullptr,
1051  static_cast<unsigned int>(Size));
1052  }
1053  }
1054 
1055  if (Status != cm_support::CM_SUCCESS) {
1056  return PI_INVALID_OPERATION;
1057  }
1058 
1059  MapBasePtr =
1060  pi_cast<char *>(cm_support::get_surface_base_addr(CmIndex->get_data()));
1061 
1062  try {
1063  *RetMem =
1064  new _pi_buffer(Context, MapBasePtr, CmBuf, CmIndex->get_data(), Size);
1065  } catch (const std::bad_alloc &) {
1066  return PI_OUT_OF_HOST_MEMORY;
1067  } catch (...) {
1068  return PI_ERROR_UNKNOWN;
1069  }
1070 
1071  std::lock_guard<std::mutex> Lock{*PiESimdSurfaceMapLock};
1072  assert(PiESimdSurfaceMap->find((*RetMem)->SurfaceIndex) ==
1073  PiESimdSurfaceMap->end() &&
1074  "Failure from CM-managed buffer creation");
1075 
1076  (*PiESimdSurfaceMap)[(*RetMem)->SurfaceIndex] = *RetMem;
1077 
1078  return PI_SUCCESS;
1079 }
1080 
1081 pi_result piMemGetInfo(pi_mem, pi_mem_info, size_t, void *, size_t *) {
1083 }
1084 
1086  if (Mem == nullptr) {
1087  return PI_INVALID_MEM_OBJECT;
1088  }
1089  ++(Mem->RefCount);
1090  return PI_SUCCESS;
1091 }
1092 
1094  if ((Mem == nullptr) || (Mem->RefCount == 0)) {
1095  return PI_INVALID_MEM_OBJECT;
1096  }
1097 
1098  if (--(Mem->RefCount) == 0) {
1099  // Removing Surface-map entry
1100  std::lock_guard<std::mutex> Lock{*PiESimdSurfaceMapLock};
1101  auto MapEntryIt = PiESimdSurfaceMap->find(Mem->SurfaceIndex);
1102  assert(MapEntryIt != PiESimdSurfaceMap->end() &&
1103  "Failure from Buffer/Image deletion");
1104  PiESimdSurfaceMap->erase(MapEntryIt);
1105  delete Mem;
1106  }
1107  return PI_SUCCESS;
1108 }
1109 
1111  int Status = cm_support::CM_FAILURE;
1112 
1113  cm_support::CmDevice *CmDevice = Context->Device->CmDevicePtr;
1114 
1116  Status = CmDevice->DestroyBufferUP(SurfacePtr.UPBufPtr);
1118  Status = CmDevice->DestroySurface(SurfacePtr.RegularBufPtr);
1120  Status = CmDevice->DestroySurface2DUP(SurfacePtr.UPImgPtr);
1122  Status = CmDevice->DestroySurface(SurfacePtr.RegularImgPtr);
1123  }
1124 
1125  assert(Status == cm_support::CM_SUCCESS &&
1126  "Surface Deletion Failure from CM_EMU");
1127 
1128  for (auto mapit = Mappings.begin(); mapit != Mappings.end();) {
1129  mapit = Mappings.erase(mapit);
1130  }
1131 }
1132 
1133 cm_support::CM_SURFACE_FORMAT
1135  using ULongPair = std::pair<unsigned long, unsigned long>;
1136  using FmtMap = std::map<ULongPair, cm_support::CM_SURFACE_FORMAT>;
1137  static const FmtMap pi2cm = {
1139  cm_support::CM_SURFACE_FORMAT_A8R8G8B8},
1140 
1142  cm_support::CM_SURFACE_FORMAT_A8R8G8B8},
1143 
1145  cm_support::CM_SURFACE_FORMAT_A8R8G8B8},
1146 
1148  cm_support::CM_SURFACE_FORMAT_R32G32B32A32F},
1149  };
1150  auto Result = pi2cm.find(
1151  {PiFormat->image_channel_data_type, PiFormat->image_channel_order});
1152  if (Result != pi2cm.end()) {
1153  return Result->second;
1154  }
1155  return cm_support::CM_SURFACE_FORMAT_UNKNOWN;
1156 }
1157 
1159  const pi_image_format *ImageFormat,
1160  const pi_image_desc *ImageDesc, void *HostPtr,
1161  pi_mem *RetImage) {
1162  if ((Flags & PI_MEM_FLAGS_ACCESS_RW) == 0) {
1163  if (PrintPiTrace) {
1164  std::cerr << "Invalid memory attribute for piMemImageCreate" << std::endl;
1165  }
1166  return PI_INVALID_OPERATION;
1167  }
1168 
1169  if (ImageFormat == nullptr || ImageDesc == nullptr)
1171 
1172  switch (ImageDesc->image_type) {
1173  case PI_MEM_TYPE_IMAGE2D:
1174  break;
1175 
1181 
1182  default:
1183  return PI_INVALID_MEM_OBJECT;
1184  }
1185 
1186  auto BytesPerPixel = 4;
1187  switch (ImageFormat->image_channel_data_type) {
1189  BytesPerPixel = 16;
1190  break;
1193  BytesPerPixel = 4;
1194  break;
1207  default:
1209  }
1210 
1211  // Flag & HostPtr argument sanity check
1212  if (!Context->checkSurfaceArgument(Flags, HostPtr)) {
1213  return PI_INVALID_OPERATION;
1214  }
1215 
1216  cm_support::CM_SURFACE_FORMAT CmSurfFormat =
1217  ConvertPiImageFormatToCmFormat(ImageFormat);
1218  if (CmSurfFormat == cm_support::CM_SURFACE_FORMAT_UNKNOWN) {
1220  }
1221 
1222  char *MapBasePtr = nullptr;
1223  cm_surface_ptr_t CmImg;
1224  cm_support::SurfaceIndex *CmIndex;
1225  int Status = cm_support::CM_SUCCESS;
1226 
1227  if (Flags & PI_MEM_FLAGS_HOST_PTR_USE) {
1229  Status = Context->Device->CmDevicePtr->CreateSurface2DUP(
1230  static_cast<unsigned int>(ImageDesc->image_width),
1231  static_cast<unsigned int>(ImageDesc->image_height), CmSurfFormat,
1232  HostPtr, CmImg.UPImgPtr);
1233  CmImg.UPImgPtr->GetIndex(CmIndex);
1234  } else {
1236  Status = Context->Device->CmDevicePtr->CreateSurface2D(
1237  static_cast<unsigned int>(ImageDesc->image_width),
1238  static_cast<unsigned int>(ImageDesc->image_height), CmSurfFormat,
1239  CmImg.RegularImgPtr);
1240  CmImg.RegularImgPtr->GetIndex(CmIndex);
1241 
1242  if (Flags & PI_MEM_FLAGS_HOST_PTR_COPY) {
1243  CmImg.RegularImgPtr->WriteSurface(
1244  reinterpret_cast<const unsigned char *>(HostPtr), nullptr,
1245  static_cast<unsigned int>(ImageDesc->image_width *
1246  ImageDesc->image_height * BytesPerPixel));
1247  }
1248  }
1249 
1250  if (Status != cm_support::CM_SUCCESS) {
1251  return PI_INVALID_OPERATION;
1252  }
1253 
1254  MapBasePtr =
1255  pi_cast<char *>(cm_support::get_surface_base_addr(CmIndex->get_data()));
1256 
1257  try {
1258  *RetImage = new _pi_image(Context, MapBasePtr, CmImg, CmIndex->get_data(),
1259  ImageDesc->image_width, ImageDesc->image_height,
1260  BytesPerPixel);
1261  } catch (const std::bad_alloc &) {
1262  return PI_OUT_OF_HOST_MEMORY;
1263  } catch (...) {
1264  return PI_ERROR_UNKNOWN;
1265  }
1266 
1267  std::lock_guard<std::mutex> Lock{*PiESimdSurfaceMapLock};
1268  assert(PiESimdSurfaceMap->find((*RetImage)->SurfaceIndex) ==
1269  PiESimdSurfaceMap->end() &&
1270  "Failure from CM-managed image creation");
1271 
1272  (*PiESimdSurfaceMap)[(*RetImage)->SurfaceIndex] = *RetImage;
1273 
1274  return PI_SUCCESS;
1275 }
1276 
1279 }
1280 
1282  pi_mem *) {
1284 }
1285 
1286 pi_result piProgramCreate(pi_context, const void *, size_t, pi_program *) {
1288 }
1289 
1291  const size_t *, const unsigned char **,
1292  size_t, const pi_device_binary_property *,
1293  pi_int32 *, pi_program *) {
1295 }
1296 
1298  const size_t *, const unsigned char **,
1299  pi_int32 *, pi_program *) {
1301 }
1302 
1304  const size_t *, pi_program *) {
1306 }
1307 
1309  size_t *) {
1311 }
1312 
1314  pi_uint32, const pi_program *,
1315  void (*)(pi_program, void *), void *, pi_program *) {
1317 }
1318 
1320  const char *, pi_uint32, const pi_program *,
1321  const char **, void (*)(pi_program, void *),
1322  void *) {
1324 }
1325 
1327  void (*)(pi_program, void *), void *) {
1329 }
1330 
1332  size_t, void *, size_t *) {
1334 }
1335 
1337 
1339 
1342 }
1343 
1345  pi_program *) {
1347 }
1348 
1351 }
1352 
1353 pi_result piKernelSetArg(pi_kernel, pi_uint32, size_t, const void *) {
1355 }
1356 
1359 }
1360 
1361 // Special version of piKernelSetArg to accept pi_sampler.
1364 }
1365 
1366 pi_result piKernelGetInfo(pi_kernel, pi_kernel_info, size_t, void *, size_t *) {
1368 }
1369 
1371  size_t, void *, size_t *) {
1373 }
1374 
1376  pi_kernel_sub_group_info, size_t,
1377  const void *, size_t, void *, size_t *) {
1379 }
1380 
1382 
1384 
1386 
1387 pi_result piEventGetInfo(pi_event, pi_event_info, size_t, void *, size_t *) {
1389 }
1390 
1392  size_t ParamValueSize, void *ParamValue,
1393  size_t *ParamValueSizeRet) {
1394  ARG_UNUSED(Event);
1395  ARG_UNUSED(ParamName);
1396  ARG_UNUSED(ParamValueSize);
1397  ARG_UNUSED(ParamValue);
1398  ARG_UNUSED(ParamValueSizeRet);
1399 
1400  if (PrintPiTrace) {
1401  std::cerr << "Warning : Profiling Not supported under PI_ESIMD_EMULATOR"
1402  << std::endl;
1403  }
1404  return PI_SUCCESS;
1405 }
1406 
1407 pi_result piEventsWait(pi_uint32 NumEvents, const pi_event *EventList) {
1408  for (int i = 0; i < (int)NumEvents; i++) {
1409  if (EventList[i]->IsDummyEvent) {
1410  // Dummy event is already completed ones done by CM. Skip
1411  // waiting.
1412  continue;
1413  }
1414  if (EventList[i]->CmEventPtr == nullptr) {
1415  return PI_INVALID_EVENT;
1416  }
1417  int Result = EventList[i]->CmEventPtr->WaitForTaskFinished();
1418  if (Result != cm_support::CM_SUCCESS) {
1419  return PI_OUT_OF_RESOURCES;
1420  }
1421  }
1422  return PI_SUCCESS;
1423 }
1424 
1426  void (*)(pi_event, pi_int32, void *), void *) {
1428 }
1429 
1431 
1433  if (Event == nullptr) {
1434  return PI_INVALID_EVENT;
1435  }
1436 
1437  ++(Event->RefCount);
1438 
1439  return PI_SUCCESS;
1440 }
1441 
1443  if (Event == nullptr || (Event->RefCount <= 0)) {
1444  return PI_INVALID_EVENT;
1445  }
1446 
1447  if (--(Event->RefCount) == 0) {
1448  if (!Event->IsDummyEvent) {
1449  if ((Event->CmEventPtr == nullptr) || (Event->OwnerQueue == nullptr)) {
1450  return PI_INVALID_EVENT;
1451  }
1452  int Result = Event->OwnerQueue->DestroyEvent(Event->CmEventPtr);
1453  if (Result != cm_support::CM_SUCCESS) {
1454  return PI_INVALID_EVENT;
1455  }
1456  }
1457  delete Event;
1458  }
1459 
1460  return PI_SUCCESS;
1461 }
1462 
1465 }
1466 
1468  pi_event *) {
1470 }
1472  pi_sampler *) {
1474 }
1475 
1477  size_t *) {
1479 }
1480 
1482 
1484 
1486  pi_event *) {
1488 }
1489 
1491  pi_event *) {
1493 }
1494 
1496  pi_bool BlockingRead, size_t Offset,
1497  size_t Size, void *Dst,
1498  pi_uint32 NumEventsInWaitList,
1499  const pi_event *EventWaitList,
1500  pi_event *Event) {
1501  ARG_UNUSED(Queue);
1502  ARG_UNUSED(EventWaitList);
1503 
1505  if (BlockingRead) {
1506  assert(false &&
1507  "ESIMD_EMULATOR support for blocking piEnqueueMemBufferRead is NYI");
1508  }
1509 
1510  assert(Offset == 0 &&
1511  "ESIMD_EMULATOR does not support buffer reading with offsets");
1512 
1513  if (NumEventsInWaitList != 0) {
1515  }
1516 
1517  _pi_buffer *buf = static_cast<_pi_buffer *>(Src);
1518 
1519  std::unique_ptr<_pi_event> RetEv{nullptr};
1520  if (Event) {
1521  RetEv = std::unique_ptr<_pi_event>(new _pi_event());
1522  RetEv->IsDummyEvent = true;
1523  }
1524 
1526  // CM does not provide 'ReadSurface' call for 'User-Provided'
1527  // Surface. memcpy is used for BufferRead PI_API call.
1528  memcpy(Dst, buf->MapHostPtr, Size);
1529  } else {
1531  int Status = buf->SurfacePtr.RegularBufPtr->ReadSurface(
1532  reinterpret_cast<unsigned char *>(Dst),
1533  nullptr, // event
1534  static_cast<uint64_t>(Size));
1535 
1536  if (Status != cm_support::CM_SUCCESS) {
1537  return PI_INVALID_MEM_OBJECT;
1538  }
1539  }
1540 
1541  if (Event) {
1542  *Event = RetEv.release();
1543  }
1544 
1545  return PI_SUCCESS;
1546 }
1547 
1550  pi_buff_rect_region, size_t, size_t,
1551  size_t, size_t, void *, pi_uint32,
1552  const pi_event *, pi_event *) {
1554 }
1555 
1557  const void *, pi_uint32, const pi_event *,
1558  pi_event *) {
1560 }
1561 
1564  pi_buff_rect_region, size_t, size_t,
1565  size_t, size_t, const void *, pi_uint32,
1566  const pi_event *, pi_event *) {
1568 }
1569 
1571  size_t, pi_uint32, const pi_event *,
1572  pi_event *) {
1574 }
1575 
1578  pi_buff_rect_region, size_t, size_t,
1579  size_t, size_t, pi_uint32,
1580  const pi_event *, pi_event *) {
1582 }
1583 
1584 pi_result piEnqueueMemBufferFill(pi_queue, pi_mem, const void *, size_t, size_t,
1585  size_t, pi_uint32, const pi_event *,
1586  pi_event *) {
1588 }
1589 
1591  pi_bool BlockingMap, pi_map_flags MapFlags,
1592  size_t Offset, size_t Size,
1593  pi_uint32 NumEventsInWaitList,
1594  const pi_event *EventWaitList, pi_event *Event,
1595  void **RetMap) {
1596  ARG_UNUSED(Queue);
1597  ARG_UNUSED(BlockingMap);
1598  ARG_UNUSED(MapFlags);
1599  ARG_UNUSED(NumEventsInWaitList);
1600  ARG_UNUSED(EventWaitList);
1601 
1602  std::unique_ptr<_pi_event> RetEv{nullptr};
1603  pi_result ret = PI_SUCCESS;
1604 
1605  if (Event) {
1606  RetEv = std::unique_ptr<_pi_event>(new _pi_event());
1607  RetEv->IsDummyEvent = true;
1608  }
1609 
1610  // Real mapping does not occur here and CPU-accessible address is
1611  // returned as the actual memory space for the buffer is located in
1612  // CPU memory and the plug-in know its base address
1613  // ('_pi_mem::MapHostPtr')
1614  *RetMap = MemObj->MapHostPtr + Offset;
1615 
1616  {
1617  std::lock_guard<std::mutex> Lock{MemObj->MappingsMutex};
1618  auto Res = MemObj->Mappings.insert({*RetMap, {Offset, Size}});
1619  // False as the second value in pair means that mapping was not inserted
1620  // because mapping already exists.
1621  if (!Res.second) {
1622  ret = PI_INVALID_VALUE;
1623  if (PrintPiTrace) {
1624  std::cerr << "piEnqueueMemBufferMap: duplicate mapping detected"
1625  << std::endl;
1626  }
1627  }
1628  }
1629 
1630  if (Event) {
1631  *Event = RetEv.release();
1632  }
1633  return ret;
1634 }
1635 
1636 pi_result piEnqueueMemUnmap(pi_queue Queue, pi_mem MemObj, void *MappedPtr,
1637  pi_uint32 NumEventsInWaitList,
1638  const pi_event *EventWaitList, pi_event *Event) {
1639  ARG_UNUSED(Queue);
1640  ARG_UNUSED(NumEventsInWaitList);
1641  ARG_UNUSED(EventWaitList);
1642 
1643  std::unique_ptr<_pi_event> RetEv{nullptr};
1644  pi_result ret = PI_SUCCESS;
1645 
1646  if (Event) {
1647  RetEv = std::unique_ptr<_pi_event>(new _pi_event());
1648  RetEv->IsDummyEvent = true;
1649  }
1650 
1651  // Real unmapping does not occur here and CPU-accessible address is
1652  // returned as the actual memory space for the buffer is located in
1653  // CPU memory and the plug-in knows its base address
1654  // ('_pi_mem::MapHostPtr')
1655  {
1656  std::lock_guard<std::mutex> Lock(MemObj->MappingsMutex);
1657  auto It = MemObj->Mappings.find(MappedPtr);
1658  if (It == MemObj->Mappings.end()) {
1659  ret = PI_INVALID_VALUE;
1660  if (PrintPiTrace) {
1661  std::cerr << "piEnqueueMemUnmap: unknown memory mapping" << std::endl;
1662  }
1663  }
1664  MemObj->Mappings.erase(It);
1665  }
1666 
1667  if (Event) {
1668  *Event = RetEv.release();
1669  }
1670 
1671  return ret;
1672 }
1673 
1674 pi_result piMemImageGetInfo(pi_mem, pi_image_info, size_t, void *, size_t *) {
1676 }
1677 
1679  pi_bool BlockingRead, pi_image_offset Origin,
1680  pi_image_region Region, size_t RowPitch,
1681  size_t SlicePitch, void *Ptr,
1682  pi_uint32 NumEventsInWaitList,
1683  const pi_event *EventWaitList,
1684  pi_event *Event) {
1685  ARG_UNUSED(CommandQueue);
1686  ARG_UNUSED(NumEventsInWaitList);
1687  ARG_UNUSED(EventWaitList);
1688 
1690  if (BlockingRead) {
1691  assert(false && "ESIMD_EMULATOR does not support Blocking Read");
1692  }
1693 
1694  // SlicePitch is for 3D image while ESIMD_EMULATOR does not
1695  // support. For 2D surfaces, SlicePitch must be 0.
1696  assert((SlicePitch == 0) && "ESIMD_EMULATOR does not support 3D-image");
1697 
1698  // CM_EMU does not support ReadSurface with offset
1699  assert(Origin->x == 0 && Origin->y == 0 && Origin->z == 0 &&
1700  "ESIMD_EMULATOR does not support 2D-image reading with offsets");
1701 
1702  _pi_image *PiImg = static_cast<_pi_image *>(Image);
1703 
1704  std::unique_ptr<_pi_event> RetEv{nullptr};
1705 
1706  if (Event) {
1707  RetEv = std::unique_ptr<_pi_event>(new _pi_event());
1708  RetEv->IsDummyEvent = true;
1709  }
1710 
1711  size_t Size = RowPitch * (Region->height);
1713  // CM does not provide 'ReadSurface' call for 'User-Provided'
1714  // Surface. memcpy is used for ImageRead PI_API call.
1715  memcpy(Ptr, PiImg->MapHostPtr, Size);
1716  } else {
1718  int Status = PiImg->SurfacePtr.RegularImgPtr->ReadSurface(
1719  reinterpret_cast<unsigned char *>(Ptr),
1720  nullptr, // event
1721  static_cast<uint64_t>(Size));
1722 
1723  if (Status != cm_support::CM_SUCCESS) {
1724  return PI_INVALID_MEM_OBJECT;
1725  }
1726  }
1727 
1728  if (Event) {
1729  *Event = RetEv.release();
1730  }
1731 
1732  return PI_SUCCESS;
1733 }
1734 
1736  pi_image_region, size_t, size_t, const void *,
1737  pi_uint32, const pi_event *, pi_event *) {
1739 }
1740 
1743  const pi_event *, pi_event *) {
1745 }
1746 
1747 pi_result piEnqueueMemImageFill(pi_queue, pi_mem, const void *, const size_t *,
1748  const size_t *, pi_uint32, const pi_event *,
1749  pi_event *) {
1751 }
1752 
1754  void *, pi_mem *) {
1756 }
1757 
1758 pi_result
1760  const size_t *GlobalWorkOffset,
1761  const size_t *GlobalWorkSize, const size_t *LocalWorkSize,
1762  pi_uint32 NumEventsInWaitList,
1763  const pi_event *EventWaitList, pi_event *Event) {
1764  ARG_UNUSED(Queue);
1765  ARG_UNUSED(NumEventsInWaitList);
1766  ARG_UNUSED(EventWaitList);
1767 
1768  const size_t LocalWorkSz[] = {1, 1, 1};
1769 
1770  if (Kernel == nullptr) {
1771  return PI_INVALID_KERNEL;
1772  }
1773 
1774  if (WorkDim > 3 || WorkDim == 0) {
1776  }
1777 
1778  if (isNull(WorkDim, LocalWorkSize)) {
1779  LocalWorkSize = LocalWorkSz;
1780  }
1781 
1782  for (pi_uint32 I = 0; I < WorkDim; I++) {
1783  if ((GlobalWorkSize[I] % LocalWorkSize[I]) != 0) {
1785  }
1786  }
1787 
1788  std::unique_ptr<_pi_event> RetEv{nullptr};
1789 
1790  if (Event) {
1791  RetEv = std::unique_ptr<_pi_event>(new _pi_event());
1792  RetEv->IsDummyEvent = true;
1793  }
1794 
1795  switch (WorkDim) {
1796  case 1:
1797  InvokeImpl<1>::invoke(Kernel, GlobalWorkOffset, GlobalWorkSize,
1798  LocalWorkSize);
1799  break;
1800  case 2:
1801  InvokeImpl<2>::invoke(Kernel, GlobalWorkOffset, GlobalWorkSize,
1802  LocalWorkSize);
1803  break;
1804  case 3:
1805  InvokeImpl<3>::invoke(Kernel, GlobalWorkOffset, GlobalWorkSize,
1806  LocalWorkSize);
1807  break;
1808  default:
1810  break;
1811  }
1812 
1813  if (Event) {
1814  *Event = RetEv.release();
1815  }
1816 
1817  return PI_SUCCESS;
1818 }
1819 
1821  pi_program, bool, pi_kernel *) {
1823 }
1824 
1827 }
1828 
1829 pi_result piEnqueueNativeKernel(pi_queue, void (*)(void *), void *, size_t,
1830  pi_uint32, const pi_mem *, const void **,
1831  pi_uint32, const pi_event *, pi_event *) {
1833 }
1834 
1836  pi_uint64 *) {
1838 }
1839 
1841  size_t, pi_uint32) {
1843 }
1844 
1846  pi_usm_mem_properties *, size_t, pi_uint32) {
1848 }
1849 
1850 pi_result piextUSMSharedAlloc(void **ResultPtr, pi_context Context,
1851  pi_device Device,
1852  pi_usm_mem_properties *Properties, size_t Size,
1853  pi_uint32 Alignment) {
1854  ARG_UNUSED(Properties);
1855  ARG_UNUSED(Alignment);
1856 
1857  if (Context == nullptr || (Device != Context->Device)) {
1858  return PI_INVALID_CONTEXT;
1859  }
1860 
1861  if (ResultPtr == nullptr) {
1862  return PI_INVALID_OPERATION;
1863  }
1864 
1865  // 'Size' must be power of two in order to prevent memory corruption
1866  // error
1867  if ((Size & (Size - 1)) != 0) {
1868  Size = sycl::detail::getNextPowerOfTwo(Size);
1869  }
1870 
1871  cm_support::CmBufferSVM *Buf = nullptr;
1872  void *SystemMemPtr = nullptr;
1873  int32_t Result = Context->Device->CmDevicePtr->CreateBufferSVM(
1874  Size, SystemMemPtr, CM_SVM_ACCESS_FLAG_DEFAULT, Buf);
1875 
1876  if (Result != cm_support::CM_SUCCESS) {
1877  return PI_OUT_OF_HOST_MEMORY;
1878  }
1879  *ResultPtr = SystemMemPtr;
1880  std::lock_guard<std::mutex> Lock(Context->Addr2CmBufferSVMLock);
1881  auto Iter = Context->Addr2CmBufferSVM.find(SystemMemPtr);
1882  if (Context->Addr2CmBufferSVM.end() != Iter) {
1883  return PI_INVALID_MEM_OBJECT;
1884  }
1885  Context->Addr2CmBufferSVM[SystemMemPtr] = Buf;
1886  return PI_SUCCESS;
1887 }
1888 
1889 pi_result piextUSMFree(pi_context Context, void *Ptr) {
1890  if (Context == nullptr) {
1891  return PI_INVALID_CONTEXT;
1892  }
1893  if (Ptr == nullptr) {
1894  return PI_INVALID_OPERATION;
1895  }
1896 
1897  std::lock_guard<std::mutex> Lock(Context->Addr2CmBufferSVMLock);
1898  cm_support::CmBufferSVM *Buf = Context->Addr2CmBufferSVM[Ptr];
1899  if (Buf == nullptr) {
1900  return PI_INVALID_MEM_OBJECT;
1901  }
1902  auto Count = Context->Addr2CmBufferSVM.erase(Ptr);
1903  if (Count != 1) {
1904  return PI_INVALID_MEM_OBJECT;
1905  }
1906  int32_t Result = Context->Device->CmDevicePtr->DestroyBufferSVM(Buf);
1907  if (cm_support::CM_SUCCESS != Result) {
1908  return PI_ERROR_UNKNOWN;
1909  }
1910  return PI_SUCCESS;
1911 }
1912 
1915 }
1916 
1918  const pi_event *, pi_event *) {
1920 }
1921 
1922 pi_result piextUSMEnqueueMemcpy(pi_queue, pi_bool, void *, const void *, size_t,
1923  pi_uint32, const pi_event *, pi_event *) {
1925 }
1926 
1928  pi_mem_advice, pi_event *) {
1930 }
1931 
1933  size_t, void *, size_t *) {
1935 }
1936 
1938  const void *) {
1940 }
1941 
1943  const void *) {
1945 }
1946 
1948  pi_uint32 RawImgSize, pi_uint32 *ImgInd) {
1951  if (RawImgSize != 1) {
1952  if (PrintPiTrace) {
1953  std::cerr
1954  << "Only single device binary image is supported in ESIMD_EMULATOR"
1955  << std::endl;
1956  }
1957  return PI_INVALID_VALUE;
1958  }
1959  *ImgInd = 0;
1960  return PI_SUCCESS;
1961 }
1962 
1965  const pi_event *, pi_event *) {
1967 }
1968 
1969 pi_result piextPluginGetOpaqueData(void *, void **OpaqueDataReturn) {
1970  *OpaqueDataReturn = reinterpret_cast<void *>(PiESimdDeviceAccess);
1971  return PI_SUCCESS;
1972 }
1973 
1975  delete reinterpret_cast<sycl::detail::ESIMDEmuPluginOpaqueData *>(
1976  PiESimdDeviceAccess->data);
1977  delete PiESimdDeviceAccess;
1978 
1979  for (auto it = PiESimdSurfaceMap->begin(); it != PiESimdSurfaceMap->end();) {
1980  auto Mem = it->second;
1981  if (Mem != nullptr) {
1982  delete Mem;
1983  } // else { /* Null-entry for SLM_BTI */ }
1984  it = PiESimdSurfaceMap->erase(it);
1985  }
1986  return PI_SUCCESS;
1987 }
1988 
1990  if (PluginInit == nullptr) {
1991  return PI_INVALID_VALUE;
1992  }
1993 
1994  size_t PluginVersionSize = sizeof(PluginInit->PluginVersion);
1995  if (strlen(_PI_H_VERSION_STRING) >= PluginVersionSize) {
1996  return PI_INVALID_VALUE;
1997  }
1998  strncpy(PluginInit->PluginVersion, _PI_H_VERSION_STRING, PluginVersionSize);
1999 
2000  PiESimdDeviceAccess = new sycl::detail::ESIMDEmuPluginOpaqueData();
2001  // 'version' to be compared with 'ESIMD_EMULATOR_DEVICE_REQUIRED_VER' defined
2002  // in device interface file
2004  PiESimdDeviceAccess->data =
2005  reinterpret_cast<void *>(new sycl::detail::ESIMDDeviceInterface());
2006 
2007  // Registering pre-defined surface index dedicated for SLM
2008  (*PiESimdSurfaceMap)[__ESIMD_DNS::SLM_BTI] = nullptr;
2009 
2010 #define _PI_API(api) \
2011  (PluginInit->PiFunctionTable).api = (decltype(&::api))(&api);
2012 #include <CL/sycl/detail/pi.def>
2013 
2014  return PI_SUCCESS;
2015 }
2016 
2017 } // extern C
piextDeviceSelectBinary
pi_result piextDeviceSelectBinary(pi_device, pi_device_binary *, pi_uint32 RawImgSize, pi_uint32 *ImgInd)
Selects the most appropriate device binary based on runtime information and the IR characteristics.
Definition: pi_esimd_emulator.cpp:1947
setErrorMessage
static void setErrorMessage(const char *message, pi_result error_code)
Definition: pi_esimd_emulator.cpp:150
PI_IMAGE_FORMAT_NOT_SUPPORTED
@ PI_IMAGE_FORMAT_NOT_SUPPORTED
Definition: pi.h:118
PI_DEVICE_INFO_HOST_UNIFIED_MEMORY
@ PI_DEVICE_INFO_HOST_UNIFIED_MEMORY
Definition: pi.h:259
cm_surface_ptr_t::TypeUserProvidedImage
@ TypeUserProvidedImage
Definition: pi_esimd_emulator.hpp:119
piextUSMSharedAlloc
pi_result piextUSMSharedAlloc(void **ResultPtr, pi_context Context, pi_device Device, pi_usm_mem_properties *Properties, size_t Size, pi_uint32 Alignment)
Allocates memory accessible on both host and device.
Definition: pi_esimd_emulator.cpp:1850
PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_CHAR
@ PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_CHAR
Definition: pi.h:209
PI_FP_ROUND_TO_ZERO
@ PI_FP_ROUND_TO_ZERO
Definition: pi.h:384
PI_DEVICE_INFO_OPENCL_C_VERSION
@ PI_DEVICE_INFO_OPENCL_C_VERSION
Definition: pi.h:279
pi_image_region_struct::height
size_t height
Definition: pi.h:884
_pi_mem
PI Mem mapping to CUDA memory allocations, both data and texture/surface.
Definition: pi_cuda.hpp:207
piPluginInit
pi_result piPluginInit(pi_plugin *PluginInit)
Definition: pi_esimd_emulator.cpp:1989
_pi_buffer::MapHostPtr
char * MapHostPtr
Definition: pi_level_zero.hpp:1095
PI_USM_ACCESS
@ PI_USM_ACCESS
Definition: pi.h:1634
PI_DEVICE_INFO_PROFILE
@ PI_DEVICE_INFO_PROFILE
Definition: pi.h:277
PI_INVALID_KERNEL
@ PI_INVALID_KERNEL
Definition: pi.h:89
piextPluginGetOpaqueData
pi_result piextPluginGetOpaqueData(void *, void **OpaqueDataReturn)
API to get Plugin internal data, opaque to SYCL RT.
Definition: pi_esimd_emulator.cpp:1969
PI_DEVICE_INFO_IMAGE3D_MAX_WIDTH
@ PI_DEVICE_INFO_IMAGE3D_MAX_WIDTH
Definition: pi.h:241
PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_CHAR
@ PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_CHAR
Definition: pi.h:223
PI_SUCCESS
@ PI_SUCCESS
Definition: pi.h:86
_pi_device::CmDevicePtr
cm_support::CmDevice * CmDevicePtr
Definition: pi_esimd_emulator.hpp:81
piProgramGetInfo
pi_result piProgramGetInfo(pi_program, pi_program_info, size_t, void *, size_t *)
Definition: pi_esimd_emulator.cpp:1308
PI_DEVICE_INFO_DOUBLE_FP_CONFIG
@ PI_DEVICE_INFO_DOUBLE_FP_CONFIG
Definition: pi.h:207
piDeviceGetInfo
pi_result piDeviceGetInfo(pi_device Device, pi_device_info ParamName, size_t ParamValueSize, void *ParamValue, size_t *ParamValueSizeRet)
Returns requested info for provided native device Return PI_DEVICE_INFO_EXTENSION_DEVICELIB_ASSERT fo...
Definition: pi_esimd_emulator.cpp:591
_pi_context_info
_pi_context_info
Definition: pi.h:340
PI_MEM_TYPE_IMAGE1D
@ PI_MEM_TYPE_IMAGE1D
Definition: pi.h:454
pi_buff_rect_offset_struct
Definition: pi.h:855
piContextCreate
pi_result piContextCreate(const pi_context_properties *Properties, pi_uint32 NumDevices, const pi_device *Devices, void(*PFnNotify)(const char *ErrInfo, const void *PrivateInfo, size_t CB, void *UserData), void *UserData, pi_context *RetContext)
Definition: pi_esimd_emulator.cpp:827
_pi_image::Width
size_t Width
Definition: pi_esimd_emulator.hpp:188
piextProgramCreateWithNativeHandle
pi_result piextProgramCreateWithNativeHandle(pi_native_handle, pi_context, bool, pi_program *)
Creates PI program object from a native handle.
Definition: pi_esimd_emulator.cpp:1344
PI_DEVICE_INFO_DRIVER_VERSION
@ PI_DEVICE_INFO_DRIVER_VERSION
Definition: pi.h:276
piProgramCreate
pi_result piProgramCreate(pi_context, const void *, size_t, pi_program *)
Definition: pi_esimd_emulator.cpp:1286
piextUSMFree
pi_result piextUSMFree(pi_context Context, void *Ptr)
Indicates that the allocated USM memory is no longer needed on the runtime side.
Definition: pi_esimd_emulator.cpp:1889
InvokeKernel
void InvokeKernel(KernelInvocationContext< NDims > *ctx)
Definition: pi_esimd_emulator.cpp:204
PI_DEVICE_INFO_GPU_EU_COUNT
@ PI_DEVICE_INFO_GPU_EU_COUNT
Definition: pi.h:307
pi_bool
pi_uint32 pi_bool
Definition: pi.h:74
PI_IMAGE_CHANNEL_TYPE_UNORM_SHORT_565
@ PI_IMAGE_CHANNEL_TYPE_UNORM_SHORT_565
Definition: pi.h:497
IDBuilder
sycl::detail::Builder IDBuilder
Definition: pi_esimd_emulator.cpp:163
PI_IMAGE_CHANNEL_TYPE_HALF_FLOAT
@ PI_IMAGE_CHANNEL_TYPE_HALF_FLOAT
Definition: pi.h:506
PI_DEVICE_INFO_PREFERRED_INTEROP_USER_SYNC
@ PI_DEVICE_INFO_PREFERRED_INTEROP_USER_SYNC
Definition: pi.h:282
T
_pi_image_format::image_channel_data_type
pi_image_channel_type image_channel_data_type
Definition: pi.h:933
ErrorMessage
thread_local char ErrorMessage[MaxMessageSize]
Definition: pi_esimd_emulator.cpp:147
piextKernelSetArgSampler
pi_result piextKernelSetArgSampler(pi_kernel, pi_uint32, const pi_sampler *)
Definition: pi_esimd_emulator.cpp:1362
sycl_get_cm_image_params
void sycl_get_cm_image_params(unsigned int IndexInput, char **BaseAddr, uint32_t *Width, uint32_t *Height, uint32_t *Bpp, std::mutex **ImgMtxLock)
Definition: pi_esimd_emulator.cpp:291
piQueueFinish
pi_result piQueueFinish(pi_queue)
Definition: pi_esimd_emulator.cpp:984
ARG_UNUSED
#define ARG_UNUSED(x)
Definition: pi_esimd_emulator.cpp:46
PI_INVALID_OPERATION
@ PI_INVALID_OPERATION
Definition: pi.h:88
type_traits.hpp
PI_DEVICE_INFO_NAME
@ PI_DEVICE_INFO_NAME
Definition: pi.h:274
piProgramCreateWithBinary
pi_result piProgramCreateWithBinary(pi_context, pi_uint32, const pi_device *, const size_t *, const unsigned char **, size_t, const pi_device_binary_property *, pi_int32 *, pi_program *)
Creates a PI program for a context and loads the given binary into it.
Definition: pi_esimd_emulator.cpp:1290
KernelInvocationContext::Func
KernelFunc< NDims > Func
Definition: pi_esimd_emulator.cpp:173
PI_DEVICE_INFO_IL_VERSION
@ PI_DEVICE_INFO_IL_VERSION
Definition: pi.h:273
_pi_mem::Mappings
std::unordered_map< void *, Mapping > Mappings
Definition: pi_esimd_emulator.hpp:155
PI_FP_ROUND_TO_NEAREST
@ PI_FP_ROUND_TO_NEAREST
Definition: pi.h:383
PI_IMAGE_CHANNEL_TYPE_UNORM_INT8
@ PI_IMAGE_CHANNEL_TYPE_UNORM_INT8
Definition: pi.h:495
piclProgramCreateWithBinary
pi_result piclProgramCreateWithBinary(pi_context, pi_uint32, const pi_device *, const size_t *, const unsigned char **, pi_int32 *, pi_program *)
Definition: pi_esimd_emulator.cpp:1297
_pi_buffer
Definition: pi_esimd_emulator.hpp:171
piextDeviceGetNativeHandle
pi_result piextDeviceGetNativeHandle(pi_device, pi_native_handle *)
Gets the native handle of a PI device object.
Definition: pi_esimd_emulator.cpp:818
_pi_platform::PiDeviceCacheMutex
std::mutex PiDeviceCacheMutex
Definition: pi_esimd_emulator.hpp:63
PI_MEM_FLAGS_HOST_PTR_COPY
constexpr pi_mem_flags PI_MEM_FLAGS_HOST_PTR_COPY
Definition: pi.h:588
piextUSMEnqueueMemcpy
pi_result piextUSMEnqueueMemcpy(pi_queue, pi_bool, void *, const void *, size_t, pi_uint32, const pi_event *, pi_event *)
USM Memcpy API.
Definition: pi_esimd_emulator.cpp:1922
piKernelSetArg
pi_result piKernelSetArg(pi_kernel, pi_uint32, size_t, const void *)
Definition: pi_esimd_emulator.cpp:1353
_pi_image
Definition: pi_esimd_emulator.hpp:180
_pi_image_format::image_channel_order
pi_image_channel_order image_channel_order
Definition: pi.h:932
PI_DEVICE_INFO_PARTITION_AFFINITY_DOMAIN
@ PI_DEVICE_INFO_PARTITION_AFFINITY_DOMAIN
Definition: pi.h:288
PI_DEVICE_INFO_IMAGE_SRGB
@ PI_DEVICE_INFO_IMAGE_SRGB
Definition: pi.h:313
_pi_image_desc::image_type
pi_mem_type image_type
Definition: pi.h:937
piEventCreate
pi_result piEventCreate(pi_context, pi_event *)
Definition: pi_esimd_emulator.cpp:1385
_pi_plugin
Definition: pi.h:1822
PI_INVALID_MEM_OBJECT
@ PI_INVALID_MEM_OBJECT
Definition: pi.h:102
RangeBuilder
Definition: pi_esimd_emulator.cpp:183
PI_DEVICE_INFO_GLOBAL_MEM_CACHE_TYPE
@ PI_DEVICE_INFO_GLOBAL_MEM_CACHE_TYPE
Definition: pi.h:249
_pi_mem_advice
_pi_mem_advice
Definition: pi.h:459
piSamplerRelease
pi_result piSamplerRelease(pi_sampler)
Definition: pi_esimd_emulator.cpp:1483
piSamplerGetInfo
pi_result piSamplerGetInfo(pi_sampler, pi_sampler_info, size_t, void *, size_t *)
Definition: pi_esimd_emulator.cpp:1476
piDeviceRelease
pi_result piDeviceRelease(pi_device Device)
Definition: pi_esimd_emulator.cpp:581
RangeBuilder< 1 >::create
static sycl::range< 1 > create(Gen G)
Definition: pi_esimd_emulator.cpp:186
libCMBatch::libCMBatch
libCMBatch(const KernelFunc< DIMS > &Kernel)
Definition: pi_esimd_emulator.cpp:245
PI_DEVICE_INFO_MAX_WORK_ITEM_SIZES
@ PI_DEVICE_INFO_MAX_WORK_ITEM_SIZES
Definition: pi.h:203
cm_surface_ptr_t
Definition: pi_esimd_emulator.hpp:110
piEnqueueEventsWaitWithBarrier
pi_result piEnqueueEventsWaitWithBarrier(pi_queue, pi_uint32, const pi_event *, pi_event *)
Definition: pi_esimd_emulator.cpp:1490
PI_DEVICE_INFO_MAX_COMPUTE_UNITS
@ PI_DEVICE_INFO_MAX_COMPUTE_UNITS
Definition: pi.h:201
PrintPiTrace
static bool PrintPiTrace
Definition: pi_esimd_emulator.cpp:113
piEventSetCallback
pi_result piEventSetCallback(pi_event, pi_int32, void(*)(pi_event, pi_int32, void *), void *)
Definition: pi_esimd_emulator.cpp:1425
piEnqueueMemUnmap
pi_result piEnqueueMemUnmap(pi_queue Queue, pi_mem MemObj, void *MappedPtr, pi_uint32 NumEventsInWaitList, const pi_event *EventWaitList, pi_event *Event)
Definition: pi_esimd_emulator.cpp:1636
_pi_result
_pi_result
Definition: pi.h:85
cl::__ESIMD_NS::barrier
__ESIMD_API void barrier()
Generic work-group barrier.
Definition: memory.hpp:891
InvokeImpl
Definition: pi_esimd_emulator.cpp:346
piProgramGetBuildInfo
pi_result piProgramGetBuildInfo(pi_program, pi_device, cl_program_build_info, size_t, void *, size_t *)
Definition: pi_esimd_emulator.cpp:1331
piextProgramGetNativeHandle
pi_result piextProgramGetNativeHandle(pi_program, pi_native_handle *)
Gets the native handle of a PI program object.
Definition: pi_esimd_emulator.cpp:1340
ErrorMessageCode
thread_local pi_result ErrorMessageCode
Definition: pi_esimd_emulator.cpp:146
pi_context_properties
intptr_t pi_context_properties
Definition: pi.h:541
pi_esimd_emulator.hpp
PI_IMAGE_CHANNEL_TYPE_SIGNED_INT8
@ PI_IMAGE_CHANNEL_TYPE_SIGNED_INT8
Definition: pi.h:500
PI_FP_INF_NAN
@ PI_FP_INF_NAN
Definition: pi.h:386
PI_INVALID_EVENT_WAIT_LIST
@ PI_INVALID_EVENT_WAIT_LIST
Definition: pi.h:105
PI_DEVICE_INFO_GPU_SLICES
@ PI_DEVICE_INFO_GPU_SLICES
Definition: pi.h:309
PI_DEVICE_INFO_MAX_MEM_ALLOC_SIZE
@ PI_DEVICE_INFO_MAX_MEM_ALLOC_SIZE
Definition: pi.h:235
_pi_device_type
_pi_device_type
Definition: pi.h:174
piextPlatformCreateWithNativeHandle
pi_result piextPlatformCreateWithNativeHandle(pi_native_handle, pi_platform *)
Creates PI platform object from a native handle.
Definition: pi_esimd_emulator.cpp:475
_pi_event::CmEventPtr
cm_support::CmEvent * CmEventPtr
Definition: pi_esimd_emulator.hpp:196
piDevicePartition
pi_result piDevicePartition(pi_device, const pi_device_partition_property *, pi_uint32, pi_device *, pi_uint32 *)
Definition: pi_esimd_emulator.cpp:813
PI_DEVICE_INFO_REFERENCE_COUNT
@ PI_DEVICE_INFO_REFERENCE_COUNT
Definition: pi.h:272
PI_DEVICE_INFO_USM_CROSS_SHARED_SUPPORT
@ PI_DEVICE_INFO_USM_CROSS_SHARED_SUPPORT
Definition: pi.h:299
PI_DEVICE_INFO_PLATFORM
@ PI_DEVICE_INFO_PLATFORM
Definition: pi.h:271
PI_DEVICE_INFO_GPU_EU_COUNT_PER_SUBSLICE
@ PI_DEVICE_INFO_GPU_EU_COUNT_PER_SUBSLICE
Definition: pi.h:311
PI_DEVICE_LOCAL_MEM_TYPE_LOCAL
@ PI_DEVICE_LOCAL_MEM_TYPE_LOCAL
Definition: pi.h:194
helpers.hpp
host_profiling_info.hpp
_pi_device::VersionStr
std::string VersionStr
Definition: pi_esimd_emulator.hpp:83
piEventGetInfo
pi_result piEventGetInfo(pi_event, pi_event_info, size_t, void *, size_t *)
Definition: pi_esimd_emulator.cpp:1387
PI_IMAGE_CHANNEL_TYPE_SNORM_INT16
@ PI_IMAGE_CHANNEL_TYPE_SNORM_INT16
Definition: pi.h:494
PI_DEVICE_INFO_ADDRESS_BITS
@ PI_DEVICE_INFO_ADDRESS_BITS
Definition: pi.h:234
_pi_image_desc::image_height
size_t image_height
Definition: pi.h:939
_pi_platform
A PI platform stores all known PI devices, in the CUDA plugin this is just a vector of available devi...
Definition: pi_cuda.hpp:63
piPluginGetLastError
pi_result piPluginGetLastError(char **message)
API to get Plugin specific warning and error messages.
Definition: pi_esimd_emulator.cpp:158
CASE_PI_UNSUPPORTED
#define CASE_PI_UNSUPPORTED(not_supported)
Definition: pi_esimd_emulator.cpp:384
PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_DOUBLE
@ PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_DOUBLE
Definition: pi.h:219
PI_DEVICE_INFO_MAX_SAMPLERS
@ PI_DEVICE_INFO_MAX_SAMPLERS
Definition: pi.h:246
pi_image_offset_struct::y
size_t y
Definition: pi.h:875
PI_DEVICE_INFO_MAX_NUM_SUB_GROUPS
@ PI_DEVICE_INFO_MAX_NUM_SUB_GROUPS
Definition: pi.h:291
PI_DEVICE_INFO_MEM_BASE_ADDR_ALIGN
@ PI_DEVICE_INFO_MEM_BASE_ADDR_ALIGN
Definition: pi.h:248
PI_DEVICE_INFO_BUILT_IN_KERNELS
@ PI_DEVICE_INFO_BUILT_IN_KERNELS
Definition: pi.h:270
PI_DEVICE_INFO_USM_DEVICE_SUPPORT
@ PI_DEVICE_INFO_USM_DEVICE_SUPPORT
Definition: pi.h:296
PI_PLATFORM_INFO_NAME
@ PI_PLATFORM_INFO_NAME
Definition: pi.h:143
PI_DEVICE_INFO_IMAGE_MAX_ARRAY_SIZE
@ PI_DEVICE_INFO_IMAGE_MAX_ARRAY_SIZE
Definition: pi.h:245
PI_FP_ROUND_TO_INF
@ PI_FP_ROUND_TO_INF
Definition: pi.h:385
_pi_device_info
_pi_device_info
Definition: pi.h:198
PI_USM_ATOMIC_ACCESS
@ PI_USM_ATOMIC_ACCESS
Definition: pi.h:1635
PI_DEVICE_INFO_MAX_CONSTANT_BUFFER_SIZE
@ PI_DEVICE_INFO_MAX_CONSTANT_BUFFER_SIZE
Definition: pi.h:254
pi_image_offset_struct::x
size_t x
Definition: pi.h:874
_pi_image_info
_pi_image_info
Definition: pi.h:391
group.hpp
libCMBatch::runIterationSpace
void runIterationSpace(const sycl::range< DIMS > &LocalSize, const sycl::range< DIMS > &GlobalSize, const sycl::id< DIMS > &GlobalOffset)
Definition: pi_esimd_emulator.cpp:248
piQueueRetain
pi_result piQueueRetain(pi_queue Queue)
Definition: pi_esimd_emulator.cpp:962
piextMemCreateWithNativeHandle
pi_result piextMemCreateWithNativeHandle(pi_native_handle, pi_context, bool, pi_mem *)
Creates PI mem object from a native handle.
Definition: pi_esimd_emulator.cpp:1281
id.hpp
isNull
static bool isNull(int NDims, const size_t *R)
Implementation for Host Kernel Launch used by piEnqueueKernelLaunch.
Definition: pi_esimd_emulator.cpp:334
PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_INT
@ PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_INT
Definition: pi.h:213
_pi_kernel
Implementation of a PI Kernel for CUDA.
Definition: pi_cuda.hpp:624
PI_DEVICE_INFO_ERROR_CORRECTION_SUPPORT
@ PI_DEVICE_INFO_ERROR_CORRECTION_SUPPORT
Definition: pi.h:258
piextKernelCreateWithNativeHandle
pi_result piextKernelCreateWithNativeHandle(pi_native_handle, pi_context, pi_program, bool, pi_kernel *)
Creates PI kernel object from a native handle.
Definition: pi_esimd_emulator.cpp:1820
PI_DEVICE_INFO_IMAGE2D_MAX_HEIGHT
@ PI_DEVICE_INFO_IMAGE2D_MAX_HEIGHT
Definition: pi.h:240
PI_INVALID_EVENT
@ PI_INVALID_EVENT
Definition: pi.h:104
piEventsWait
pi_result piEventsWait(pi_uint32 NumEvents, const pi_event *EventList)
Definition: pi_esimd_emulator.cpp:1407
PI_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE
constexpr pi_queue_properties PI_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE
Definition: pi.h:624
PiPlatformCachePopulated
static bool PiPlatformCachePopulated
Definition: pi_level_zero.cpp:440
piextQueueGetNativeHandle
pi_result piextQueueGetNativeHandle(pi_queue, pi_native_handle *)
Gets the native handle of a PI queue object.
Definition: pi_esimd_emulator.cpp:998
PI_DEVICE_INFO_GLOBAL_MEM_CACHE_SIZE
@ PI_DEVICE_INFO_GLOBAL_MEM_CACHE_SIZE
Definition: pi.h:252
PI_DEVICE_EXEC_CAPABILITIES_KERNEL
constexpr pi_device_exec_capabilities PI_DEVICE_EXEC_CAPABILITIES_KERNEL
Definition: pi.h:544
piextDeviceCreateWithNativeHandle
pi_result piextDeviceCreateWithNativeHandle(pi_native_handle, pi_platform, pi_device *)
Creates PI device object from a native handle.
Definition: pi_esimd_emulator.cpp:822
cl::sycl::detail::memcpy
void memcpy(void *Dst, const void *Src, std::size_t Size)
piMemBufferPartition
pi_result piMemBufferPartition(pi_mem, pi_mem_flags, pi_buffer_create_type, void *, pi_mem *)
Definition: pi_esimd_emulator.cpp:1753
cm_surface_ptr_t::RegularBufPtr
cm_support::CmBuffer * RegularBufPtr
Definition: pi_esimd_emulator.hpp:124
_pi_mem::Context
pi_context Context
Definition: pi_esimd_emulator.hpp:132
piEnqueueNativeKernel
pi_result piEnqueueNativeKernel(pi_queue, void(*)(void *), void *, size_t, pi_uint32, const pi_mem *, const void **, pi_uint32, const pi_event *, pi_event *)
Definition: pi_esimd_emulator.cpp:1829
_pi_queue_info
_pi_queue_info
Definition: pi.h:350
_pi_mem::~_pi_mem
~_pi_mem()
Definition: pi_cuda.hpp:348
_pi_buffer_create_type
_pi_buffer_create_type
Definition: pi.h:510
PI_ERROR_UNKNOWN
@ PI_ERROR_UNKNOWN
Definition: pi.h:131
piEnqueueMemImageRead
pi_result piEnqueueMemImageRead(pi_queue CommandQueue, pi_mem Image, pi_bool BlockingRead, pi_image_offset Origin, pi_image_region Region, size_t RowPitch, size_t SlicePitch, void *Ptr, pi_uint32 NumEventsInWaitList, const pi_event *EventWaitList, pi_event *Event)
Definition: pi_esimd_emulator.cpp:1678
piKernelGetSubGroupInfo
pi_result piKernelGetSubGroupInfo(pi_kernel, pi_device, pi_kernel_sub_group_info, size_t, const void *, size_t, void *, size_t *)
API to query information from the sub-group from a kernel.
Definition: pi_esimd_emulator.cpp:1375
piextUSMEnqueueMemAdvise
pi_result piextUSMEnqueueMemAdvise(pi_queue, const void *, size_t, pi_mem_advice, pi_event *)
USM Memadvise API.
Definition: pi_esimd_emulator.cpp:1927
export.hpp
PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_HALF
@ PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_HALF
Definition: pi.h:232
piDevicesGet
pi_result piDevicesGet(pi_platform Platform, pi_device_type DeviceType, pi_uint32 NumEntries, pi_device *Devices, pi_uint32 *NumDevices)
Definition: pi_esimd_emulator.cpp:479
cm_surface_ptr_t::RegularImgPtr
cm_support::CmSurface2D * RegularImgPtr
Definition: pi_esimd_emulator.hpp:126
piEnqueueMemBufferReadRect
pi_result piEnqueueMemBufferReadRect(pi_queue, pi_mem, pi_bool, pi_buff_rect_offset, pi_buff_rect_offset, pi_buff_rect_region, size_t, size_t, size_t, size_t, void *, pi_uint32, const pi_event *, pi_event *)
Definition: pi_esimd_emulator.cpp:1548
PI_DEVICE_INFO_IMAGE_MAX_BUFFER_SIZE
@ PI_DEVICE_INFO_IMAGE_MAX_BUFFER_SIZE
Definition: pi.h:244
piEnqueueMemBufferWrite
pi_result piEnqueueMemBufferWrite(pi_queue, pi_mem, pi_bool, size_t, size_t, const void *, pi_uint32, const pi_event *, pi_event *)
Definition: pi_esimd_emulator.cpp:1556
PI_DEVICE_INFO_LOCAL_MEM_SIZE
@ PI_DEVICE_INFO_LOCAL_MEM_SIZE
Definition: pi.h:257
piKernelSetExecInfo
pi_result piKernelSetExecInfo(pi_kernel, pi_kernel_exec_info, size_t, const void *)
API to set attributes controlling kernel execution.
Definition: pi_esimd_emulator.cpp:1937
piSamplerCreate
pi_result piSamplerCreate(pi_context, const pi_sampler_properties *, pi_sampler *)
Definition: pi_esimd_emulator.cpp:1471
piEnqueueMemBufferMap
pi_result piEnqueueMemBufferMap(pi_queue Queue, pi_mem MemObj, pi_bool BlockingMap, pi_map_flags MapFlags, size_t Offset, size_t Size, pi_uint32 NumEventsInWaitList, const pi_event *EventWaitList, pi_event *Event, void **RetMap)
Definition: pi_esimd_emulator.cpp:1590
_pi_queue
PI queue mapping on to CUstream objects.
Definition: pi_cuda.hpp:378
_pi_context::Addr2CmBufferSVMLock
std::mutex Addr2CmBufferSVMLock
Definition: pi_esimd_emulator.hpp:96
PiPlatformCacheLock
static std::mutex * PiPlatformCacheLock
Definition: pi_esimd_emulator.cpp:125
PI_IMAGE_CHANNEL_TYPE_SIGNED_INT16
@ PI_IMAGE_CHANNEL_TYPE_SIGNED_INT16
Definition: pi.h:501
piextKernelSetArgMemObj
pi_result piextKernelSetArgMemObj(pi_kernel, pi_uint32, const pi_mem *)
Definition: pi_esimd_emulator.cpp:1357
PI_IMAGE_CHANNEL_TYPE_SIGNED_INT32
@ PI_IMAGE_CHANNEL_TYPE_SIGNED_INT32
Definition: pi.h:502
_pi_mem::MapHostPtr
char * MapHostPtr
Definition: pi_esimd_emulator.hpp:135
PI_DEVICE_TYPE_GPU
@ PI_DEVICE_TYPE_GPU
A PI device that is a GPU.
Definition: pi.h:182
pi_uint32
uint32_t pi_uint32
Definition: pi.h:72
piProgramCompile
pi_result piProgramCompile(pi_program, pi_uint32, const pi_device *, const char *, pi_uint32, const pi_program *, const char **, void(*)(pi_program, void *), void *)
Definition: pi_esimd_emulator.cpp:1319
cm_surface_ptr_t::tag
SurfaceType tag
Definition: pi_esimd_emulator.hpp:121
kernel.hpp
pi_buff_rect_region_struct
Definition: pi.h:864
pi_device_binary_struct
This struct is a record of the device binary information.
Definition: pi.h:798
cl::__ESIMD_NS::fence
__ESIMD_API void fence()
esimd::fence sets the memory read/write order.
Definition: memory.hpp:878
piextMemGetNativeHandle
pi_result piextMemGetNativeHandle(pi_mem, pi_native_handle *)
Gets the native handle of a PI mem object.
Definition: pi_esimd_emulator.cpp:1277
KernelInvocationContext
Definition: pi_esimd_emulator.cpp:172
piextContextGetNativeHandle
pi_result piextContextGetNativeHandle(pi_context, pi_native_handle *)
Gets the native handle of a PI context object.
Definition: pi_esimd_emulator.cpp:868
PI_DEVICE_INFO_LINKER_AVAILABLE
@ PI_DEVICE_INFO_LINKER_AVAILABLE
Definition: pi.h:265
PI_DEVICE_INFO_EXECUTION_CAPABILITIES
@ PI_DEVICE_INFO_EXECUTION_CAPABILITIES
Definition: pi.h:266
piPlatformGetInfo
pi_result piPlatformGetInfo(pi_platform Platform, pi_platform_info ParamName, size_t ParamValueSize, void *ParamValue, size_t *ParamValueSizeRet)
Definition: pi_esimd_emulator.cpp:439
PiESimdSurfaceMap
static std::unordered_map< unsigned int, _pi_mem * > * PiESimdSurfaceMap
Definition: pi_esimd_emulator.cpp:128
PI_DEVICE_INFO_PROFILING_TIMER_RESOLUTION
@ PI_DEVICE_INFO_PROFILING_TIMER_RESOLUTION
Definition: pi.h:260
_pi_kernel_exec_info
_pi_kernel_exec_info
Definition: pi.h:1325
_pi_mem::SurfacePtr
cm_surface_ptr_t SurfacePtr
Definition: pi_esimd_emulator.hpp:159
PI_MEM_TYPE_IMAGE2D
@ PI_MEM_TYPE_IMAGE2D
Definition: pi.h:451
_pi_platform::PiDeviceCache
std::unique_ptr< _pi_device > PiDeviceCache
Definition: pi_esimd_emulator.hpp:62
piContextGetInfo
pi_result piContextGetInfo(pi_context, pi_context_info, size_t, void *, size_t *)
Definition: pi_esimd_emulator.cpp:858
range.hpp
cl::sycl::detail::pi::die
void die(const char *Message)
Definition: pi.cpp:537
PI_DEVICE_INFO_HALF_FP_CONFIG
@ PI_DEVICE_INFO_HALF_FP_CONFIG
Definition: pi.h:206
piEnqueueMemBufferWriteRect
pi_result piEnqueueMemBufferWriteRect(pi_queue, pi_mem, pi_bool, pi_buff_rect_offset, pi_buff_rect_offset, pi_buff_rect_region, size_t, size_t, size_t, size_t, const void *, pi_uint32, const pi_event *, pi_event *)
Definition: pi_esimd_emulator.cpp:1562
pi_mem_flags
pi_bitfield pi_mem_flags
Definition: pi.h:582
PI_IMAGE_CHANNEL_TYPE_FLOAT
@ PI_IMAGE_CHANNEL_TYPE_FLOAT
Definition: pi.h:507
piextProgramSetSpecializationConstant
pi_result piextProgramSetSpecializationConstant(pi_program, pi_uint32, size_t, const void *)
Sets a specialization constant to a specific value.
Definition: pi_esimd_emulator.cpp:1942
PI_USM_CONCURRENT_ACCESS
@ PI_USM_CONCURRENT_ACCESS
Definition: pi.h:1636
_pi_platform::DeviceCachePopulated
bool DeviceCachePopulated
Definition: pi_esimd_emulator.hpp:64
PI_DEVICE_INFO_VENDOR
@ PI_DEVICE_INFO_VENDOR
Definition: pi.h:275
pi_context_extended_deleter
void(* pi_context_extended_deleter)(void *user_data)
Definition: pi.h:1091
PI_MEM_TYPE_IMAGE1D_ARRAY
@ PI_MEM_TYPE_IMAGE1D_ARRAY
Definition: pi.h:455
PI_INVALID_QUEUE_PROPERTIES
@ PI_INVALID_QUEUE_PROPERTIES
Definition: pi.h:90
piextGetDeviceFunctionPointer
pi_result piextGetDeviceFunctionPointer(pi_device, pi_program, const char *, pi_uint64 *)
Retrieves a device function pointer to a user-defined function.
Definition: pi_esimd_emulator.cpp:1835
piMemGetInfo
pi_result piMemGetInfo(pi_mem, pi_mem_info, size_t, void *, size_t *)
Definition: pi_esimd_emulator.cpp:1081
piextUSMEnqueuePrefetch
pi_result piextUSMEnqueuePrefetch(pi_queue, const void *, size_t, pi_usm_migration_flags, pi_uint32, const pi_event *, pi_event *)
Hint to migrate memory to the device.
Definition: pi_esimd_emulator.cpp:1963
cl::__ESIMD_NS::detail::SLM_BTI
static constexpr SurfaceIndex SLM_BTI
Definition: common.hpp:116
InvokeImpl::invoke
static void invoke(pi_kernel Kernel, const size_t *GlobalWorkOffset, const size_t *GlobalWorkSize, const size_t *LocalWorkSize)
Definition: pi_esimd_emulator.cpp:357
PI_DEVICE_INFO_USM_SYSTEM_SHARED_SUPPORT
@ PI_DEVICE_INFO_USM_SYSTEM_SHARED_SUPPORT
Definition: pi.h:301
PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_SHORT
@ PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_SHORT
Definition: pi.h:211
cl::__ESIMD_NS::SurfaceIndex
unsigned int SurfaceIndex
Surface index type.
Definition: common.hpp:105
_pi_device::Platform
pi_platform Platform
Definition: pi_esimd_emulator.hpp:78
piextUSMGetMemAllocInfo
pi_result piextUSMGetMemAllocInfo(pi_context, const void *, pi_mem_alloc_info, size_t, void *, size_t *)
API to query information about USM allocated pointers Valid Queries: PI_MEM_ALLOC_TYPE returns host/d...
Definition: pi_esimd_emulator.cpp:1932
ESIMDEmuPluginDataVersion
#define ESIMDEmuPluginDataVersion
Definition: pi_esimd_emulator.cpp:135
piKernelGetInfo
pi_result piKernelGetInfo(pi_kernel, pi_kernel_info, size_t, void *, size_t *)
Definition: pi_esimd_emulator.cpp:1366
PI_DEVICE_INFO_GLOBAL_MEM_CACHELINE_SIZE
@ PI_DEVICE_INFO_GLOBAL_MEM_CACHELINE_SIZE
Definition: pi.h:250
PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_LONG
@ PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_LONG
Definition: pi.h:227
piextKernelSetArgPointer
pi_result piextKernelSetArgPointer(pi_kernel, pi_uint32, size_t, const void *)
Sets up pointer arguments for CL kernels.
Definition: pi_esimd_emulator.cpp:1913
piextEventGetNativeHandle
pi_result piextEventGetNativeHandle(pi_event, pi_native_handle *)
Gets the native handle of a PI event object.
Definition: pi_esimd_emulator.cpp:1463
_pi_platform::CmEmuVersion
std::string CmEmuVersion
Definition: pi_esimd_emulator.hpp:70
piextContextSetExtendedDeleter
pi_result piextContextSetExtendedDeleter(pi_context, pi_context_extended_deleter, void *)
Definition: pi_esimd_emulator.cpp:863
PI_DEVICE_INFO_USM_HOST_SUPPORT
@ PI_DEVICE_INFO_USM_HOST_SUPPORT
Definition: pi.h:295
cl::__ESIMD_ENS::split_barrier
__ESIMD_API void split_barrier()
Generic work-group split barrier.
Definition: memory.hpp:28
PI_INVALID_QUEUE
@ PI_INVALID_QUEUE
Definition: pi.h:96
PI_MEM_FLAGS_HOST_PTR_USE
constexpr pi_mem_flags PI_MEM_FLAGS_HOST_PTR_USE
Definition: pi.h:587
PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_FLOAT
@ PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_FLOAT
Definition: pi.h:228
piEnqueueMemBufferCopyRect
pi_result piEnqueueMemBufferCopyRect(pi_queue, pi_mem, pi_mem, pi_buff_rect_offset, pi_buff_rect_offset, pi_buff_rect_region, size_t, size_t, size_t, size_t, pi_uint32, const pi_event *, pi_event *)
Definition: pi_esimd_emulator.cpp:1576
piTearDown
pi_result piTearDown(void *)
API to notify that the plugin should clean up its resources.
Definition: pi_esimd_emulator.cpp:1974
ESIMDEmuPluginInterfaceVersion
#define ESIMDEmuPluginInterfaceVersion
Definition: pi_esimd_emulator.cpp:139
piEnqueueEventsWait
pi_result piEnqueueEventsWait(pi_queue, pi_uint32, const pi_event *, pi_event *)
Definition: pi_esimd_emulator.cpp:1485
PiESimdDeviceAccess
static sycl::detail::ESIMDEmuPluginOpaqueData * PiESimdDeviceAccess
Definition: pi_esimd_emulator.cpp:120
PI_DEVICE_INFO_PARENT_DEVICE
@ PI_DEVICE_INFO_PARENT_DEVICE
Definition: pi.h:284
PI_PLATFORM_INFO_VERSION
@ PI_PLATFORM_INFO_VERSION
Definition: pi.h:146
PI_IMAGE_CHANNEL_TYPE_UNSIGNED_INT16
@ PI_IMAGE_CHANNEL_TYPE_UNSIGNED_INT16
Definition: pi.h:504
pi_uint64
uint64_t pi_uint64
Definition: pi.h:73
PI_DEVICE_INFO_PARTITION_TYPE
@ PI_DEVICE_INFO_PARTITION_TYPE
Definition: pi.h:290
_pi_event_info
_pi_event_info
Definition: pi.h:409
piMemBufferCreate
pi_result piMemBufferCreate(pi_context Context, pi_mem_flags Flags, size_t Size, void *HostPtr, pi_mem *RetMem, const pi_mem_properties *properties)
Definition: pi_esimd_emulator.cpp:1007
_pi_image::BytesPerPixel
size_t BytesPerPixel
Definition: pi_esimd_emulator.hpp:190
_pi_sampler_info
_pi_sampler_info
Definition: pi.h:517
_pi_device_binary_property_struct
Definition: pi.h:701
pi_mem_properties
pi_bitfield pi_mem_properties
Definition: pi.h:600
_pi_program
Implementation of PI Program on CUDA Module object.
Definition: pi_cuda.hpp:569
PI_FP_DENORM
@ PI_FP_DENORM
Definition: pi.h:387
_pi_sampler
Implementation of samplers for CUDA.
Definition: pi_cuda.hpp:795
PI_DEVICE_INFO_PRINTF_BUFFER_SIZE
@ PI_DEVICE_INFO_PRINTF_BUFFER_SIZE
Definition: pi.h:281
PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_DOUBLE
@ PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_DOUBLE
Definition: pi.h:230
pi_device_exec_capabilities
pi_bitfield pi_device_exec_capabilities
Definition: pi.h:543
PI_PLATFORM_INFO_PROFILE
@ PI_PLATFORM_INFO_PROFILE
Definition: pi.h:144
PI_DEVICE_INFO_MAX_CONSTANT_ARGS
@ PI_DEVICE_INFO_MAX_CONSTANT_ARGS
Definition: pi.h:255
PI_DEVICE_INFO_PCI_ADDRESS
@ PI_DEVICE_INFO_PCI_ADDRESS
Definition: pi.h:306
pi_native_handle
uintptr_t pi_native_handle
Definition: pi.h:76
PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_SHORT
@ PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_SHORT
Definition: pi.h:224
cl::sycl::detail::getNextPowerOfTwo
constexpr size_t getNextPowerOfTwo(size_t Var)
Definition: common.hpp:324
PI_DEVICE_INFO_GLOBAL_MEM_SIZE
@ PI_DEVICE_INFO_GLOBAL_MEM_SIZE
Definition: pi.h:253
InvokeImpl::get_range
static sycl::range< NDims > get_range(const size_t *Array)
Definition: pi_esimd_emulator.cpp:348
pi_sampler_properties
pi_bitfield pi_sampler_properties
Definition: pi.h:549
piKernelRetain
pi_result piKernelRetain(pi_kernel)
Definition: pi_esimd_emulator.cpp:1381
PI_DEVICE_MEM_CACHE_TYPE_READ_WRITE_CACHE
@ PI_DEVICE_MEM_CACHE_TYPE_READ_WRITE_CACHE
Definition: pi.h:190
PI_MEM_FLAGS_ACCESS_RW
constexpr pi_mem_flags PI_MEM_FLAGS_ACCESS_RW
Definition: pi.h:584
accessor_impl.hpp
_pi_event::release
pi_result release()
Definition: pi_cuda.cpp:534
cm_surface_ptr_t::TypeRegularImage
@ TypeRegularImage
Definition: pi_esimd_emulator.hpp:118
piclProgramCreateWithSource
pi_result piclProgramCreateWithSource(pi_context, pi_uint32, const char **, const size_t *, pi_program *)
Definition: pi_esimd_emulator.cpp:1303
_pi_image::Height
size_t Height
Definition: pi_esimd_emulator.hpp:189
_pi_image_format
Definition: pi.h:931
PI_IMAGE_CHANNEL_TYPE_UNORM_INT_101010
@ PI_IMAGE_CHANNEL_TYPE_UNORM_INT_101010
Definition: pi.h:499
PI_PLATFORM_INFO_VENDOR
@ PI_PLATFORM_INFO_VENDOR
Definition: pi.h:145
_pi_kernel_info
_pi_kernel_info
Definition: pi.h:359
piKernelRelease
pi_result piKernelRelease(pi_kernel)
Definition: pi_esimd_emulator.cpp:1383
PI_INVALID_CONTEXT
@ PI_INVALID_CONTEXT
Definition: pi.h:92
PI_DEVICE_INFO_USM_SINGLE_SHARED_SUPPORT
@ PI_DEVICE_INFO_USM_SINGLE_SHARED_SUPPORT
Definition: pi.h:297
PI_IMAGE_CHANNEL_ORDER_RGBA
@ PI_IMAGE_CHANNEL_ORDER_RGBA
Definition: pi.h:480
cm_surface_ptr_t::UPBufPtr
cm_support::CmBufferUP * UPBufPtr
Definition: pi_esimd_emulator.hpp:125
PI_DEVICE_INFO_QUEUE_PROPERTIES
@ PI_DEVICE_INFO_QUEUE_PROPERTIES
Definition: pi.h:208
PI_DEVICE_INFO_PARTITION_MAX_SUB_DEVICES
@ PI_DEVICE_INFO_PARTITION_MAX_SUB_DEVICES
Definition: pi.h:286
PI_EXT_ONEAPI_DEVICE_INFO_MAX_WORK_GROUPS_3D
@ PI_EXT_ONEAPI_DEVICE_INFO_MAX_WORK_GROUPS_3D
Definition: pi.h:324
piextPlatformGetNativeHandle
pi_result piextPlatformGetNativeHandle(pi_platform, pi_native_handle *)
Gets the native handle of a PI platform object.
Definition: pi_esimd_emulator.cpp:471
KernelFunc
std::function< void(const sycl::nd_item< NDims > &)> KernelFunc
Definition: pi_esimd_emulator.cpp:166
PI_MEM_TYPE_IMAGE1D_BUFFER
@ PI_MEM_TYPE_IMAGE1D_BUFFER
Definition: pi.h:456
PI_INVALID_IMAGE_FORMAT_DESCRIPTOR
@ PI_INVALID_IMAGE_FORMAT_DESCRIPTOR
Definition: pi.h:117
piProgramBuild
pi_result piProgramBuild(pi_program, pi_uint32, const pi_device *, const char *, void(*)(pi_program, void *), void *)
Definition: pi_esimd_emulator.cpp:1326
KernelInvocationContext::GlobalSize
const sycl::range< NDims > & GlobalSize
Definition: pi_esimd_emulator.cpp:175
_pi_image_desc::image_width
size_t image_width
Definition: pi.h:938
piextUSMEnqueueMemset
pi_result piextUSMEnqueueMemset(pi_queue, void *, pi_int32, size_t, pi_uint32, const pi_event *, pi_event *)
USM Memset API.
Definition: pi_esimd_emulator.cpp:1917
PI_DEVICE_INFO_SINGLE_FP_CONFIG
@ PI_DEVICE_INFO_SINGLE_FP_CONFIG
Definition: pi.h:205
cl::sycl::info::platform::version
@ version
PI_INVALID_VALUE
@ PI_INVALID_VALUE
Definition: pi.h:91
piextUSMDeviceAlloc
pi_result piextUSMDeviceAlloc(void **, pi_context, pi_device, pi_usm_mem_properties *, size_t, pi_uint32)
Allocates device memory.
Definition: pi_esimd_emulator.cpp:1845
PI_DEVICE_INFO_BUILD_ON_SUBDEVICE
@ PI_DEVICE_INFO_BUILD_ON_SUBDEVICE
Definition: pi.h:315
_pi_context::Addr2CmBufferSVM
std::unordered_map< void *, cm_support::CmBufferSVM * > Addr2CmBufferSVM
Definition: pi_esimd_emulator.hpp:94
PI_DEVICE_INFO_MAX_WORK_ITEM_DIMENSIONS
@ PI_DEVICE_INFO_MAX_WORK_ITEM_DIMENSIONS
Definition: pi.h:202
PI_DEVICE_INFO_LOCAL_MEM_TYPE
@ PI_DEVICE_INFO_LOCAL_MEM_TYPE
Definition: pi.h:256
PI_DEVICE_INFO_IMAGE3D_MAX_DEPTH
@ PI_DEVICE_INFO_IMAGE3D_MAX_DEPTH
Definition: pi.h:243
PI_DEVICE_INFO_COMPILER_AVAILABLE
@ PI_DEVICE_INFO_COMPILER_AVAILABLE
Definition: pi.h:264
_PI_H_VERSION_STRING
#define _PI_H_VERSION_STRING
Definition: pi.h:55
PiESimdSurfaceMapLock
static std::mutex * PiESimdSurfaceMapLock
Definition: pi_esimd_emulator.cpp:131
pi_image_offset_struct::z
size_t z
Definition: pi.h:876
backend_types.hpp
PI_DEVICE_INFO_VERSION
@ PI_DEVICE_INFO_VERSION
Definition: pi.h:278
PI_DEVICE_INFO_MAX_MEM_BANDWIDTH
@ PI_DEVICE_INFO_MAX_MEM_BANDWIDTH
Definition: pi.h:312
PI_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8
@ PI_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8
Definition: pi.h:503
PI_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32
@ PI_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32
Definition: pi.h:505
piEnqueueMemBufferFill
pi_result piEnqueueMemBufferFill(pi_queue, pi_mem, const void *, size_t, size_t, size_t, pi_uint32, const pi_event *, pi_event *)
Definition: pi_esimd_emulator.cpp:1584
pi_queue_properties
pi_bitfield pi_queue_properties
Definition: pi.h:623
PI_EXT_ONEAPI_DEVICE_INFO_MAX_WORK_GROUPS_2D
@ PI_EXT_ONEAPI_DEVICE_INFO_MAX_WORK_GROUPS_2D
Definition: pi.h:323
PI_DEVICE_INFO_ATOMIC_64
@ PI_DEVICE_INFO_ATOMIC_64
Definition: pi.h:316
PI_DEVICE_INFO_MAX_WRITE_IMAGE_ARGS
@ PI_DEVICE_INFO_MAX_WRITE_IMAGE_ARGS
Definition: pi.h:238
_pi_platform_info
_pi_platform_info
Definition: pi.h:141
MaxMessageSize
constexpr size_t MaxMessageSize
Definition: pi_esimd_emulator.cpp:145
PI_INVALID_WORK_GROUP_SIZE
@ PI_INVALID_WORK_GROUP_SIZE
Definition: pi.h:108
piEnqueueMemBufferCopy
pi_result piEnqueueMemBufferCopy(pi_queue, pi_mem, pi_mem, size_t, size_t, size_t, pi_uint32, const pi_event *, pi_event *)
Definition: pi_esimd_emulator.cpp:1570
piEnqueueMemBufferRead
pi_result piEnqueueMemBufferRead(pi_queue Queue, pi_mem Src, pi_bool BlockingRead, size_t Offset, size_t Size, void *Dst, pi_uint32 NumEventsInWaitList, const pi_event *EventWaitList, pi_event *Event)
Definition: pi_esimd_emulator.cpp:1495
PI_IMAGE_CHANNEL_TYPE_UNORM_SHORT_555
@ PI_IMAGE_CHANNEL_TYPE_UNORM_SHORT_555
Definition: pi.h:498
PI_IMAGE_CHANNEL_ORDER_ARGB
@ PI_IMAGE_CHANNEL_ORDER_ARGB
Definition: pi.h:482
_pi_image_desc
Definition: pi.h:936
PI_DEVICE_INFO_ENDIAN_LITTLE
@ PI_DEVICE_INFO_ENDIAN_LITTLE
Definition: pi.h:262
piQueueCreate
pi_result piQueueCreate(pi_context Context, pi_device Device, pi_queue_properties Properties, pi_queue *Queue)
Definition: pi_esimd_emulator.cpp:930
_pi_queue::CmQueuePtr
cm_support::CmQueue * CmQueuePtr
Definition: pi_esimd_emulator.hpp:107
piContextRelease
pi_result piContextRelease(pi_context Context)
Definition: pi_esimd_emulator.cpp:888
piQueueFlush
pi_result piQueueFlush(pi_queue)
Definition: pi_esimd_emulator.cpp:991
kernel_desc.hpp
_pi_event
PI Event mapping to CUevent.
Definition: pi_cuda.hpp:458
piQueueRelease
pi_result piQueueRelease(pi_queue Queue)
Definition: pi_esimd_emulator.cpp:970
DIE_NO_IMPLEMENTATION
#define DIE_NO_IMPLEMENTATION
Definition: pi_esimd_emulator.cpp:368
PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_INT
@ PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_INT
Definition: pi.h:226
PI_DEVICE_INFO_MAX_READ_IMAGE_ARGS
@ PI_DEVICE_INFO_MAX_READ_IMAGE_ARGS
Definition: pi.h:237
_pi_mem::MappingsMutex
std::mutex MappingsMutex
Definition: pi_esimd_emulator.hpp:157
PI_DEVICE_INFO_GPU_SUBSLICES_PER_SLICE
@ PI_DEVICE_INFO_GPU_SUBSLICES_PER_SLICE
Definition: pi.h:310
PI_OUT_OF_RESOURCES
@ PI_OUT_OF_RESOURCES
Definition: pi.h:103
piKernelCreate
pi_result piKernelCreate(pi_program, const char *, pi_kernel *)
Definition: pi_esimd_emulator.cpp:1349
piEventRelease
pi_result piEventRelease(pi_event Event)
Definition: pi_esimd_emulator.cpp:1442
ESimdEmuVersionString
static char ESimdEmuVersionString[32]
Definition: pi_esimd_emulator.cpp:142
PI_MEM_TYPE_IMAGE2D_ARRAY
@ PI_MEM_TYPE_IMAGE2D_ARRAY
Definition: pi.h:453
piMemRetain
pi_result piMemRetain(pi_mem Mem)
Definition: pi_esimd_emulator.cpp:1085
piKernelGetGroupInfo
pi_result piKernelGetGroupInfo(pi_kernel, pi_device, pi_kernel_group_info, size_t, void *, size_t *)
Definition: pi_esimd_emulator.cpp:1370
_pi_mem_info
_pi_mem_info
Definition: pi.h:952
pi_image_offset_struct
Definition: pi.h:873
_pi_plugin::PluginVersion
char PluginVersion[10]
Definition: pi.h:1832
piMemRelease
pi_result piMemRelease(pi_mem Mem)
Definition: pi_esimd_emulator.cpp:1093
PiPlatformCache
static pi_platform PiPlatformCache
Definition: pi_esimd_emulator.cpp:123
cm_surface_ptr_t::TypeRegularBuffer
@ TypeRegularBuffer
Definition: pi_esimd_emulator.hpp:116
PI_INVALID_PLATFORM
@ PI_INVALID_PLATFORM
Definition: pi.h:93
PI_OUT_OF_HOST_MEMORY
@ PI_OUT_OF_HOST_MEMORY
Definition: pi.h:97
PI_EXT_ONEAPI_DEVICE_INFO_MAX_WORK_GROUPS_1D
@ PI_EXT_ONEAPI_DEVICE_INFO_MAX_WORK_GROUPS_1D
Definition: pi.h:322
PI_DEVICE_INFO_VENDOR_ID
@ PI_DEVICE_INFO_VENDOR_ID
Definition: pi.h:200
common.hpp
piextEventCreateWithNativeHandle
pi_result piextEventCreateWithNativeHandle(pi_native_handle, pi_context, bool, pi_event *)
Creates PI event object from a native handle.
Definition: pi_esimd_emulator.cpp:1467
pi_image_region_struct
Definition: pi.h:882
PI_FP_CORRECTLY_ROUNDED_DIVIDE_SQRT
@ PI_FP_CORRECTLY_ROUNDED_DIVIDE_SQRT
Definition: pi.h:382
piDeviceRetain
pi_result piDeviceRetain(pi_device Device)
Definition: pi_esimd_emulator.cpp:571
PI_DEVICE_INFO_EXTENSIONS
@ PI_DEVICE_INFO_EXTENSIONS
Definition: pi.h:280
piextUSMHostAlloc
pi_result piextUSMHostAlloc(void **, pi_context, pi_usm_mem_properties *, size_t, pi_uint32)
Allocates host memory accessible by the device.
Definition: pi_esimd_emulator.cpp:1840
piProgramRetain
pi_result piProgramRetain(pi_program)
Definition: pi_esimd_emulator.cpp:1336
KernelInvocationContext::LocalSize
const sycl::range< NDims > & LocalSize
Definition: pi_esimd_emulator.cpp:174
PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_FLOAT
@ PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_FLOAT
Definition: pi.h:217
cm_surface_ptr_t::UPImgPtr
cm_support::CmSurface2DUP * UPImgPtr
Definition: pi_esimd_emulator.hpp:127
_pi_mem::SurfaceIndex
unsigned int SurfaceIndex
Definition: pi_esimd_emulator.hpp:141
_pi_usm_migration_flags
_pi_usm_migration_flags
Definition: pi.h:1658
KernelInvocationContext::GlobalOffset
const sycl::id< NDims > & GlobalOffset
Definition: pi_esimd_emulator.cpp:176
cm_surface_ptr_t::TypeUserProvidedBuffer
@ TypeUserProvidedBuffer
Definition: pi_esimd_emulator.hpp:117
common.hpp
piextContextCreateWithNativeHandle
pi_result piextContextCreateWithNativeHandle(pi_native_handle, pi_uint32, const pi_device *, bool, pi_context *)
Creates PI context object from a native handle.
Definition: pi_esimd_emulator.cpp:872
piEnqueueKernelLaunch
pi_result piEnqueueKernelLaunch(pi_queue Queue, pi_kernel Kernel, pi_uint32 WorkDim, const size_t *GlobalWorkOffset, const size_t *GlobalWorkSize, const size_t *LocalWorkSize, pi_uint32 NumEventsInWaitList, const pi_event *EventWaitList, pi_event *Event)
Definition: pi_esimd_emulator.cpp:1759
PI_FP_FMA
@ PI_FP_FMA
Definition: pi.h:388
sycl_get_cm_surface_index
unsigned int sycl_get_cm_surface_index(void *PtrInput)
Definition: pi_esimd_emulator.cpp:266
_pi_platform::populateDeviceCacheIfNeeded
pi_result populateDeviceCacheIfNeeded()
Definition: pi_esimd_emulator.cpp:523
_pi_event::IsDummyEvent
bool IsDummyEvent
Definition: pi_esimd_emulator.hpp:199
nd_item.hpp
PI_DEVICE_INFO_TYPE
@ PI_DEVICE_INFO_TYPE
Definition: pi.h:199
piMemImageGetInfo
pi_result piMemImageGetInfo(pi_mem, pi_image_info, size_t, void *, size_t *)
Definition: pi_esimd_emulator.cpp:1674
PI_DEVICE_INFO_MAX_WORK_GROUP_SIZE
@ PI_DEVICE_INFO_MAX_WORK_GROUP_SIZE
Definition: pi.h:204
piPlatformsGet
pi_result piPlatformsGet(pi_uint32 NumEntries, pi_platform *Platforms, pi_uint32 *NumPlatforms)
Definition: pi_esimd_emulator.cpp:394
_pi_program_info
_pi_program_info
Definition: pi.h:328
_pi_profiling_info
_pi_profiling_info
Definition: pi.h:571
RangeBuilder< 3 >::create
static sycl::range< 3 > create(Gen G)
Definition: pi_esimd_emulator.cpp:196
pi_device_affinity_domain
pi_bitfield pi_device_affinity_domain
Definition: pi.h:671
CONTINUE_NO_IMPLEMENTATION
#define CONTINUE_NO_IMPLEMENTATION
Definition: pi_esimd_emulator.cpp:376
PI_DEVICE_INFO_IMAGE_SUPPORT
@ PI_DEVICE_INFO_IMAGE_SUPPORT
Definition: pi.h:236
piMemImageCreate
pi_result piMemImageCreate(pi_context Context, pi_mem_flags Flags, const pi_image_format *ImageFormat, const pi_image_desc *ImageDesc, void *HostPtr, pi_mem *RetImage)
Definition: pi_esimd_emulator.cpp:1158
PI_QUEUE_ON_DEVICE
constexpr pi_queue_properties PI_QUEUE_ON_DEVICE
Definition: pi.h:628
PI_DEVICE_INFO_MAX_PARAMETER_SIZE
@ PI_DEVICE_INFO_MAX_PARAMETER_SIZE
Definition: pi.h:247
libCMBatch
Definition: pi_esimd_emulator.cpp:239
_pi_kernel_group_info
_pi_kernel_group_info
Definition: pi.h:368
PI_DEVICE_INFO_SUB_GROUP_SIZES_INTEL
@ PI_DEVICE_INFO_SUB_GROUP_SIZES_INTEL
Definition: pi.h:294
PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_HALF
@ PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_HALF
Definition: pi.h:221
piEnqueueMemImageCopy
pi_result piEnqueueMemImageCopy(pi_queue, pi_mem, pi_mem, pi_image_offset, pi_image_offset, pi_image_region, pi_uint32, const pi_event *, pi_event *)
Definition: pi_esimd_emulator.cpp:1741
_pi_mem::SurfaceLock
std::mutex SurfaceLock
Definition: pi_esimd_emulator.hpp:137
piProgramRelease
pi_result piProgramRelease(pi_program)
Definition: pi_esimd_emulator.cpp:1338
piContextRetain
pi_result piContextRetain(pi_context Context)
Definition: pi_esimd_emulator.cpp:878
_pi_mem_alloc_info
_pi_mem_alloc_info
Definition: pi.h:1641
PI_DEVICE_INFO_MAX_CLOCK_FREQUENCY
@ PI_DEVICE_INFO_MAX_CLOCK_FREQUENCY
Definition: pi.h:233
piEventSetStatus
pi_result piEventSetStatus(pi_event, pi_int32)
Definition: pi_esimd_emulator.cpp:1430
_pi_context::Device
pi_device Device
Definition: pi_esimd_emulator.hpp:90
PI_DEVICE_INFO_IMAGE2D_MAX_WIDTH
@ PI_DEVICE_INFO_IMAGE2D_MAX_WIDTH
Definition: pi.h:239
PI_DEVICE_INFO_AVAILABLE
@ PI_DEVICE_INFO_AVAILABLE
Definition: pi.h:263
sycl_get_cm_buffer_params
void sycl_get_cm_buffer_params(unsigned int IndexInput, char **BaseAddr, uint32_t *Width, std::mutex **BufMtxLock)
Definition: pi_esimd_emulator.cpp:274
pi_usm_mem_properties
pi_bitfield pi_usm_mem_properties
Definition: pi.h:607
pi_map_flags
pi_bitfield pi_map_flags
Definition: pi.h:592
piEventGetProfilingInfo
pi_result piEventGetProfilingInfo(pi_event Event, pi_profiling_info ParamName, size_t ParamValueSize, void *ParamValue, size_t *ParamValueSizeRet)
Definition: pi_esimd_emulator.cpp:1391
PI_DEVICE_INFO_IMAGE3D_MAX_HEIGHT
@ PI_DEVICE_INFO_IMAGE3D_MAX_HEIGHT
Definition: pi.h:242
PI_DEVICE_INFO_GPU_EU_SIMD_WIDTH
@ PI_DEVICE_INFO_GPU_EU_SIMD_WIDTH
Definition: pi.h:308
PI_MEM_TYPE_IMAGE3D
@ PI_MEM_TYPE_IMAGE3D
Definition: pi.h:452
_pi_event::OwnerQueue
cm_support::CmQueue * OwnerQueue
Definition: pi_esimd_emulator.hpp:197
piEnqueueMemImageWrite
pi_result piEnqueueMemImageWrite(pi_queue, pi_mem, pi_bool, pi_image_offset, pi_image_region, size_t, size_t, const void *, pi_uint32, const pi_event *, pi_event *)
Definition: pi_esimd_emulator.cpp:1735
piQueueGetInfo
pi_result piQueueGetInfo(pi_queue, pi_queue_info, size_t, void *, size_t *)
Definition: pi_esimd_emulator.cpp:958
piSamplerRetain
pi_result piSamplerRetain(pi_sampler)
Definition: pi_esimd_emulator.cpp:1481
PI_USM_CONCURRENT_ATOMIC_ACCESS
@ PI_USM_CONCURRENT_ATOMIC_ACCESS
Definition: pi.h:1637
RangeBuilder< 2 >::create
static sycl::range< 2 > create(Gen G)
Definition: pi_esimd_emulator.cpp:191
piProgramLink
pi_result piProgramLink(pi_context, pi_uint32, const pi_device *, const char *, pi_uint32, const pi_program *, void(*)(pi_program, void *), void *, pi_program *)
Definition: pi_esimd_emulator.cpp:1313
PI_IMAGE_CHANNEL_TYPE_SNORM_INT8
@ PI_IMAGE_CHANNEL_TYPE_SNORM_INT8
Definition: pi.h:493
pi_device_partition_property
intptr_t pi_device_partition_property
Definition: pi.h:663
PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_LONG
@ PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_LONG
Definition: pi.h:215
ConvertPiImageFormatToCmFormat
cm_support::CM_SURFACE_FORMAT ConvertPiImageFormatToCmFormat(const pi_image_format *PiFormat)
Definition: pi_esimd_emulator.cpp:1134
PI_INVALID_DEVICE
@ PI_INVALID_DEVICE
Definition: pi.h:94
piextKernelGetNativeHandle
pi_result piextKernelGetNativeHandle(pi_kernel, pi_native_handle *)
Gets the native handle of a PI kernel object.
Definition: pi_esimd_emulator.cpp:1825
piEnqueueMemImageFill
pi_result piEnqueueMemImageFill(pi_queue, pi_mem, const void *, const size_t *, const size_t *, pi_uint32, const pi_event *, pi_event *)
Definition: pi_esimd_emulator.cpp:1747
pi_int32
int32_t pi_int32
Definition: pi.h:71
_pi_context
PI context mapping to a CUDA context object.
Definition: pi_cuda.hpp:150
piEventRetain
pi_result piEventRetain(pi_event Event)
Definition: pi_esimd_emulator.cpp:1432
PI_IMAGE_CHANNEL_TYPE_UNORM_INT16
@ PI_IMAGE_CHANNEL_TYPE_UNORM_INT16
Definition: pi.h:496
_pi_device
PI device mapping to a CUdevice.
Definition: pi_cuda.hpp:73
PI_DEVICE_INFO_PARTITION_PROPERTIES
@ PI_DEVICE_INFO_PARTITION_PROPERTIES
Definition: pi.h:285
_pi_kernel_sub_group_info
_pi_kernel_sub_group_info
Definition: pi.h:401
PI_PLATFORM_INFO_EXTENSIONS
@ PI_PLATFORM_INFO_EXTENSIONS
Definition: pi.h:142
_pi_context::checkSurfaceArgument
bool checkSurfaceArgument(pi_mem_flags Flags, void *HostPtr)
Definition: pi_esimd_emulator.cpp:906
piextQueueCreateWithNativeHandle
pi_result piextQueueCreateWithNativeHandle(pi_native_handle, pi_context, pi_device, bool, pi_queue *)
Creates PI queue object from a native handle.
Definition: pi_esimd_emulator.cpp:1002
PI_DEVICE_INFO_SUB_GROUP_INDEPENDENT_FORWARD_PROGRESS
@ PI_DEVICE_INFO_SUB_GROUP_INDEPENDENT_FORWARD_PROGRESS
Definition: pi.h:292
_pi_buffer::Size
size_t Size
Definition: pi_esimd_emulator.hpp:177
PI_DEVICE_INFO_ATOMIC_MEMORY_ORDER_CAPABILITIES
@ PI_DEVICE_INFO_ATOMIC_MEMORY_ORDER_CAPABILITIES
Definition: pi.h:317
PI_EXT_ONEAPI_DEVICE_INFO_MAX_GLOBAL_WORK_GROUPS
@ PI_EXT_ONEAPI_DEVICE_INFO_MAX_GLOBAL_WORK_GROUPS
Definition: pi.h:321