diff --git a/modules/cpu/common/DeviceRTImpl_ispc.cpp b/modules/cpu/common/DeviceRTImpl_ispc.cpp index 0423091cd..746337374 100644 --- a/modules/cpu/common/DeviceRTImpl_ispc.cpp +++ b/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 diff --git a/modules/cpu/common/DeviceRTImpl_ispc.h b/modules/cpu/common/DeviceRTImpl_ispc.h index d9ba75b60..909c9fc33 100644 --- a/modules/cpu/common/DeviceRTImpl_ispc.h +++ b/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; diff --git a/modules/cpu/common/DeviceRTImpl_sycl.cpp b/modules/cpu/common/DeviceRTImpl_sycl.cpp index 8ac92a7b4..1bec6b627 100644 --- a/modules/cpu/common/DeviceRTImpl_sycl.cpp +++ b/modules/cpu/common/DeviceRTImpl_sycl.cpp @@ -91,7 +91,15 @@ DeviceImpl::DeviceImpl(void *devicePtr, void *contextPtr, bool debug) << device.get_info() << " 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(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(sampledHandle.raw_handle); +} + +void DeviceImpl::freeSampledImageHandle(void *handle) { + syclexp::sampled_image_handle sampledHandle; + sampledHandle.raw_handle = reinterpret_cast(handle); + syclexp::destroy_image_handle(sampledHandle, queue); +} + } // namespace devicert } // namespace ospray diff --git a/modules/cpu/common/DeviceRTImpl_sycl.h b/modules/cpu/common/DeviceRTImpl_sycl.h index 95b6b1de5..d6ee081ad 100644 --- a/modules/cpu/common/DeviceRTImpl_sycl.h +++ b/modules/cpu/common/DeviceRTImpl_sycl.h @@ -6,7 +6,7 @@ #include "common/DeviceRT.h" #include - +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 imageMemCache; + sycl::device device; sycl::context context; sycl::queue queue; diff --git a/modules/cpu/texture/Texture2D.cpp b/modules/cpu/texture/Texture2D.cpp index 655014139..59d77ae51 100644 --- a/modules/cpu/texture/Texture2D.cpp +++ b/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 diff --git a/modules/cpu/texture/Texture2D.ispc b/modules/cpu/texture/Texture2D.ispc index c183a6fab..13b16c2d2 100644 --- a/modules/cpu/texture/Texture2D.ispc +++ b/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( + 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) diff --git a/ospray/common/DeviceRT.h b/ospray/common/DeviceRT.h index b920117a0..4aff5ba74 100644 --- a/ospray/common/DeviceRT.h +++ b/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;