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 
17 #include <detail/accessor_impl.hpp>
18 #include <sycl/backend_types.hpp>
19 #include <sycl/detail/common.hpp>
20 #include <sycl/detail/export.hpp>
21 #include <sycl/detail/helpers.hpp>
25 #include <sycl/ext/intel/esimd/common.hpp> // SLM_BTI
26 #include <sycl/group.hpp>
27 #include <sycl/id.hpp>
28 #include <sycl/kernel.hpp>
29 #include <sycl/nd_item.hpp>
30 #include <sycl/range.hpp>
31 
32 #include <esimdemu_support.h>
33 
34 #include <cstdarg>
35 #include <cstdio>
36 #include <cstring>
37 #include <functional>
38 #include <iostream>
39 #include <map>
40 #include <memory>
41 #include <string>
42 #include <string_view>
43 #include <thread>
44 #include <utility>
45 
46 #include "pi_esimd_emulator.hpp"
47 
48 #define ARG_UNUSED(x) (void)x
49 
50 namespace {
51 
52 // Helper functions for unified 'Return' type declaration - imported
53 // from pi_level_zero.cpp
54 template <typename T, typename Assign>
55 pi_result getInfoImpl(size_t ParamValueSize, void *ParamValue,
56  size_t *ParamValueSizeRet, T Value, size_t ValueSize,
57  Assign &&AssignFunc) {
58  if (ParamValue != nullptr) {
59  if (ParamValueSize < ValueSize) {
60  return PI_ERROR_INVALID_VALUE;
61  }
62  AssignFunc(ParamValue, Value, ValueSize);
63  }
64  if (ParamValueSizeRet != nullptr) {
65  *ParamValueSizeRet = ValueSize;
66  }
67  return PI_SUCCESS;
68 }
69 
70 template <typename T>
71 pi_result getInfo(size_t ParamValueSize, void *ParamValue,
72  size_t *ParamValueSizeRet, T Value) {
73  auto assignment = [](void *ParamValue, T Value, size_t ValueSize) {
74  ARG_UNUSED(ValueSize);
75  *static_cast<T *>(ParamValue) = Value;
76  };
77  return getInfoImpl(ParamValueSize, ParamValue, ParamValueSizeRet, Value,
78  sizeof(T), assignment);
79 }
80 
81 template <typename T>
82 pi_result getInfoArray(size_t ArrayLength, size_t ParamValueSize,
83  void *ParamValue, size_t *ParamValueSizeRet, T *Value) {
84  return getInfoImpl(ParamValueSize, ParamValue, ParamValueSizeRet, Value,
85  ArrayLength * sizeof(T), memcpy);
86 }
87 
88 template <>
89 pi_result getInfo<const char *>(size_t ParamValueSize, void *ParamValue,
90  size_t *ParamValueSizeRet, const char *Value) {
91  return getInfoArray(strlen(Value) + 1, ParamValueSize, ParamValue,
92  ParamValueSizeRet, Value);
93 }
94 
95 class ReturnHelper {
96 public:
97  ReturnHelper(size_t ArgParamValueSize, void *ArgParamValue,
98  size_t *ArgParamValueSizeRet)
99  : ParamValueSize(ArgParamValueSize), ParamValue(ArgParamValue),
100  ParamValueSizeRet(ArgParamValueSizeRet) {}
101 
102  template <class T> pi_result operator()(const T &t) {
103  return getInfo(ParamValueSize, ParamValue, ParamValueSizeRet, t);
104  }
105 
106 private:
107  size_t ParamValueSize;
108  void *ParamValue;
109  size_t *ParamValueSizeRet;
110 };
111 
112 } // anonymous namespace
113 
114 // Controls PI level tracing prints.
115 static bool PrintPiTrace = false;
116 
117 static void PiTrace(std::string TraceString) {
118  if (PrintPiTrace) {
119  std::cout << TraceString << std::endl;
120  }
121 }
122 
123 // Global variables used in PI_esimd_emulator
124 // Note we only create a simple pointer variables such that C++ RT won't
125 // deallocate them automatically at the end of the main program.
126 // The heap memory allocated for this global variable reclaimed only when
127 // Sycl RT calls piTearDown().
128 static sycl::detail::ESIMDEmuPluginOpaqueData *PiESimdDeviceAccess;
129 
130 // Single-entry cache for piPlatformsGet call.
132 // TODO/FIXME : Memory leak. Handle with 'piTearDown'.
133 static std::mutex *PiPlatformCacheLock = new std::mutex;
134 
135 // Mapping between surface index and CM-managed surface
136 static std::unordered_map<unsigned int, _pi_mem *> *PiESimdSurfaceMap =
137  new std::unordered_map<unsigned int, _pi_mem *>;
138 // TODO/FIXME : Memory leak. Handle with 'piTearDown'.
139 static std::mutex *PiESimdSurfaceMapLock = new std::mutex;
140 
141 // To be compared with ESIMD_EMULATOR_PLUGIN_OPAQUE_DATA_VERSION in device
142 // interface header file
143 #define ESIMDEmuPluginDataVersion 0
144 
145 // To be compared with ESIMD_DEVICE_INTERFACE_VERSION in device
146 // interface header file
147 #define ESIMDEmuPluginInterfaceVersion 1
148 
149 // For PI_DEVICE_INFO_DRIVER_VERSION info
150 static char ESimdEmuVersionString[32];
151 
152 // Global variables for PI_ERROR_PLUGIN_SPECIFIC_ERROR
153 constexpr size_t MaxMessageSize = 256;
154 thread_local pi_result ErrorMessageCode = PI_SUCCESS;
155 thread_local char ErrorMessage[MaxMessageSize];
156 
157 // Utility function for setting a message and warning
158 [[maybe_unused]] static void setErrorMessage(const char *message,
159  pi_result error_code) {
160  assert(strlen(message) <= MaxMessageSize);
161  strcpy(ErrorMessage, message);
162  ErrorMessageCode = error_code;
163 }
164 
165 // Returns plugin specific error and warning messages
167  *message = &ErrorMessage[0];
168  return ErrorMessageCode;
169 }
170 
171 // Returns plugin specific backend option.
172 // Current support is only for optimization options.
173 // Return empty string for esimd emulator.
174 // TODO: Determine correct string to be passed.
175 pi_result piPluginGetBackendOption(pi_platform, const char *frontend_option,
176  const char **backend_option) {
177  using namespace std::literals;
178  if (frontend_option == nullptr)
179  return PI_ERROR_INVALID_VALUE;
180  if (frontend_option == "-O0"sv || frontend_option == "-O1"sv ||
181  frontend_option == "-O2"sv || frontend_option == "-O3"sv ||
182  frontend_option == ""sv) {
183  *backend_option = "";
184  return PI_SUCCESS;
185  }
186  return PI_ERROR_INVALID_VALUE;
187 }
188 
189 using IDBuilder = sycl::detail::Builder;
190 
191 template <int NDims>
192 using KernelFunc = std::function<void(const sycl::nd_item<NDims> &)>;
193 
194 // Struct to wrap dimension info and lambda function to be invoked by
195 // CM Kernel launcher that only accepts raw function pointer for
196 // kernel execution. Function instances of 'InvokeKernel' un-wrap
197 // this struct instance and invoke lambda function ('Func')
198 template <int NDims> struct KernelInvocationContext {
200  const sycl::range<NDims> &LocalSize;
201  const sycl::range<NDims> &GlobalSize;
202  const sycl::id<NDims> &GlobalOffset;
203 };
204 
205 // A helper structure to create multi-dimensional range when
206 // dimensionality is given as a template parameter. `create` function
207 // in specializations accepts a template `Gen` function which
208 // generates range extent for a dimension given as an argument.
209 template <int NDims> struct RangeBuilder;
210 
211 template <> struct RangeBuilder<1> {
212  template <typename Gen> static sycl::range<1> create(Gen G) {
213  return sycl::range<1>{G(0)};
214  }
215 };
216 template <> struct RangeBuilder<2> {
217  template <typename Gen> static sycl::range<2> create(Gen G) {
218  return sycl::range<2>{G(0), G(1)};
219  }
220 };
221 template <> struct RangeBuilder<3> {
222  template <typename Gen> static sycl::range<3> create(Gen G) {
223  return sycl::range<3>{G(0), G(1), G(2)};
224  }
225 };
226 
227 // Function template to generate entry point of kernel execution as
228 // raw function pointer. CM kernel launcher executes one instance of
229 // this function per 'NDims'
230 template <int NDims> void InvokeKernel(KernelInvocationContext<NDims> *ctx) {
231 
232  sycl::range<NDims> GroupSize{
233  sycl::detail::InitializedVal<NDims, sycl::range>::template get<0>()};
234 
235  for (int i = 0; i < NDims; ++i) {
236  GroupSize[i] = ctx->GlobalSize[i] / ctx->LocalSize[i];
237  }
238 
239  const sycl::id<NDims> LocalID = RangeBuilder<NDims>::create(
240  [](int i) { return cm_support::get_thread_idx(i); });
241 
242  const sycl::id<NDims> GroupID = RangeBuilder<NDims>::create(
243  [](int i) { return cm_support::get_group_idx(i); });
244 
245  const sycl::group<NDims> Group = IDBuilder::createGroup<NDims>(
246  ctx->GlobalSize, ctx->LocalSize, GroupSize, GroupID);
247 
248  const sycl::id<NDims> GlobalID =
249  GroupID * ctx->LocalSize + LocalID + ctx->GlobalOffset;
250 
251  const sycl::item<NDims, /*Offset=*/true> GlobalItem =
252  IDBuilder::createItem<NDims, true>(ctx->GlobalSize, GlobalID,
253  ctx->GlobalOffset);
254 
255  const sycl::item<NDims, /*Offset=*/false> LocalItem =
256  IDBuilder::createItem<NDims, false>(ctx->LocalSize, LocalID);
257 
258  const sycl::nd_item<NDims> NDItem =
259  IDBuilder::createNDItem<NDims>(GlobalItem, LocalItem, Group);
260 
261  ctx->Func(NDItem);
262 }
263 
264 // Interface for lauching kernels using libcm from CM EMU project.
265 template <int DIMS> class libCMBatch {
266 private:
267  const KernelFunc<DIMS> &MKernel;
268  std::vector<uint32_t> GroupDim, SpaceDim;
269 
270 public:
272  : MKernel(Kernel), GroupDim{1, 1, 1}, SpaceDim{1, 1, 1} {}
273 
274  void runIterationSpace(const sycl::range<DIMS> &LocalSize,
275  const sycl::range<DIMS> &GlobalSize,
276  const sycl::id<DIMS> &GlobalOffset) {
277 
278  for (int I = 0; I < DIMS; I++) {
279  SpaceDim[I] = (uint32_t)LocalSize[I];
280  GroupDim[I] = (uint32_t)(GlobalSize[I] / LocalSize[I]);
281  }
282 
283  const auto InvokeKernelArg = KernelInvocationContext<DIMS>{
284  MKernel, LocalSize, GlobalSize, GlobalOffset};
285 
286  EsimdemuKernel{reinterpret_cast<fptrVoid>(InvokeKernel<DIMS>),
287  GroupDim.data(), SpaceDim.data()}
288  .launchMT(sizeof(InvokeKernelArg), &InvokeKernelArg);
289  }
290 };
291 
292 unsigned int sycl_get_cm_surface_index(void *PtrInput) {
293  _pi_mem *Surface = static_cast<_pi_mem *>(PtrInput);
294 
295  return Surface->SurfaceIndex;
296 }
297 
298 // Function to provide image info for kernel compilation using surface
299 // index without dependency on '_pi_image' definition
300 void sycl_get_cm_buffer_params(unsigned int IndexInput, char **BaseAddr,
301  uint32_t *Width, std::mutex **BufMtxLock) {
302  std::lock_guard<std::mutex> Lock{*PiESimdSurfaceMapLock};
303  auto MemIter = PiESimdSurfaceMap->find(IndexInput);
304 
305  assert(MemIter != PiESimdSurfaceMap->end() && "Invalid Surface Index");
306 
307  _pi_buffer *Buf = static_cast<_pi_buffer *>(MemIter->second);
308 
309  *BaseAddr = Buf->MapHostPtr;
310  *Width = static_cast<uint32_t>(Buf->Size);
311 
312  *BufMtxLock = &(Buf->SurfaceLock);
313 }
314 
315 // Function to provide image info for kernel compilation using surface
316 // index without dependency on '_pi_image' definition
317 void sycl_get_cm_image_params(unsigned int IndexInput, char **BaseAddr,
318  uint32_t *Width, uint32_t *Height, uint32_t *Bpp,
319  std::mutex **ImgMtxLock) {
320  std::lock_guard<std::mutex> Lock{*PiESimdSurfaceMapLock};
321  auto MemIter = PiESimdSurfaceMap->find(IndexInput);
322  assert(MemIter != PiESimdSurfaceMap->end() && "Invalid Surface Index");
323 
324  _pi_image *Img = static_cast<_pi_image *>(MemIter->second);
325 
326  *BaseAddr = Img->MapHostPtr;
327 
328  *Bpp = static_cast<uint32_t>(Img->BytesPerPixel);
329  *Width = static_cast<uint32_t>(Img->Width) * (*Bpp);
330  *Height = static_cast<uint32_t>(Img->Height);
331 
332  *ImgMtxLock = &(Img->SurfaceLock);
333 }
334 
337 sycl::detail::ESIMDDeviceInterface::ESIMDDeviceInterface() {
339  reserved = nullptr;
340 
341  /* From 'esimd_emulator_functions_v1.h' : Start */
342  cm_barrier_ptr = cm_support::barrier;
343  cm_sbarrier_ptr = cm_support::split_barrier;
344  cm_fence_ptr = cm_support::fence;
345 
346  sycl_get_surface_base_addr_ptr = cm_support::get_surface_base_addr;
347  __cm_emu_get_slm_ptr = cm_support::get_slm_base;
348  cm_slm_init_ptr = cm_support::init_slm;
349 
350  sycl_get_cm_surface_index_ptr = sycl_get_cm_surface_index;
351  sycl_get_cm_buffer_params_ptr = sycl_get_cm_buffer_params;
352  sycl_get_cm_image_params_ptr = sycl_get_cm_image_params;
353 
354  /* From 'esimd_emulator_functions_v1.h' : End */
355 }
356 
359 
360 static bool isNull(int NDims, const size_t *R) {
361  return ((0 == R[0]) && (NDims < 2 || 0 == R[1]) && (NDims < 3 || 0 == R[2]));
362 }
363 
364 // NDims is the number of dimensions in the ND-range. Kernels are
365 // normalized in the handler so that all kernels take an sycl::nd_item
366 // as argument (see StoreLambda in sycl/handler.hpp). For kernels
367 // whose workgroup size (LocalWorkSize) is unspecified, InvokeImpl
368 // sets LocalWorkSize to {1, 1, 1}, i.e. each workgroup contains just
369 // one work item. CM emulator will run several workgroups in parallel
370 // depending on environment settings.
371 
372 template <int NDims> struct InvokeImpl {
373 
374  static sycl::range<NDims> get_range(const size_t *Array) {
375  if constexpr (NDims == 1)
376  return sycl::range<NDims>{Array[0]};
377  else if constexpr (NDims == 2)
378  return sycl::range<NDims>{Array[0], Array[1]};
379  else if constexpr (NDims == 3)
380  return sycl::range<NDims>{Array[0], Array[1], Array[2]};
381  }
382 
383  static void invoke(pi_kernel Kernel, const size_t *GlobalWorkOffset,
384  const size_t *GlobalWorkSize,
385  const size_t *LocalWorkSize) {
386  libCMBatch<NDims>{*reinterpret_cast<KernelFunc<NDims> *>(Kernel)}
387  .runIterationSpace(get_range(LocalWorkSize), get_range(GlobalWorkSize),
388  sycl::id<NDims>{get_range(GlobalWorkOffset)});
389  }
390 };
391 
392 extern "C" {
393 
394 #define DIE_NO_IMPLEMENTATION \
395  if (PrintPiTrace) { \
396  std::cerr << "Not Implemented : " << __FUNCTION__ \
397  << " - File : " << __FILE__; \
398  std::cerr << " / Line : " << __LINE__ << std::endl; \
399  } \
400  return PI_ERROR_INVALID_OPERATION;
401 
402 #define CONTINUE_NO_IMPLEMENTATION \
403  if (PrintPiTrace) { \
404  std::cerr << "Warning : Not Implemented : " << __FUNCTION__ \
405  << " - File : " << __FILE__; \
406  std::cerr << " / Line : " << __LINE__ << std::endl; \
407  } \
408  return PI_SUCCESS;
409 
410 #define CASE_PI_UNSUPPORTED(not_supported) \
411  case not_supported: \
412  if (PrintPiTrace) { \
413  std::cerr << std::endl \
414  << "Unsupported PI case : " << #not_supported << " in " \
415  << __FUNCTION__ << ":" << __LINE__ << "(" << __FILE__ << ")" \
416  << std::endl; \
417  } \
418  return PI_ERROR_INVALID_OPERATION;
419 
421  pi_uint32 *NumPlatforms) {
422  static bool PiPlatformCachePopulated = false;
423  static const char *PiTraceEnv = std::getenv("SYCL_PI_TRACE");
424  static const int PiTraceValue = PiTraceEnv ? std::stoi(PiTraceEnv) : 0;
425 
426  if (PiTraceValue == -1) { // Means print all PI traces
427  PrintPiTrace = true;
428  }
429 
430  if (NumPlatforms) {
431  *NumPlatforms = 1;
432  }
433 
434  if (NumEntries == 0) {
436  if (Platforms != nullptr) {
437  PiTrace("Invalid Arguments for piPlatformsGet of "
438  "esimd_emulator (Platforms!=nullptr) "
439  "while querying number of platforms");
440  return PI_ERROR_INVALID_VALUE;
441  }
442  return PI_SUCCESS;
443  }
444 
445  if (Platforms == nullptr && NumPlatforms == nullptr) {
446  return PI_ERROR_INVALID_VALUE;
447  }
448 
449  std::lock_guard<std::mutex> Lock{*PiPlatformCacheLock};
452  PiPlatformCache->CmEmuVersion = std::string("0.0.1");
454  }
455 
456  if (Platforms && NumEntries > 0) {
457  *Platforms = PiPlatformCache;
458  }
459 
460  return PI_SUCCESS;
461 }
462 
464  size_t ParamValueSize, void *ParamValue,
465  size_t *ParamValueSizeRet) {
466  if (Platform == nullptr) {
467  return PI_ERROR_INVALID_PLATFORM;
468  }
469  ReturnHelper ReturnValue(ParamValueSize, ParamValue, ParamValueSizeRet);
470 
471  switch (ParamName) {
473  return ReturnValue("Intel(R) ESIMD_EMULATOR/GPU");
474 
476  return ReturnValue("Intel(R) Corporation");
477 
479  return ReturnValue(Platform->CmEmuVersion.c_str());
480 
482  return ReturnValue("FULL_PROFILE");
483 
485  return ReturnValue("");
486 
488  return getInfo<pi_platform_backend>(ParamValueSize, ParamValue,
489  ParamValueSizeRet,
491 
492  default:
493  // TODO: implement other parameters
494  die("Unsupported ParamName in piPlatformGetInfo");
495  }
496 
497  return PI_SUCCESS;
498 }
499 
502 }
503 
506 }
507 
509  pi_uint32 NumEntries, pi_device *Devices,
510  pi_uint32 *NumDevices) {
511  if (Platform == nullptr) {
512  return PI_ERROR_INVALID_PLATFORM;
513  }
514 
515  pi_result Res = Platform->populateDeviceCacheIfNeeded();
516  if (Res != PI_SUCCESS) {
517  return Res;
518  }
519 
520  // CM has single-root-GPU-device without sub-device support.
521  pi_uint32 DeviceCount = (DeviceType & PI_DEVICE_TYPE_GPU) ? 1 : 0;
522 
523  if (NumDevices) {
524  *NumDevices = DeviceCount;
525  }
526 
527  if (NumEntries == 0) {
529  if (Devices != nullptr) {
530  PiTrace("Invalid Arguments for piDevicesGet of esimd_emultor "
531  "(Devices!=nullptr) while querying number of platforms");
532  return PI_ERROR_INVALID_VALUE;
533  }
534  return PI_SUCCESS;
535  }
536 
537  if (DeviceCount == 0) {
539  return PI_SUCCESS;
540  }
541 
542  if (Devices) {
543  *Devices = Platform->PiDeviceCache.get();
544  }
545  return PI_SUCCESS;
546 }
547 
548 // Check the device cache and load it if necessary.
550  std::lock_guard<std::mutex> Lock(PiDeviceCacheMutex);
551 
552  if (DeviceCachePopulated) {
553  return PI_SUCCESS;
554  }
555  cm_support::CmDevice *CmDevice = nullptr;
556  // TODO FIXME Implement proper version checking and reporting:
557  // - version passed to cm_support::CreateCmDevice
558  // - CmEmuVersion
559  // - PluginVersion
560  // - ESIMDEmuPluginOpaqueData::version
561  //
562  // PI_DEVICE_INFO_DRIVER_VERSION could report the ESIMDDeviceInterface
563  // version, PI_PLATFORM_INFO_VERSION - the underlying libCM library version.
564  unsigned int Version = 0;
565 
566  int Result = cm_support::CreateCmDevice(CmDevice, Version);
567 
568  if (Result != cm_support::CM_SUCCESS) {
569  return PI_ERROR_INVALID_DEVICE;
570  }
571 
572  // CM Device version info consists of two decimal numbers - major
573  // and minor. Minor is single-digit. Version info is encoded into a
574  // unsigned integer value = 100 * major + minor. Second from right
575  // digit in decimal must be zero as it is used as 'dot'
576  // REF - $CM_EMU/common/cm_version_defs.h - 'CURRENT_CM_VERSION'
577  // e.g. CM version 7.3 => Device version = 703
578 
579  if (((Version / 10) % 10) != 0) {
580  PiTrace("Invalid Arguments for piPlatformsGet of "
581  "esimd_emulator (Platforms!=nullptr) "
582  "while querying number of platforms");
583  return PI_ERROR_INVALID_DEVICE;
584  }
585 
586  std::ostringstream StrFormat;
587  StrFormat << (int)(Version / 100) << "." << (int)(Version % 10);
588 
589  std::unique_ptr<_pi_device> Device(
590  new _pi_device(this, CmDevice, StrFormat.str()));
591  PiDeviceCache = std::move(Device);
592  DeviceCachePopulated = true;
593  return PI_SUCCESS;
594 }
595 
597  if (Device == nullptr) {
598  return PI_ERROR_INVALID_DEVICE;
599  }
600 
601  // CM supports only single device, which is root-device. 'Retain' is
602  // No-op.
603  return PI_SUCCESS;
604 }
605 
607  if (Device == nullptr) {
608  return PI_ERROR_INVALID_DEVICE;
609  }
610 
611  // CM supports only single device, which is root-device. 'Release'
612  // is No-op.
613  return PI_SUCCESS;
614 }
615 
617  size_t ParamValueSize, void *ParamValue,
618  size_t *ParamValueSizeRet) {
619  ReturnHelper ReturnValue(ParamValueSize, ParamValue, ParamValueSizeRet);
620 
621  switch (ParamName) {
622  case PI_DEVICE_INFO_TYPE:
623  return ReturnValue(PI_DEVICE_TYPE_GPU);
625  return ReturnValue(pi_device{0});
627  return ReturnValue(Device->Platform);
628  case PI_DEVICE_INFO_NAME:
629  return ReturnValue("ESIMD_EMULATOR");
631  return ReturnValue(pi_bool{true});
639  return ReturnValue(ESimdEmuVersionString);
641  return ReturnValue("Intel(R) Corporation");
643  return ReturnValue(size_t{8192});
645  return ReturnValue(size_t{8192});
647  return ReturnValue(pi_bool{1});
649  // TODO : Populate return string accordingly - e.g. cl_khr_fp16,
650  // cl_khr_fp64, cl_khr_int64_base_atomics,
651  // cl_khr_int64_extended_atomics
652  return ReturnValue("cl_khr_fp64");
654  return ReturnValue(Device->VersionStr.c_str());
655  case PI_DEVICE_INFO_BUILD_ON_SUBDEVICE: // emulator doesn't support partition
656  return ReturnValue(pi_bool{true});
658  return ReturnValue(pi_bool{false});
660  return ReturnValue(pi_bool{false});
662  return ReturnValue(pi_uint32{256});
664  return ReturnValue(pi_uint32{0});
666  return ReturnValue(pi_device_partition_property{0});
668  // '0x8086' : 'Intel HD graphics vendor ID'
669  return ReturnValue(pi_uint32{0x8086});
671  // Default SLM_MAX_SIZE from CM_EMU
672  return ReturnValue(pi_uint32{65536});
674  return ReturnValue(size_t{256});
676  // Imported from level_zero
677  return ReturnValue(pi_uint32{8});
681  // Default minimum values required by the SYCL specification.
682  return ReturnValue(size_t{2048});
684  return ReturnValue(pi_uint32{3});
686  return ReturnValue(pi_device_partition_property{0});
688  return ReturnValue("");
690  return ReturnValue(pi_queue_properties{PI_QUEUE_FLAG_ON_DEVICE});
692  struct {
693  size_t Arr[3];
694  } MaxGroupSize = {{256, 256, 1}};
695  return ReturnValue(MaxGroupSize);
696  }
711  return ReturnValue(pi_uint32{1});
712 
713  // Imported from level_zero
719  pi_uint64 Supported = 0;
720  // TODO[1.0]: how to query for USM support now?
721  if (true) {
722  // TODO: Use ze_memory_access_capabilities_t
723  Supported = PI_USM_ACCESS | PI_USM_ATOMIC_ACCESS |
725  }
726  return ReturnValue(Supported);
727  }
729  return ReturnValue(
730  pi_uint32{sizeof(void *) * std::numeric_limits<unsigned char>::digits});
732  return ReturnValue(pi_uint32{1000});
734  return ReturnValue(pi_bool{true});
736  return ReturnValue(pi_bool{true});
740  return ReturnValue(pi_uint32{0});
743  return ReturnValue(size_t{0x80000000});
746  return ReturnValue(size_t{0});
749  return ReturnValue(pi_uint32{16});
752  return ReturnValue(size_t{32});
757  uint64_t FPValue = PI_FP_CORRECTLY_ROUNDED_DIVIDE_SQRT |
760  PI_FP_FMA;
761  return ReturnValue(pi_uint64{FPValue});
762  }
764  return ReturnValue(PI_DEVICE_MEM_CACHE_TYPE_READ_WRITE_CACHE);
766  // TODO : CHECK
767  return ReturnValue(pi_uint32{64});
769  // TODO : CHECK
770  return ReturnValue(pi_uint64{0});
772  // TODO : CHECK
773  return ReturnValue(pi_uint64{0});
775  // TODO : CHECK
776  return ReturnValue(pi_uint64{0});
778  // TODO : CHECK
779  return ReturnValue(pi_uint32{64});
781  // TODO : CHECK
782  return ReturnValue(PI_DEVICE_LOCAL_MEM_TYPE_LOCAL);
784  return ReturnValue(pi_bool{false});
786  // TODO : CHECK
787  return ReturnValue(size_t{0});
789  // TODO : CHECK
790  return ReturnValue("");
792  // TODO : CHECK
793  return ReturnValue(size_t{1024});
795  return ReturnValue(pi_bool{false});
797  return ReturnValue(pi_device_affinity_domain{0});
799  // TODO : CHECK
800  return ReturnValue(pi_uint64{0});
802  // TODO : CHECK
803  return ReturnValue(
806  return ReturnValue("FULL_PROFILE");
808  // TODO : CHECK
809  return ReturnValue(pi_uint32{0});
811  return ReturnValue(size_t{1});
813  return ReturnValue(pi_int32{1});
815  return ReturnValue(pi_uint32{1}); // Minimum required by SYCL 2020 spec
817  // The mem-channel buffer property is not supported on the ESIMD emulator.
818  return ReturnValue(pi_bool{false});
820  // The sRGB images are not supported on the ESIMD emulator.
821  return ReturnValue(pi_bool{false});
823  return ReturnValue(pi_bool{false});
824 
827 
828  // Intel-specific extensions
845 
846  default:
848  }
849  return PI_SUCCESS;
850 }
851 
853  pi_uint32, pi_device *, pi_uint32 *) {
855 }
856 
859 }
860 
862  pi_device *) {
864 }
865 
867  pi_uint32 NumDevices, const pi_device *Devices,
868  void (*PFnNotify)(const char *ErrInfo,
869  const void *PrivateInfo, size_t CB,
870  void *UserData),
871  void *UserData, pi_context *RetContext) {
872  ARG_UNUSED(Properties);
873  ARG_UNUSED(PFnNotify);
874  ARG_UNUSED(UserData);
875 
876  if (NumDevices != 1) {
877  return PI_ERROR_INVALID_VALUE;
878  }
879  if (Devices == nullptr) {
880  return PI_ERROR_INVALID_DEVICE;
881  }
882  if (RetContext == nullptr) {
883  return PI_ERROR_INVALID_VALUE;
884  }
885 
886  try {
888  *RetContext = new _pi_context(Devices[0]);
889  } catch (const std::bad_alloc &) {
890  return PI_ERROR_OUT_OF_HOST_MEMORY;
891  } catch (...) {
892  return PI_ERROR_UNKNOWN;
893  }
894  return PI_SUCCESS;
895 }
896 
898  size_t *) {
900 }
901 
903  pi_context_extended_deleter, void *) {
905 }
906 
909 }
910 
912  const pi_device *, bool,
913  pi_context *) {
915 }
916 
918  if (Context == nullptr) {
919  return PI_ERROR_INVALID_CONTEXT;
920  }
921 
922  ++(Context->RefCount);
923 
924  return PI_SUCCESS;
925 }
926 
928  if (Context == nullptr || (Context->RefCount <= 0)) {
929  return PI_ERROR_INVALID_CONTEXT;
930  }
931 
932  if (--(Context->RefCount) == 0) {
935  std::lock_guard<std::mutex> Lock(Context->Addr2CmBufferSVMLock);
936  for (auto &Entry : Context->Addr2CmBufferSVM) {
937  Context->Device->CmDevicePtr->DestroyBufferSVM(Entry.second);
938  }
939  delete Context;
940  }
941 
942  return PI_SUCCESS;
943 }
944 
947  if (HostPtr == nullptr) {
948  PiTrace("HostPtr argument is required for "
949  "PI_MEM_FLAGS_HOST_PTR_USE/COPY");
950  return false;
951  }
952  // COPY and USE are mutually exclusive
955  PiTrace("PI_MEM_FLAGS_HOST_PTR_USE and _COPY cannot be used together");
956  return false;
957  }
958  }
959  return true;
960 }
961 
963  pi_queue_properties *Properties, pi_queue *Queue) {
964  assert(Properties);
965  // Expect flags mask to be passed first.
966  assert(Properties[0] == PI_QUEUE_FLAGS);
967  if (Properties[0] != PI_QUEUE_FLAGS)
968  return PI_ERROR_INVALID_VALUE;
969  pi_queue_properties Flags = Properties[1];
970  // Extra data isn't supported yet.
971  assert(Properties[2] == 0);
972  if (Properties[2] != 0)
973  return PI_ERROR_INVALID_VALUE;
974  return piQueueCreate(Context, Device, Flags, Queue);
975 }
976 
978  pi_queue_properties Properties, pi_queue *Queue) {
979  ARG_UNUSED(Device);
980 
982  // TODO : Support Out-of-order Queue
983  *Queue = nullptr;
984  return PI_ERROR_INVALID_QUEUE_PROPERTIES;
985  }
986 
987  cm_support::CmQueue *CmQueue = nullptr;
988 
989  int Result = Context->Device->CmDevicePtr->CreateQueue(CmQueue);
990  if (Result != cm_support::CM_SUCCESS) {
991  return PI_ERROR_INVALID_CONTEXT;
992  }
993 
994  try {
995  *Queue = new _pi_queue(Context, CmQueue);
996  } catch (const std::bad_alloc &) {
997  return PI_ERROR_OUT_OF_HOST_MEMORY;
998  } catch (...) {
999  return PI_ERROR_UNKNOWN;
1000  }
1001 
1002  return PI_SUCCESS;
1003 }
1004 
1005 pi_result piQueueGetInfo(pi_queue, pi_queue_info, size_t, void *, size_t *) {
1007 }
1008 
1010  if (Queue == nullptr) {
1011  return PI_ERROR_INVALID_QUEUE;
1012  }
1013  ++(Queue->RefCount);
1014  return PI_SUCCESS;
1015 }
1016 
1018  if ((Queue == nullptr) || (Queue->CmQueuePtr == nullptr)) {
1019  return PI_ERROR_INVALID_QUEUE;
1020  }
1021 
1022  if (--(Queue->RefCount) == 0) {
1023  // CM's 'DestoryQueue' is no-op
1024  // Queue->Context->Device->CmDevicePTr->DestroyQueue(Queue->CmQueuePtr);
1025  delete Queue;
1026  }
1027 
1028  return PI_SUCCESS;
1029 }
1030 
1032  // No-op as enqueued commands with ESIMD_EMULATOR plugin are blocking
1033  // ones that do not return until their completion - kernel execution
1034  // and memory read.
1036 }
1037 
1039  // No-op as enqueued commands with ESIMD_EMULATOR plugin are blocking
1040  // ones that do not return until their completion - kernel execution
1041  // and memory read.
1043 }
1044 
1047 }
1048 
1050  pi_context, pi_device, bool,
1053 }
1054 
1056  void *HostPtr, pi_mem *RetMem,
1057  const pi_mem_properties *properties) {
1058  ARG_UNUSED(properties);
1059 
1060  if ((Flags & PI_MEM_FLAGS_ACCESS_RW) == 0) {
1061  PiTrace("Invalid memory attribute for piMemBufferCreate");
1062  return PI_ERROR_INVALID_OPERATION;
1063  }
1064 
1065  if (Context == nullptr) {
1066  return PI_ERROR_INVALID_CONTEXT;
1067  }
1068  if (RetMem == nullptr) {
1069  return PI_ERROR_INVALID_VALUE;
1070  }
1071 
1072  // Flag & HostPtr argument sanity check
1073  if (!Context->checkSurfaceArgument(Flags, HostPtr)) {
1074  return PI_ERROR_INVALID_OPERATION;
1075  }
1076 
1077  char *MapBasePtr = nullptr;
1078  cm_surface_ptr_t CmBuf;
1079  cm_support::SurfaceIndex *CmIndex = nullptr;
1080  int Status = cm_support::CM_FAILURE;
1081 
1082  if (Flags & PI_MEM_FLAGS_HOST_PTR_USE) {
1084  Status = Context->Device->CmDevicePtr->CreateBufferUP(
1085  static_cast<unsigned int>(Size), HostPtr, CmBuf.UPBufPtr);
1086  CmBuf.UPBufPtr->GetIndex(CmIndex);
1087  } else {
1089  Status = Context->Device->CmDevicePtr->CreateBuffer(
1090  static_cast<unsigned int>(Size), CmBuf.RegularBufPtr);
1091  CmBuf.RegularBufPtr->GetIndex(CmIndex);
1092 
1093  if (Flags & PI_MEM_FLAGS_HOST_PTR_COPY) {
1094  CmBuf.RegularBufPtr->WriteSurface(
1095  reinterpret_cast<const unsigned char *>(HostPtr), nullptr,
1096  static_cast<unsigned int>(Size));
1097  }
1098  }
1099 
1100  if (Status != cm_support::CM_SUCCESS) {
1101  return PI_ERROR_INVALID_OPERATION;
1102  }
1103 
1104  MapBasePtr =
1105  pi_cast<char *>(cm_support::get_surface_base_addr(CmIndex->get_data()));
1106 
1107  try {
1108  *RetMem =
1109  new _pi_buffer(Context, MapBasePtr, CmBuf, CmIndex->get_data(), Size);
1110  } catch (const std::bad_alloc &) {
1111  return PI_ERROR_OUT_OF_HOST_MEMORY;
1112  } catch (...) {
1113  return PI_ERROR_UNKNOWN;
1114  }
1115 
1116  std::lock_guard<std::mutex> Lock{*PiESimdSurfaceMapLock};
1117  if (PiESimdSurfaceMap->find((*RetMem)->SurfaceIndex) !=
1118  PiESimdSurfaceMap->end()) {
1119  PiTrace("Failure from CM-managed buffer creation");
1120  return PI_ERROR_INVALID_MEM_OBJECT;
1121  }
1122 
1123  (*PiESimdSurfaceMap)[(*RetMem)->SurfaceIndex] = *RetMem;
1124 
1125  return PI_SUCCESS;
1126 }
1127 
1128 pi_result piMemGetInfo(pi_mem, pi_mem_info, size_t, void *, size_t *) {
1130 }
1131 
1133  if (Mem == nullptr) {
1134  return PI_ERROR_INVALID_MEM_OBJECT;
1135  }
1136  ++(Mem->RefCount);
1137  return PI_SUCCESS;
1138 }
1139 
1141  if ((Mem == nullptr) || (Mem->RefCount == 0)) {
1142  return PI_ERROR_INVALID_MEM_OBJECT;
1143  }
1144 
1145  if (--(Mem->RefCount) == 0) {
1146  // Removing Surface-map entry
1147  std::lock_guard<std::mutex> Lock{*PiESimdSurfaceMapLock};
1148  auto MapEntryIt = PiESimdSurfaceMap->find(Mem->SurfaceIndex);
1149  if (MapEntryIt == PiESimdSurfaceMap->end()) {
1150  PiTrace("Failure from Buffer/Image deletion");
1151  return PI_ERROR_INVALID_MEM_OBJECT;
1152  }
1153  PiESimdSurfaceMap->erase(MapEntryIt);
1154  delete Mem;
1155  }
1156  return PI_SUCCESS;
1157 }
1158 
1160  int Status = cm_support::CM_FAILURE;
1161 
1162  cm_support::CmDevice *CmDevice = Context->Device->CmDevicePtr;
1163 
1165  Status = CmDevice->DestroyBufferUP(SurfacePtr.UPBufPtr);
1167  Status = CmDevice->DestroySurface(SurfacePtr.RegularBufPtr);
1169  Status = CmDevice->DestroySurface2DUP(SurfacePtr.UPImgPtr);
1171  Status = CmDevice->DestroySurface(SurfacePtr.RegularImgPtr);
1172  }
1173 
1174  sycl::detail::pi::assertion(Status == cm_support::CM_SUCCESS &&
1175  "Surface Deletion Failure from CM_EMU");
1176 
1177  for (auto mapit = Mappings.begin(); mapit != Mappings.end();) {
1178  mapit = Mappings.erase(mapit);
1179  }
1180 }
1181 
1182 cm_support::CM_SURFACE_FORMAT
1184  using ULongPair = std::pair<unsigned long, unsigned long>;
1185  using FmtMap = std::map<ULongPair, cm_support::CM_SURFACE_FORMAT>;
1186  static const FmtMap pi2cm = {
1188  cm_support::CM_SURFACE_FORMAT_A8R8G8B8},
1189 
1191  cm_support::CM_SURFACE_FORMAT_A8R8G8B8},
1192 
1194  cm_support::CM_SURFACE_FORMAT_A8R8G8B8},
1195 
1197  cm_support::CM_SURFACE_FORMAT_R32G32B32A32F},
1198  };
1199  auto Result = pi2cm.find(
1200  {PiFormat->image_channel_data_type, PiFormat->image_channel_order});
1201  if (Result != pi2cm.end()) {
1202  return Result->second;
1203  }
1204  return cm_support::CM_SURFACE_FORMAT_UNKNOWN;
1205 }
1206 
1208  const pi_image_format *ImageFormat,
1209  const pi_image_desc *ImageDesc, void *HostPtr,
1210  pi_mem *RetImage) {
1211  if ((Flags & PI_MEM_FLAGS_ACCESS_RW) == 0) {
1212  PiTrace("Invalid memory attribute for piMemImageCreate");
1213  return PI_ERROR_INVALID_OPERATION;
1214  }
1215 
1216  if (ImageFormat == nullptr || ImageDesc == nullptr)
1217  return PI_ERROR_INVALID_IMAGE_FORMAT_DESCRIPTOR;
1218 
1219  switch (ImageDesc->image_type) {
1220  case PI_MEM_TYPE_IMAGE2D:
1221  break;
1222 
1228 
1229  default:
1230  return PI_ERROR_INVALID_MEM_OBJECT;
1231  }
1232 
1233  auto BytesPerPixel = 4;
1234  switch (ImageFormat->image_channel_data_type) {
1236  BytesPerPixel = 16;
1237  break;
1240  BytesPerPixel = 4;
1241  break;
1254  default:
1255  return PI_ERROR_IMAGE_FORMAT_NOT_SUPPORTED;
1256  }
1257 
1258  // Flag & HostPtr argument sanity check
1259  if (!Context->checkSurfaceArgument(Flags, HostPtr)) {
1260  return PI_ERROR_INVALID_OPERATION;
1261  }
1262 
1263  cm_support::CM_SURFACE_FORMAT CmSurfFormat =
1264  ConvertPiImageFormatToCmFormat(ImageFormat);
1265  if (CmSurfFormat == cm_support::CM_SURFACE_FORMAT_UNKNOWN) {
1266  return PI_ERROR_IMAGE_FORMAT_NOT_SUPPORTED;
1267  }
1268 
1269  char *MapBasePtr = nullptr;
1270  cm_surface_ptr_t CmImg;
1271  cm_support::SurfaceIndex *CmIndex = nullptr;
1272  int Status = cm_support::CM_SUCCESS;
1273 
1274  if (Flags & PI_MEM_FLAGS_HOST_PTR_USE) {
1276  Status = Context->Device->CmDevicePtr->CreateSurface2DUP(
1277  static_cast<unsigned int>(ImageDesc->image_width),
1278  static_cast<unsigned int>(ImageDesc->image_height), CmSurfFormat,
1279  HostPtr, CmImg.UPImgPtr);
1280  CmImg.UPImgPtr->GetIndex(CmIndex);
1281  } else {
1283  Status = Context->Device->CmDevicePtr->CreateSurface2D(
1284  static_cast<unsigned int>(ImageDesc->image_width),
1285  static_cast<unsigned int>(ImageDesc->image_height), CmSurfFormat,
1286  CmImg.RegularImgPtr);
1287  CmImg.RegularImgPtr->GetIndex(CmIndex);
1288 
1289  if (Flags & PI_MEM_FLAGS_HOST_PTR_COPY) {
1290  CmImg.RegularImgPtr->WriteSurface(
1291  reinterpret_cast<const unsigned char *>(HostPtr), nullptr,
1292  static_cast<unsigned int>(ImageDesc->image_width *
1293  ImageDesc->image_height * BytesPerPixel));
1294  }
1295  }
1296 
1297  if (Status != cm_support::CM_SUCCESS) {
1298  return PI_ERROR_INVALID_OPERATION;
1299  }
1300 
1301  MapBasePtr =
1302  pi_cast<char *>(cm_support::get_surface_base_addr(CmIndex->get_data()));
1303 
1304  try {
1305  *RetImage = new _pi_image(Context, MapBasePtr, CmImg, CmIndex->get_data(),
1306  ImageDesc->image_width, ImageDesc->image_height,
1307  BytesPerPixel);
1308  } catch (const std::bad_alloc &) {
1309  return PI_ERROR_OUT_OF_HOST_MEMORY;
1310  } catch (...) {
1311  return PI_ERROR_UNKNOWN;
1312  }
1313 
1314  std::lock_guard<std::mutex> Lock{*PiESimdSurfaceMapLock};
1315  if (PiESimdSurfaceMap->find((*RetImage)->SurfaceIndex) !=
1316  PiESimdSurfaceMap->end()) {
1317  PiTrace("Failure from CM-managed image creation");
1318  return PI_ERROR_INVALID_VALUE;
1319  }
1320 
1321  (*PiESimdSurfaceMap)[(*RetImage)->SurfaceIndex] = *RetImage;
1322 
1323  return PI_SUCCESS;
1324 }
1325 
1328 }
1329 
1331  pi_mem *) {
1333 }
1334 
1336  bool, const pi_image_format *,
1337  const pi_image_desc *, pi_mem *) {
1339 }
1340 
1341 pi_result piProgramCreate(pi_context, const void *, size_t, pi_program *) {
1343 }
1344 
1346  const size_t *, const unsigned char **,
1347  size_t, const pi_device_binary_property *,
1348  pi_int32 *, pi_program *) {
1350 }
1351 
1353  const size_t *, const unsigned char **,
1354  pi_int32 *, pi_program *) {
1356 }
1357 
1359  const size_t *, pi_program *) {
1361 }
1362 
1364  size_t *) {
1366 }
1367 
1369  pi_uint32, const pi_program *,
1370  void (*)(pi_program, void *), void *, pi_program *) {
1372 }
1373 
1375  const char *, pi_uint32, const pi_program *,
1376  const char **, void (*)(pi_program, void *),
1377  void *) {
1379 }
1380 
1382  void (*)(pi_program, void *), void *) {
1384 }
1385 
1387  size_t, void *, size_t *) {
1389 }
1390 
1392 
1394 
1397 }
1398 
1400  pi_program *) {
1402 }
1403 
1406 }
1407 
1408 pi_result piKernelSetArg(pi_kernel, pi_uint32, size_t, const void *) {
1410 }
1411 
1414 }
1415 
1416 // Special version of piKernelSetArg to accept pi_sampler.
1419 }
1420 
1421 pi_result piKernelGetInfo(pi_kernel, pi_kernel_info, size_t, void *, size_t *) {
1423 }
1424 
1426  size_t, void *, size_t *) {
1428 }
1429 
1431  pi_kernel_sub_group_info, size_t,
1432  const void *, size_t, void *, size_t *) {
1434 }
1435 
1437 
1439 
1441 
1443  size_t ParamValueSize, void *ParamValue,
1444  size_t *ParamValueSizeRet) {
1445  if (ParamName != PI_EVENT_INFO_COMMAND_EXECUTION_STATUS) {
1447  }
1448 
1449  auto CheckAndFillStatus = [&](const cm_support::CM_STATUS &State) {
1450  pi_int32 Result = PI_EVENT_RUNNING;
1451  if (State == cm_support::CM_STATUS_FINISHED)
1452  Result = PI_EVENT_COMPLETE;
1453  if (ParamValue) {
1454  if (ParamValueSize < sizeof(Result))
1455  return PI_ERROR_INVALID_VALUE;
1456  *static_cast<pi_int32 *>(ParamValue) = Result;
1457  }
1458  if (ParamValueSizeRet) {
1459  *ParamValueSizeRet = sizeof(Result);
1460  }
1461  return PI_SUCCESS;
1462  };
1463  // Dummy event is already completed ones done by CM.
1464  if (Event->IsDummyEvent)
1465  return CheckAndFillStatus(cm_support::CM_STATUS_FINISHED);
1466 
1467  if (Event->CmEventPtr == nullptr)
1468  return PI_ERROR_INVALID_EVENT;
1469 
1470  cm_support::CM_STATUS Status;
1471  int32_t Result = Event->CmEventPtr->GetStatus(Status);
1472  if (Result != cm_support::CM_SUCCESS)
1473  return PI_ERROR_COMMAND_EXECUTION_FAILURE;
1474 
1475  return CheckAndFillStatus(Status);
1476 }
1477 
1479  size_t ParamValueSize, void *ParamValue,
1480  size_t *ParamValueSizeRet) {
1481  ARG_UNUSED(Event);
1482  ARG_UNUSED(ParamName);
1483  ARG_UNUSED(ParamValueSize);
1484  ARG_UNUSED(ParamValue);
1485  ARG_UNUSED(ParamValueSizeRet);
1486 
1487  PiTrace("Warning : Profiling Not supported under PI_ESIMD_EMULATOR");
1488  return PI_SUCCESS;
1489 }
1490 
1491 pi_result piEventsWait(pi_uint32 NumEvents, const pi_event *EventList) {
1492  for (int i = 0; i < (int)NumEvents; i++) {
1493  if (EventList[i]->IsDummyEvent) {
1494  // Dummy event is already completed ones done by CM. Skip
1495  // waiting.
1496  continue;
1497  }
1498  if (EventList[i]->CmEventPtr == nullptr) {
1499  return PI_ERROR_INVALID_EVENT;
1500  }
1501  int Result = EventList[i]->CmEventPtr->WaitForTaskFinished();
1502  if (Result != cm_support::CM_SUCCESS) {
1503  return PI_ERROR_OUT_OF_RESOURCES;
1504  }
1505  }
1506  return PI_SUCCESS;
1507 }
1508 
1510  void (*)(pi_event, pi_int32, void *), void *) {
1512 }
1513 
1515 
1517  if (Event == nullptr) {
1518  return PI_ERROR_INVALID_EVENT;
1519  }
1520 
1521  ++(Event->RefCount);
1522 
1523  return PI_SUCCESS;
1524 }
1525 
1527  if (Event == nullptr || (Event->RefCount <= 0)) {
1528  return PI_ERROR_INVALID_EVENT;
1529  }
1530 
1531  if (--(Event->RefCount) == 0) {
1532  if (!Event->IsDummyEvent) {
1533  if ((Event->CmEventPtr == nullptr) || (Event->OwnerQueue == nullptr)) {
1534  return PI_ERROR_INVALID_EVENT;
1535  }
1536  int Result = Event->OwnerQueue->DestroyEvent(Event->CmEventPtr);
1537  if (Result != cm_support::CM_SUCCESS) {
1538  return PI_ERROR_INVALID_EVENT;
1539  }
1540  }
1541  delete Event;
1542  }
1543 
1544  return PI_SUCCESS;
1545 }
1546 
1549 }
1550 
1552  pi_event *) {
1554 }
1556  pi_sampler *) {
1558 }
1559 
1561  size_t *) {
1563 }
1564 
1566 
1568 
1570  pi_event *) {
1572 }
1573 
1575  pi_event *) {
1577 }
1578 
1580  pi_bool BlockingRead, size_t Offset,
1581  size_t Size, void *Dst,
1582  pi_uint32 NumEventsInWaitList,
1583  const pi_event *EventWaitList,
1584  pi_event *Event) {
1585  ARG_UNUSED(Queue);
1586  ARG_UNUSED(EventWaitList);
1587 
1589  if (BlockingRead) {
1590  PiTrace(
1591  "ESIMD_EMULATOR support for blocking piEnqueueMemBufferRead is NYI");
1592  return PI_ERROR_INVALID_OPERATION;
1593  }
1594 
1595  if (Offset != 0) {
1596  PiTrace("ESIMD_EMULATOR does not support buffer reading with offsets");
1597  return PI_ERROR_INVALID_ARG_VALUE;
1598  }
1599 
1600  if (NumEventsInWaitList != 0) {
1601  return PI_ERROR_INVALID_EVENT_WAIT_LIST;
1602  }
1603 
1604  _pi_buffer *buf = static_cast<_pi_buffer *>(Src);
1605 
1606  std::unique_ptr<_pi_event> RetEv{nullptr};
1607  if (Event) {
1608  RetEv = std::unique_ptr<_pi_event>(new _pi_event());
1609  RetEv->IsDummyEvent = true;
1610  }
1611 
1613  // CM does not provide 'ReadSurface' call for 'User-Provided'
1614  // Surface. memcpy is used for BufferRead PI_API call.
1615  memcpy(Dst, buf->MapHostPtr, Size);
1616  } else {
1618  return PI_ERROR_INVALID_MEM_OBJECT;
1619  }
1620  int Status = buf->SurfacePtr.RegularBufPtr->ReadSurface(
1621  reinterpret_cast<unsigned char *>(Dst),
1622  nullptr, // event
1623  static_cast<uint64_t>(Size));
1624 
1625  if (Status != cm_support::CM_SUCCESS) {
1626  return PI_ERROR_INVALID_MEM_OBJECT;
1627  }
1628  }
1629 
1630  if (Event) {
1631  *Event = RetEv.release();
1632  }
1633 
1634  return PI_SUCCESS;
1635 }
1636 
1639  pi_buff_rect_region, size_t, size_t,
1640  size_t, size_t, void *, pi_uint32,
1641  const pi_event *, pi_event *) {
1643 }
1644 
1646  const void *, pi_uint32, const pi_event *,
1647  pi_event *) {
1649 }
1650 
1653  pi_buff_rect_region, size_t, size_t,
1654  size_t, size_t, const void *, pi_uint32,
1655  const pi_event *, pi_event *) {
1657 }
1658 
1660  size_t, pi_uint32, const pi_event *,
1661  pi_event *) {
1663 }
1664 
1667  pi_buff_rect_region, size_t, size_t,
1668  size_t, size_t, pi_uint32,
1669  const pi_event *, pi_event *) {
1671 }
1672 
1673 pi_result piEnqueueMemBufferFill(pi_queue, pi_mem, const void *, size_t, size_t,
1674  size_t, pi_uint32, const pi_event *,
1675  pi_event *) {
1677 }
1678 
1680  pi_bool BlockingMap, pi_map_flags MapFlags,
1681  size_t Offset, size_t Size,
1682  pi_uint32 NumEventsInWaitList,
1683  const pi_event *EventWaitList, pi_event *Event,
1684  void **RetMap) {
1685  ARG_UNUSED(Queue);
1686  ARG_UNUSED(BlockingMap);
1687  ARG_UNUSED(MapFlags);
1688  ARG_UNUSED(NumEventsInWaitList);
1689  ARG_UNUSED(EventWaitList);
1690 
1691  std::unique_ptr<_pi_event> RetEv{nullptr};
1692  pi_result ret = PI_SUCCESS;
1693 
1694  if (Event) {
1695  RetEv = std::unique_ptr<_pi_event>(new _pi_event());
1696  RetEv->IsDummyEvent = true;
1697  }
1698 
1699  // Real mapping does not occur here and CPU-accessible address is
1700  // returned as the actual memory space for the buffer is located in
1701  // CPU memory and the plug-in know its base address
1702  // ('_pi_mem::MapHostPtr')
1703  *RetMap = MemObj->MapHostPtr + Offset;
1704 
1705  {
1706  std::lock_guard<std::mutex> Lock{MemObj->MappingsMutex};
1707  auto Res = MemObj->Mappings.insert({*RetMap, {Offset, Size}});
1708  // False as the second value in pair means that mapping was not inserted
1709  // because mapping already exists.
1710  if (!Res.second) {
1711  ret = PI_ERROR_INVALID_VALUE;
1712  PiTrace("piEnqueueMemBufferMap: duplicate mapping detected");
1713  }
1714  }
1715 
1716  if (Event) {
1717  *Event = RetEv.release();
1718  }
1719  return ret;
1720 }
1721 
1722 pi_result piEnqueueMemUnmap(pi_queue Queue, pi_mem MemObj, void *MappedPtr,
1723  pi_uint32 NumEventsInWaitList,
1724  const pi_event *EventWaitList, pi_event *Event) {
1725  ARG_UNUSED(Queue);
1726  ARG_UNUSED(NumEventsInWaitList);
1727  ARG_UNUSED(EventWaitList);
1728 
1729  std::unique_ptr<_pi_event> RetEv{nullptr};
1730  pi_result ret = PI_SUCCESS;
1731 
1732  if (Event) {
1733  RetEv = std::unique_ptr<_pi_event>(new _pi_event());
1734  RetEv->IsDummyEvent = true;
1735  }
1736 
1737  // Real unmapping does not occur here and CPU-accessible address is
1738  // returned as the actual memory space for the buffer is located in
1739  // CPU memory and the plug-in knows its base address
1740  // ('_pi_mem::MapHostPtr')
1741  {
1742  std::lock_guard<std::mutex> Lock(MemObj->MappingsMutex);
1743  auto It = MemObj->Mappings.find(MappedPtr);
1744  if (It == MemObj->Mappings.end()) {
1745  ret = PI_ERROR_INVALID_VALUE;
1746  PiTrace("piEnqueueMemUnmap: unknown memory mapping");
1747  }
1748  MemObj->Mappings.erase(It);
1749  }
1750 
1751  if (Event) {
1752  *Event = RetEv.release();
1753  }
1754 
1755  return ret;
1756 }
1757 
1758 pi_result piMemImageGetInfo(pi_mem, pi_image_info, size_t, void *, size_t *) {
1760 }
1761 
1763  pi_bool BlockingRead, pi_image_offset Origin,
1764  pi_image_region Region, size_t RowPitch,
1765  size_t SlicePitch, void *Ptr,
1766  pi_uint32 NumEventsInWaitList,
1767  const pi_event *EventWaitList,
1768  pi_event *Event) {
1769  ARG_UNUSED(CommandQueue);
1770  ARG_UNUSED(NumEventsInWaitList);
1771  ARG_UNUSED(EventWaitList);
1772 
1774  if (BlockingRead) {
1775  PiTrace("ESIMD_EMULATOR support for blocking piEnqueueMemImageRead is NYI");
1776  return PI_ERROR_INVALID_OPERATION;
1777  }
1778 
1779  // SlicePitch is for 3D image while ESIMD_EMULATOR does not
1780  // support. For 2D surfaces, SlicePitch must be 0.
1781  if (SlicePitch != 0) {
1782  PiTrace("ESIMD_EMULATOR does not support 3D-image");
1783  return PI_ERROR_INVALID_ARG_VALUE;
1784  }
1785 
1786  // CM_EMU does not support ReadSurface with offset
1787  if (Origin->x != 0 || Origin->y != 0 || Origin->z != 0) {
1788  PiTrace("ESIMD_EMULATOR does not support 2D-image reading with offsets");
1789  return PI_ERROR_INVALID_ARG_VALUE;
1790  }
1791 
1792  _pi_image *PiImg = static_cast<_pi_image *>(Image);
1793 
1794  std::unique_ptr<_pi_event> RetEv{nullptr};
1795 
1796  if (Event) {
1797  RetEv = std::unique_ptr<_pi_event>(new _pi_event());
1798  RetEv->IsDummyEvent = true;
1799  }
1800 
1801  size_t Size = RowPitch * (Region->height);
1803  // CM does not provide 'ReadSurface' call for 'User-Provided'
1804  // Surface. memcpy is used for ImageRead PI_API call.
1805  memcpy(Ptr, PiImg->MapHostPtr, Size);
1806  } else {
1808  return PI_ERROR_INVALID_MEM_OBJECT;
1809  }
1810  int Status = PiImg->SurfacePtr.RegularImgPtr->ReadSurface(
1811  reinterpret_cast<unsigned char *>(Ptr),
1812  nullptr, // event
1813  static_cast<uint64_t>(Size));
1814 
1815  if (Status != cm_support::CM_SUCCESS) {
1816  return PI_ERROR_INVALID_MEM_OBJECT;
1817  }
1818  }
1819 
1820  if (Event) {
1821  *Event = RetEv.release();
1822  }
1823 
1824  return PI_SUCCESS;
1825 }
1826 
1828  pi_image_region, size_t, size_t, const void *,
1829  pi_uint32, const pi_event *, pi_event *) {
1831 }
1832 
1835  const pi_event *, pi_event *) {
1837 }
1838 
1839 pi_result piEnqueueMemImageFill(pi_queue, pi_mem, const void *, const size_t *,
1840  const size_t *, pi_uint32, const pi_event *,
1841  pi_event *) {
1843 }
1844 
1846  void *, pi_mem *) {
1848 }
1849 
1850 pi_result
1852  const size_t *GlobalWorkOffset,
1853  const size_t *GlobalWorkSize, const size_t *LocalWorkSize,
1854  pi_uint32 NumEventsInWaitList,
1855  const pi_event *EventWaitList, pi_event *Event) {
1856  ARG_UNUSED(Queue);
1857  ARG_UNUSED(NumEventsInWaitList);
1858  ARG_UNUSED(EventWaitList);
1859 
1860  const size_t LocalWorkSz[] = {1, 1, 1};
1861 
1862  if (Kernel == nullptr) {
1863  return PI_ERROR_INVALID_KERNEL;
1864  }
1865 
1866  if (WorkDim > 3 || WorkDim == 0) {
1867  return PI_ERROR_INVALID_WORK_GROUP_SIZE;
1868  }
1869 
1870  if (isNull(WorkDim, LocalWorkSize)) {
1871  LocalWorkSize = LocalWorkSz;
1872  }
1873 
1874  for (pi_uint32 I = 0; I < WorkDim; I++) {
1875  if ((GlobalWorkSize[I] % LocalWorkSize[I]) != 0) {
1876  return PI_ERROR_INVALID_WORK_GROUP_SIZE;
1877  }
1878  }
1879 
1880  std::unique_ptr<_pi_event> RetEv{nullptr};
1881 
1882  if (Event) {
1883  RetEv = std::unique_ptr<_pi_event>(new _pi_event());
1884  RetEv->IsDummyEvent = true;
1885  }
1886 
1887  switch (WorkDim) {
1888  case 1:
1889  InvokeImpl<1>::invoke(Kernel, GlobalWorkOffset, GlobalWorkSize,
1890  LocalWorkSize);
1891  break;
1892  case 2:
1893  InvokeImpl<2>::invoke(Kernel, GlobalWorkOffset, GlobalWorkSize,
1894  LocalWorkSize);
1895  break;
1896  case 3:
1897  InvokeImpl<3>::invoke(Kernel, GlobalWorkOffset, GlobalWorkSize,
1898  LocalWorkSize);
1899  break;
1900  default:
1902  break;
1903  }
1904 
1905  if (Event) {
1906  *Event = RetEv.release();
1907  }
1908 
1909  return PI_SUCCESS;
1910 }
1911 
1913  pi_program, bool, pi_kernel *) {
1915 }
1916 
1919 }
1920 
1921 pi_result piEnqueueNativeKernel(pi_queue, void (*)(void *), void *, size_t,
1922  pi_uint32, const pi_mem *, const void **,
1923  pi_uint32, const pi_event *, pi_event *) {
1925 }
1926 
1928  pi_uint64 *) {
1930 }
1931 
1933  size_t, pi_uint32) {
1935 }
1936 
1938  pi_usm_mem_properties *, size_t, pi_uint32) {
1940 }
1941 
1942 pi_result piextUSMSharedAlloc(void **ResultPtr, pi_context Context,
1943  pi_device Device,
1944  pi_usm_mem_properties *Properties, size_t Size,
1945  pi_uint32 Alignment) {
1946  ARG_UNUSED(Properties);
1948 
1949  if (Context == nullptr || (Device != Context->Device)) {
1950  return PI_ERROR_INVALID_CONTEXT;
1951  }
1952 
1953  if (ResultPtr == nullptr) {
1954  return PI_ERROR_INVALID_OPERATION;
1955  }
1956 
1957  // 'Size' must be power of two in order to prevent memory corruption
1958  // error
1959  if ((Size & (Size - 1)) != 0) {
1960  Size = sycl::detail::getNextPowerOfTwo(Size);
1961  }
1962 
1963  cm_support::CmBufferSVM *Buf = nullptr;
1964  void *SystemMemPtr = nullptr;
1965  int32_t Result = Context->Device->CmDevicePtr->CreateBufferSVM(
1966  Size, SystemMemPtr, CM_SVM_ACCESS_FLAG_DEFAULT, Buf);
1967 
1968  if (Result != cm_support::CM_SUCCESS) {
1969  return PI_ERROR_OUT_OF_HOST_MEMORY;
1970  }
1971  *ResultPtr = SystemMemPtr;
1972  std::lock_guard<std::mutex> Lock(Context->Addr2CmBufferSVMLock);
1973  auto Iter = Context->Addr2CmBufferSVM.find(SystemMemPtr);
1974  if (Context->Addr2CmBufferSVM.end() != Iter) {
1975  return PI_ERROR_INVALID_MEM_OBJECT;
1976  }
1977  Context->Addr2CmBufferSVM[SystemMemPtr] = Buf;
1978  return PI_SUCCESS;
1979 }
1980 
1981 pi_result piextUSMFree(pi_context Context, void *Ptr) {
1982  if (Context == nullptr) {
1983  return PI_ERROR_INVALID_CONTEXT;
1984  }
1985  if (Ptr == nullptr) {
1986  return PI_ERROR_INVALID_OPERATION;
1987  }
1988 
1989  std::lock_guard<std::mutex> Lock(Context->Addr2CmBufferSVMLock);
1990  cm_support::CmBufferSVM *Buf = Context->Addr2CmBufferSVM[Ptr];
1991  if (Buf == nullptr) {
1992  return PI_ERROR_INVALID_MEM_OBJECT;
1993  }
1994  auto Count = Context->Addr2CmBufferSVM.erase(Ptr);
1995  if (Count != 1) {
1996  return PI_ERROR_INVALID_MEM_OBJECT;
1997  }
1998  int32_t Result = Context->Device->CmDevicePtr->DestroyBufferSVM(Buf);
1999  if (cm_support::CM_SUCCESS != Result) {
2000  return PI_ERROR_UNKNOWN;
2001  }
2002  return PI_SUCCESS;
2003 }
2004 
2007 }
2008 
2010  const pi_event *, pi_event *) {
2012 }
2013 
2014 pi_result piextUSMEnqueueMemcpy(pi_queue, pi_bool, void *, const void *, size_t,
2015  pi_uint32, const pi_event *, pi_event *) {
2017 }
2018 
2020  pi_mem_advice, pi_event *) {
2022 }
2023 
2024 pi_result piextUSMEnqueueFill2D(pi_queue, void *, size_t, size_t, const void *,
2025  size_t, size_t, pi_uint32, const pi_event *,
2026  pi_event *) {
2028 }
2029 
2030 pi_result piextUSMEnqueueMemset2D(pi_queue, void *, size_t, int, size_t, size_t,
2031  pi_uint32, const pi_event *, pi_event *) {
2033 }
2034 
2036  const void *, size_t, size_t, size_t,
2037  pi_uint32, const pi_event *, pi_event *) {
2039 }
2040 
2042  size_t, void *, size_t *) {
2044 }
2045 
2048  void *, size_t, pi_uint32, const pi_event *,
2049  pi_event *) {
2051 }
2052 
2054  void *, size_t, pi_uint32, const pi_event *,
2055  pi_event *) {
2057 }
2058 
2060  const void *) {
2062 }
2063 
2065  const void *) {
2067 }
2068 
2070  pi_uint32 RawImgSize, pi_uint32 *ImgInd) {
2073  if (RawImgSize != 1) {
2074  PiTrace("Only single device binary image is supported in ESIMD_EMULATOR");
2075  return PI_ERROR_INVALID_VALUE;
2076  }
2077  *ImgInd = 0;
2078  return PI_SUCCESS;
2079 }
2080 
2083  const pi_event *, pi_event *) {
2085 }
2086 
2088  const char *, pi_bool, size_t,
2089  size_t, const void *, pi_uint32,
2090  const pi_event *, pi_event *) {
2092 }
2093 
2095  const char *, pi_bool, size_t,
2096  size_t, void *, pi_uint32,
2097  const pi_event *, pi_event *) {
2099 }
2100 
2101 pi_result piextPluginGetOpaqueData(void *, void **OpaqueDataReturn) {
2102  *OpaqueDataReturn = reinterpret_cast<void *>(PiESimdDeviceAccess);
2103  return PI_SUCCESS;
2104 }
2105 
2106 // Windows: dynamically loaded plugins might have been unloaded already
2107 // when this is called. Sycl RT holds onto the PI plugin so it can be
2108 // called safely. But this is not transitive. If the PI plugin in turn
2109 // dynamically loaded a different DLL, that may have been unloaded.
2111  delete reinterpret_cast<sycl::detail::ESIMDEmuPluginOpaqueData *>(
2112  PiESimdDeviceAccess->data);
2113  delete PiESimdDeviceAccess;
2114 
2115  for (auto it = PiESimdSurfaceMap->begin(); it != PiESimdSurfaceMap->end();) {
2116  auto Mem = it->second;
2117  if (Mem != nullptr) {
2118  delete Mem;
2119  } // else { /* Null-entry for SLM_BTI */ }
2120  it = PiESimdSurfaceMap->erase(it);
2121  }
2122  return PI_SUCCESS;
2123 }
2124 
2125 pi_result piGetDeviceAndHostTimer(pi_device, uint64_t *, uint64_t *) {
2126  PiTrace(
2127  "Warning : Querying device clock not supported under PI_ESIMD_EMULATOR");
2128  return PI_SUCCESS;
2129 }
2131 
2133  if (PluginInit == nullptr) {
2134  return PI_ERROR_INVALID_VALUE;
2135  }
2136 
2137  // Check that the major version matches in PiVersion and SupportedVersion
2139 
2140  size_t PluginVersionSize = sizeof(PluginInit->PluginVersion);
2141  if (strlen(_PI_H_VERSION_STRING) >= PluginVersionSize) {
2142  return PI_ERROR_INVALID_VALUE;
2143  }
2144  strncpy(PluginInit->PluginVersion, SupportedVersion, PluginVersionSize);
2145 
2146  PiESimdDeviceAccess = new sycl::detail::ESIMDEmuPluginOpaqueData();
2147  // 'version' to be compared with 'ESIMD_EMULATOR_DEVICE_REQUIRED_VER' defined
2148  // in device interface file
2150  PiESimdDeviceAccess->data =
2151  reinterpret_cast<void *>(new sycl::detail::ESIMDDeviceInterface());
2152 
2153  // Registering pre-defined surface index dedicated for SLM
2154  (*PiESimdSurfaceMap)[__ESIMD_DNS::SLM_BTI] = nullptr;
2155 
2156 #define _PI_API(api) \
2157  (PluginInit->PiFunctionTable).api = (decltype(&::api))(&api);
2158 #include <sycl/detail/pi.def>
2159 
2160  return PI_SUCCESS;
2161 }
2162 
2163 #ifdef _WIN32
2164 #define __SYCL_PLUGIN_DLL_NAME "pi_esimd_emulator.dll"
2165 #include "../common_win_pi_trace/common_win_pi_trace.hpp"
2166 #undef __SYCL_PLUGIN_DLL_NAME
2167 #endif
2168 
2169 } // extern C
sycl::_V1::ext::intel::esimd::fence
__ESIMD_API void fence()
esimd::fence sets the memory read/write order.
Definition: memory.hpp:1628
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:2069
setErrorMessage
static void setErrorMessage(const char *message, pi_result error_code)
Definition: pi_esimd_emulator.cpp:158
PiPlatformCachePopulated
bool PiPlatformCachePopulated
Definition: ur.cpp:29
PI_DEVICE_INFO_HOST_UNIFIED_MEMORY
@ PI_DEVICE_INFO_HOST_UNIFIED_MEMORY
Definition: pi.h:282
cm_surface_ptr_t::TypeUserProvidedImage
@ TypeUserProvidedImage
Definition: pi_esimd_emulator.hpp:127
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:1942
PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_CHAR
@ PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_CHAR
Definition: pi.h:243
PI_DEVICE_INFO_OPENCL_C_VERSION
@ PI_DEVICE_INFO_OPENCL_C_VERSION
Definition: pi.h:300
pi_image_region_struct::height
size_t height
Definition: pi.h:957
_pi_mem
PI Mem mapping to CUDA memory allocations, both data and texture/surface.
Definition: pi_cuda.hpp:224
SupportedVersion
const char SupportedVersion[]
Definition: pi_esimd_emulator.cpp:2130
piPluginInit
pi_result piPluginInit(pi_plugin *PluginInit)
Definition: pi_esimd_emulator.cpp:2132
PI_USM_ACCESS
@ PI_USM_ACCESS
Definition: pi.h:1740
PI_DEVICE_INFO_PROFILE
@ PI_DEVICE_INFO_PROFILE
Definition: pi.h:298
piextPluginGetOpaqueData
pi_result piextPluginGetOpaqueData(void *, void **OpaqueDataReturn)
API to get Plugin internal data, opaque to SYCL RT.
Definition: pi_esimd_emulator.cpp:2101
PI_DEVICE_INFO_IMAGE3D_MAX_WIDTH
@ PI_DEVICE_INFO_IMAGE3D_MAX_WIDTH
Definition: pi.h:265
PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_CHAR
@ PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_CHAR
Definition: pi.h:250
_pi_device::CmDevicePtr
cm_support::CmDevice * CmDevicePtr
Definition: pi_esimd_emulator.hpp:89
piProgramGetInfo
pi_result piProgramGetInfo(pi_program, pi_program_info, size_t, void *, size_t *)
Definition: pi_esimd_emulator.cpp:1363
PI_DEVICE_INFO_DOUBLE_FP_CONFIG
@ PI_DEVICE_INFO_DOUBLE_FP_CONFIG
Definition: pi.h:241
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:616
_pi_context_info
_pi_context_info
Definition: pi.h:373
PI_MEM_TYPE_IMAGE1D
@ PI_MEM_TYPE_IMAGE1D
Definition: pi.h:486
pi_buff_rect_offset_struct
Definition: pi.h:928
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:866
_pi_image::Width
size_t Width
Definition: pi_esimd_emulator.hpp:196
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:1399
PI_DEVICE_INFO_DRIVER_VERSION
@ PI_DEVICE_INFO_DRIVER_VERSION
Definition: pi.h:297
piProgramCreate
pi_result piProgramCreate(pi_context, const void *, size_t, pi_program *)
Definition: pi_esimd_emulator.cpp:1341
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:1981
InvokeKernel
void InvokeKernel(KernelInvocationContext< NDims > *ctx)
Definition: pi_esimd_emulator.cpp:230
PI_DEVICE_INFO_GPU_EU_COUNT
@ PI_DEVICE_INFO_GPU_EU_COUNT
Definition: pi.h:322
pi_bool
pi_uint32 pi_bool
Definition: pi.h:144
PI_IMAGE_CHANNEL_TYPE_UNORM_SHORT_565
@ PI_IMAGE_CHANNEL_TYPE_UNORM_SHORT_565
Definition: pi.h:530
IDBuilder
sycl::detail::Builder IDBuilder
Definition: pi_esimd_emulator.cpp:189
PI_IMAGE_CHANNEL_TYPE_HALF_FLOAT
@ PI_IMAGE_CHANNEL_TYPE_HALF_FLOAT
Definition: pi.h:539
piextEnqueueReadHostPipe
pi_result piextEnqueueReadHostPipe(pi_queue, pi_program, const char *, pi_bool, void *, size_t, pi_uint32, const pi_event *, pi_event *)
Host Pipes.
Definition: pi_esimd_emulator.cpp:2047
PI_DEVICE_INFO_PREFERRED_INTEROP_USER_SYNC
@ PI_DEVICE_INFO_PREFERRED_INTEROP_USER_SYNC
Definition: pi.h:303
T
_pi_image_format::image_channel_data_type
pi_image_channel_type image_channel_data_type
Definition: pi.h:1006
ErrorMessage
thread_local char ErrorMessage[MaxMessageSize]
Definition: pi_esimd_emulator.cpp:155
piextKernelSetArgSampler
pi_result piextKernelSetArgSampler(pi_kernel, pi_uint32, const pi_sampler *)
Definition: pi_esimd_emulator.cpp:1417
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:317
piQueueFinish
pi_result piQueueFinish(pi_queue)
Definition: pi_esimd_emulator.cpp:1031
ARG_UNUSED
#define ARG_UNUSED(x)
Definition: pi_esimd_emulator.cpp:48
type_traits.hpp
PI_DEVICE_INFO_NAME
@ PI_DEVICE_INFO_NAME
Definition: pi.h:295
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:1345
KernelInvocationContext::Func
KernelFunc< NDims > Func
Definition: pi_esimd_emulator.cpp:199
PI_DEVICE_INFO_IL_VERSION
@ PI_DEVICE_INFO_IL_VERSION
Definition: pi.h:294
_pi_mem::Mappings
std::unordered_map< void *, Mapping > Mappings
Definition: pi_esimd_emulator.hpp:163
PI_QUEUE_FLAG_ON_DEVICE
constexpr pi_queue_properties PI_QUEUE_FLAG_ON_DEVICE
Definition: pi.h:652
PI_IMAGE_CHANNEL_TYPE_UNORM_INT8
@ PI_IMAGE_CHANNEL_TYPE_UNORM_INT8
Definition: pi.h:528
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:1352
_pi_buffer
Definition: pi_esimd_emulator.hpp:179
piextDeviceGetNativeHandle
pi_result piextDeviceGetNativeHandle(pi_device, pi_native_handle *)
Gets the native handle of a PI device object.
Definition: pi_esimd_emulator.cpp:857
_pi_platform::PiDeviceCacheMutex
std::mutex PiDeviceCacheMutex
Definition: pi_esimd_emulator.hpp:71
PI_MEM_FLAGS_HOST_PTR_COPY
constexpr pi_mem_flags PI_MEM_FLAGS_HOST_PTR_COPY
Definition: pi.h:617
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:2014
piKernelSetArg
pi_result piKernelSetArg(pi_kernel, pi_uint32, size_t, const void *)
Definition: pi_esimd_emulator.cpp:1408
common.hpp
_pi_image
Definition: pi_esimd_emulator.hpp:188
_pi_image_format::image_channel_order
pi_image_channel_order image_channel_order
Definition: pi.h:1005
PI_DEVICE_INFO_PARTITION_AFFINITY_DOMAIN
@ PI_DEVICE_INFO_PARTITION_AFFINITY_DOMAIN
Definition: pi.h:307
PI_DEVICE_INFO_IMAGE_SRGB
@ PI_DEVICE_INFO_IMAGE_SRGB
Definition: pi.h:328
_pi_image_desc::image_type
pi_mem_type image_type
Definition: pi.h:1010
piEventCreate
pi_result piEventCreate(pi_context, pi_event *)
Create PI event object in a signalled/completed state.
Definition: pi_esimd_emulator.cpp:1440
_pi_plugin
Definition: pi.h:2103
RangeBuilder
Definition: pi_esimd_emulator.cpp:209
piextUSMEnqueueMemcpy2D
pi_result piextUSMEnqueueMemcpy2D(pi_queue, pi_bool, void *, size_t, const void *, size_t, size_t, size_t, pi_uint32, const pi_event *, pi_event *)
USM 2D Memcpy API.
Definition: pi_esimd_emulator.cpp:2035
PI_DEVICE_INFO_GLOBAL_MEM_CACHE_TYPE
@ PI_DEVICE_INFO_GLOBAL_MEM_CACHE_TYPE
Definition: pi.h:273
sycl::_V1::ext::intel::experimental::esimd::split_barrier
__ESIMD_API void split_barrier()
Generic work-group split barrier.
Definition: memory.hpp:28
_pi_mem_advice
_pi_mem_advice
Definition: pi.h:491
piSamplerRelease
pi_result piSamplerRelease(pi_sampler)
Definition: pi_esimd_emulator.cpp:1567
piSamplerGetInfo
pi_result piSamplerGetInfo(pi_sampler, pi_sampler_info, size_t, void *, size_t *)
Definition: pi_esimd_emulator.cpp:1560
piDeviceRelease
pi_result piDeviceRelease(pi_device Device)
Definition: pi_esimd_emulator.cpp:606
RangeBuilder< 1 >::create
static sycl::range< 1 > create(Gen G)
Definition: pi_esimd_emulator.cpp:212
libCMBatch::libCMBatch
libCMBatch(const KernelFunc< DIMS > &Kernel)
Definition: pi_esimd_emulator.cpp:271
PI_DEVICE_INFO_MAX_WORK_ITEM_SIZES
@ PI_DEVICE_INFO_MAX_WORK_ITEM_SIZES
Definition: pi.h:237
cm_surface_ptr_t
Definition: pi_esimd_emulator.hpp:118
piEnqueueEventsWaitWithBarrier
pi_result piEnqueueEventsWaitWithBarrier(pi_queue, pi_uint32, const pi_event *, pi_event *)
Definition: pi_esimd_emulator.cpp:1574
PI_DEVICE_INFO_MAX_COMPUTE_UNITS
@ PI_DEVICE_INFO_MAX_COMPUTE_UNITS
Definition: pi.h:235
PrintPiTrace
static bool PrintPiTrace
Definition: pi_esimd_emulator.cpp:115
piEventSetCallback
pi_result piEventSetCallback(pi_event, pi_int32, void(*)(pi_event, pi_int32, void *), void *)
Definition: pi_esimd_emulator.cpp:1509
_pi_plugin::PluginVersion
char PluginVersion[20]
Definition: pi.h:2113
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:1722
_pi_result
_pi_result
Definition: pi.h:153
sycl::_V1::detail::memcpy
void memcpy(void *Dst, const void *Src, size_t Size)
Definition: memcpy.hpp:16
InvokeImpl
Definition: pi_esimd_emulator.cpp:372
piextProgramGetNativeHandle
pi_result piextProgramGetNativeHandle(pi_program, pi_native_handle *)
Gets the native handle of a PI program object.
Definition: pi_esimd_emulator.cpp:1395
ErrorMessageCode
thread_local pi_result ErrorMessageCode
Definition: pi_esimd_emulator.cpp:154
pi_context_properties
intptr_t pi_context_properties
Definition: pi.h:572
pi_esimd_emulator.hpp
PI_IMAGE_CHANNEL_TYPE_SIGNED_INT8
@ PI_IMAGE_CHANNEL_TYPE_SIGNED_INT8
Definition: pi.h:533
piextEnqueueDeviceGlobalVariableRead
pi_result piextEnqueueDeviceGlobalVariableRead(pi_queue, pi_program, const char *, pi_bool, size_t, size_t, void *, pi_uint32, const pi_event *, pi_event *)
API reading data from a device global variable to host.
Definition: pi_esimd_emulator.cpp:2094
piGetDeviceAndHostTimer
pi_result piGetDeviceAndHostTimer(pi_device, uint64_t *, uint64_t *)
Queries device for it's global timestamp in nanoseconds, and updates HostTime with the value of the h...
Definition: pi_esimd_emulator.cpp:2125
PI_DEVICE_INFO_GPU_SLICES
@ PI_DEVICE_INFO_GPU_SLICES
Definition: pi.h:324
PI_DEVICE_INFO_MAX_MEM_ALLOC_SIZE
@ PI_DEVICE_INFO_MAX_MEM_ALLOC_SIZE
Definition: pi.h:259
_pi_device_type
_pi_device_type
Definition: pi.h:201
piextPlatformCreateWithNativeHandle
pi_result piextPlatformCreateWithNativeHandle(pi_native_handle, pi_platform *)
Creates PI platform object from a native handle.
Definition: pi_esimd_emulator.cpp:504
_pi_event::CmEventPtr
cm_support::CmEvent * CmEventPtr
Definition: pi_esimd_emulator.hpp:204
piDevicePartition
pi_result piDevicePartition(pi_device, const pi_device_partition_property *, pi_uint32, pi_device *, pi_uint32 *)
Definition: pi_esimd_emulator.cpp:852
PI_DEVICE_INFO_REFERENCE_COUNT
@ PI_DEVICE_INFO_REFERENCE_COUNT
Definition: pi.h:293
PI_DEVICE_INFO_USM_CROSS_SHARED_SUPPORT
@ PI_DEVICE_INFO_USM_CROSS_SHARED_SUPPORT
Definition: pi.h:315
PI_DEVICE_INFO_PLATFORM
@ PI_DEVICE_INFO_PLATFORM
Definition: pi.h:292
PI_DEVICE_INFO_GPU_EU_COUNT_PER_SUBSLICE
@ PI_DEVICE_INFO_GPU_EU_COUNT_PER_SUBSLICE
Definition: pi.h:326
PI_DEVICE_LOCAL_MEM_TYPE_LOCAL
@ PI_DEVICE_LOCAL_MEM_TYPE_LOCAL
Definition: pi.h:228
PI_DEVICE_INFO_DEVICE_ID
@ PI_DEVICE_INFO_DEVICE_ID
Definition: pi.h:320
helpers.hpp
host_profiling_info.hpp
PI_IMAGE_CHANNEL_TYPE_SNORM_INT16
@ PI_IMAGE_CHANNEL_TYPE_SNORM_INT16
Definition: pi.h:527
PI_DEVICE_INFO_ADDRESS_BITS
@ PI_DEVICE_INFO_ADDRESS_BITS
Definition: pi.h:258
PI_EVENT_RUNNING
@ PI_EVENT_RUNNING
Definition: pi.h:163
_pi_image_desc::image_height
size_t image_height
Definition: pi.h:1012
_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:74
piPluginGetLastError
pi_result piPluginGetLastError(char **message)
API to get Plugin specific warning and error messages.
Definition: pi_esimd_emulator.cpp:166
CASE_PI_UNSUPPORTED
#define CASE_PI_UNSUPPORTED(not_supported)
Definition: pi_esimd_emulator.cpp:410
PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_DOUBLE
@ PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_DOUBLE
Definition: pi.h:248
PI_DEVICE_INFO_MAX_SAMPLERS
@ PI_DEVICE_INFO_MAX_SAMPLERS
Definition: pi.h:270
pi_image_offset_struct::y
size_t y
Definition: pi.h:948
piextUSMEnqueueFill2D
pi_result piextUSMEnqueueFill2D(pi_queue, void *, size_t, size_t, const void *, size_t, size_t, pi_uint32, const pi_event *, pi_event *)
USM 2D fill API.
Definition: pi_esimd_emulator.cpp:2024
PI_DEVICE_INFO_MAX_NUM_SUB_GROUPS
@ PI_DEVICE_INFO_MAX_NUM_SUB_GROUPS
Definition: pi.h:309
PI_DEVICE_INFO_MEM_BASE_ADDR_ALIGN
@ PI_DEVICE_INFO_MEM_BASE_ADDR_ALIGN
Definition: pi.h:272
piextEnqueueWriteHostPipe
pi_result piextEnqueueWriteHostPipe(pi_queue, pi_program, const char *, pi_bool, void *, size_t, pi_uint32, const pi_event *, pi_event *)
Write to pipe of a given name.
Definition: pi_esimd_emulator.cpp:2053
PI_DEVICE_INFO_BUILT_IN_KERNELS
@ PI_DEVICE_INFO_BUILT_IN_KERNELS
Definition: pi.h:291
PI_DEVICE_INFO_USM_DEVICE_SUPPORT
@ PI_DEVICE_INFO_USM_DEVICE_SUPPORT
Definition: pi.h:313
PI_PLATFORM_INFO_NAME
@ PI_PLATFORM_INFO_NAME
Definition: pi.h:170
PI_FP_ROUND_TO_NEAREST
static constexpr pi_device_fp_config PI_FP_ROUND_TO_NEAREST
Definition: pi.h:732
PI_DEVICE_INFO_IMAGE_MAX_ARRAY_SIZE
@ PI_DEVICE_INFO_IMAGE_MAX_ARRAY_SIZE
Definition: pi.h:269
_pi_device_info
_pi_device_info
Definition: pi.h:232
PI_USM_ATOMIC_ACCESS
@ PI_USM_ATOMIC_ACCESS
Definition: pi.h:1741
PI_DEVICE_INFO_MAX_CONSTANT_BUFFER_SIZE
@ PI_DEVICE_INFO_MAX_CONSTANT_BUFFER_SIZE
Definition: pi.h:277
pi_image_offset_struct::x
size_t x
Definition: pi.h:947
_pi_image_info
_pi_image_info
Definition: pi.h:422
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:274
piQueueRetain
pi_result piQueueRetain(pi_queue Queue)
Definition: pi_esimd_emulator.cpp:1009
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:1330
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:360
PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_INT
@ PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_INT
Definition: pi.h:245
_pi_kernel
Implementation of a PI Kernel for CUDA.
Definition: pi_cuda.hpp:816
PI_DEVICE_INFO_ERROR_CORRECTION_SUPPORT
@ PI_DEVICE_INFO_ERROR_CORRECTION_SUPPORT
Definition: pi.h:281
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:1912
PI_DEVICE_INFO_IMAGE2D_MAX_HEIGHT
@ PI_DEVICE_INFO_IMAGE2D_MAX_HEIGHT
Definition: pi.h:264
piEventsWait
pi_result piEventsWait(pi_uint32 NumEvents, const pi_event *EventList)
Definition: pi_esimd_emulator.cpp:1491
_pi_plugin::PiVersion
char PiVersion[20]
Definition: pi.h:2111
PI_DEVICE_INFO_GLOBAL_MEM_CACHE_SIZE
@ PI_DEVICE_INFO_GLOBAL_MEM_CACHE_SIZE
Definition: pi.h:275
PI_DEVICE_EXEC_CAPABILITIES_KERNEL
constexpr pi_device_exec_capabilities PI_DEVICE_EXEC_CAPABILITIES_KERNEL
Definition: pi.h:575
piextDeviceCreateWithNativeHandle
pi_result piextDeviceCreateWithNativeHandle(pi_native_handle, pi_platform, pi_device *)
Creates PI device object from a native handle.
Definition: pi_esimd_emulator.cpp:861
piMemBufferPartition
pi_result piMemBufferPartition(pi_mem, pi_mem_flags, pi_buffer_create_type, void *, pi_mem *)
Definition: pi_esimd_emulator.cpp:1845
cm_surface_ptr_t::RegularBufPtr
cm_support::CmBuffer * RegularBufPtr
Definition: pi_esimd_emulator.hpp:132
piextQueueCreateWithNativeHandle
pi_result piextQueueCreateWithNativeHandle(pi_native_handle, int32_t, pi_context, pi_device, bool, pi_queue_properties *, pi_queue *)
Creates PI queue object from a native handle.
Definition: pi_esimd_emulator.cpp:1049
_pi_mem::Context
pi_context Context
Definition: pi_esimd_emulator.hpp:140
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:1921
PI_EXT_DEVICE_INFO_ATOMIC_MEMORY_SCOPE_CAPABILITIES
@ PI_EXT_DEVICE_INFO_ATOMIC_MEMORY_SCOPE_CAPABILITIES
Definition: pi.h:343
_pi_queue_info
_pi_queue_info
Definition: pi.h:390
_pi_mem::~_pi_mem
~_pi_mem()
Definition: pi_cuda.hpp:365
_pi_buffer_create_type
_pi_buffer_create_type
Definition: pi.h:543
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:1762
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:1430
piextUSMEnqueueMemAdvise
pi_result piextUSMEnqueueMemAdvise(pi_queue, const void *, size_t, pi_mem_advice, pi_event *)
USM Memadvise API.
Definition: pi_esimd_emulator.cpp:2019
sycl::_V1::ext::intel::esimd::barrier
__ESIMD_API void barrier()
Generic work-group barrier.
Definition: memory.hpp:1641
export.hpp
PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_HALF
@ PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_HALF
Definition: pi.h:256
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:508
cm_surface_ptr_t::RegularImgPtr
cm_support::CmSurface2D * RegularImgPtr
Definition: pi_esimd_emulator.hpp:134
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:1637
PI_DEVICE_INFO_IMAGE_MAX_BUFFER_SIZE
@ PI_DEVICE_INFO_IMAGE_MAX_BUFFER_SIZE
Definition: pi.h:268
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:1645
PI_DEVICE_INFO_LOCAL_MEM_SIZE
@ PI_DEVICE_INFO_LOCAL_MEM_SIZE
Definition: pi.h:280
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:2059
piSamplerCreate
pi_result piSamplerCreate(pi_context, const pi_sampler_properties *, pi_sampler *)
Definition: pi_esimd_emulator.cpp:1555
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:1679
_pi_queue
PI queue mapping on to CUstream objects.
Definition: pi_cuda.hpp:395
_pi_context::Addr2CmBufferSVMLock
std::mutex Addr2CmBufferSVMLock
Definition: pi_esimd_emulator.hpp:104
PiPlatformCacheLock
static std::mutex * PiPlatformCacheLock
Definition: pi_esimd_emulator.cpp:133
PI_IMAGE_CHANNEL_TYPE_SIGNED_INT16
@ PI_IMAGE_CHANNEL_TYPE_SIGNED_INT16
Definition: pi.h:534
piextKernelSetArgMemObj
pi_result piextKernelSetArgMemObj(pi_kernel, pi_uint32, const pi_mem *)
Definition: pi_esimd_emulator.cpp:1412
PI_IMAGE_CHANNEL_TYPE_SIGNED_INT32
@ PI_IMAGE_CHANNEL_TYPE_SIGNED_INT32
Definition: pi.h:535
_pi_mem::MapHostPtr
char * MapHostPtr
Definition: pi_esimd_emulator.hpp:143
PI_DEVICE_TYPE_GPU
@ PI_DEVICE_TYPE_GPU
A PI device that is a GPU.
Definition: pi.h:206
pi_uint32
uint32_t pi_uint32
Definition: pi.h:142
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:1374
cm_surface_ptr_t::tag
SurfaceType tag
Definition: pi_esimd_emulator.hpp:129
kernel.hpp
pi_buff_rect_region_struct
Definition: pi.h:937
pi_device_binary_struct
This struct is a record of the device binary information.
Definition: pi.h:871
sycl::_V1::ext::intel::esimd::SurfaceIndex
unsigned int SurfaceIndex
Surface index type.
Definition: common.hpp:64
piextMemGetNativeHandle
pi_result piextMemGetNativeHandle(pi_mem, pi_native_handle *)
Gets the native handle of a PI mem object.
Definition: pi_esimd_emulator.cpp:1326
KernelInvocationContext
Definition: pi_esimd_emulator.cpp:198
piextContextGetNativeHandle
pi_result piextContextGetNativeHandle(pi_context, pi_native_handle *)
Gets the native handle of a PI context object.
Definition: pi_esimd_emulator.cpp:907
PI_DEVICE_INFO_LINKER_AVAILABLE
@ PI_DEVICE_INFO_LINKER_AVAILABLE
Definition: pi.h:287
PI_DEVICE_INFO_EXECUTION_CAPABILITIES
@ PI_DEVICE_INFO_EXECUTION_CAPABILITIES
Definition: pi.h:288
piPlatformGetInfo
pi_result piPlatformGetInfo(pi_platform Platform, pi_platform_info ParamName, size_t ParamValueSize, void *ParamValue, size_t *ParamValueSizeRet)
Definition: pi_esimd_emulator.cpp:463
PiESimdSurfaceMap
static std::unordered_map< unsigned int, _pi_mem * > * PiESimdSurfaceMap
Definition: pi_esimd_emulator.cpp:136
PI_DEVICE_INFO_PROFILING_TIMER_RESOLUTION
@ PI_DEVICE_INFO_PROFILING_TIMER_RESOLUTION
Definition: pi.h:283
_pi_kernel_exec_info
_pi_kernel_exec_info
Definition: pi.h:1426
_pi_mem::SurfacePtr
cm_surface_ptr_t SurfacePtr
Definition: pi_esimd_emulator.hpp:167
PI_MEM_TYPE_IMAGE2D
@ PI_MEM_TYPE_IMAGE2D
Definition: pi.h:483
_pi_platform::PiDeviceCache
std::unique_ptr< _pi_device > PiDeviceCache
Definition: pi_esimd_emulator.hpp:70
piContextGetInfo
pi_result piContextGetInfo(pi_context, pi_context_info, size_t, void *, size_t *)
Definition: pi_esimd_emulator.cpp:897
piextMemImageCreateWithNativeHandle
pi_result piextMemImageCreateWithNativeHandle(pi_native_handle, pi_context, bool, const pi_image_format *, const pi_image_desc *, pi_mem *)
Creates PI image object from a native handle.
Definition: pi_esimd_emulator.cpp:1335
range.hpp
PI_DEVICE_INFO_HALF_FP_CONFIG
@ PI_DEVICE_INFO_HALF_FP_CONFIG
Definition: pi.h:240
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:1651
PI_FP_DENORM
static constexpr pi_device_fp_config PI_FP_DENORM
Definition: pi.h:730
pi_mem_flags
pi_bitfield pi_mem_flags
Definition: pi.h:611
PI_IMAGE_CHANNEL_TYPE_FLOAT
@ PI_IMAGE_CHANNEL_TYPE_FLOAT
Definition: pi.h:540
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:2064
PI_USM_CONCURRENT_ACCESS
@ PI_USM_CONCURRENT_ACCESS
Definition: pi.h:1742
_pi_platform::DeviceCachePopulated
bool DeviceCachePopulated
Definition: pi_esimd_emulator.hpp:72
PI_DEVICE_INFO_VENDOR
@ PI_DEVICE_INFO_VENDOR
Definition: pi.h:296
pi_context_extended_deleter
void(* pi_context_extended_deleter)(void *user_data)
Definition: pi.h:1161
PI_MEM_TYPE_IMAGE1D_ARRAY
@ PI_MEM_TYPE_IMAGE1D_ARRAY
Definition: pi.h:487
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:1927
piMemGetInfo
pi_result piMemGetInfo(pi_mem, pi_mem_info, size_t, void *, size_t *)
Definition: pi_esimd_emulator.cpp:1128
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:2081
piPluginGetBackendOption
pi_result piPluginGetBackendOption(pi_platform, const char *frontend_option, const char **backend_option)
API to get backend specific option.
Definition: pi_esimd_emulator.cpp:175
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:383
PI_DEVICE_INFO_USM_SYSTEM_SHARED_SUPPORT
@ PI_DEVICE_INFO_USM_SYSTEM_SHARED_SUPPORT
Definition: pi.h:316
PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_SHORT
@ PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_SHORT
Definition: pi.h:244
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:2041
ESIMDEmuPluginDataVersion
#define ESIMDEmuPluginDataVersion
Definition: pi_esimd_emulator.cpp:143
piKernelGetInfo
pi_result piKernelGetInfo(pi_kernel, pi_kernel_info, size_t, void *, size_t *)
Definition: pi_esimd_emulator.cpp:1421
PI_DEVICE_INFO_GLOBAL_MEM_CACHELINE_SIZE
@ PI_DEVICE_INFO_GLOBAL_MEM_CACHELINE_SIZE
Definition: pi.h:274
PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_LONG
@ PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_LONG
Definition: pi.h:253
piextKernelSetArgPointer
pi_result piextKernelSetArgPointer(pi_kernel, pi_uint32, size_t, const void *)
Sets up pointer arguments for CL kernels.
Definition: pi_esimd_emulator.cpp:2005
piextEventGetNativeHandle
pi_result piextEventGetNativeHandle(pi_event, pi_native_handle *)
Gets the native handle of a PI event object.
Definition: pi_esimd_emulator.cpp:1547
_pi_platform::CmEmuVersion
std::string CmEmuVersion
Definition: pi_esimd_emulator.hpp:78
common.hpp
piextContextSetExtendedDeleter
pi_result piextContextSetExtendedDeleter(pi_context, pi_context_extended_deleter, void *)
Definition: pi_esimd_emulator.cpp:902
PI_DEVICE_INFO_USM_HOST_SUPPORT
@ PI_DEVICE_INFO_USM_HOST_SUPPORT
Definition: pi.h:312
PI_MEM_FLAGS_HOST_PTR_USE
constexpr pi_mem_flags PI_MEM_FLAGS_HOST_PTR_USE
Definition: pi.h:616
PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_FLOAT
@ PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_FLOAT
Definition: pi.h:254
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:1665
piTearDown
pi_result piTearDown(void *)
API to notify that the plugin should clean up its resources.
Definition: pi_esimd_emulator.cpp:2110
ESIMDEmuPluginInterfaceVersion
#define ESIMDEmuPluginInterfaceVersion
Definition: pi_esimd_emulator.cpp:147
sycl::_V1::ext::oneapi::experimental::detail::Alignment
@ Alignment
Definition: property.hpp:189
piEnqueueEventsWait
pi_result piEnqueueEventsWait(pi_queue, pi_uint32, const pi_event *, pi_event *)
Definition: pi_esimd_emulator.cpp:1569
PiESimdDeviceAccess
static sycl::detail::ESIMDEmuPluginOpaqueData * PiESimdDeviceAccess
Definition: pi_esimd_emulator.cpp:128
PI_DEVICE_INFO_PARENT_DEVICE
@ PI_DEVICE_INFO_PARENT_DEVICE
Definition: pi.h:304
PI_PLATFORM_INFO_VERSION
@ PI_PLATFORM_INFO_VERSION
Definition: pi.h:173
PI_IMAGE_CHANNEL_TYPE_UNSIGNED_INT16
@ PI_IMAGE_CHANNEL_TYPE_UNSIGNED_INT16
Definition: pi.h:537
PI_FP_ROUND_TO_INF
static constexpr pi_device_fp_config PI_FP_ROUND_TO_INF
Definition: pi.h:734
pi_uint64
uint64_t pi_uint64
Definition: pi.h:143
PI_DEVICE_INFO_PARTITION_TYPE
@ PI_DEVICE_INFO_PARTITION_TYPE
Definition: pi.h:308
sycl::_V1::ext::intel::esimd::detail::SLM_BTI
static constexpr SurfaceIndex SLM_BTI
Definition: common.hpp:98
_pi_event_info
_pi_event_info
Definition: pi.h:439
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:1055
_pi_image::BytesPerPixel
size_t BytesPerPixel
Definition: pi_esimd_emulator.hpp:198
_pi_sampler_info
_pi_sampler_info
Definition: pi.h:548
_pi_device_binary_property_struct
Definition: pi.h:768
pi_mem_properties
pi_bitfield pi_mem_properties
Definition: pi.h:627
pi_device_exec_capabilities
pi_bitfield pi_device_exec_capabilities
Definition: pi.h:574
_pi_program
Implementation of PI Program on CUDA Module object.
Definition: pi_cuda.hpp:760
_pi_sampler
Implementation of samplers for CUDA.
Definition: pi_cuda.hpp:987
PI_DEVICE_INFO_PRINTF_BUFFER_SIZE
@ PI_DEVICE_INFO_PRINTF_BUFFER_SIZE
Definition: pi.h:302
PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_DOUBLE
@ PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_DOUBLE
Definition: pi.h:255
_PI_ESIMD_PLUGIN_VERSION_STRING
#define _PI_ESIMD_PLUGIN_VERSION_STRING
Definition: pi_esimd_emulator.hpp:31
PI_PLATFORM_INFO_PROFILE
@ PI_PLATFORM_INFO_PROFILE
Definition: pi.h:171
getInfo< const char * >
ur_result_t getInfo< const char * >(size_t param_value_size, void *param_value, size_t *param_value_size_ret, const char *value)
Definition: ur.hpp:258
PI_DEVICE_INFO_MAX_CONSTANT_ARGS
@ PI_DEVICE_INFO_MAX_CONSTANT_ARGS
Definition: pi.h:278
PI_DEVICE_INFO_PCI_ADDRESS
@ PI_DEVICE_INFO_PCI_ADDRESS
Definition: pi.h:321
pi_native_handle
uintptr_t pi_native_handle
Definition: pi.h:146
PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_SHORT
@ PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_SHORT
Definition: pi.h:251
PI_DEVICE_INFO_GLOBAL_MEM_SIZE
@ PI_DEVICE_INFO_GLOBAL_MEM_SIZE
Definition: pi.h:276
InvokeImpl::get_range
static sycl::range< NDims > get_range(const size_t *Array)
Definition: pi_esimd_emulator.cpp:374
pi_sampler_properties
pi_bitfield pi_sampler_properties
Definition: pi.h:580
piKernelRetain
pi_result piKernelRetain(pi_kernel)
Definition: pi_esimd_emulator.cpp:1436
PI_DEVICE_MEM_CACHE_TYPE_READ_WRITE_CACHE
@ PI_DEVICE_MEM_CACHE_TYPE_READ_WRITE_CACHE
Definition: pi.h:224
PI_MEM_FLAGS_ACCESS_RW
constexpr pi_mem_flags PI_MEM_FLAGS_ACCESS_RW
Definition: pi.h:613
accessor_impl.hpp
_pi_event::release
pi_result release()
Definition: pi_cuda.cpp:738
cm_surface_ptr_t::TypeRegularImage
@ TypeRegularImage
Definition: pi_esimd_emulator.hpp:126
piclProgramCreateWithSource
pi_result piclProgramCreateWithSource(pi_context, pi_uint32, const char **, const size_t *, pi_program *)
Definition: pi_esimd_emulator.cpp:1358
_pi_image::Height
size_t Height
Definition: pi_esimd_emulator.hpp:197
_pi_image_format
Definition: pi.h:1004
PI_EVENT_COMPLETE
@ PI_EVENT_COMPLETE
Definition: pi.h:162
PI_IMAGE_CHANNEL_TYPE_UNORM_INT_101010
@ PI_IMAGE_CHANNEL_TYPE_UNORM_INT_101010
Definition: pi.h:532
getInfoArray
ur_result_t getInfoArray(size_t array_length, size_t param_value_size, void *param_value, size_t *param_value_size_ret, const T *value)
Definition: ur.hpp:235
PI_PLATFORM_INFO_VENDOR
@ PI_PLATFORM_INFO_VENDOR
Definition: pi.h:172
_pi_kernel_info
_pi_kernel_info
Definition: pi.h:402
piKernelRelease
pi_result piKernelRelease(pi_kernel)
Definition: pi_esimd_emulator.cpp:1438
PI_DEVICE_INFO_USM_SINGLE_SHARED_SUPPORT
@ PI_DEVICE_INFO_USM_SINGLE_SHARED_SUPPORT
Definition: pi.h:314
sycl::_V1::detail::pi::die
void die(const char *Message)
Definition: pi.cpp:531
PI_IMAGE_CHANNEL_ORDER_RGBA
@ PI_IMAGE_CHANNEL_ORDER_RGBA
Definition: pi.h:513
USMAllocationForceResidencyType::Device
@ Device
cm_surface_ptr_t::UPBufPtr
cm_support::CmBufferUP * UPBufPtr
Definition: pi_esimd_emulator.hpp:133
PI_DEVICE_INFO_QUEUE_PROPERTIES
@ PI_DEVICE_INFO_QUEUE_PROPERTIES
Definition: pi.h:242
PI_DEVICE_INFO_PARTITION_MAX_SUB_DEVICES
@ PI_DEVICE_INFO_PARTITION_MAX_SUB_DEVICES
Definition: pi.h:306
piEventGetInfo
pi_result piEventGetInfo(pi_event Event, pi_event_info ParamName, size_t ParamValueSize, void *ParamValue, size_t *ParamValueSizeRet)
Definition: pi_esimd_emulator.cpp:1442
sycl::_V1::detail::getNextPowerOfTwo
constexpr size_t getNextPowerOfTwo(size_t Var)
Definition: common.hpp:429
getInfoImpl
ur_result_t getInfoImpl(size_t param_value_size, void *param_value, size_t *param_value_size_ret, T value, size_t value_size, Assign &&assign_func)
Definition: ur.hpp:201
PI_EXT_ONEAPI_DEVICE_INFO_MAX_WORK_GROUPS_3D
@ PI_EXT_ONEAPI_DEVICE_INFO_MAX_WORK_GROUPS_3D
Definition: pi.h:351
piextPlatformGetNativeHandle
pi_result piextPlatformGetNativeHandle(pi_platform, pi_native_handle *)
Gets the native handle of a PI platform object.
Definition: pi_esimd_emulator.cpp:500
KernelFunc
std::function< void(const sycl::nd_item< NDims > &)> KernelFunc
Definition: pi_esimd_emulator.cpp:192
PI_MEM_TYPE_IMAGE1D_BUFFER
@ PI_MEM_TYPE_IMAGE1D_BUFFER
Definition: pi.h:488
piProgramBuild
pi_result piProgramBuild(pi_program, pi_uint32, const pi_device *, const char *, void(*)(pi_program, void *), void *)
Definition: pi_esimd_emulator.cpp:1381
KernelInvocationContext::GlobalSize
const sycl::range< NDims > & GlobalSize
Definition: pi_esimd_emulator.cpp:201
_pi_image_desc::image_width
size_t image_width
Definition: pi.h:1011
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:2009
PI_DEVICE_INFO_SINGLE_FP_CONFIG
@ PI_DEVICE_INFO_SINGLE_FP_CONFIG
Definition: pi.h:239
_PI_PLUGIN_VERSION_CHECK
#define _PI_PLUGIN_VERSION_CHECK(PI_API_VERSION, PI_PLUGIN_VERSION)
Definition: pi.h:119
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:1937
PI_DEVICE_INFO_BUILD_ON_SUBDEVICE
@ PI_DEVICE_INFO_BUILD_ON_SUBDEVICE
Definition: pi.h:330
_pi_context::Addr2CmBufferSVM
std::unordered_map< void *, cm_support::CmBufferSVM * > Addr2CmBufferSVM
Definition: pi_esimd_emulator.hpp:102
PI_DEVICE_INFO_MAX_WORK_ITEM_DIMENSIONS
@ PI_DEVICE_INFO_MAX_WORK_ITEM_DIMENSIONS
Definition: pi.h:236
PI_DEVICE_INFO_LOCAL_MEM_TYPE
@ PI_DEVICE_INFO_LOCAL_MEM_TYPE
Definition: pi.h:279
PI_DEVICE_INFO_IMAGE3D_MAX_DEPTH
@ PI_DEVICE_INFO_IMAGE3D_MAX_DEPTH
Definition: pi.h:267
PI_DEVICE_INFO_COMPILER_AVAILABLE
@ PI_DEVICE_INFO_COMPILER_AVAILABLE
Definition: pi.h:286
_PI_H_VERSION_STRING
#define _PI_H_VERSION_STRING
Definition: pi.h:113
PiESimdSurfaceMapLock
static std::mutex * PiESimdSurfaceMapLock
Definition: pi_esimd_emulator.cpp:139
PI_FP_FMA
static constexpr pi_device_fp_config PI_FP_FMA
Definition: pi.h:735
pi_image_offset_struct::z
size_t z
Definition: pi.h:949
backend_types.hpp
PI_DEVICE_INFO_VERSION
@ PI_DEVICE_INFO_VERSION
Definition: pi.h:299
PI_EXT_PLATFORM_BACKEND_ESIMD
@ PI_EXT_PLATFORM_BACKEND_ESIMD
The backend is ESIMD.
Definition: pi.h:218
PI_DEVICE_INFO_MAX_MEM_BANDWIDTH
@ PI_DEVICE_INFO_MAX_MEM_BANDWIDTH
Definition: pi.h:327
PI_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8
@ PI_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8
Definition: pi.h:536
PI_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32
@ PI_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32
Definition: pi.h:538
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:1673
pi_queue_properties
pi_bitfield pi_queue_properties
Definition: pi.h:646
PI_EXT_ONEAPI_DEVICE_INFO_MAX_WORK_GROUPS_2D
@ PI_EXT_ONEAPI_DEVICE_INFO_MAX_WORK_GROUPS_2D
Definition: pi.h:350
PI_DEVICE_INFO_ATOMIC_64
@ PI_DEVICE_INFO_ATOMIC_64
Definition: pi.h:341
PI_DEVICE_INFO_MAX_WRITE_IMAGE_ARGS
@ PI_DEVICE_INFO_MAX_WRITE_IMAGE_ARGS
Definition: pi.h:262
_pi_platform_info
_pi_platform_info
Definition: pi.h:168
MaxMessageSize
constexpr size_t MaxMessageSize
Definition: pi_esimd_emulator.cpp:153
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:1659
piProgramGetBuildInfo
pi_result piProgramGetBuildInfo(pi_program, pi_device, pi_program_build_info, size_t, void *, size_t *)
Definition: pi_esimd_emulator.cpp:1386
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:1579
PI_IMAGE_CHANNEL_TYPE_UNORM_SHORT_555
@ PI_IMAGE_CHANNEL_TYPE_UNORM_SHORT_555
Definition: pi.h:531
PI_IMAGE_CHANNEL_ORDER_ARGB
@ PI_IMAGE_CHANNEL_ORDER_ARGB
Definition: pi.h:515
_pi_image_desc
Definition: pi.h:1009
PI_DEVICE_INFO_ENDIAN_LITTLE
@ PI_DEVICE_INFO_ENDIAN_LITTLE
Definition: pi.h:284
piQueueCreate
pi_result piQueueCreate(pi_context Context, pi_device Device, pi_queue_properties Properties, pi_queue *Queue)
Definition: pi_esimd_emulator.cpp:977
_pi_queue::CmQueuePtr
cm_support::CmQueue * CmQueuePtr
Definition: pi_esimd_emulator.hpp:115
piContextRelease
pi_result piContextRelease(pi_context Context)
Definition: pi_esimd_emulator.cpp:927
PI_EXT_PLATFORM_INFO_BACKEND
@ PI_EXT_PLATFORM_INFO_BACKEND
Definition: pi.h:174
piextQueueGetNativeHandle
pi_result piextQueueGetNativeHandle(pi_queue, pi_native_handle *, int32_t *)
Gets the native handle of a PI queue object.
Definition: pi_esimd_emulator.cpp:1045
piQueueFlush
pi_result piQueueFlush(pi_queue)
Definition: pi_esimd_emulator.cpp:1038
kernel_desc.hpp
_pi_event
PI Event mapping to CUevent.
Definition: pi_cuda.hpp:632
piQueueRelease
pi_result piQueueRelease(pi_queue Queue)
Definition: pi_esimd_emulator.cpp:1017
DIE_NO_IMPLEMENTATION
#define DIE_NO_IMPLEMENTATION
Definition: pi_esimd_emulator.cpp:394
PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_INT
@ PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_INT
Definition: pi.h:252
PI_DEVICE_INFO_MAX_READ_IMAGE_ARGS
@ PI_DEVICE_INFO_MAX_READ_IMAGE_ARGS
Definition: pi.h:261
_pi_mem::MappingsMutex
std::mutex MappingsMutex
Definition: pi_esimd_emulator.hpp:165
PI_FP_INF_NAN
static constexpr pi_device_fp_config PI_FP_INF_NAN
Definition: pi.h:731
PI_DEVICE_INFO_GPU_SUBSLICES_PER_SLICE
@ PI_DEVICE_INFO_GPU_SUBSLICES_PER_SLICE
Definition: pi.h:325
piKernelCreate
pi_result piKernelCreate(pi_program, const char *, pi_kernel *)
Definition: pi_esimd_emulator.cpp:1404
piEventRelease
pi_result piEventRelease(pi_event Event)
Definition: pi_esimd_emulator.cpp:1526
ESimdEmuVersionString
static char ESimdEmuVersionString[32]
Definition: pi_esimd_emulator.cpp:150
PI_MEM_TYPE_IMAGE2D_ARRAY
@ PI_MEM_TYPE_IMAGE2D_ARRAY
Definition: pi.h:485
piMemRetain
pi_result piMemRetain(pi_mem Mem)
Definition: pi_esimd_emulator.cpp:1132
PI_EXT_INTEL_DEVICE_INFO_MAX_COMPUTE_QUEUE_INDICES
@ PI_EXT_INTEL_DEVICE_INFO_MAX_COMPUTE_QUEUE_INDICES
Definition: pi.h:340
piKernelGetGroupInfo
pi_result piKernelGetGroupInfo(pi_kernel, pi_device, pi_kernel_group_info, size_t, void *, size_t *)
Definition: pi_esimd_emulator.cpp:1425
_pi_mem_info
_pi_mem_info
Definition: pi.h:1025
pi_image_offset_struct
Definition: pi.h:946
piMemRelease
pi_result piMemRelease(pi_mem Mem)
Definition: pi_esimd_emulator.cpp:1140
sycl::_V1::detail::pi::assertion
void assertion(bool Condition, const char *Message=nullptr)
Definition: pi.cpp:536
PiPlatformCache
static pi_platform PiPlatformCache
Definition: pi_esimd_emulator.cpp:131
cm_surface_ptr_t::TypeRegularBuffer
@ TypeRegularBuffer
Definition: pi_esimd_emulator.hpp:124
PI_EXT_ONEAPI_DEVICE_INFO_MAX_WORK_GROUPS_1D
@ PI_EXT_ONEAPI_DEVICE_INFO_MAX_WORK_GROUPS_1D
Definition: pi.h:349
PI_DEVICE_INFO_VENDOR_ID
@ PI_DEVICE_INFO_VENDOR_ID
Definition: pi.h:234
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:1551
pi_image_region_struct
Definition: pi.h:955
piDeviceRetain
pi_result piDeviceRetain(pi_device Device)
Definition: pi_esimd_emulator.cpp:596
PI_DEVICE_INFO_EXTENSIONS
@ PI_DEVICE_INFO_EXTENSIONS
Definition: pi.h:301
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:1932
piProgramRetain
pi_result piProgramRetain(pi_program)
Definition: pi_esimd_emulator.cpp:1391
getInfo
ur_result_t getInfo(size_t param_value_size, void *param_value, size_t *param_value_size_ret, T value)
Definition: ur.hpp:222
KernelInvocationContext::LocalSize
const sycl::range< NDims > & LocalSize
Definition: pi_esimd_emulator.cpp:200
PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_FLOAT
@ PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_FLOAT
Definition: pi.h:247
cm_surface_ptr_t::UPImgPtr
cm_support::CmSurface2DUP * UPImgPtr
Definition: pi_esimd_emulator.hpp:135
_pi_mem::SurfaceIndex
unsigned int SurfaceIndex
Definition: pi_esimd_emulator.hpp:149
_pi_usm_migration_flags
_pi_usm_migration_flags
Definition: pi.h:1763
KernelInvocationContext::GlobalOffset
const sycl::id< NDims > & GlobalOffset
Definition: pi_esimd_emulator.cpp:202
cm_surface_ptr_t::TypeUserProvidedBuffer
@ TypeUserProvidedBuffer
Definition: pi_esimd_emulator.hpp:125
piextQueueCreate
pi_result piextQueueCreate(pi_context Context, pi_device Device, pi_queue_properties *Properties, pi_queue *Queue)
Definition: pi_esimd_emulator.cpp:962
PI_QUEUE_FLAG_OUT_OF_ORDER_EXEC_MODE_ENABLE
constexpr pi_queue_properties PI_QUEUE_FLAG_OUT_OF_ORDER_EXEC_MODE_ENABLE
Definition: pi.h:650
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:911
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:1851
PI_FP_CORRECTLY_ROUNDED_DIVIDE_SQRT
static constexpr pi_device_fp_config PI_FP_CORRECTLY_ROUNDED_DIVIDE_SQRT
Definition: pi.h:737
sycl_get_cm_surface_index
unsigned int sycl_get_cm_surface_index(void *PtrInput)
Definition: pi_esimd_emulator.cpp:292
_pi_platform::populateDeviceCacheIfNeeded
pi_result populateDeviceCacheIfNeeded()
Definition: pi_esimd_emulator.cpp:549
_pi_event::IsDummyEvent
bool IsDummyEvent
Definition: pi_esimd_emulator.hpp:207
nd_item.hpp
PI_DEVICE_INFO_TYPE
@ PI_DEVICE_INFO_TYPE
Definition: pi.h:233
piMemImageGetInfo
pi_result piMemImageGetInfo(pi_mem, pi_image_info, size_t, void *, size_t *)
Definition: pi_esimd_emulator.cpp:1758
PI_DEVICE_INFO_MAX_WORK_GROUP_SIZE
@ PI_DEVICE_INFO_MAX_WORK_GROUP_SIZE
Definition: pi.h:238
piPlatformsGet
pi_result piPlatformsGet(pi_uint32 NumEntries, pi_platform *Platforms, pi_uint32 *NumPlatforms)
Definition: pi_esimd_emulator.cpp:420
_pi_program_info
_pi_program_info
Definition: pi.h:361
_pi_profiling_info
_pi_profiling_info
Definition: pi.h:600
RangeBuilder< 3 >::create
static sycl::range< 3 > create(Gen G)
Definition: pi_esimd_emulator.cpp:222
pi_device_affinity_domain
pi_bitfield pi_device_affinity_domain
Definition: pi.h:714
CONTINUE_NO_IMPLEMENTATION
#define CONTINUE_NO_IMPLEMENTATION
Definition: pi_esimd_emulator.cpp:402
PI_DEVICE_INFO_IMAGE_SUPPORT
@ PI_DEVICE_INFO_IMAGE_SUPPORT
Definition: pi.h:260
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:1207
PI_EXT_DEVICE_INFO_ATOMIC_FENCE_ORDER_CAPABILITIES
@ PI_EXT_DEVICE_INFO_ATOMIC_FENCE_ORDER_CAPABILITIES
Definition: pi.h:354
PI_EXT_INTEL_DEVICE_INFO_MEM_CHANNEL_SUPPORT
@ PI_EXT_INTEL_DEVICE_INFO_MEM_CHANNEL_SUPPORT
Definition: pi.h:356
PI_DEVICE_INFO_MAX_PARAMETER_SIZE
@ PI_DEVICE_INFO_MAX_PARAMETER_SIZE
Definition: pi.h:271
PI_EVENT_INFO_COMMAND_EXECUTION_STATUS
@ PI_EVENT_INFO_COMMAND_EXECUTION_STATUS
Definition: pi.h:443
libCMBatch
Definition: pi_esimd_emulator.cpp:265
_pi_kernel_group_info
_pi_kernel_group_info
Definition: pi.h:411
PI_DEVICE_INFO_SUB_GROUP_SIZES_INTEL
@ PI_DEVICE_INFO_SUB_GROUP_SIZES_INTEL
Definition: pi.h:311
PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_HALF
@ PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_HALF
Definition: pi.h:249
PI_EXT_DEVICE_INFO_ATOMIC_MEMORY_ORDER_CAPABILITIES
@ PI_EXT_DEVICE_INFO_ATOMIC_MEMORY_ORDER_CAPABILITIES
Definition: pi.h:342
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:1833
_pi_mem::SurfaceLock
std::mutex SurfaceLock
Definition: pi_esimd_emulator.hpp:145
piProgramRelease
pi_result piProgramRelease(pi_program)
Definition: pi_esimd_emulator.cpp:1393
std::cout
__SYCL_EXTERN_STREAM_ATTRS ostream cout
Linked to standard output.
piContextRetain
pi_result piContextRetain(pi_context Context)
Definition: pi_esimd_emulator.cpp:917
_pi_mem_alloc_info
_pi_mem_alloc_info
Definition: pi.h:1746
PI_DEVICE_INFO_MAX_CLOCK_FREQUENCY
@ PI_DEVICE_INFO_MAX_CLOCK_FREQUENCY
Definition: pi.h:257
piEventSetStatus
pi_result piEventSetStatus(pi_event, pi_int32)
Definition: pi_esimd_emulator.cpp:1514
_pi_context::Device
pi_device Device
Definition: pi_esimd_emulator.hpp:98
PI_DEVICE_INFO_IMAGE2D_MAX_WIDTH
@ PI_DEVICE_INFO_IMAGE2D_MAX_WIDTH
Definition: pi.h:263
PI_DEVICE_INFO_AVAILABLE
@ PI_DEVICE_INFO_AVAILABLE
Definition: pi.h:285
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:300
pi_usm_mem_properties
pi_bitfield pi_usm_mem_properties
Definition: pi.h:633
pi_map_flags
pi_bitfield pi_map_flags
Definition: pi.h:621
piEventGetProfilingInfo
pi_result piEventGetProfilingInfo(pi_event Event, pi_profiling_info ParamName, size_t ParamValueSize, void *ParamValue, size_t *ParamValueSizeRet)
Definition: pi_esimd_emulator.cpp:1478
piextEnqueueDeviceGlobalVariableWrite
pi_result piextEnqueueDeviceGlobalVariableWrite(pi_queue, pi_program, const char *, pi_bool, size_t, size_t, const void *, pi_uint32, const pi_event *, pi_event *)
Device global variable.
Definition: pi_esimd_emulator.cpp:2087
PI_DEVICE_INFO_IMAGE3D_MAX_HEIGHT
@ PI_DEVICE_INFO_IMAGE3D_MAX_HEIGHT
Definition: pi.h:266
PI_DEVICE_INFO_GPU_EU_SIMD_WIDTH
@ PI_DEVICE_INFO_GPU_EU_SIMD_WIDTH
Definition: pi.h:323
_pi_program_build_info
_pi_program_build_info
Definition: pi.h:177
PI_MEM_TYPE_IMAGE3D
@ PI_MEM_TYPE_IMAGE3D
Definition: pi.h:484
_pi_event::OwnerQueue
cm_support::CmQueue * OwnerQueue
Definition: pi_esimd_emulator.hpp:205
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:1827
piQueueGetInfo
pi_result piQueueGetInfo(pi_queue, pi_queue_info, size_t, void *, size_t *)
Definition: pi_esimd_emulator.cpp:1005
piSamplerRetain
pi_result piSamplerRetain(pi_sampler)
Definition: pi_esimd_emulator.cpp:1565
PI_USM_CONCURRENT_ATOMIC_ACCESS
@ PI_USM_CONCURRENT_ATOMIC_ACCESS
Definition: pi.h:1743
RangeBuilder< 2 >::create
static sycl::range< 2 > create(Gen G)
Definition: pi_esimd_emulator.cpp:217
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:1368
PI_IMAGE_CHANNEL_TYPE_SNORM_INT8
@ PI_IMAGE_CHANNEL_TYPE_SNORM_INT8
Definition: pi.h:526
PiTrace
static void PiTrace(std::string TraceString)
Definition: pi_esimd_emulator.cpp:117
pi_device_partition_property
intptr_t pi_device_partition_property
Definition: pi.h:701
PI_FP_ROUND_TO_ZERO
static constexpr pi_device_fp_config PI_FP_ROUND_TO_ZERO
Definition: pi.h:733
PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_LONG
@ PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_LONG
Definition: pi.h:246
ConvertPiImageFormatToCmFormat
cm_support::CM_SURFACE_FORMAT ConvertPiImageFormatToCmFormat(const pi_image_format *PiFormat)
Definition: pi_esimd_emulator.cpp:1183
piextKernelGetNativeHandle
pi_result piextKernelGetNativeHandle(pi_kernel, pi_native_handle *)
Gets the native handle of a PI kernel object.
Definition: pi_esimd_emulator.cpp:1917
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:1839
pi_int32
int32_t pi_int32
Definition: pi.h:141
_pi_context
PI context mapping to a CUDA context object.
Definition: pi_cuda.hpp:170
piextUSMEnqueueMemset2D
pi_result piextUSMEnqueueMemset2D(pi_queue, void *, size_t, int, size_t, size_t, pi_uint32, const pi_event *, pi_event *)
USM 2D Memset API.
Definition: pi_esimd_emulator.cpp:2030
piEventRetain
pi_result piEventRetain(pi_event Event)
Definition: pi_esimd_emulator.cpp:1516
PI_IMAGE_CHANNEL_TYPE_UNORM_INT16
@ PI_IMAGE_CHANNEL_TYPE_UNORM_INT16
Definition: pi.h:529
PI_QUEUE_FLAGS
constexpr pi_queue_properties PI_QUEUE_FLAGS
Definition: pi.h:647
_pi_device
PI device mapping to a CUdevice.
Definition: pi_cuda.hpp:83
PI_DEVICE_INFO_PARTITION_PROPERTIES
@ PI_DEVICE_INFO_PARTITION_PROPERTIES
Definition: pi.h:305
_pi_kernel_sub_group_info
_pi_kernel_sub_group_info
Definition: pi.h:432
PI_PLATFORM_INFO_EXTENSIONS
@ PI_PLATFORM_INFO_EXTENSIONS
Definition: pi.h:169
_pi_context::checkSurfaceArgument
bool checkSurfaceArgument(pi_mem_flags Flags, void *HostPtr)
Definition: pi_esimd_emulator.cpp:945
PI_EXT_DEVICE_INFO_ATOMIC_FENCE_SCOPE_CAPABILITIES
@ PI_EXT_DEVICE_INFO_ATOMIC_FENCE_SCOPE_CAPABILITIES
Definition: pi.h:355
PI_DEVICE_INFO_SUB_GROUP_INDEPENDENT_FORWARD_PROGRESS
@ PI_DEVICE_INFO_SUB_GROUP_INDEPENDENT_FORWARD_PROGRESS
Definition: pi.h:310
_pi_buffer::Size
size_t Size
Definition: pi_esimd_emulator.hpp:185
PI_EXT_ONEAPI_DEVICE_INFO_MAX_GLOBAL_WORK_GROUPS
@ PI_EXT_ONEAPI_DEVICE_INFO_MAX_GLOBAL_WORK_GROUPS
Definition: pi.h:348