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

llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td

@@ -50,6 +50,17 @@ def AspectExt_intel_memory_clock_rate : Aspect<"ext_intel_memory_clock_rate">;
 def AspectExt_intel_memory_bus_width : Aspect<"ext_intel_memory_bus_width">;
 def AspectEmulated : Aspect<"emulated">;
 def AspectExt_intel_legacy_image : Aspect<"ext_intel_legacy_image">;
+def AspectExt_oneapi_bindless_images : Aspect<"ext_oneapi_bindless_images">;
+def AspectExt_oneapi_bindless_images_shared_usm : Aspect<"ext_oneapi_bindless_images_shared_usm">;
+def AspectExt_oneapi_bindless_images_1d_usm : Aspect<"ext_oneapi_bindless_images_1d_usm">;
+def AspectExt_oneapi_bindless_images_2d_usm : Aspect<"ext_oneapi_bindless_images_2d_usm">;
+def AspectExt_oneapi_interop_memory_import : Aspect<"ext_oneapi_interop_memory_import">;
+def AspectExt_oneapi_interop_memory_export : Aspect<"ext_oneapi_interop_memory_export">;
+def AspectExt_oneapi_interop_semaphore_import : Aspect<"ext_oneapi_interop_semaphore_import">;
+def AspectExt_oneapi_interop_semaphore_export : Aspect<"ext_oneapi_interop_semaphore_export">;
+def AspectExt_oneapi_mipmap : Aspect<"ext_oneapi_mipmap">;
+def AspectExt_oneapi_mipmap_anisotropy : Aspect<"ext_oneapi_mipmap_anisotropy">;
+def AspectExt_oneapi_mipmap_level_reference : Aspect<"ext_oneapi_mipmap_level_reference">;
 // Deprecated aspects
 def AspectInt64_base_atomics : Aspect<"int64_base_atomics">;
 def AspectInt64_extended_atomics : Aspect<"int64_extended_atomics">;
@@ -94,7 +105,11 @@ def : TargetInfo<"__TestAspectList",
     AspectExt_oneapi_native_assert, AspectHost_debuggable, AspectExt_intel_gpu_hw_threads_per_eu,
     AspectExt_oneapi_cuda_async_barrier, AspectExt_oneapi_bfloat16_math_functions, AspectExt_intel_free_memory,
     AspectExt_intel_device_id, AspectExt_intel_memory_clock_rate, AspectExt_intel_memory_bus_width, AspectEmulated,
-    AspectExt_intel_legacy_image],
+    AspectExt_intel_legacy_image, AspectExt_oneapi_bindless_images,
+    AspectExt_oneapi_bindless_images_shared_usm, AspectExt_oneapi_bindless_images_1d_usm, AspectExt_oneapi_bindless_images_2d_usm,
+    AspectExt_oneapi_interop_memory_import, AspectExt_oneapi_interop_memory_export,
+    AspectExt_oneapi_interop_semaphore_import, AspectExt_oneapi_interop_semaphore_export,
+    AspectExt_oneapi_mipmap, AspectExt_oneapi_mipmap_anisotropy, AspectExt_oneapi_mipmap_level_reference],
     []>;
 // This definition serves the only purpose of testing whether the deprecated aspect list defined in here and in SYCL RT
 // match.

sycl/include/sycl/detail/cg.hpp

