[SYCL][Bindless][2/4] Add experimental implementation of SYCL bindless images extension

sycl/include/sycl/detail/pi.def

@@ -174,4 +174,28 @@ _PI_API(piextCommandBufferMemBufferRead)
 _PI_API(piextCommandBufferMemBufferReadRect)
 _PI_API(piextEnqueueCommandBuffer)
 
+_PI_API(piextUSMPitchedAlloc)
+
+// Bindless Images
+_PI_API(piextMemUnsampledImageHandleDestroy)
+_PI_API(piextMemSampledImageHandleDestroy)
+_PI_API(piextBindlessImageSamplerCreate)
+_PI_API(piextMemImageAllocate)
+_PI_API(piextMemImageFree)
+_PI_API(piextMemUnsampledImageCreate)
+_PI_API(piextMemSampledImageCreate)
+_PI_API(piextMemImageCopy)
+_PI_API(piextMemImageGetInfo)
+_PI_API(piextMemMipmapGetLevel)
+_PI_API(piextMemMipmapFree)
+
+// Interop
+_PI_API(piextMemImportOpaqueFD)
+_PI_API(piextMemReleaseInterop)
+_PI_API(piextMemMapExternalArray)
+_PI_API(piextImportExternalSemaphoreOpaqueFD)
+_PI_API(piextDestroyExternalSemaphore)
+_PI_API(piextWaitExternalSemaphore)
+_PI_API(piextSignalExternalSemaphore)
+
 #undef _PI_API

sycl/include/sycl/detail/pi.hpp

@@ -156,6 +156,13 @@ using PiExtSyncPoint = ::pi_ext_sync_point;
 using PiExtCommandBuffer = ::pi_ext_command_buffer;
 using PiExtCommandBufferDesc = ::pi_ext_command_buffer_desc;
 using PiPeerAttr = ::pi_peer_attr;
