DPC++ Runtime
Runtime libraries for oneAPI DPC++
|
|
Go to the documentation of this file.
32 #include <esimdemu_support.h>
42 #include <string_view>
48 #define ARG_UNUSED(x) (void)x
54 template <
typename T,
typename Assign>
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;
62 AssignFunc(ParamValue, Value, ValueSize);
64 if (ParamValueSizeRet !=
nullptr) {
65 *ParamValueSizeRet = ValueSize;
72 size_t *ParamValueSizeRet,
T Value) {
73 auto assignment = [](
void *ParamValue,
T Value,
size_t ValueSize) {
75 *
static_cast<T *
>(ParamValue) = Value;
77 return getInfoImpl(ParamValueSize, ParamValue, ParamValueSizeRet, Value,
78 sizeof(
T), assignment);
83 void *ParamValue,
size_t *ParamValueSizeRet,
T *Value) {
84 return getInfoImpl(ParamValueSize, ParamValue, ParamValueSizeRet, Value,
85 ArrayLength *
sizeof(
T),
memcpy);
90 size_t *ParamValueSizeRet,
const char *Value) {
91 return getInfoArray(strlen(Value) + 1, ParamValueSize, ParamValue,
92 ParamValueSizeRet, Value);
97 ReturnHelper(
size_t ArgParamValueSize,
void *ArgParamValue,
98 size_t *ArgParamValueSizeRet)
99 : ParamValueSize(ArgParamValueSize), ParamValue(ArgParamValue),
100 ParamValueSizeRet(ArgParamValueSizeRet) {}
102 template <
class T>
pi_result operator()(
const T &t) {
103 return getInfo(ParamValueSize, ParamValue, ParamValueSizeRet, t);
107 size_t ParamValueSize;
109 size_t *ParamValueSizeRet;
117 static void PiTrace(std::string TraceString) {
137 new std::unordered_map<unsigned int, _pi_mem *>;
143 #define ESIMDEmuPluginDataVersion 0
147 #define ESIMDEmuPluginInterfaceVersion 1
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 =
"";
186 return PI_ERROR_INVALID_VALUE;
192 using KernelFunc = std::function<void(
const sycl::nd_item<NDims> &)>;
212 template <
typename Gen>
static sycl::range<1>
create(Gen G) {
213 return sycl::range<1>{G(0)};
217 template <
typename Gen>
static sycl::range<2>
create(Gen G) {
218 return sycl::range<2>{G(0), G(1)};
222 template <
typename Gen>
static sycl::range<3>
create(Gen G) {
223 return sycl::range<3>{G(0), G(1), G(2)};
232 sycl::range<NDims> GroupSize{
233 sycl::detail::InitializedVal<NDims, sycl::range>::template get<0>()};
235 for (
int i = 0; i < NDims; ++i) {
240 [](
int i) {
return cm_support::get_thread_idx(i); });
243 [](
int i) {
return cm_support::get_group_idx(i); });
245 const sycl::group<NDims> Group = IDBuilder::createGroup<NDims>(
248 const sycl::id<NDims> GlobalID =
251 const sycl::item<NDims,
true> GlobalItem =
252 IDBuilder::createItem<NDims, true>(ctx->
GlobalSize, GlobalID,
255 const sycl::item<NDims,
false> LocalItem =
256 IDBuilder::createItem<NDims, false>(ctx->
LocalSize, LocalID);
258 const sycl::nd_item<NDims> NDItem =
259 IDBuilder::createNDItem<NDims>(GlobalItem, LocalItem, Group);
268 std::vector<uint32_t> GroupDim, SpaceDim;
272 : MKernel(Kernel), GroupDim{1, 1, 1}, SpaceDim{1, 1, 1} {}
275 const sycl::range<DIMS> &GlobalSize,
276 const sycl::id<DIMS> &GlobalOffset) {
278 for (
int I = 0; I < DIMS; I++) {
279 SpaceDim[I] = (uint32_t)LocalSize[I];
280 GroupDim[I] = (uint32_t)(GlobalSize[I] / LocalSize[I]);
284 MKernel, LocalSize, GlobalSize, GlobalOffset};
286 EsimdemuKernel{
reinterpret_cast<fptrVoid
>(InvokeKernel<DIMS>),
287 GroupDim.data(), SpaceDim.data()}
288 .launchMT(
sizeof(InvokeKernelArg), &InvokeKernelArg);
301 uint32_t *Width, std::mutex **BufMtxLock) {
310 *Width =
static_cast<uint32_t
>(Buf->
Size);
318 uint32_t *Width, uint32_t *Height, uint32_t *Bpp,
319 std::mutex **ImgMtxLock) {
329 *Width =
static_cast<uint32_t
>(Img->
Width) * (*Bpp);
330 *Height =
static_cast<uint32_t
>(Img->
Height);
337 sycl::detail::ESIMDDeviceInterface::ESIMDDeviceInterface() {
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;
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]));
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]};
384 const size_t *GlobalWorkSize,
385 const size_t *LocalWorkSize) {
388 sycl::id<NDims>{
get_range(GlobalWorkOffset)});
394 #define DIE_NO_IMPLEMENTATION \
395 if (PrintPiTrace) { \
396 std::cerr << "Not Implemented : " << __FUNCTION__ \
397 << " - File : " << __FILE__; \
398 std::cerr << " / Line : " << __LINE__ << std::endl; \
400 return PI_ERROR_INVALID_OPERATION;
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; \
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__ << ")" \
418 return PI_ERROR_INVALID_OPERATION;
423 static const char *PiTraceEnv = std::getenv(
"SYCL_PI_TRACE");
424 static const int PiTraceValue = PiTraceEnv ? std::stoi(PiTraceEnv) : 0;
426 if (PiTraceValue == -1) {
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;
445 if (Platforms ==
nullptr && NumPlatforms ==
nullptr) {
446 return PI_ERROR_INVALID_VALUE;
456 if (Platforms && NumEntries > 0) {
464 size_t ParamValueSize,
void *ParamValue,
465 size_t *ParamValueSizeRet) {
466 if (Platform ==
nullptr) {
467 return PI_ERROR_INVALID_PLATFORM;
469 ReturnHelper ReturnValue(ParamValueSize, ParamValue, ParamValueSizeRet);
473 return ReturnValue(
"Intel(R) ESIMD_EMULATOR/GPU");
476 return ReturnValue(
"Intel(R) Corporation");
482 return ReturnValue(
"FULL_PROFILE");
485 return ReturnValue(
"");
488 return getInfo<pi_platform_backend>(ParamValueSize, ParamValue,
494 die(
"Unsupported ParamName in piPlatformGetInfo");
511 if (Platform ==
nullptr) {
512 return PI_ERROR_INVALID_PLATFORM;
516 if (Res != PI_SUCCESS) {
524 *NumDevices = DeviceCount;
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;
537 if (DeviceCount == 0) {
555 cm_support::CmDevice *CmDevice =
nullptr;
564 unsigned int Version = 0;
566 int Result = cm_support::CreateCmDevice(CmDevice, Version);
568 if (Result != cm_support::CM_SUCCESS) {
569 return PI_ERROR_INVALID_DEVICE;
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;
586 std::ostringstream StrFormat;
587 StrFormat << (int)(Version / 100) <<
"." << (int)(Version % 10);
589 std::unique_ptr<_pi_device>
Device(
590 new _pi_device(
this, CmDevice, StrFormat.str()));
597 if (Device ==
nullptr) {
598 return PI_ERROR_INVALID_DEVICE;
607 if (Device ==
nullptr) {
608 return PI_ERROR_INVALID_DEVICE;
617 size_t ParamValueSize,
void *ParamValue,
618 size_t *ParamValueSizeRet) {
619 ReturnHelper ReturnValue(ParamValueSize, ParamValue, ParamValueSizeRet);
627 return ReturnValue(Device->Platform);
629 return ReturnValue(
"ESIMD_EMULATOR");
631 return ReturnValue(
pi_bool{
true});
641 return ReturnValue(
"Intel(R) Corporation");
643 return ReturnValue(
size_t{8192});
645 return ReturnValue(
size_t{8192});
647 return ReturnValue(
pi_bool{1});
652 return ReturnValue(
"cl_khr_fp64");
654 return ReturnValue(Device->VersionStr.c_str());
656 return ReturnValue(
pi_bool{
true});
658 return ReturnValue(
pi_bool{
false});
660 return ReturnValue(
pi_bool{
false});
674 return ReturnValue(
size_t{256});
682 return ReturnValue(
size_t{2048});
688 return ReturnValue(
"");
694 } MaxGroupSize = {{256, 256, 1}};
695 return ReturnValue(MaxGroupSize);
726 return ReturnValue(Supported);
730 pi_uint32{
sizeof(
void *) * std::numeric_limits<unsigned char>::digits});
734 return ReturnValue(
pi_bool{
true});
736 return ReturnValue(
pi_bool{
true});
743 return ReturnValue(
size_t{0x80000000});
746 return ReturnValue(
size_t{0});
752 return ReturnValue(
size_t{32});
784 return ReturnValue(
pi_bool{
false});
787 return ReturnValue(
size_t{0});
790 return ReturnValue(
"");
793 return ReturnValue(
size_t{1024});
795 return ReturnValue(
pi_bool{
false});
806 return ReturnValue(
"FULL_PROFILE");
811 return ReturnValue(
size_t{1});
818 return ReturnValue(
pi_bool{
false});
821 return ReturnValue(
pi_bool{
false});
823 return ReturnValue(
pi_bool{
false});
868 void (*PFnNotify)(
const char *ErrInfo,
869 const void *PrivateInfo,
size_t CB,
876 if (NumDevices != 1) {
877 return PI_ERROR_INVALID_VALUE;
879 if (Devices ==
nullptr) {
880 return PI_ERROR_INVALID_DEVICE;
882 if (RetContext ==
nullptr) {
883 return PI_ERROR_INVALID_VALUE;
889 }
catch (
const std::bad_alloc &) {
890 return PI_ERROR_OUT_OF_HOST_MEMORY;
892 return PI_ERROR_UNKNOWN;
918 if (Context ==
nullptr) {
919 return PI_ERROR_INVALID_CONTEXT;
922 ++(Context->RefCount);
928 if (Context ==
nullptr || (Context->RefCount <= 0)) {
929 return PI_ERROR_INVALID_CONTEXT;
932 if (--(Context->RefCount) == 0) {
947 if (HostPtr ==
nullptr) {
948 PiTrace(
"HostPtr argument is required for "
949 "PI_MEM_FLAGS_HOST_PTR_USE/COPY");
955 PiTrace(
"PI_MEM_FLAGS_HOST_PTR_USE and _COPY cannot be used together");
968 return PI_ERROR_INVALID_VALUE;
971 assert(Properties[2] == 0);
972 if (Properties[2] != 0)
973 return PI_ERROR_INVALID_VALUE;
984 return PI_ERROR_INVALID_QUEUE_PROPERTIES;
987 cm_support::CmQueue *CmQueue =
nullptr;
990 if (Result != cm_support::CM_SUCCESS) {
991 return PI_ERROR_INVALID_CONTEXT;
995 *Queue =
new _pi_queue(Context, CmQueue);
996 }
catch (
const std::bad_alloc &) {
997 return PI_ERROR_OUT_OF_HOST_MEMORY;
999 return PI_ERROR_UNKNOWN;
1010 if (Queue ==
nullptr) {
1011 return PI_ERROR_INVALID_QUEUE;
1013 ++(Queue->RefCount);
1018 if ((Queue ==
nullptr) || (Queue->
CmQueuePtr ==
nullptr)) {
1019 return PI_ERROR_INVALID_QUEUE;
1022 if (--(Queue->RefCount) == 0) {
1056 void *HostPtr,
pi_mem *RetMem,
1061 PiTrace(
"Invalid memory attribute for piMemBufferCreate");
1062 return PI_ERROR_INVALID_OPERATION;
1065 if (Context ==
nullptr) {
1066 return PI_ERROR_INVALID_CONTEXT;
1068 if (RetMem ==
nullptr) {
1069 return PI_ERROR_INVALID_VALUE;
1074 return PI_ERROR_INVALID_OPERATION;
1077 char *MapBasePtr =
nullptr;
1080 int Status = cm_support::CM_FAILURE;
1085 static_cast<unsigned int>(Size), HostPtr, CmBuf.
UPBufPtr);
1095 reinterpret_cast<const unsigned char *
>(HostPtr),
nullptr,
1096 static_cast<unsigned int>(Size));
1100 if (Status != cm_support::CM_SUCCESS) {
1101 return PI_ERROR_INVALID_OPERATION;
1105 pi_cast<char *>(cm_support::get_surface_base_addr(CmIndex->get_data()));
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;
1113 return PI_ERROR_UNKNOWN;
1119 PiTrace(
"Failure from CM-managed buffer creation");
1120 return PI_ERROR_INVALID_MEM_OBJECT;
1123 (*PiESimdSurfaceMap)[(*RetMem)->SurfaceIndex] = *RetMem;
1133 if (Mem ==
nullptr) {
1134 return PI_ERROR_INVALID_MEM_OBJECT;
1141 if ((Mem ==
nullptr) || (Mem->RefCount == 0)) {
1142 return PI_ERROR_INVALID_MEM_OBJECT;
1145 if (--(Mem->RefCount) == 0) {
1150 PiTrace(
"Failure from Buffer/Image deletion");
1151 return PI_ERROR_INVALID_MEM_OBJECT;
1160 int Status = cm_support::CM_FAILURE;
1175 "Surface Deletion Failure from CM_EMU");
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},
1191 cm_support::CM_SURFACE_FORMAT_A8R8G8B8},
1194 cm_support::CM_SURFACE_FORMAT_A8R8G8B8},
1197 cm_support::CM_SURFACE_FORMAT_R32G32B32A32F},
1199 auto Result = pi2cm.find(
1201 if (Result != pi2cm.end()) {
1202 return Result->second;
1204 return cm_support::CM_SURFACE_FORMAT_UNKNOWN;
1212 PiTrace(
"Invalid memory attribute for piMemImageCreate");
1213 return PI_ERROR_INVALID_OPERATION;
1216 if (ImageFormat ==
nullptr || ImageDesc ==
nullptr)
1217 return PI_ERROR_INVALID_IMAGE_FORMAT_DESCRIPTOR;
1230 return PI_ERROR_INVALID_MEM_OBJECT;
1233 auto BytesPerPixel = 4;
1255 return PI_ERROR_IMAGE_FORMAT_NOT_SUPPORTED;
1260 return PI_ERROR_INVALID_OPERATION;
1263 cm_support::CM_SURFACE_FORMAT CmSurfFormat =
1265 if (CmSurfFormat == cm_support::CM_SURFACE_FORMAT_UNKNOWN) {
1266 return PI_ERROR_IMAGE_FORMAT_NOT_SUPPORTED;
1269 char *MapBasePtr =
nullptr;
1272 int Status = cm_support::CM_SUCCESS;
1277 static_cast<unsigned int>(ImageDesc->
image_width),
1278 static_cast<unsigned int>(ImageDesc->
image_height), CmSurfFormat,
1284 static_cast<unsigned int>(ImageDesc->
image_width),
1285 static_cast<unsigned int>(ImageDesc->
image_height), CmSurfFormat,
1291 reinterpret_cast<const unsigned char *
>(HostPtr),
nullptr,
1292 static_cast<unsigned int>(ImageDesc->
image_width *
1297 if (Status != cm_support::CM_SUCCESS) {
1298 return PI_ERROR_INVALID_OPERATION;
1302 pi_cast<char *>(cm_support::get_surface_base_addr(CmIndex->get_data()));
1305 *RetImage =
new _pi_image(Context, MapBasePtr, CmImg, CmIndex->get_data(),
1308 }
catch (
const std::bad_alloc &) {
1309 return PI_ERROR_OUT_OF_HOST_MEMORY;
1311 return PI_ERROR_UNKNOWN;
1317 PiTrace(
"Failure from CM-managed image creation");
1318 return PI_ERROR_INVALID_VALUE;
1321 (*PiESimdSurfaceMap)[(*RetImage)->SurfaceIndex] = *RetImage;
1346 const size_t *,
const unsigned char **,
1353 const size_t *,
const unsigned char **,
1387 size_t,
void *,
size_t *) {
1426 size_t,
void *,
size_t *) {
1432 const void *,
size_t,
void *,
size_t *) {
1443 size_t ParamValueSize,
void *ParamValue,
1444 size_t *ParamValueSizeRet) {
1449 auto CheckAndFillStatus = [&](
const cm_support::CM_STATUS &State) {
1451 if (State == cm_support::CM_STATUS_FINISHED)
1454 if (ParamValueSize <
sizeof(Result))
1455 return PI_ERROR_INVALID_VALUE;
1456 *
static_cast<pi_int32 *
>(ParamValue) = Result;
1458 if (ParamValueSizeRet) {
1459 *ParamValueSizeRet =
sizeof(Result);
1465 return CheckAndFillStatus(cm_support::CM_STATUS_FINISHED);
1468 return PI_ERROR_INVALID_EVENT;
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;
1475 return CheckAndFillStatus(Status);
1479 size_t ParamValueSize,
void *ParamValue,
1480 size_t *ParamValueSizeRet) {
1487 PiTrace(
"Warning : Profiling Not supported under PI_ESIMD_EMULATOR");
1492 for (
int i = 0; i < (int)NumEvents; i++) {
1493 if (EventList[i]->IsDummyEvent) {
1498 if (EventList[i]->CmEventPtr ==
nullptr) {
1499 return PI_ERROR_INVALID_EVENT;
1501 int Result = EventList[i]->
CmEventPtr->WaitForTaskFinished();
1502 if (Result != cm_support::CM_SUCCESS) {
1503 return PI_ERROR_OUT_OF_RESOURCES;
1517 if (Event ==
nullptr) {
1518 return PI_ERROR_INVALID_EVENT;
1521 ++(Event->RefCount);
1527 if (Event ==
nullptr || (Event->RefCount <= 0)) {
1528 return PI_ERROR_INVALID_EVENT;
1531 if (--(Event->RefCount) == 0) {
1534 return PI_ERROR_INVALID_EVENT;
1537 if (Result != cm_support::CM_SUCCESS) {
1538 return PI_ERROR_INVALID_EVENT;
1580 pi_bool BlockingRead,
size_t Offset,
1581 size_t Size,
void *Dst,
1591 "ESIMD_EMULATOR support for blocking piEnqueueMemBufferRead is NYI");
1592 return PI_ERROR_INVALID_OPERATION;
1596 PiTrace(
"ESIMD_EMULATOR does not support buffer reading with offsets");
1597 return PI_ERROR_INVALID_ARG_VALUE;
1600 if (NumEventsInWaitList != 0) {
1601 return PI_ERROR_INVALID_EVENT_WAIT_LIST;
1606 std::unique_ptr<_pi_event> RetEv{
nullptr};
1608 RetEv = std::unique_ptr<_pi_event>(
new _pi_event());
1609 RetEv->IsDummyEvent =
true;
1618 return PI_ERROR_INVALID_MEM_OBJECT;
1621 reinterpret_cast<unsigned char *
>(Dst),
1623 static_cast<uint64_t
>(Size));
1625 if (Status != cm_support::CM_SUCCESS) {
1626 return PI_ERROR_INVALID_MEM_OBJECT;
1654 size_t,
size_t,
const void *,
pi_uint32,
1681 size_t Offset,
size_t Size,
1691 std::unique_ptr<_pi_event> RetEv{
nullptr};
1695 RetEv = std::unique_ptr<_pi_event>(
new _pi_event());
1696 RetEv->IsDummyEvent =
true;
1707 auto Res = MemObj->
Mappings.insert({*RetMap, {Offset, Size}});
1711 ret = PI_ERROR_INVALID_VALUE;
1712 PiTrace(
"piEnqueueMemBufferMap: duplicate mapping detected");
1729 std::unique_ptr<_pi_event> RetEv{
nullptr};
1733 RetEv = std::unique_ptr<_pi_event>(
new _pi_event());
1734 RetEv->IsDummyEvent =
true;
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");
1765 size_t SlicePitch,
void *Ptr,
1775 PiTrace(
"ESIMD_EMULATOR support for blocking piEnqueueMemImageRead is NYI");
1776 return PI_ERROR_INVALID_OPERATION;
1781 if (SlicePitch != 0) {
1782 PiTrace(
"ESIMD_EMULATOR does not support 3D-image");
1783 return PI_ERROR_INVALID_ARG_VALUE;
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;
1794 std::unique_ptr<_pi_event> RetEv{
nullptr};
1797 RetEv = std::unique_ptr<_pi_event>(
new _pi_event());
1798 RetEv->IsDummyEvent =
true;
1801 size_t Size = RowPitch * (Region->
height);
1808 return PI_ERROR_INVALID_MEM_OBJECT;
1811 reinterpret_cast<unsigned char *
>(Ptr),
1813 static_cast<uint64_t
>(Size));
1815 if (Status != cm_support::CM_SUCCESS) {
1816 return PI_ERROR_INVALID_MEM_OBJECT;
1852 const size_t *GlobalWorkOffset,
1853 const size_t *GlobalWorkSize,
const size_t *LocalWorkSize,
1860 const size_t LocalWorkSz[] = {1, 1, 1};
1862 if (Kernel ==
nullptr) {
1863 return PI_ERROR_INVALID_KERNEL;
1866 if (WorkDim > 3 || WorkDim == 0) {
1867 return PI_ERROR_INVALID_WORK_GROUP_SIZE;
1870 if (
isNull(WorkDim, LocalWorkSize)) {
1871 LocalWorkSize = LocalWorkSz;
1874 for (
pi_uint32 I = 0; I < WorkDim; I++) {
1875 if ((GlobalWorkSize[I] % LocalWorkSize[I]) != 0) {
1876 return PI_ERROR_INVALID_WORK_GROUP_SIZE;
1880 std::unique_ptr<_pi_event> RetEv{
nullptr};
1883 RetEv = std::unique_ptr<_pi_event>(
new _pi_event());
1884 RetEv->IsDummyEvent =
true;
1949 if (Context ==
nullptr || (Device != Context->
Device)) {
1950 return PI_ERROR_INVALID_CONTEXT;
1953 if (ResultPtr ==
nullptr) {
1954 return PI_ERROR_INVALID_OPERATION;
1959 if ((Size & (Size - 1)) != 0) {
1963 cm_support::CmBufferSVM *Buf =
nullptr;
1964 void *SystemMemPtr =
nullptr;
1966 Size, SystemMemPtr, CM_SVM_ACCESS_FLAG_DEFAULT, Buf);
1968 if (Result != cm_support::CM_SUCCESS) {
1969 return PI_ERROR_OUT_OF_HOST_MEMORY;
1971 *ResultPtr = SystemMemPtr;
1975 return PI_ERROR_INVALID_MEM_OBJECT;
1982 if (Context ==
nullptr) {
1983 return PI_ERROR_INVALID_CONTEXT;
1985 if (Ptr ==
nullptr) {
1986 return PI_ERROR_INVALID_OPERATION;
1991 if (Buf ==
nullptr) {
1992 return PI_ERROR_INVALID_MEM_OBJECT;
1996 return PI_ERROR_INVALID_MEM_OBJECT;
1999 if (cm_support::CM_SUCCESS != Result) {
2000 return PI_ERROR_UNKNOWN;
2036 const void *,
size_t,
size_t,
size_t,
2042 size_t,
void *,
size_t *) {
2073 if (RawImgSize != 1) {
2074 PiTrace(
"Only single device binary image is supported in ESIMD_EMULATOR");
2075 return PI_ERROR_INVALID_VALUE;
2088 const char *,
pi_bool,
size_t,
2095 const char *,
pi_bool,
size_t,
2111 delete reinterpret_cast<sycl::detail::ESIMDEmuPluginOpaqueData *
>(
2116 auto Mem = it->second;
2117 if (Mem !=
nullptr) {
2127 "Warning : Querying device clock not supported under PI_ESIMD_EMULATOR");
2133 if (PluginInit ==
nullptr) {
2134 return PI_ERROR_INVALID_VALUE;
2140 size_t PluginVersionSize =
sizeof(PluginInit->
PluginVersion);
2142 return PI_ERROR_INVALID_VALUE;
2151 reinterpret_cast<void *
>(
new sycl::detail::ESIMDDeviceInterface());
2156 #define _PI_API(api) \
2157 (PluginInit->PiFunctionTable).api = (decltype(&::api))(&api);
2158 #include <sycl/detail/pi.def>
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
__ESIMD_API void fence()
esimd::fence sets the memory read/write order.
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.
static void setErrorMessage(const char *message, pi_result error_code)
bool PiPlatformCachePopulated
@ PI_DEVICE_INFO_HOST_UNIFIED_MEMORY
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.
@ PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_CHAR
@ PI_DEVICE_INFO_OPENCL_C_VERSION
PI Mem mapping to CUDA memory allocations, both data and texture/surface.
const char SupportedVersion[]
pi_result piPluginInit(pi_plugin *PluginInit)
pi_result piextPluginGetOpaqueData(void *, void **OpaqueDataReturn)
API to get Plugin internal data, opaque to SYCL RT.
@ PI_DEVICE_INFO_IMAGE3D_MAX_WIDTH
@ PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_CHAR
cm_support::CmDevice * CmDevicePtr
pi_result piProgramGetInfo(pi_program, pi_program_info, size_t, void *, size_t *)
@ PI_DEVICE_INFO_DOUBLE_FP_CONFIG
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...
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)
pi_result piextProgramCreateWithNativeHandle(pi_native_handle, pi_context, bool, pi_program *)
Creates PI program object from a native handle.
@ PI_DEVICE_INFO_DRIVER_VERSION
pi_result piProgramCreate(pi_context, const void *, size_t, pi_program *)
pi_result piextUSMFree(pi_context Context, void *Ptr)
Indicates that the allocated USM memory is no longer needed on the runtime side.
void InvokeKernel(KernelInvocationContext< NDims > *ctx)
@ PI_DEVICE_INFO_GPU_EU_COUNT
@ PI_IMAGE_CHANNEL_TYPE_UNORM_SHORT_565
sycl::detail::Builder IDBuilder
@ PI_IMAGE_CHANNEL_TYPE_HALF_FLOAT
pi_result piextEnqueueReadHostPipe(pi_queue, pi_program, const char *, pi_bool, void *, size_t, pi_uint32, const pi_event *, pi_event *)
Host Pipes.
@ PI_DEVICE_INFO_PREFERRED_INTEROP_USER_SYNC
thread_local char ErrorMessage[MaxMessageSize]
pi_result piextKernelSetArgSampler(pi_kernel, pi_uint32, const pi_sampler *)
void sycl_get_cm_image_params(unsigned int IndexInput, char **BaseAddr, uint32_t *Width, uint32_t *Height, uint32_t *Bpp, std::mutex **ImgMtxLock)
pi_result piQueueFinish(pi_queue)
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.
@ PI_DEVICE_INFO_IL_VERSION
std::unordered_map< void *, Mapping > Mappings
constexpr pi_queue_properties PI_QUEUE_FLAG_ON_DEVICE
@ PI_IMAGE_CHANNEL_TYPE_UNORM_INT8
pi_result piclProgramCreateWithBinary(pi_context, pi_uint32, const pi_device *, const size_t *, const unsigned char **, pi_int32 *, pi_program *)
pi_result piextDeviceGetNativeHandle(pi_device, pi_native_handle *)
Gets the native handle of a PI device object.
constexpr pi_mem_flags PI_MEM_FLAGS_HOST_PTR_COPY
pi_result piextUSMEnqueueMemcpy(pi_queue, pi_bool, void *, const void *, size_t, pi_uint32, const pi_event *, pi_event *)
USM Memcpy API.
pi_result piKernelSetArg(pi_kernel, pi_uint32, size_t, const void *)
@ PI_DEVICE_INFO_PARTITION_AFFINITY_DOMAIN
@ PI_DEVICE_INFO_IMAGE_SRGB
pi_result piEventCreate(pi_context, pi_event *)
Create PI event object in a signalled/completed state.
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.
@ PI_DEVICE_INFO_GLOBAL_MEM_CACHE_TYPE
__ESIMD_API void split_barrier()
Generic work-group split barrier.
pi_result piSamplerRelease(pi_sampler)
pi_result piSamplerGetInfo(pi_sampler, pi_sampler_info, size_t, void *, size_t *)
pi_result piDeviceRelease(pi_device Device)
static sycl::range< 1 > create(Gen G)
libCMBatch(const KernelFunc< DIMS > &Kernel)
@ PI_DEVICE_INFO_MAX_WORK_ITEM_SIZES
pi_result piEnqueueEventsWaitWithBarrier(pi_queue, pi_uint32, const pi_event *, pi_event *)
@ PI_DEVICE_INFO_MAX_COMPUTE_UNITS
pi_result piEventSetCallback(pi_event, pi_int32, void(*)(pi_event, pi_int32, void *), void *)
pi_result piEnqueueMemUnmap(pi_queue Queue, pi_mem MemObj, void *MappedPtr, pi_uint32 NumEventsInWaitList, const pi_event *EventWaitList, pi_event *Event)
void memcpy(void *Dst, const void *Src, size_t Size)
pi_result piextProgramGetNativeHandle(pi_program, pi_native_handle *)
Gets the native handle of a PI program object.
thread_local pi_result ErrorMessageCode
intptr_t pi_context_properties
@ PI_IMAGE_CHANNEL_TYPE_SIGNED_INT8
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.
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...
@ PI_DEVICE_INFO_GPU_SLICES
@ PI_DEVICE_INFO_MAX_MEM_ALLOC_SIZE
pi_result piextPlatformCreateWithNativeHandle(pi_native_handle, pi_platform *)
Creates PI platform object from a native handle.
cm_support::CmEvent * CmEventPtr
pi_result piDevicePartition(pi_device, const pi_device_partition_property *, pi_uint32, pi_device *, pi_uint32 *)
@ PI_DEVICE_INFO_REFERENCE_COUNT
@ PI_DEVICE_INFO_USM_CROSS_SHARED_SUPPORT
@ PI_DEVICE_INFO_PLATFORM
@ PI_DEVICE_INFO_GPU_EU_COUNT_PER_SUBSLICE
@ PI_DEVICE_LOCAL_MEM_TYPE_LOCAL
@ PI_DEVICE_INFO_DEVICE_ID
@ PI_IMAGE_CHANNEL_TYPE_SNORM_INT16
@ PI_DEVICE_INFO_ADDRESS_BITS
pi_result piPluginGetLastError(char **message)
API to get Plugin specific warning and error messages.
#define CASE_PI_UNSUPPORTED(not_supported)
@ PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_DOUBLE
@ PI_DEVICE_INFO_MAX_SAMPLERS
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.
@ PI_DEVICE_INFO_MAX_NUM_SUB_GROUPS
@ PI_DEVICE_INFO_MEM_BASE_ADDR_ALIGN
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.
@ PI_DEVICE_INFO_BUILT_IN_KERNELS
@ PI_DEVICE_INFO_USM_DEVICE_SUPPORT
static constexpr pi_device_fp_config PI_FP_ROUND_TO_NEAREST
@ PI_DEVICE_INFO_IMAGE_MAX_ARRAY_SIZE
@ PI_DEVICE_INFO_MAX_CONSTANT_BUFFER_SIZE
void runIterationSpace(const sycl::range< DIMS > &LocalSize, const sycl::range< DIMS > &GlobalSize, const sycl::id< DIMS > &GlobalOffset)
pi_result piQueueRetain(pi_queue Queue)
pi_result piextMemCreateWithNativeHandle(pi_native_handle, pi_context, bool, pi_mem *)
Creates PI mem object from a native handle.
static bool isNull(int NDims, const size_t *R)
Implementation for Host Kernel Launch used by piEnqueueKernelLaunch.
@ PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_INT
Implementation of a PI Kernel for CUDA.
@ PI_DEVICE_INFO_ERROR_CORRECTION_SUPPORT
pi_result piextKernelCreateWithNativeHandle(pi_native_handle, pi_context, pi_program, bool, pi_kernel *)
Creates PI kernel object from a native handle.
@ PI_DEVICE_INFO_IMAGE2D_MAX_HEIGHT
pi_result piEventsWait(pi_uint32 NumEvents, const pi_event *EventList)
@ PI_DEVICE_INFO_GLOBAL_MEM_CACHE_SIZE
constexpr pi_device_exec_capabilities PI_DEVICE_EXEC_CAPABILITIES_KERNEL
pi_result piextDeviceCreateWithNativeHandle(pi_native_handle, pi_platform, pi_device *)
Creates PI device object from a native handle.
pi_result piMemBufferPartition(pi_mem, pi_mem_flags, pi_buffer_create_type, void *, pi_mem *)
cm_support::CmBuffer * RegularBufPtr
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.
pi_result piEnqueueNativeKernel(pi_queue, void(*)(void *), void *, size_t, pi_uint32, const pi_mem *, const void **, pi_uint32, const pi_event *, pi_event *)
@ PI_EXT_DEVICE_INFO_ATOMIC_MEMORY_SCOPE_CAPABILITIES
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)
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.
pi_result piextUSMEnqueueMemAdvise(pi_queue, const void *, size_t, pi_mem_advice, pi_event *)
USM Memadvise API.
__ESIMD_API void barrier()
Generic work-group barrier.
@ PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_HALF
pi_result piDevicesGet(pi_platform Platform, pi_device_type DeviceType, pi_uint32 NumEntries, pi_device *Devices, pi_uint32 *NumDevices)
cm_support::CmSurface2D * RegularImgPtr
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 *)
@ PI_DEVICE_INFO_IMAGE_MAX_BUFFER_SIZE
pi_result piEnqueueMemBufferWrite(pi_queue, pi_mem, pi_bool, size_t, size_t, const void *, pi_uint32, const pi_event *, pi_event *)
@ PI_DEVICE_INFO_LOCAL_MEM_SIZE
pi_result piKernelSetExecInfo(pi_kernel, pi_kernel_exec_info, size_t, const void *)
API to set attributes controlling kernel execution.
pi_result piSamplerCreate(pi_context, const pi_sampler_properties *, pi_sampler *)
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)
PI queue mapping on to CUstream objects.
std::mutex Addr2CmBufferSVMLock
static std::mutex * PiPlatformCacheLock
@ PI_IMAGE_CHANNEL_TYPE_SIGNED_INT16
pi_result piextKernelSetArgMemObj(pi_kernel, pi_uint32, const pi_mem *)
@ PI_IMAGE_CHANNEL_TYPE_SIGNED_INT32
@ PI_DEVICE_TYPE_GPU
A PI device that is a GPU.
pi_result piProgramCompile(pi_program, pi_uint32, const pi_device *, const char *, pi_uint32, const pi_program *, const char **, void(*)(pi_program, void *), void *)
This struct is a record of the device binary information.
unsigned int SurfaceIndex
Surface index type.
pi_result piextMemGetNativeHandle(pi_mem, pi_native_handle *)
Gets the native handle of a PI mem object.
pi_result piextContextGetNativeHandle(pi_context, pi_native_handle *)
Gets the native handle of a PI context object.
@ PI_DEVICE_INFO_LINKER_AVAILABLE
@ PI_DEVICE_INFO_EXECUTION_CAPABILITIES
pi_result piPlatformGetInfo(pi_platform Platform, pi_platform_info ParamName, size_t ParamValueSize, void *ParamValue, size_t *ParamValueSizeRet)
static std::unordered_map< unsigned int, _pi_mem * > * PiESimdSurfaceMap
@ PI_DEVICE_INFO_PROFILING_TIMER_RESOLUTION
cm_surface_ptr_t SurfacePtr
pi_result piContextGetInfo(pi_context, pi_context_info, size_t, void *, size_t *)
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.
@ PI_DEVICE_INFO_HALF_FP_CONFIG
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 *)
static constexpr pi_device_fp_config PI_FP_DENORM
@ PI_IMAGE_CHANNEL_TYPE_FLOAT
pi_result piextProgramSetSpecializationConstant(pi_program, pi_uint32, size_t, const void *)
Sets a specialization constant to a specific value.
@ PI_USM_CONCURRENT_ACCESS
void(* pi_context_extended_deleter)(void *user_data)
@ PI_MEM_TYPE_IMAGE1D_ARRAY
pi_result piextGetDeviceFunctionPointer(pi_device, pi_program, const char *, pi_uint64 *)
Retrieves a device function pointer to a user-defined function.
pi_result piMemGetInfo(pi_mem, pi_mem_info, size_t, void *, size_t *)
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.
pi_result piPluginGetBackendOption(pi_platform, const char *frontend_option, const char **backend_option)
API to get backend specific option.
static void invoke(pi_kernel Kernel, const size_t *GlobalWorkOffset, const size_t *GlobalWorkSize, const size_t *LocalWorkSize)
@ PI_DEVICE_INFO_USM_SYSTEM_SHARED_SUPPORT
@ PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_SHORT
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...
#define ESIMDEmuPluginDataVersion
pi_result piKernelGetInfo(pi_kernel, pi_kernel_info, size_t, void *, size_t *)
@ PI_DEVICE_INFO_GLOBAL_MEM_CACHELINE_SIZE
@ PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_LONG
pi_result piextKernelSetArgPointer(pi_kernel, pi_uint32, size_t, const void *)
Sets up pointer arguments for CL kernels.
pi_result piextEventGetNativeHandle(pi_event, pi_native_handle *)
Gets the native handle of a PI event object.
pi_result piextContextSetExtendedDeleter(pi_context, pi_context_extended_deleter, void *)
@ PI_DEVICE_INFO_USM_HOST_SUPPORT
constexpr pi_mem_flags PI_MEM_FLAGS_HOST_PTR_USE
@ PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_FLOAT
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 *)
pi_result piTearDown(void *)
API to notify that the plugin should clean up its resources.
#define ESIMDEmuPluginInterfaceVersion
pi_result piEnqueueEventsWait(pi_queue, pi_uint32, const pi_event *, pi_event *)
static sycl::detail::ESIMDEmuPluginOpaqueData * PiESimdDeviceAccess
@ PI_DEVICE_INFO_PARENT_DEVICE
@ PI_PLATFORM_INFO_VERSION
@ PI_IMAGE_CHANNEL_TYPE_UNSIGNED_INT16
static constexpr pi_device_fp_config PI_FP_ROUND_TO_INF
@ PI_DEVICE_INFO_PARTITION_TYPE
static constexpr SurfaceIndex SLM_BTI
pi_result piMemBufferCreate(pi_context Context, pi_mem_flags Flags, size_t Size, void *HostPtr, pi_mem *RetMem, const pi_mem_properties *properties)
pi_bitfield pi_mem_properties
pi_bitfield pi_device_exec_capabilities
Implementation of PI Program on CUDA Module object.
Implementation of samplers for CUDA.
@ PI_DEVICE_INFO_PRINTF_BUFFER_SIZE
@ PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_DOUBLE
#define _PI_ESIMD_PLUGIN_VERSION_STRING
@ PI_PLATFORM_INFO_PROFILE
ur_result_t getInfo< const char * >(size_t param_value_size, void *param_value, size_t *param_value_size_ret, const char *value)
@ PI_DEVICE_INFO_MAX_CONSTANT_ARGS
@ PI_DEVICE_INFO_PCI_ADDRESS
uintptr_t pi_native_handle
@ PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_SHORT
@ PI_DEVICE_INFO_GLOBAL_MEM_SIZE
static sycl::range< NDims > get_range(const size_t *Array)
pi_bitfield pi_sampler_properties
pi_result piKernelRetain(pi_kernel)
@ PI_DEVICE_MEM_CACHE_TYPE_READ_WRITE_CACHE
constexpr pi_mem_flags PI_MEM_FLAGS_ACCESS_RW
pi_result piclProgramCreateWithSource(pi_context, pi_uint32, const char **, const size_t *, pi_program *)
@ PI_IMAGE_CHANNEL_TYPE_UNORM_INT_101010
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)
@ PI_PLATFORM_INFO_VENDOR
pi_result piKernelRelease(pi_kernel)
@ PI_DEVICE_INFO_USM_SINGLE_SHARED_SUPPORT
void die(const char *Message)
@ PI_IMAGE_CHANNEL_ORDER_RGBA
cm_support::CmBufferUP * UPBufPtr
@ PI_DEVICE_INFO_QUEUE_PROPERTIES
@ PI_DEVICE_INFO_PARTITION_MAX_SUB_DEVICES
pi_result piEventGetInfo(pi_event Event, pi_event_info ParamName, size_t ParamValueSize, void *ParamValue, size_t *ParamValueSizeRet)
constexpr size_t getNextPowerOfTwo(size_t Var)
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)
@ PI_EXT_ONEAPI_DEVICE_INFO_MAX_WORK_GROUPS_3D
pi_result piextPlatformGetNativeHandle(pi_platform, pi_native_handle *)
Gets the native handle of a PI platform object.
std::function< void(const sycl::nd_item< NDims > &)> KernelFunc
@ PI_MEM_TYPE_IMAGE1D_BUFFER
pi_result piProgramBuild(pi_program, pi_uint32, const pi_device *, const char *, void(*)(pi_program, void *), void *)
const sycl::range< NDims > & GlobalSize
pi_result piextUSMEnqueueMemset(pi_queue, void *, pi_int32, size_t, pi_uint32, const pi_event *, pi_event *)
USM Memset API.
@ PI_DEVICE_INFO_SINGLE_FP_CONFIG
#define _PI_PLUGIN_VERSION_CHECK(PI_API_VERSION, PI_PLUGIN_VERSION)
pi_result piextUSMDeviceAlloc(void **, pi_context, pi_device, pi_usm_mem_properties *, size_t, pi_uint32)
Allocates device memory.
@ PI_DEVICE_INFO_BUILD_ON_SUBDEVICE
std::unordered_map< void *, cm_support::CmBufferSVM * > Addr2CmBufferSVM
@ PI_DEVICE_INFO_MAX_WORK_ITEM_DIMENSIONS
@ PI_DEVICE_INFO_LOCAL_MEM_TYPE
@ PI_DEVICE_INFO_IMAGE3D_MAX_DEPTH
@ PI_DEVICE_INFO_COMPILER_AVAILABLE
#define _PI_H_VERSION_STRING
static std::mutex * PiESimdSurfaceMapLock
static constexpr pi_device_fp_config PI_FP_FMA
@ PI_EXT_PLATFORM_BACKEND_ESIMD
The backend is ESIMD.
@ PI_DEVICE_INFO_MAX_MEM_BANDWIDTH
@ PI_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8
@ PI_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32
pi_result piEnqueueMemBufferFill(pi_queue, pi_mem, const void *, size_t, size_t, size_t, pi_uint32, const pi_event *, pi_event *)
pi_bitfield pi_queue_properties
@ PI_EXT_ONEAPI_DEVICE_INFO_MAX_WORK_GROUPS_2D
@ PI_DEVICE_INFO_ATOMIC_64
@ PI_DEVICE_INFO_MAX_WRITE_IMAGE_ARGS
constexpr size_t MaxMessageSize
pi_result piEnqueueMemBufferCopy(pi_queue, pi_mem, pi_mem, size_t, size_t, size_t, pi_uint32, const pi_event *, pi_event *)
pi_result piProgramGetBuildInfo(pi_program, pi_device, pi_program_build_info, size_t, void *, size_t *)
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)
@ PI_IMAGE_CHANNEL_TYPE_UNORM_SHORT_555
@ PI_IMAGE_CHANNEL_ORDER_ARGB
@ PI_DEVICE_INFO_ENDIAN_LITTLE
pi_result piQueueCreate(pi_context Context, pi_device Device, pi_queue_properties Properties, pi_queue *Queue)
cm_support::CmQueue * CmQueuePtr
pi_result piContextRelease(pi_context Context)
@ PI_EXT_PLATFORM_INFO_BACKEND
pi_result piextQueueGetNativeHandle(pi_queue, pi_native_handle *, int32_t *)
Gets the native handle of a PI queue object.
pi_result piQueueFlush(pi_queue)
PI Event mapping to CUevent.
pi_result piQueueRelease(pi_queue Queue)
#define DIE_NO_IMPLEMENTATION
@ PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_INT
@ PI_DEVICE_INFO_MAX_READ_IMAGE_ARGS
static constexpr pi_device_fp_config PI_FP_INF_NAN
@ PI_DEVICE_INFO_GPU_SUBSLICES_PER_SLICE
pi_result piKernelCreate(pi_program, const char *, pi_kernel *)
pi_result piEventRelease(pi_event Event)
static char ESimdEmuVersionString[32]
@ PI_MEM_TYPE_IMAGE2D_ARRAY
pi_result piMemRetain(pi_mem Mem)
@ PI_EXT_INTEL_DEVICE_INFO_MAX_COMPUTE_QUEUE_INDICES
pi_result piKernelGetGroupInfo(pi_kernel, pi_device, pi_kernel_group_info, size_t, void *, size_t *)
pi_result piMemRelease(pi_mem Mem)
void assertion(bool Condition, const char *Message=nullptr)
static pi_platform PiPlatformCache
@ PI_EXT_ONEAPI_DEVICE_INFO_MAX_WORK_GROUPS_1D
@ PI_DEVICE_INFO_VENDOR_ID
pi_result piextEventCreateWithNativeHandle(pi_native_handle, pi_context, bool, pi_event *)
Creates PI event object from a native handle.
pi_result piDeviceRetain(pi_device Device)
@ PI_DEVICE_INFO_EXTENSIONS
pi_result piextUSMHostAlloc(void **, pi_context, pi_usm_mem_properties *, size_t, pi_uint32)
Allocates host memory accessible by the device.
pi_result piProgramRetain(pi_program)
ur_result_t getInfo(size_t param_value_size, void *param_value, size_t *param_value_size_ret, T value)
const sycl::range< NDims > & LocalSize
@ PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_FLOAT
cm_support::CmSurface2DUP * UPImgPtr
unsigned int SurfaceIndex
const sycl::id< NDims > & GlobalOffset
pi_result piextQueueCreate(pi_context Context, pi_device Device, pi_queue_properties *Properties, pi_queue *Queue)
constexpr pi_queue_properties PI_QUEUE_FLAG_OUT_OF_ORDER_EXEC_MODE_ENABLE
pi_result piextContextCreateWithNativeHandle(pi_native_handle, pi_uint32, const pi_device *, bool, pi_context *)
Creates PI context object from a native handle.
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)
static constexpr pi_device_fp_config PI_FP_CORRECTLY_ROUNDED_DIVIDE_SQRT
unsigned int sycl_get_cm_surface_index(void *PtrInput)
pi_result piMemImageGetInfo(pi_mem, pi_image_info, size_t, void *, size_t *)
@ PI_DEVICE_INFO_MAX_WORK_GROUP_SIZE
pi_result piPlatformsGet(pi_uint32 NumEntries, pi_platform *Platforms, pi_uint32 *NumPlatforms)
static sycl::range< 3 > create(Gen G)
pi_bitfield pi_device_affinity_domain
#define CONTINUE_NO_IMPLEMENTATION
@ PI_DEVICE_INFO_IMAGE_SUPPORT
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)
@ PI_EXT_DEVICE_INFO_ATOMIC_FENCE_ORDER_CAPABILITIES
@ PI_EXT_INTEL_DEVICE_INFO_MEM_CHANNEL_SUPPORT
@ PI_DEVICE_INFO_MAX_PARAMETER_SIZE
@ PI_EVENT_INFO_COMMAND_EXECUTION_STATUS
@ PI_DEVICE_INFO_SUB_GROUP_SIZES_INTEL
@ PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_HALF
@ PI_EXT_DEVICE_INFO_ATOMIC_MEMORY_ORDER_CAPABILITIES
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 *)
pi_result piProgramRelease(pi_program)
__SYCL_EXTERN_STREAM_ATTRS ostream cout
Linked to standard output.
pi_result piContextRetain(pi_context Context)
@ PI_DEVICE_INFO_MAX_CLOCK_FREQUENCY
pi_result piEventSetStatus(pi_event, pi_int32)
@ PI_DEVICE_INFO_IMAGE2D_MAX_WIDTH
@ PI_DEVICE_INFO_AVAILABLE
void sycl_get_cm_buffer_params(unsigned int IndexInput, char **BaseAddr, uint32_t *Width, std::mutex **BufMtxLock)
pi_bitfield pi_usm_mem_properties
pi_result piEventGetProfilingInfo(pi_event Event, pi_profiling_info ParamName, size_t ParamValueSize, void *ParamValue, size_t *ParamValueSizeRet)
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.
@ PI_DEVICE_INFO_IMAGE3D_MAX_HEIGHT
@ PI_DEVICE_INFO_GPU_EU_SIMD_WIDTH
cm_support::CmQueue * OwnerQueue
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 *)
pi_result piQueueGetInfo(pi_queue, pi_queue_info, size_t, void *, size_t *)
pi_result piSamplerRetain(pi_sampler)
@ PI_USM_CONCURRENT_ATOMIC_ACCESS
static sycl::range< 2 > create(Gen G)
pi_result piProgramLink(pi_context, pi_uint32, const pi_device *, const char *, pi_uint32, const pi_program *, void(*)(pi_program, void *), void *, pi_program *)
@ PI_IMAGE_CHANNEL_TYPE_SNORM_INT8
static void PiTrace(std::string TraceString)
intptr_t pi_device_partition_property
static constexpr pi_device_fp_config PI_FP_ROUND_TO_ZERO
@ PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_LONG
cm_support::CM_SURFACE_FORMAT ConvertPiImageFormatToCmFormat(const pi_image_format *PiFormat)
pi_result piextKernelGetNativeHandle(pi_kernel, pi_native_handle *)
Gets the native handle of a PI kernel object.
pi_result piEnqueueMemImageFill(pi_queue, pi_mem, const void *, const size_t *, const size_t *, pi_uint32, const pi_event *, pi_event *)
PI context mapping to a CUDA context object.
pi_result piextUSMEnqueueMemset2D(pi_queue, void *, size_t, int, size_t, size_t, pi_uint32, const pi_event *, pi_event *)
USM 2D Memset API.
pi_result piEventRetain(pi_event Event)
@ PI_IMAGE_CHANNEL_TYPE_UNORM_INT16
constexpr pi_queue_properties PI_QUEUE_FLAGS
PI device mapping to a CUdevice.
@ PI_DEVICE_INFO_PARTITION_PROPERTIES
_pi_kernel_sub_group_info
@ PI_PLATFORM_INFO_EXTENSIONS
bool checkSurfaceArgument(pi_mem_flags Flags, void *HostPtr)
@ PI_EXT_DEVICE_INFO_ATOMIC_FENCE_SCOPE_CAPABILITIES
@ PI_DEVICE_INFO_SUB_GROUP_INDEPENDENT_FORWARD_PROGRESS
@ PI_EXT_ONEAPI_DEVICE_INFO_MAX_GLOBAL_WORK_GROUPS