Adding SYCL bindless image support

modules/cpu/common/DeviceRTImpl_ispc.cpp

@@ -381,6 +381,30 @@ void *DeviceImpl::getSyclQueuePtr()
   // SYCL not used
   return nullptr;
 }
+void *createImageMemHandle(void **hostData,
+    const size_t width,
+    const size_t height,
+    const unsigned int numLevels,
+    const OSPTextureFormat format) 
+{
+  return nullptr;
+}
+
+void freeImageMemHandle(void *handle)  
+{
+
+}
+void *createSampledImageHandle(
+    void *imgMemHandlePtr,
+    const OSPTextureFilter filter,
+    const vec2ui wrapMode) 
+{
+  return nullptr;
+}
 
+void freeSampledImageHandle(void *handle)  
+{
+
+}
 } // namespace devicert
 } // namespace ospray

modules/cpu/common/DeviceRTImpl_ispc.h

@@ -104,7 +104,19 @@ struct OSPRAY_SDK_INTERFACE DeviceImpl : public Device
   void *getSyclDevicePtr() override;
   void *getSyclContextPtr() override;
   void *getSyclQueuePtr() override;
+  void *createImageMemHandle(void ** hostData,
+    const size_t width,
+    const size_t height,
+    const unsigned int numLevels,
+    const OSPTextureFormat format) override;
 
+  void freeImageMemHandle(void *handle) override;
+
+  void *createSampledImageHandle(void *imgMemHandle,
+      const OSPTextureFilter filter,
+      const vec2ui wrapMode) override;
+
+  void freeSampledImageHandle(void *handle) override;
  private:
   // Inner command classes
   class Command;

modules/cpu/common/DeviceRTImpl_sycl.cpp

@@ -91,7 +91,15 @@ DeviceImpl::DeviceImpl(void *devicePtr, void *contextPtr, bool debug)
                               << device.get_info<sycl::info::device::name>()
                               << " device (provided externally)";
 }
-
+DeviceImpl::~DeviceImpl()
+{
+  for (auto &entry : imageMemCache) 
+  {
+      syclexp::free_image_mem(
+            entry.second.memHandle, syclexp::image_type::mipmap, queue);
+  }
+  imageMemCache.clear();
+}
 void *DeviceImpl::deviceMalloc(std::size_t size)
 {
   return sycl::malloc_device(size, queue);
@@ -213,5 +221,168 @@ void *DeviceImpl::getSyclQueuePtr()
   return &queue;
 }
 