+using PiImageHandle = ::pi_image_handle;
+using PiImageMemHandle = ::pi_image_mem_handle;
+using PiImageCopyFlags = ::pi_image_copy_flags;
+using PiInteropMemHandle = ::pi_interop_mem_handle;
+using PiInteropSemaphoreHandle = ::pi_interop_semaphore_handle;
+using PiImageOffset = ::pi_image_offset_struct;
+using PiImageRegion = ::pi_image_region_struct;
 
 __SYCL_EXPORT void contextSetExtendedDeleter(const sycl::context &constext,
                                              pi_context_extended_deleter func,

sycl/plugins/cuda/CMakeLists.txt

@@ -64,6 +64,8 @@ add_sycl_plugin(cuda
     "../unified_runtime/ur/adapters/cuda/enqueue.cpp"
     "../unified_runtime/ur/adapters/cuda/event.cpp"
     "../unified_runtime/ur/adapters/cuda/event.hpp"
+    "../unified_runtime/ur/adapters/cuda/image.cpp"
+    "../unified_runtime/ur/adapters/cuda/image.hpp"
     "../unified_runtime/ur/adapters/cuda/kernel.cpp"
     "../unified_runtime/ur/adapters/cuda/kernel.hpp"
     "../unified_runtime/ur/adapters/cuda/memory.cpp"

sycl/plugins/cuda/pi_cuda.cpp

@@ -158,6 +158,7 @@ pi_result piPluginInit(pi_plugin *PluginInit) {
   _PI_CL(piextUSMHostAlloc, pi2ur::piextUSMHostAlloc)
   _PI_CL(piextUSMDeviceAlloc, pi2ur::piextUSMDeviceAlloc)
   _PI_CL(piextUSMSharedAlloc, pi2ur::piextUSMSharedAlloc)
+  _PI_CL(piextUSMPitchedAlloc, pi2ur::piextUSMPitchedAlloc)
   _PI_CL(piextUSMFree, pi2ur::piextUSMFree)
   _PI_CL(piextUSMEnqueueMemset, pi2ur::piextUSMEnqueueMemset)
   _PI_CL(piextUSMEnqueueMemcpy, pi2ur::piextUSMEnqueueMemcpy)
@@ -196,10 +197,38 @@ pi_result piPluginInit(pi_plugin *PluginInit) {
   _PI_CL(piextCommandBufferMemBufferCopyRect,
          pi2ur::piextCommandBufferMemBufferCopyRect)
   _PI_CL(piextEnqueueCommandBuffer, pi2ur::piextEnqueueCommandBuffer)
+
   // Peer to Peer
   _PI_CL(piextEnablePeerAccess, pi2ur::piextEnablePeerAccess)
   _PI_CL(piextDisablePeerAccess, pi2ur::piextDisablePeerAccess)
   _PI_CL(piextPeerAccessGetInfo, pi2ur::piextPeerAccessGetInfo)
+
+  // Bindless Images
+  _PI_CL(piextMemUnsampledImageHandleDestroy,
+         pi2ur::piextMemUnsampledImageHandleDestroy)
+  _PI_CL(piextMemSampledImageHandleDestroy,
+         pi2ur::piextMemSampledImageHandleDestroy)
+  _PI_CL(piextMemImageAllocate, pi2ur::piextMemImageAllocate)
+  _PI_CL(piextMemImageFree, pi2ur::piextMemImageFree)
+  _PI_CL(piextMemUnsampledImageCreate, pi2ur::piextMemUnsampledImageCreate)
+  _PI_CL(piextMemSampledImageCreate, pi2ur::piextMemSampledImageCreate)
+  _PI_CL(piextBindlessImageSamplerCreate,
+         pi2ur::piextBindlessImageSamplerCreate)
+  _PI_CL(piextMemImageCopy, pi2ur::piextMemImageCopy)
+  _PI_CL(piextMemImageGetInfo, pi2ur::piextMemImageGetInfo)
+
+  _PI_CL(piextMemMipmapGetLevel, pi2ur::piextMemMipmapGetLevel)
+  _PI_CL(piextMemMipmapFree, pi2ur::piextMemMipmapFree)
+
+  _PI_CL(piextMemImportOpaqueFD, pi2ur::piextMemImportOpaqueFD)
+  _PI_CL(piextMemReleaseInterop, pi2ur::piextMemReleaseInterop)
+  _PI_CL(piextMemMapExternalArray, pi2ur::piextMemMapExternalArray)
+  _PI_CL(piextImportExternalSemaphoreOpaqueFD,
+         pi2ur::piextImportExternalSemaphoreOpaqueFD)
+  _PI_CL(piextDestroyExternalSemaphore, pi2ur::piextDestroyExternalSemaphore)
+  _PI_CL(piextWaitExternalSemaphore, pi2ur::piextWaitExternalSemaphore)
+  _PI_CL(piextSignalExternalSemaphore, pi2ur::piextSignalExternalSemaphore)
+
 #undef _PI_CL
 
   return PI_SUCCESS;

sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp

@@ -1826,6 +1826,12 @@ pi_result piEnqueueMemImageWrite(pi_queue, pi_mem, pi_bool, pi_image_offset,
   DIE_NO_IMPLEMENTATION;
 }
 
+pi_result piextBindlessImageSamplerCreate(pi_context,
+                                          const pi_sampler_properties *, float,
+                                          float, float, pi_sampler *) {
+  DIE_NO_IMPLEMENTATION;
+}
+
 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 *) {
@@ -1968,6 +1974,12 @@ pi_result piextUSMSharedAlloc(void **ResultPtr, pi_context Context,
   return PI_SUCCESS;
 }
 
+pi_result piextUSMPitchedAlloc(void **, size_t *, pi_context, pi_device,
+                               pi_usm_mem_properties *, size_t, size_t,
+                               unsigned int) {
+  DIE_NO_IMPLEMENTATION;
+}
+
 pi_result piextUSMFree(pi_context Context, void *Ptr) {
   if (Context == nullptr) {
     return PI_ERROR_INVALID_CONTEXT;
@@ -2267,6 +2279,96 @@ pi_result piextPeerAccessGetInfo(pi_device command_device,
   return ReturnValue(pi_int32{0});
 }
 
+pi_result piextMemUnsampledImageHandleDestroy(pi_context, pi_device,
+                                              pi_image_handle) {
+  DIE_NO_IMPLEMENTATION;
+}
+
+pi_result piextMemSampledImageHandleDestroy(pi_context, pi_device,
+                                            pi_image_handle) {
+  DIE_NO_IMPLEMENTATION;
+}
+
+pi_result piextMemImageAllocate(pi_context, pi_device, pi_image_format *,
+                                pi_image_desc *, pi_image_mem_handle *) {
+  DIE_NO_IMPLEMENTATION;
+}
+
+pi_result piextMemMipmapGetLevel(pi_context, pi_device, pi_image_mem_handle,
+                                 unsigned int, pi_image_mem_handle *) {
+  DIE_NO_IMPLEMENTATION;
+}
+
+pi_result piextMemImageFree(pi_context, pi_device, pi_image_mem_handle) {
+  DIE_NO_IMPLEMENTATION;
+}
+
+pi_result piextMemMipmapFree(pi_context, pi_device, pi_image_mem_handle) {
+  DIE_NO_IMPLEMENTATION;
+}
+
+pi_result piextMemUnsampledImageCreate(pi_context, pi_device,
+                                       pi_image_mem_handle, pi_image_format *,
+                                       pi_image_desc *, pi_mem *,
+                                       pi_image_handle *) {
+  DIE_NO_IMPLEMENTATION;
+}
+
+pi_result piextMemSampledImageCreate(pi_context, pi_device, pi_image_mem_handle,
+                                     pi_image_format *, pi_image_desc *,
+                                     pi_sampler, pi_mem *, pi_image_handle *) {
+  DIE_NO_IMPLEMENTATION;
+}
+
+pi_result piextMemImageCopy(pi_queue, void *, void *, const pi_image_format *,
+                            const pi_image_desc *, const pi_image_copy_flags,
+                            pi_image_offset, pi_image_offset, pi_image_region,
+                            pi_image_region, pi_uint32, const pi_event *,
+                            pi_event *) {
+  DIE_NO_IMPLEMENTATION;
+}
+
+pi_result piextMemImageGetInfo(const pi_image_mem_handle, pi_image_info, void *,
+                               size_t *) {
+  DIE_NO_IMPLEMENTATION;
+}
+
+pi_result piextMemImportOpaqueFD(pi_context, pi_device, size_t, int,
+                                 pi_interop_mem_handle *) {
+  DIE_NO_IMPLEMENTATION;
+}
+
+pi_result piextMemMapExternalArray(pi_context, pi_device, pi_image_format *,
+                                   pi_image_desc *, pi_interop_mem_handle,
+                                   pi_image_mem_handle *) {
+  DIE_NO_IMPLEMENTATION;
+}
+
+pi_result piextMemReleaseInterop(pi_context, pi_device, pi_interop_mem_handle) {
+  DIE_NO_IMPLEMENTATION;
+}
+
+pi_result piextImportExternalSemaphoreOpaqueFD(pi_context, pi_device, int,
+                                               pi_interop_semaphore_handle *) {
+  DIE_NO_IMPLEMENTATION;
+}
+
+pi_result piextDestroyExternalSemaphore(pi_context, pi_device,
+                                        pi_interop_semaphore_handle) {
+  DIE_NO_IMPLEMENTATION;
+}
+
+pi_result piextWaitExternalSemaphore(pi_queue, pi_interop_semaphore_handle,
+                                     pi_uint32, const pi_event *, pi_event *) {
+  DIE_NO_IMPLEMENTATION;
+}
+
+pi_result piextSignalExternalSemaphore(pi_queue, pi_interop_semaphore_handle,
+                                       pi_uint32, const pi_event *,
+                                       pi_event *) {
+  DIE_NO_IMPLEMENTATION;
+}
+
 #ifdef _WIN32
 #define __SYCL_PLUGIN_DLL_NAME "pi_esimd_emulator.dll"
 #include "../common_win_pi_trace/common_win_pi_trace.hpp"

sycl/plugins/level_zero/CMakeLists.txt

@@ -107,8 +107,9 @@ add_sycl_plugin(level_zero
     "../unified_runtime/ur/adapters/level_zero/context.hpp"
     "../unified_runtime/ur/adapters/level_zero/device.hpp"
     "../unified_runtime/ur/adapters/level_zero/event.hpp"
-    "../unified_runtime/ur/adapters/level_zero/memory.hpp"
+    "../unified_runtime/ur/adapters/level_zero/image.hpp"
     "../unified_runtime/ur/adapters/level_zero/kernel.hpp"
+    "../unified_runtime/ur/adapters/level_zero/memory.hpp"
     "../unified_runtime/ur/adapters/level_zero/platform.hpp"
     "../unified_runtime/ur/adapters/level_zero/program.hpp"
     "../unified_runtime/ur/adapters/level_zero/queue.hpp"
@@ -120,8 +121,9 @@ add_sycl_plugin(level_zero
     "../unified_runtime/ur/adapters/level_zero/context.cpp"
     "../unified_runtime/ur/adapters/level_zero/device.cpp"
     "../unified_runtime/ur/adapters/level_zero/event.cpp"
-    "../unified_runtime/ur/adapters/level_zero/memory.cpp"
+    "../unified_runtime/ur/adapters/level_zero/image.cpp"
     "../unified_runtime/ur/adapters/level_zero/kernel.cpp"
+    "../unified_runtime/ur/adapters/level_zero/memory.cpp"
     "../unified_runtime/ur/adapters/level_zero/platform.cpp"
     "../unified_runtime/ur/adapters/level_zero/program.cpp"
     "../unified_runtime/ur/adapters/level_zero/queue.cpp"