@@ -73,6 +73,9 @@ class CG {
     CopyFromDeviceGlobal = 20,
     ReadWriteHostPipe = 21,
     ExecCommandBuffer = 22,
+    CopyImage = 23,
+    SemaphoreWait = 24,
+    SemaphoreSignal = 25,
   };
 
   struct StorageInitHelper {
@@ -496,6 +499,77 @@ class CGCopyFromDeviceGlobal : public CG {
   size_t getNumBytes() { return MNumBytes; }
   size_t getOffset() { return MOffset; }
 };
+/// "Copy Image" command group class.
+class CGCopyImage : public CG {
+  void *MSrc;
+  void *MDst;
+  sycl::detail::pi::PiMemImageDesc MImageDesc;
+  sycl::detail::pi::PiMemImageFormat MImageFormat;
+  sycl::detail::pi::PiImageCopyFlags MImageCopyFlags;
+  sycl::detail::pi::PiImageOffset MSrcOffset;
+  sycl::detail::pi::PiImageOffset MDstOffset;
+  sycl::detail::pi::PiImageRegion MHostExtent;
+  sycl::detail::pi::PiImageRegion MCopyExtent;
+
+public:
+  CGCopyImage(void *Src, void *Dst, sycl::detail::pi::PiMemImageDesc ImageDesc,
+              sycl::detail::pi::PiMemImageFormat ImageFormat,
+              sycl::detail::pi::PiImageCopyFlags ImageCopyFlags,
+              sycl::detail::pi::PiImageOffset SrcOffset,
+              sycl::detail::pi::PiImageOffset DstOffset,
+              sycl::detail::pi::PiImageRegion HostExtent,
+              sycl::detail::pi::PiImageRegion CopyExtent,
+              CG::StorageInitHelper CGData, detail::code_location loc = {})
+      : CG(CopyImage, std::move(CGData), std::move(loc)), MSrc(Src), MDst(Dst),
+        MImageDesc(ImageDesc), MImageFormat(ImageFormat),
+        MImageCopyFlags(ImageCopyFlags), MSrcOffset(SrcOffset),
+        MDstOffset(DstOffset), MHostExtent(HostExtent),
+        MCopyExtent(CopyExtent) {}
+
+  void *getSrc() const { return MSrc; }
+  void *getDst() const { return MDst; }
+  sycl::detail::pi::PiMemImageDesc getDesc() const { return MImageDesc; }
+  sycl::detail::pi::PiMemImageFormat getFormat() const { return MImageFormat; }
+  sycl::detail::pi::PiImageCopyFlags getCopyFlags() const {
+    return MImageCopyFlags;
+  }
+  sycl::detail::pi::PiImageOffset getSrcOffset() const { return MSrcOffset; }
+  sycl::detail::pi::PiImageOffset getDstOffset() const { return MDstOffset; }
+  sycl::detail::pi::PiImageRegion getHostExtent() const { return MHostExtent; }
+  sycl::detail::pi::PiImageRegion getCopyExtent() const { return MCopyExtent; }
+};
+
+/// "Semaphore Wait" command group class.
+class CGSemaphoreWait : public CG {
+  sycl::detail::pi::PiInteropSemaphoreHandle MInteropSemaphoreHandle;
+
+public:
+  CGSemaphoreWait(
+      sycl::detail::pi::PiInteropSemaphoreHandle InteropSemaphoreHandle,
+      CG::StorageInitHelper CGData, detail::code_location loc = {})
+      : CG(SemaphoreWait, std::move(CGData), std::move(loc)),
+        MInteropSemaphoreHandle(InteropSemaphoreHandle) {}
+
+  sycl::detail::pi::PiInteropSemaphoreHandle getInteropSemaphoreHandle() const {
+    return MInteropSemaphoreHandle;
+  }
+};
+
+/// "Semaphore Signal" command group class.
+class CGSemaphoreSignal : public CG {
+  sycl::detail::pi::PiInteropSemaphoreHandle MInteropSemaphoreHandle;
+
+public:
+  CGSemaphoreSignal(
+      sycl::detail::pi::PiInteropSemaphoreHandle InteropSemaphoreHandle,
+      CG::StorageInitHelper CGData, detail::code_location loc = {})
+      : CG(SemaphoreSignal, std::move(CGData), std::move(loc)),
+        MInteropSemaphoreHandle(InteropSemaphoreHandle) {}
+
+  sycl::detail::pi::PiInteropSemaphoreHandle getInteropSemaphoreHandle() const {
+    return MInteropSemaphoreHandle;
+  }
+};
 
 /// "Execute command-buffer" command group class.
 class CGExecCommandBuffer : public CG {

sycl/include/sycl/device_aspect_macros.hpp

@@ -218,6 +218,61 @@
 #define __SYCL_ALL_DEVICES_HAVE_ext_intel_legacy_image__ 0
 #endif
 
+#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_bindless_images__
+// __SYCL_ASPECT(ext_oneapi_bindless_images, 42)
+#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_bindless_images__ 0
+#endif
+
+#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_bindless_images_shared_usm__
+//__SYCL_ASPECT(ext_oneapi_bindless_images_shared_usm, 43)
+#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_bindless_images_shared_usm__ 0
+#endif
+
+#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_bindless_images_1d_usm__
+//__SYCL_ASPECT(ext_oneapi_bindless_images_1d_usm, 44)
+#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_bindless_images_1d_usm__ 0
+#endif
+
+#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_bindless_images_2d_usm__
+//__SYCL_ASPECT(ext_oneapi_bindless_images_2d_usm, 45)
+#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_bindless_images_2d_usm__ 0
+#endif
+
+#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_interop_memory_import__
+//__SYCL_ASPECT(ext_oneapi_interop_memory_import, 46)
+#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_interop_memory_import__ 0
+#endif
+
+#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_interop_memory_export__
+//__SYCL_ASPECT(ext_oneapi_interop_memory_export, 47)
+#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_interop_memory_export__ 0
+#endif
+
+#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_interop_semaphore_import__
+//__SYCL_ASPECT(ext_oneapi_interop_semaphore_import, 48)
+#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_interop_semaphore_import__ 0
+#endif
+
+#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_interop_semaphore_export__
+//__SYCL_ASPECT(ext_oneapi_interop_semaphore_export, 49)
+#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_interop_semaphore_export__ 0
+#endif
+
+#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_mipmap__
+//__SYCL_ASPECT(ext_oneapi_mipmap, 50)
+#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_mipmap__ 0
+#endif
+
+#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_mipmap_anisotropy__
+//__SYCL_ASPECT(ext_oneapi_mipmap_anisotropy, 51)
+#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_mipmap_anisotropy__ 0
+#endif
+
+#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_mipmap_level_reference__
+//__SYCL_ASPECT(ext_oneapi_mipmap_level_reference, 52)
+#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_mipmap_level_reference__ 0
+#endif
+
 #ifndef __SYCL_ANY_DEVICE_HAS_host__
 // __SYCL_ASPECT(host, 0)
 #define __SYCL_ANY_DEVICE_HAS_host__ 0
@@ -427,3 +482,58 @@
 // __SYCL_ASPECT(ext_intel_legacy_image, 41)
 #define __SYCL_ANY_DEVICE_HAS_ext_intel_legacy_image__ 0
 #endif
+
+#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_bindless_images__
+// __SYCL_ASPECT(ext_oneapi_bindless_images, 42)
+#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_bindless_images__ 0
+#endif
+
+#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_bindless_images_shared_usm__
+//__SYCL_ASPECT(ext_oneapi_bindless_images_shared_usm, 43)
+#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_bindless_images_shared_usm__ 0
+#endif
+
+#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_bindless_images_1d_usm__
+//__SYCL_ASPECT(ext_oneapi_bindless_images_1d_usm, 44)
+#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_bindless_images_1d_usm__ 0
+#endif
+
+#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_bindless_images_2d_usm__
+//__SYCL_ASPECT(ext_oneapi_bindless_images_2d_usm, 45)
+#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_bindless_images_2d_usm__ 0
+#endif
+
+#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_interop_memory_import__
+//__SYCL_ASPECT(ext_oneapi_interop_memory_import, 46)
+#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_interop_memory_import__ 0
+#endif
+
+#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_interop_memory_export__
+//__SYCL_ASPECT(ext_oneapi_interop_memory_export, 47)
+#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_interop_memory_export__ 0
+#endif
+
+#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_interop_semaphore_import__
+//__SYCL_ASPECT(ext_oneapi_interop_semaphore_import, 48)
+#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_interop_semaphore_import__ 0
+#endif
+
+#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_interop_semaphore_export__
+//__SYCL_ASPECT(ext_oneapi_interop_semaphore_export, 49)
+#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_interop_semaphore_export__ 0
+#endif
+
+#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_mipmap__
+//__SYCL_ASPECT(ext_oneapi_mipmap, 50)
+#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_mipmap__ 0
+#endif
+
+#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_mipmap_anisotropy__
+//__SYCL_ASPECT(ext_oneapi_mipmap_anisotropy, 51)
+#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_mipmap_anisotropy__ 0
+#endif
+
+#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_mipmap_level_reference__
+//__SYCL_ASPECT(ext_oneapi_mipmap_level_reference, 52)
+#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_mipmap_level_reference__ 0
+#endif