+void *DeviceImpl::createImageMemHandle(void **hostData,
+    const size_t width,
+    const size_t height,
+    const unsigned int numLevels,
+    const OSPTextureFormat format)
+{
+  // Determine number of channels and channel data type based on the texture format
+  size_t numChannels = 0;
+  syclexp::image_channel_type channelType;
+
+  switch (format) {
+    case OSP_TEXTURE_RGBA8:
+    case OSP_TEXTURE_SRGBA:
+      numChannels = 4;
+      channelType = syclexp::image_channel_type::unorm_int8;
+    break;
+    case OSP_TEXTURE_RGBA32F:
+      numChannels = 4;
+      channelType = syclexp::image_channel_type::fp32;
+      break;
+    case OSP_TEXTURE_RGBA16:
+      numChannels = 4;
+      channelType = syclexp::image_channel_type::unorm_int16;
+      break;
+    case OSP_TEXTURE_RGBA16F:
+      numChannels = 4;
+      channelType = syclexp::image_channel_type::fp16;
+      break;
+    case OSP_TEXTURE_RGB8:
+    case OSP_TEXTURE_SRGB:
+      numChannels = 3;
+      channelType = syclexp::image_channel_type::unorm_int8;
+      break;
+    case OSP_TEXTURE_RGB32F:
+      numChannels = 3;
+      channelType = syclexp::image_channel_type::fp32;
+      break;
+    case OSP_TEXTURE_RGB16:
+      numChannels = 3;
+      channelType = syclexp::image_channel_type::unorm_int16;
+      break;
+    case OSP_TEXTURE_RGB16F:
+      numChannels = 3;
+      channelType = syclexp::image_channel_type::fp16;
+      break;
+    case OSP_TEXTURE_RA8:
+    case OSP_TEXTURE_LA8:
+      numChannels = 2;
+      channelType = syclexp::image_channel_type::unorm_int8;
+      break;
+    case OSP_TEXTURE_RA32F:
+      numChannels = 2;
+      channelType = syclexp::image_channel_type::fp32;
+      break;
+    case OSP_TEXTURE_RA16:
+      numChannels = 2;
+      channelType = syclexp::image_channel_type::unorm_int16;
+      break;
+    case OSP_TEXTURE_RA16F:
+      numChannels = 2;
+      channelType = syclexp::image_channel_type::fp16;
+      break;
+    case OSP_TEXTURE_R8:
+    case OSP_TEXTURE_L8:
+      numChannels = 1;
+      channelType = syclexp::image_channel_type::unorm_int8;
+      break;
+    case OSP_TEXTURE_R32F:
+      numChannels = 1;
+      channelType = syclexp::image_channel_type::fp32;
+      break;
+    case OSP_TEXTURE_R16:
+      numChannels = 1;
+      channelType = syclexp::image_channel_type::unorm_int16;
+      break;
+    case OSP_TEXTURE_R16F:
+      numChannels = 1;
+      channelType = syclexp::image_channel_type::fp16;
+      break;
+    default:
+      throw std::runtime_error("Unsupported texture format for bindless images");
+  }
+  ImageMemEntry imgMemEntry;
+  // Construct the image descriptor.
+  syclexp::image_descriptor imgDesc(
+    {width, height},     // Dimensions
+    numChannels,         // Channel count
+    channelType,          // Channel data type
+    image_type::mipmap,   // Image type (using mipmap type to support multiple levels)
+    numLevels            // Number of mipmap levels   
+  );
+
+  syclexp::image_mem_handle memHandle = syclexp::alloc_image_mem(imgDesc, queue); 
+  for (size_t i = 0; i < numLevels; i++) {
+      syclexp::image_mem_handle levelHandle =
+          syclexp::get_mip_level_mem_handle(memHandle, i, queue);
+      syclexp::image_descriptor levelDesc = imgDesc.get_mip_level_desc(i);
+      queue.ext_oneapi_copy(hostData[i], levelHandle, levelDesc);
+  }
+
+  imgMemEntry.desc = imgDesc;
+  imgMemEntry.memHandle = memHandle;
+  imageMemCache[memHandle.raw_handle] = imgMemEntry;
+  queue.wait_and_throw();
+  return key;
+}
+
+void DeviceImpl::freeImageMemHandle(void *handle)
+{
+  syclexp::image_mem_handle memHandle;
+  memHandle.raw_handle = handle;
+  syclexp::free_image_mem(memHandle, syclexp::image_type::mipmap, queue);
+  imageMemCache.erase(handle);
+}
+
+void *DeviceImpl::createSampledImageHandle(
+    void *imgMemHandlePtr, const OSPTextureFilter filter, const vec2ui wrapMode)
+{
+    sycl::addressing_mode addressingMode;
+    switch (wrapMode.x) {
+    case OSP_TEXTURE_WRAP_REPEAT:
+        addressingMode = sycl::addressing_mode::repeat;
+        break;
+    case OSP_TEXTURE_WRAP_MIRRORED_REPEAT:
+        addressingMode = sycl::addressing_mode::mirrored_repeat;
+        break;
+    case OSP_TEXTURE_WRAP_CLAMP_TO_EDGE:
+        addressingMode = sycl::addressing_mode::clamp_to_edge;
+        break;
+    default:
+        addressingMode = sycl::addressing_mode::repeat;
+    }
+
+    sycl::filtering_mode filteringMode = (filter == OSP_TEXTURE_FILTER_NEAREST)
+        ? sycl::filtering_mode::nearest
+        : sycl::filtering_mode::linear;
+
+    syclexp::bindless_image_sampler sampler(
+        addressingMode,
+        sycl::coordinate_normalization_mode::normalized,
+        filteringMode,
+        filteringMode,
+        0.f,
+        static_cast<float>(32),
+        0.f);
+    //Get the image descriptor for this image handle
+    syclexp::image_descriptor imgDesc = imageMemCache[imgMemHandlePtr].desc;
+    //Rebuild the image handle from the pointer
+    syclexp::image_mem_handle memHandle;
+    memHandle.raw_handle = imgMemHandlePtr;
+
+    syclexp::sampled_image_handle sampledHandle =
+        syclexp::create_image(memHandle, sampler, imgDesc, queue);
+
+    return reinterpret_cast<void *>(sampledHandle.raw_handle);
+}
+
+void DeviceImpl::freeSampledImageHandle(void *handle) {
+    syclexp::sampled_image_handle sampledHandle;
+    sampledHandle.raw_handle = reinterpret_cast<syclexp::sampled_image_handle::raw_handle_type>(handle);
+    syclexp::destroy_image_handle(sampledHandle, queue);
+}
+
 } // namespace devicert
 } // namespace ospray

modules/cpu/common/DeviceRTImpl_sycl.h

@@ -6,7 +6,7 @@
 #include "common/DeviceRT.h"
 
 #include <sycl/sycl.hpp>
