diff --git a/PBR/Render/include/sycl_renderer.hpp b/PBR/Render/include/sycl_renderer.hpp index 04b1dc4..e7f2b78 100644 --- a/PBR/Render/include/sycl_renderer.hpp +++ b/PBR/Render/include/sycl_renderer.hpp @@ -16,13 +16,15 @@ #include "PBR/Render/brdf/cook_torrance.hpp" #include "PBR/HitData/hit_data.hpp" #include "PBR/Render/shared/core_renderer.hpp" +#include "PBR/TextureManager/general_gpu_texture.hpp" namespace syclexp = sycl::ext::oneapi::experimental; class SYCL_Renderer : public IComputeRenderer { private: sycl::ext::oneapi::experimental::sampled_image_handle* m_textureHandles = nullptr; + GPUTexture* m_bufferTextures = nullptr; // For devices with no bindless image supportGPUTexture int m_numTextures = 0; sycl::queue m_queue; - + bool m_useBindlessImage = true; public: SYCL_Renderer(){ @@ -72,11 +74,31 @@ class SYCL_Renderer : public IComputeRenderer { std::cout << " Uploaded " << m_numTextures << " texture handles to GPU" << std::endl; } } + void setGPUBufferTextures(const std::vector& textures) { + if (m_bufferTextures) { + sycl::free(m_bufferTextures,m_queue); + m_bufferTextures = nullptr; + } + m_useBindlessImage = false; + m_numTextures = textures.size(); + + m_bufferTextures = sycl::malloc_device( + m_numTextures, m_queue + ); + m_queue.memcpy(m_bufferTextures, textures.data(), + m_numTextures * sizeof(GPUTexture)).wait(); + + std::cout << " Uploaded " << m_numTextures << "buffer textures to GPU" << std::endl; + } ~SYCL_Renderer() { if (m_textureHandles) { sycl::free(m_textureHandles, m_queue); m_textureHandles = nullptr; } + if(m_bufferTextures){ + sycl::free(m_bufferTextures, m_queue); + m_bufferTextures = nullptr; + } } }; diff --git a/PBR/Render/src/sycl_renderer.cpp b/PBR/Render/src/sycl_renderer.cpp index 1aa2cc7..22800b4 100644 --- a/PBR/Render/src/sycl_renderer.cpp +++ b/PBR/Render/src/sycl_renderer.cpp @@ -168,6 +168,8 @@ std::vector SYCL_Renderer::RenderScene( int numNodes = nodes.size(); int numLights = lightsList.size(); int numTextures = m_numTextures; + + //Checking Texture properties auto textureHandles = m_textureHandles; // Progressive loop OUTSIDE diff --git a/PBR/Space/space.cpp b/PBR/Space/space.cpp index 5a5fb0d..75981a5 100644 --- a/PBR/Space/space.cpp +++ b/PBR/Space/space.cpp @@ -124,8 +124,16 @@ void Space::sendTexturesToRender() if (syclRenderer && m_textureManager) { // Set textures on renderer auto syclTexMgr = dynamic_cast(m_textureManager.get()); - if (syclTexMgr) - syclRenderer->setSyclTextureHandles(syclTexMgr->getImageHandles()); + if (syclTexMgr){ + if(syclTexMgr->hasBindlessSupport()) + { + syclRenderer->setSyclTextureHandles(syclTexMgr->getImageHandles()); + } + else{ + syclRenderer->setGPUBufferTextures(syclTexMgr->getBufferTextures()); + } + } + } break; } diff --git a/PBR/TextureManager/general_gpu_texture.hpp b/PBR/TextureManager/general_gpu_texture.hpp new file mode 100644 index 0000000..0703f19 --- /dev/null +++ b/PBR/TextureManager/general_gpu_texture.hpp @@ -0,0 +1,10 @@ +#if !defined(_GENERAL_GPU_BUFFER_H_) +#define _GENERAL_GPU_BUFFER_H_ + +struct GPUTexture { + float* data; // Device pointer to RGBA float data + int width; + int height; +}; + +#endif // _GENERAL_GPU_BUFFER_H_ diff --git a/PBR/TextureManager/sampler2d_texture.hpp b/PBR/TextureManager/sampler2d_texture.hpp index c77e1ec..7536e99 100644 --- a/PBR/TextureManager/sampler2d_texture.hpp +++ b/PBR/TextureManager/sampler2d_texture.hpp @@ -2,6 +2,7 @@ #define _SAMPLER_2D_TEXTURE_HPP_ #include "gpu/include/fgt_cpu_device.hpp" #include "texture_types.hpp" +#include "general_gpu_texture.hpp" // Portable sampling function fgt_device_gpu inline fungt::Vec3 sampleTexture2D( const TextureDeviceObject& texture, @@ -24,5 +25,51 @@ fgt_device_gpu inline fungt::Vec3 sampleTexture2D( return fungt::Vec3(1.0f, 0.0f, 1.0f); // Magenta error color #endif } +// Buffer path - manual bilinear filter for devices without bindless image support +fgt_device_gpu inline fungt::Vec3 sampleTexture2D( + const GPUTexture& texture, + float u, + float v +) { + int w = texture.width; + int h = texture.height; + + u = u < 0.0f ? 0.0f : (u > 1.0f ? 1.0f : u); + v = v < 0.0f ? 0.0f : (v > 1.0f ? 1.0f : v); + + float fx = u * (w - 1); + float fy = v * (h - 1); + + int x0 = (int)fx; + int y0 = (int)fy; + int x1 = x0 + 1 < w ? x0 + 1 : x0; + int y1 = y0 + 1 < h ? y0 + 1 : y0; + + float tx = fx - (float)x0; + float ty = fy - (float)y0; + + const float* d = texture.data; +#define FETCH(x, y, c) d[((y) * w + (x)) * 4 + (c)] + + float r = (1.0f - tx) * (1.0f - ty) * FETCH(x0, y0, 0) + + tx * (1.0f - ty) * FETCH(x1, y0, 0) + + (1.0f - tx) * ty * FETCH(x0, y1, 0) + + tx * ty * FETCH(x1, y1, 0); + + float g = (1.0f - tx) * (1.0f - ty) * FETCH(x0, y0, 1) + + tx * (1.0f - ty) * FETCH(x1, y0, 1) + + (1.0f - tx) * ty * FETCH(x0, y1, 1) + + tx * ty * FETCH(x1, y1, 1); + + float b = (1.0f - tx) * (1.0f - ty) * FETCH(x0, y0, 2) + + tx * (1.0f - ty) * FETCH(x1, y0, 2) + + (1.0f - tx) * ty * FETCH(x0, y1, 2) + + tx * ty * FETCH(x1, y1, 2); + +#undef FETCH + + return fungt::Vec3(r, g, b); +} + #endif // _SAMPLER_2D_TEXTURE_HPP_ diff --git a/PBR/TextureManager/sycl_texture.cpp b/PBR/TextureManager/sycl_texture.cpp index 3d85bc0..0ea65b8 100644 --- a/PBR/TextureManager/sycl_texture.cpp +++ b/PBR/TextureManager/sycl_texture.cpp @@ -4,10 +4,13 @@ SYCLTexture::SYCLTexture(sycl::queue& queue) :m_queue{&queue}{ + auto currDevice = m_queue->get_device(); std::cout << "SYCLTexture: Initialized with queue for device: " - << queue.get_device().get_info() + << currDevice.get_info() << std::endl; - + + m_useBindlessImages = currDevice.has(sycl::aspect::ext_oneapi_bindless_images); + std::cout << "Bindless Image support: " << (m_useBindlessImages ? "YES" : "NO") << std::endl; } SYCLTexture::~SYCLTexture() { cleanup(); @@ -29,40 +32,17 @@ std::cout << "SYCLTexture: Loaded " << path << " (" << width << "x" << height << ", " << channels << " channels)" << std::endl; - try{ - const unsigned int numChannels = 4; - const auto channelType = sycl::image_channel_type::unorm_int8; - syclexp::image_descriptor desc( - { static_cast(width), static_cast(height) }, - numChannels, - channelType - ); - syclexp::image_mem imgMem(desc, *m_queue); - - auto cpyToDeviceEvent = m_queue->ext_oneapi_copy( - data, //Source - imgMem.get_handle(), //Destination - desc //Image descriptor - ); - cpyToDeviceEvent.wait_and_throw(); - syclexp::bindless_image_sampler sampler( - sycl::addressing_mode::repeat, - sycl::coordinate_normalization_mode::normalized, - sycl::filtering_mode::linear - ); - syclexp::sampled_image_handle imgHandle = - syclexp::create_image(imgMem, sampler,desc, *m_queue); - - SYCLTextureData texData; - texData.imgHandle = imgHandle; - texData.imgMem = std::move(imgMem); - texData.width = width; - texData.height = height; - texData.path = path; - int index = textures.size(); - textures.push_back(std::move(texData)); - pathToIndex[path] = index; + + + int index = -1; + + if (m_useBindlessImages) { + index = loadBindlessTexture(data, width, height, path); + } + else { + index = loadBufferTexture(data, width, height, path); + } stbi_image_free(data); @@ -97,16 +77,24 @@ // Wait for all operations to complete m_queue->wait_and_throw(); - for (size_t i = 0; i < textures.size(); i++) { - auto& tex = textures[i]; - std::cout << "SYCLTexture: Destroying texture " << i << std::endl; - - try { - syclexp::destroy_image_handle(tex.imgHandle, *m_queue); + if(m_useBindlessImages){ + for (size_t i = 0; i < textures.size(); i++) { + auto& tex = textures[i]; + std::cout << "SYCLTexture: Destroying texture " << i << std::endl; + + try { + syclexp::destroy_image_handle(tex.imgHandle, *m_queue); + } + catch (const sycl::exception& e) { + std::cerr << "SYCLTexture: Error destroying texture " << i + << ": " << e.what() << std::endl; + } } - catch (const sycl::exception& e) { - std::cerr << "SYCLTexture: Error destroying texture " << i - << ": " << e.what() << std::endl; + } + else{ + for (auto& tex : m_bufferTextures) { + std::cout << "SYCLTexture: Destroying buffer texture " << i << std::endl; + sycl::free(tex.deviceData, *m_queue); } } @@ -119,3 +107,66 @@ textures.clear(); pathToIndex.clear(); } + + int SYCLTexture::loadBindlessTexture(unsigned char* data, int width, int height, const std::string& path) + { + const unsigned int numChannels = 4; + const auto channelType = sycl::image_channel_type::unorm_int8; + syclexp::image_descriptor desc( + { static_cast(width), static_cast(height) }, + numChannels, + channelType + ); + syclexp::image_mem imgMem(desc, *m_queue); + + auto cpyToDeviceEvent = m_queue->ext_oneapi_copy( + data, //Source + imgMem.get_handle(), //Destination + desc //Image descriptor + ); + cpyToDeviceEvent.wait_and_throw(); + syclexp::bindless_image_sampler sampler( + sycl::addressing_mode::repeat, + sycl::coordinate_normalization_mode::normalized, + sycl::filtering_mode::linear + ); + syclexp::sampled_image_handle imgHandle = + syclexp::create_image(imgMem, sampler, desc, *m_queue); + + SYCLTextureData texData; + texData.imgHandle = imgHandle; + texData.imgMem = std::move(imgMem); + texData.width = width; + texData.height = height; + texData.path = path; + int index = textures.size(); + textures.push_back(std::move(texData)); + pathToIndex[path] = index; + + return index; + } + + int SYCLTexture::loadBufferTexture(unsigned char* data, int w, int h, const std::string& path) + { + std::vector floatData(w * h * 4); + for (int i = 0; i < w * h * 4; i++) { + floatData[i] = data[i] / 255.0f; + } + + float* devData = sycl::malloc_device(w * h * 4, *m_queue); //m_queue is of type pointer + m_queue->memcpy(devData, floatData.data(), w * h * 4 * sizeof(float)).wait(); + + + BufferTextureData tex; + tex.deviceData = devData; + tex.width = w; + tex.height = h; + tex.path = path; + + + int index = m_bufferTextures.size(); + m_bufferTextures.push_back(tex); + + std::cout << "Loaded buffer texture " << index << " (" << path << ")" << std::endl; + return index; + } diff --git a/PBR/TextureManager/sycl_texture.hpp b/PBR/TextureManager/sycl_texture.hpp index bfb65aa..858abd3 100644 --- a/PBR/TextureManager/sycl_texture.hpp +++ b/PBR/TextureManager/sycl_texture.hpp @@ -2,6 +2,7 @@ #define _SYCL_TEXTURE_HPP_ #include "idevice_texture.hpp" +#include "general_gpu_texture.hpp" #include #include #include @@ -17,13 +18,21 @@ struct SYCLTextureData { int width, height; std::string path; }; +//General buffers for textures +struct BufferTextureData { + float* deviceData; // Device pointer + int width; + int height; + std::string path; +}; class SYCLTexture : public IDeviceTexture { private: std::vector textures; std::map pathToIndex; sycl::queue* m_queue; - + bool m_useBindlessImages = true; + std::vector m_bufferTextures; // Add to class public: SYCLTexture(sycl::queue& queue); ~SYCLTexture(); @@ -31,8 +40,10 @@ class SYCLTexture : public IDeviceTexture { int loadTexture(const std::string& path) override; int getTextureCount() const override{}; void cleanup() override; - - // MATCHING CUDA PATTERN - return host-side handles! + int loadBindlessTexture(unsigned char* data, int w, int h, const std::string& path); + int loadBufferTexture(unsigned char* data, int w, int h, const std::string& path); + bool hasBindlessSupport(){return m_useBindlessImages; } + // return host-side handles! std::vector getImageHandles() { std::vector handles; for (const auto& tex : textures) { @@ -41,6 +52,22 @@ class SYCLTexture : public IDeviceTexture { std::cout<< "HANDLES SIZE : " < getBufferTextures() { + if (m_useBindlessImages) { + throw std::runtime_error("Use bindless handles instead!"); + } + + std::vector result; + for (const auto& tex : m_bufferTextures) { + GPUTexture gpuTex; + gpuTex.data = tex.deviceData; + gpuTex.width = tex.width; + gpuTex.height = tex.height; + result.push_back(gpuTex); + } + return result; + } + }; #endif // _SYCL_TEXTURE_HPP_ \ No newline at end of file diff --git a/PBR/TextureManager/texture_types.hpp b/PBR/TextureManager/texture_types.hpp index dcdcd48..2d334a3 100644 --- a/PBR/TextureManager/texture_types.hpp +++ b/PBR/TextureManager/texture_types.hpp @@ -8,8 +8,8 @@ using TextureDeviceObject = cudaTextureObject_t; #elif defined(FUNGT_USE_SYCL) && !defined(__CUDACC__) #include -#include // ← ADD THIS! -namespace syclexp = sycl::ext::oneapi::experimental; // ← FIX THIS! +#include +namespace syclexp = sycl::ext::oneapi::experimental; using TextureDeviceObject = syclexp::sampled_image_handle; #define TEXTURE_BACKEND_SYCL