[SYCL][Bindless][3/4] Add experimental implementation of SYCL bindles…

…s images extension (intel#10454)

This commit stands as the third commit of four to make code review
easier, mostly covering the changes made to the user-facing SYCL API for
the [bindless images extension
proposal](intel#9842).

### Overview

The bindless images extension provides a new interface for allocating,
creating, and accessing images in SYCL. Image memory allocation is
seperated from image handle creation, and image handles can be passed to
kernels without requesting access through accessors. This approach
provides much more flexibility to the user, as well as enabling programs
to implement features that were impossible to implement using standard
SYCL images, such as a texture atlas. In addition to providing a new
interface for images, this extension also provides initial experimental
support for importing external memory into SYCL.

### Following Split PRs

- [4/4] Add tests

### Authors

Co-authored-by: Isaac Ault <isaac.ault@codeplay.com>
Co-authored-by: Hugh Bird <hugh.bird@codeplay.com>
Co-authored-by: Duncan Brawley <duncan.brawley@codeplay.com>
Co-authored-by: Przemek Malon <przemek.malon@codeplay.com>
Co-authored-by: Chedy Najjar <chedy.najjar@codeplay.com>
Co-authored-by: Sean Stirling <sean.stirling@codeplay.com>
Co-authored-by: Peter Zuzek <peter@codeplay.com>

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