-
+namespace syclexp = sycl::ext::oneapi::experimental;
 namespace ospray {
 namespace devicert {
 
@@ -43,6 +43,7 @@ struct OSPRAY_SDK_INTERFACE DeviceImpl : public Device
   DeviceImpl(bool debug);
   DeviceImpl(uint32_t deviceId, bool debug);
   DeviceImpl(void *devicePtr, void *contextPtr, bool debug);
+  ~DeviveImpl();
 
   // Allocate device memory
   void *deviceMalloc(std::size_t size) override;
@@ -89,7 +90,28 @@ struct OSPRAY_SDK_INTERFACE DeviceImpl : public Device
   void *getSyclContextPtr() override;
   void *getSyclQueuePtr() override;
 
+  void *createImageMemHandle(void ** hostData,
+    const size_t width,
+    const size_t height,
+    const unsigned int numLevels,
+    const OSPTextureFormat format) override;
+
+  void freeImageMemHandle(void *handle) override;
+
+  void *createSampledImageHandle(void *imgMemHandle,
+      const OSPTextureFilter filter,
+      const vec2ui wrapMode) override;
+
+  void freeSampledImageHandle(void *handle) override;
+
  private:
+
+ struct ImageMemEntry {
+    syclexp::image_mem_handle memHandle;
+    syclexp::image_descriptor desc;
+  };
+  std::unordered_map<void *, ImageMemEntry> imageMemCache;
+
   sycl::device device;
   sycl::context context;
   sycl::queue queue;

modules/cpu/texture/Texture2D.cpp

@@ -141,6 +141,15 @@ void Texture2D::commit()
   // Initialize ispc shared structure
   getSh()->set(
       size, dataPtr.data(), dataPtr.size() - 1, format, filter, wrapMode);
+
+  void *imgMemHandle = getISPCDevice().getDRTDevice().createImageMemHandle(
+    dataPtr.data(), size.x, size.y, dataPtr.size(), format);
+
+  void *sampledHandle = getISPCDevice().getDRTDevice().createSampledImageHandle(
+    imgMemHandle, filter, wrapMode);
+
+  getSh()->data[0] = sampledHandle;
+
 }
 
 } // namespace ospray

modules/cpu/texture/Texture2D.ispc

@@ -5,7 +5,10 @@
 #include "Texture2DFormats.ih"
 
 OSPRAY_BEGIN_ISPC_NAMESPACE
-
+#ifdef OSPRAY_TARGET_SYCL
+using half = sycl::half;
+namespace syclexp = sycl::ext::oneapi::experimental;
+#endif
 // TODO tiling
 
 // Texture coordinate utilities
@@ -136,6 +139,17 @@ inline vec4f Texture2D_bilinear_l(
 SYCL_EXTERNAL __noinline vec4f Texture2D_get(
     const Texture2D *uniform self, const vec2f &st, const float pixelFootprint)
 {
+ #ifdef OSPRAY_TARGET_SYCL
+  syclexp::sampled_image_handle handle;
+  handle.raw_handle = (syclexp::sampled_image_handle::raw_handle_type)self->data[0];
+
+  sycl::float4 res = syclexp::sample_mipmap<sycl::float4>(
+      handle,
+      sycl::float2{st.x, st.y},
+      0.f);
+
+  return make_vec4f(res.x(), res.y(), res.z(), res.w());
+#else
   int mipLevel = 0;
   float frac = 0.f;
   const uniform bool filter_nearest = self->filter & OSP_TEXTURE_FILTER_NEAREST;
@@ -145,15 +159,14 @@ SYCL_EXTERNAL __noinline vec4f Texture2D_get(
     mipLevel = min((int)lambda, self->maxLevel - 1);
     frac = min(lambda - mipLevel, 1.f);
   }
-
   vec4f t0 = filter_nearest ? Texture2D_nearest_l(self, st, mipLevel)
                             : Texture2D_bilinear_l(self, st, mipLevel);
   if (frac == 0.f)
     return t0;
-
   vec4f t1 = filter_nearest ? Texture2D_nearest_l(self, st, mipLevel + 1)
                             : Texture2D_bilinear_l(self, st, mipLevel + 1);
   return lerp(frac, t0, t1);
+#endif
 }
 
 export void *uniform Texture2D_get_addr(const uniform uint32 type)

ospray/common/DeviceRT.h

@@ -187,7 +187,20 @@ struct OSPRAY_CORE_INTERFACE Device
   virtual void *getSyclDevicePtr() = 0;
   virtual void *getSyclContextPtr() = 0;
   virtual void *getSyclQueuePtr() = 0;
-
+  // For bindless image support:
+  virtual void *createImageMemHandle(void **hostData,
+    const size_t width,
+    const size_t height,
+    const unsigned int numLevels,
+    const OSPTextureFormat format
+  ) = 0;
+
+  //Receives what comes from createImageMemHandle and creates a handle for the sampled image
+  virtual void *createSampledImageHandle(void *imgMemHandle,
+    const OSPTextureFilter filter,
+    const vec2ui wrapMode) = 0;
+  virtual void freeImageMemHandle(void *handle) = 0;
+  virtual void freeSampledImageHandle(void *handle) = 0;
   // Check if device is in debug mode
   bool isDebug() const;