From b6ec82cc7d75a81c0efdd71fce37d6b8a4cf1480 Mon Sep 17 00:00:00 2001 From: juanchuletas Date: Mon, 17 Aug 2026 23:53:55 -0600 Subject: [PATCH 1/7] Adding OpenCL backend --- PBR/TextureManager/idevice_texture.hpp | 1 - PBR/TextureManager/ocl_texture.hpp | 28 ++++++++++++++++++++++++++ 2 files changed, 28 insertions(+), 1 deletion(-) create mode 100644 PBR/TextureManager/ocl_texture.hpp diff --git a/PBR/TextureManager/idevice_texture.hpp b/PBR/TextureManager/idevice_texture.hpp index 60d1972..db68049 100644 --- a/PBR/TextureManager/idevice_texture.hpp +++ b/PBR/TextureManager/idevice_texture.hpp @@ -1,6 +1,5 @@ #if !defined(_IDEVICE_TEXTURE_HPP) #define _IDEVICE_TEXTURE_HPP -//#include "texture_types.hpp" #include "Vector/vector3.hpp" #include #include diff --git a/PBR/TextureManager/ocl_texture.hpp b/PBR/TextureManager/ocl_texture.hpp new file mode 100644 index 0000000..684b833 --- /dev/null +++ b/PBR/TextureManager/ocl_texture.hpp @@ -0,0 +1,28 @@ +#if !defined(_OPENCL_TEXTURE_HPP_) +#define _OPENCL_TEXTURE_HPP_ +#include +#include +#include +#include "idevice_texture.hpp" +#define CL_MEM_BINDLESS_IMAGE_INTEL 0x10060 +#define CL_IMAGE_BINDLESS_HANDLE_INTEL 0x10061 + +class OpenCLTexture +{ + uint64_t m_handleId; + cl_image_format m_imageFormat{}; + m_image_format.image_channel_order = CL_RGBA; + m_image_format.image_channel_data_type = CL_UNORM_INT8; + + public: + OCLTexture(); + + + + +}; + + + + +#endif // _OPENCL_TEXTURE_HPP_ From 9da05c22e6051442b13ecba62c1c9069ce16aa1c Mon Sep 17 00:00:00 2001 From: juanchuletas Date: Tue, 18 Aug 2026 12:27:46 -0600 Subject: [PATCH 2/7] Refactor: texture addition in Space class --- PBR/Render/include/cpu_renderer.hpp | 3 + PBR/Render/include/cuda_renderer.hpp | 56 +++++++++--------- PBR/Render/include/icompute_renderer.hpp | 2 + PBR/Render/include/sycl_renderer.hpp | 64 +++++++++++++-------- PBR/Render/src/cuda_renderer.cu | 4 +- PBR/Render/src/sycl_renderer.cpp | 3 + PBR/Space/space.cpp | 72 ++---------------------- PBR/Space/space.hpp | 11 ---- PBR/TextureManager/cuda_texture.cpp | 3 +- PBR/TextureManager/cuda_texture.hpp | 3 + PBR/TextureManager/ocl_texture.hpp | 51 ++++++++++++----- PBR/TextureManager/sycl_texture.cpp | 2 + PBR/TextureManager/sycl_texture.hpp | 7 ++- Samples/hdr_test/main.cpp | 19 +++---- Samples/iwocl/CMakeLists.txt | 4 ++ 15 files changed, 144 insertions(+), 160 deletions(-) diff --git a/PBR/Render/include/cpu_renderer.hpp b/PBR/Render/include/cpu_renderer.hpp index 34901ce..5a6218c 100644 --- a/PBR/Render/include/cpu_renderer.hpp +++ b/PBR/Render/include/cpu_renderer.hpp @@ -5,11 +5,14 @@ #include "../Ray/ray.hpp" #include "../HitData/hit_data.hpp" #include "PBR/Intersection/intersection.hpp" +#include "PBR/TextureManager/cpu_texture.hpp" class CPU_Renderer : public IComputeRenderer{ + CPUTexture m_textures; public: + IDeviceTexture& textures() override { return m_textures; } fungt::Vec3 shadeNormal(const fungt::Vec3& normal) { // Convert from [-1,1] to [0,1] //return 0.5f * (normal + fungt::Vec3(1.0f, 1.0f, 1.0f)); diff --git a/PBR/Render/include/cuda_renderer.hpp b/PBR/Render/include/cuda_renderer.hpp index 71669f2..ce63a0a 100644 --- a/PBR/Render/include/cuda_renderer.hpp +++ b/PBR/Render/include/cuda_renderer.hpp @@ -22,16 +22,45 @@ using cudaTextureObject_t = unsigned long long; #include "PBR/BVH/bvh_node.hpp" #include "PBR/Light/light.hpp" +#include "PBR/TextureManager/cuda_texture.hpp" #define CUDA_CHECK(err) do { cudaError_t e = (err); if (e != cudaSuccess) { \ std::cerr << "CUDA error: " << cudaGetErrorString(e) << " at " << __FILE__ << ":" << __LINE__ << std::endl; exit(1); }} while(0) class CUDA_Renderer : public IComputeRenderer{ - + CUDATexture m_textures; cudaTextureObject_t* m_textureObj= nullptr; int m_numTextures = 0; + + void prepareTextures() { + if (m_textures.handlesAreDirty()) { + setCudaTextureObjects(m_textures.getTextureObjects()); + m_textures.markHandlesClean(); + } + } + + void setCudaTextureObjects(const std::vector& textureObj) { + std::cout<<"*** SETTING CUDA TEXTURE OBJECTS*** "<0){ + CUDA_CHECK(cudaMalloc(&m_textureObj, + m_numTextures * sizeof(cudaTextureObject_t))); + CUDA_CHECK(cudaMemcpy(m_textureObj, + textureObj.data(), + m_numTextures * sizeof(cudaTextureObject_t), + cudaMemcpyHostToDevice)); + std::cout << " Uploaded " << m_numTextures << " textures to GPU" << std::endl; + } + } + public: CUDA_Renderer() = default; + IDeviceTexture& textures() override { return m_textures; } std::vector RenderScene( int width, @@ -44,31 +73,6 @@ class CUDA_Renderer : public IComputeRenderer{ int samplesPerPixel, int sampleOffset ); - void setCudaTextureObjects(const std::vector& textureObj) { - std::cout<<"*** SETTING CUDA TEXTURE OBJECTS*** "<0){ - // Allocate GPU memory - CUDA_CHECK(cudaMalloc(&m_textureObj, - m_numTextures * sizeof(cudaTextureObject_t))); - - // Copy to GPU - CUDA_CHECK(cudaMemcpy(m_textureObj, - textureObj.data(), - m_numTextures * sizeof(cudaTextureObject_t), - cudaMemcpyHostToDevice)); - - std::cout << " Uploaded " << m_numTextures << " textures to GPU" << std::endl; - - } - - } ~CUDA_Renderer(){ // Cleanup device textures if (m_textureObj) { diff --git a/PBR/Render/include/icompute_renderer.hpp b/PBR/Render/include/icompute_renderer.hpp index 7d0576a..9f396f9 100644 --- a/PBR/Render/include/icompute_renderer.hpp +++ b/PBR/Render/include/icompute_renderer.hpp @@ -5,6 +5,7 @@ #include "Vector/vector3.hpp" #include "PBR/BVH/bvh_node.hpp" #include "PBR/Light/light.hpp" +#include "PBR/TextureManager/idevice_texture.hpp" // Forward declarations class PBRCamera; @@ -13,6 +14,7 @@ class IComputeRenderer{ public: virtual ~IComputeRenderer() = default; + virtual IDeviceTexture& textures() = 0; virtual std::vector RenderScene( int width, int height, const std::vector &triangleList, diff --git a/PBR/Render/include/sycl_renderer.hpp b/PBR/Render/include/sycl_renderer.hpp index d75713a..48265cf 100644 --- a/PBR/Render/include/sycl_renderer.hpp +++ b/PBR/Render/include/sycl_renderer.hpp @@ -2,6 +2,8 @@ #define SYCL_RENDERER_HPP #include // MUST be first! #include +#include +#include #include #include #include @@ -16,40 +18,30 @@ #include "PBR/Render/brdf/cook_torrance.hpp" #include "PBR/HitData/hit_data.hpp" #include "PBR/Render/shared/core_renderer.hpp" +#include "PBR/TextureManager/sycl_texture.hpp" namespace syclexp = sycl::ext::oneapi::experimental; class SYCL_Renderer : public IComputeRenderer { private: sycl::ext::oneapi::experimental::sampled_image_handle* m_textureHandles = nullptr; int m_numTextures = 0; sycl::queue m_queue; + std::unique_ptr m_textures; -public: - - SYCL_Renderer(){ - + void prepareTextures() { + if (!m_textures) { + throw std::runtime_error("SYCL textures requested before queue initialization"); + } + if (m_textures->handlesAreDirty()) { + setSyclTextureHandles(m_textures->getImageHandles()); + m_textures->markHandlesClean(); + } } - std::vector RenderScene( - int width, - int height, - const std::vector& triangleList, - const std::vector& nodes, - const std::vector& lightsList, - const std::vector& emissiveTriIndices, - const PBRCamera& camera, - int samplesPerPixel, - int sampleOffset - ) override; - void createQueue(const std::string &name, flib::vendor v = flib::vendor::INTEL, - flib::device dev = flib::device::GPU, - flib::backend b = flib::backend::LEVEL_ZERO); - sycl::queue &getQueue(); - // MATCHING CUDA PATTERN! + void setSyclTextureHandles( const std::vector& handles ) { std::cout << "*** SETTING SYCL TEXTURE HANDLES ***" << std::endl; - // Free old handles if (m_textureHandles) { sycl::free(m_textureHandles, m_queue); m_textureHandles = nullptr; @@ -59,12 +51,9 @@ class SYCL_Renderer : public IComputeRenderer { std::cout << "*** NUM SYCL TEXTURE HANDLES *** " << m_numTextures << std::endl; if (m_numTextures > 0) { - // Allocate GPU memory m_textureHandles = sycl::malloc_device( m_numTextures, m_queue ); - - // Copy to GPU m_queue.memcpy( m_textureHandles, handles.data(), @@ -74,6 +63,31 @@ class SYCL_Renderer : public IComputeRenderer { std::cout << " Uploaded " << m_numTextures << " texture handles to GPU" << std::endl; } } + +public: + + SYCL_Renderer() = default; + std::vector RenderScene( + int width, + int height, + const std::vector& triangleList, + const std::vector& nodes, + const std::vector& lightsList, + const std::vector& emissiveTriIndices, + const PBRCamera& camera, + int samplesPerPixel, + int sampleOffset + ) override; + void createQueue(const std::string &name, flib::vendor v = flib::vendor::INTEL, + flib::device dev = flib::device::GPU, + flib::backend b = flib::backend::LEVEL_ZERO); + sycl::queue &getQueue(); + IDeviceTexture& textures() override { + if (!m_textures) { + throw std::runtime_error("SYCL texture manager requested before queue initialization"); + } + return *m_textures; + } ~SYCL_Renderer() { if (m_textureHandles) { sycl::free(m_textureHandles, m_queue); @@ -82,4 +96,4 @@ class SYCL_Renderer : public IComputeRenderer { } }; -#endif // SYCL_RENDERER_HPP \ No newline at end of file +#endif // SYCL_RENDERER_HPP diff --git a/PBR/Render/src/cuda_renderer.cu b/PBR/Render/src/cuda_renderer.cu index 380bc4e..52fc292 100644 --- a/PBR/Render/src/cuda_renderer.cu +++ b/PBR/Render/src/cuda_renderer.cu @@ -202,6 +202,8 @@ std::vector CUDA_Renderer::RenderScene( int samplesPerPixel, int sampleOffset ) { + prepareTextures(); + std::vector framebuffer; const int imageSize = width * height; framebuffer.resize(imageSize); @@ -291,4 +293,4 @@ std::vector CUDA_Renderer::RenderScene( return framebuffer; -} \ No newline at end of file +} diff --git a/PBR/Render/src/sycl_renderer.cpp b/PBR/Render/src/sycl_renderer.cpp index 219358f..004ccd2 100644 --- a/PBR/Render/src/sycl_renderer.cpp +++ b/PBR/Render/src/sycl_renderer.cpp @@ -140,6 +140,8 @@ std::vector SYCL_Renderer::RenderScene( int sampleOffset ) { + prepareTextures(); + int imageSize = width * height; std::vector framebuffer(imageSize); std::cout << "SYCL_Renderer: Rendering " << width << "x" << height @@ -270,6 +272,7 @@ void SYCL_Renderer::createQueue(const std::string& name, flib::vendor vendor , flib::sycl_handler::get_device_info(); // Prints current device info m_queue = flib::sycl_handler::get_queue(name); + m_textures = std::make_unique(m_queue); } catch (const std::exception& e) diff --git a/PBR/Space/space.cpp b/PBR/Space/space.cpp index 645773a..4c24bdc 100644 --- a/PBR/Space/space.cpp +++ b/PBR/Space/space.cpp @@ -3,11 +3,9 @@ // Conditional includes - only include backend headers when enabled #ifdef FUNGT_USE_CUDA #include "PBR/Render/include/cuda_renderer.hpp" -#include "PBR/TextureManager/cuda_texture.hpp" #endif #ifdef FUNGT_USE_SYCL #include "PBR/Render/include/sycl_renderer.hpp" -#include "PBR/TextureManager/sycl_texture.hpp" #endif #define STB_IMAGE_WRITE_IMPLEMENTATION #include "../../vendor/stb_image/stb_image_write.h" @@ -28,7 +26,6 @@ Space::Space(){ /* code */ std::cout << "Using CPU to render scene" << std::endl; m_computeRenderer = std::make_unique(); - m_textureManager = std::make_shared(); break; } #ifdef FUNGT_USE_CUDA @@ -97,57 +94,6 @@ std::vector Space::Render(const int width, const int height,int sam return frameBuffer; } -void Space::sendTexturesToRender() -{ - if (!m_computeRenderer) { - std::cout << "Error in sendTexturesToRender : computeRenderer pointer is null" <(m_computeRenderer.get()); - if (cudaRenderer && m_textureManager) { - auto cudaTexMgr = dynamic_cast(m_textureManager.get()); - if (cudaTexMgr) - cudaRenderer->setCudaTextureObjects(cudaTexMgr->getTextureObjects()); - } - break; - } -#endif - -#ifdef FUNGT_USE_SYCL - case Compute::Backend::SYCL_CUDA: - case Compute::Backend::SYCL: - { - std::cout << "Sending SYCL Textures" << std::endl; - SYCL_Renderer* syclRenderer = dynamic_cast(m_computeRenderer.get()); - if (syclRenderer && m_textureManager) { - // Set textures on renderer - auto syclTexMgr = dynamic_cast(m_textureManager.get()); - if (syclTexMgr) - syclRenderer->setSyclTextureHandles(syclTexMgr->getImageHandles()); - } - break; - } -#endif - - default: - throw std::runtime_error("Unknown Compute API!"); - } - -} - void Space::InitComputeRenderBackend() { if (!m_computeRenderer) { @@ -159,15 +105,12 @@ void Space::InitComputeRenderBackend() { case Compute::Backend::CPU: std::cout << "Initializing CPU backend" << std::endl; - m_textureManager = std::make_shared(); break; #ifdef FUNGT_USE_CUDA case Compute::Backend::CUDA: { std::cout << "Initializing CUDA backend" << std::endl; - m_textureManager = std::make_shared(); - break; } #endif @@ -179,7 +122,6 @@ void Space::InitComputeRenderBackend() SYCL_Renderer* syclRenderer = dynamic_cast(m_computeRenderer.get()); if (syclRenderer) { syclRenderer->createQueue("intel_queue", flib::vendor::INTEL, flib::device::GPU, flib::backend::LEVEL_ZERO); - m_textureManager = std::make_shared(syclRenderer->getQueue()); } break; } @@ -189,7 +131,6 @@ void Space::InitComputeRenderBackend() SYCL_Renderer* syclRenderer = dynamic_cast(m_computeRenderer.get()); if (syclRenderer) { syclRenderer->createQueue("nvidia_queue", flib::vendor::NVIDIA, flib::device::GPU, flib::backend::CUDA); - m_textureManager = std::make_shared(syclRenderer->getQueue()); } break; } @@ -238,10 +179,10 @@ void Space::LoadModelToRender(const SimpleModel& Simplemodel) global_material.reflectance = 0.05f; global_material.emission = 0.0f; global_material.baseColorTexIdx = -1; - if (!textures.empty() && m_textureManager != nullptr) { + if (!textures.empty()) { std::string texPath = textures[0].getPath(); std::cout << " Loading texture: " << texPath << std::endl; - global_material.baseColorTexIdx = m_textureManager->loadTexture(texPath); + global_material.baseColorTexIdx = m_computeRenderer->textures().loadTexture(texPath); std::cout << " Assigned texture index: " << global_material.baseColorTexIdx << std::endl; } else { @@ -304,8 +245,6 @@ void Space::LoadModelToRender(const SimpleModel& Simplemodel) std::cout << " Metallic: " << global_material.metallic << std::endl; } - - sendTexturesToRender(); } void Space::LoadGeometryToRender(const SimpleGeometry& geometry) { auto primitive = geometry.getPrimitive(); @@ -324,10 +263,10 @@ void Space::LoadGeometryToRender(const SimpleGeometry& geometry) { std::cout << "Loading geometry with " << indices.size() / 3 << " triangles" << std::endl; int baseColorTexId = -1; - if (geometry.isTexturized() && m_textureManager != nullptr) { + if (geometry.isTexturized()) { std::string texPath = constPrimitive->texture.getPath(); std::cout << " Loading texture: " << texPath << std::endl; - baseColorTexId = m_textureManager->loadTexture(texPath); + baseColorTexId = m_computeRenderer->textures().loadTexture(texPath); } else { std::cout << " No texture (using base color)" << std::endl; @@ -375,7 +314,6 @@ void Space::LoadGeometryToRender(const SimpleGeometry& geometry) { } std::cout << "Geometry loaded. Total triangles in scene: " << m_triangles.size() << std::endl; - sendTexturesToRender(); } void Space::loadLightsFromScene(const std::vector& sceneLights) { @@ -474,5 +412,3 @@ void Space::setSamples(int numOfSamples) { m_samplesPerPixel = numOfSamples; } - - diff --git a/PBR/Space/space.hpp b/PBR/Space/space.hpp index 2e48ff5..0410a84 100644 --- a/PBR/Space/space.hpp +++ b/PBR/Space/space.hpp @@ -15,9 +15,6 @@ #include "PBR/Render/include/icompute_renderer.hpp" #include "PBR/Render/include/cpu_renderer.hpp" -#include "PBR/TextureManager/idevice_texture.hpp" - -#include "PBR/TextureManager/cpu_texture.hpp" #include "PBR/BVH/bvh_builder.hpp" #include "SimpleGeometry/simple_geometry.hpp" #include @@ -29,13 +26,10 @@ class Space { std::unique_ptr m_computeRenderer; std::vector m_lights; int m_samplesPerPixel = 16; - std::shared_ptr m_textureManager; std::vector m_bvh_nodes; std::vector m_bvh_indices; std::vector m_emissiveTriIndices; - void sendTexturesToRender(); - public: Space(); Space(std::vector& triangleList); @@ -70,11 +64,6 @@ class Space { void LoadTrianglesToRender(const std::vector& triangles) { m_triangles = triangles; } - std::shared_ptr getTextureManager() const { - return m_textureManager; - } - - }; diff --git a/PBR/TextureManager/cuda_texture.cpp b/PBR/TextureManager/cuda_texture.cpp index 57ded6d..c2c3610 100644 --- a/PBR/TextureManager/cuda_texture.cpp +++ b/PBR/TextureManager/cuda_texture.cpp @@ -84,6 +84,7 @@ int CUDATexture::loadTexture(const std::string& path) { int idx = textures.size(); textures.push_back(cudaTex); pathToIndex[path] = idx; + m_handlesDirty = true; std::cout << " textures size : " << textures.size() << std::endl; std::cout << " [CUDA] Texture index: " << idx << std::endl; return idx; @@ -135,6 +136,6 @@ void CUDATexture::cleanup() { textures.clear(); pathToIndex.clear(); + m_handlesDirty = true; std::cout << " [CUDA] Cleanup complete!" << std::endl; } - diff --git a/PBR/TextureManager/cuda_texture.hpp b/PBR/TextureManager/cuda_texture.hpp index 6ad4714..d14788a 100644 --- a/PBR/TextureManager/cuda_texture.hpp +++ b/PBR/TextureManager/cuda_texture.hpp @@ -20,6 +20,7 @@ class CUDATexture : public IDeviceTexture { private: std::vector textures; std::map pathToIndex; // Cache: path -> index + bool m_handlesDirty = false; public: CUDATexture(); @@ -29,6 +30,8 @@ class CUDATexture : public IDeviceTexture { int getTextureCount() const override; void cleanup() override; std::vector getTextureObjects(); + bool handlesAreDirty() const { return m_handlesDirty; } + void markHandlesClean() { m_handlesDirty = false; } }; diff --git a/PBR/TextureManager/ocl_texture.hpp b/PBR/TextureManager/ocl_texture.hpp index 684b833..d6d9f9e 100644 --- a/PBR/TextureManager/ocl_texture.hpp +++ b/PBR/TextureManager/ocl_texture.hpp @@ -1,27 +1,48 @@ #if !defined(_OPENCL_TEXTURE_HPP_) #define _OPENCL_TEXTURE_HPP_ -#include -#include -#include #include "idevice_texture.hpp" +#include +#include +#include +#include + #define CL_MEM_BINDLESS_IMAGE_INTEL 0x10060 #define CL_IMAGE_BINDLESS_HANDLE_INTEL 0x10061 -class OpenCLTexture -{ - uint64_t m_handleId; - cl_image_format m_imageFormat{}; - m_image_format.image_channel_order = CL_RGBA; - m_image_format.image_channel_data_type = CL_UNORM_INT8; - - public: - OCLTexture(); - - - +struct OpenCLTextureData { + cl_mem image; + uint64_t bindlessHandle; // only meaningful if isBindless + bool isBindless; + int width, height; + std::string path; +}; +class OpenCLTexture : public IDeviceTexture { +private: + std::vector textures; + std::map pathToIndex; + + cl_context context; + cl_platform_id platform; + bool useBindless; // decided once at construction, e.g. from driver capability check + + typedef cl_mem(CL_API_CALL* clCreateImageWithPropertiesINTEL_fn)( + cl_context, const cl_bitfield*, cl_mem_flags, + const cl_image_format*, const cl_image_desc*, void*, cl_int*); + clCreateImageWithPropertiesINTEL_fn createBindlessImage = nullptr; + +public: + OpenCLTexture(cl_context context, cl_platform_id platform, bool preferBindless); + ~OpenCLTexture(); + + int loadTexture(const std::string& path) override; + int getTextureCount() const override; + void cleanup() override; + std::vector getTextureObjects(); // analogous to CUDA's getTextureObjects + bool isBindlessMode() const { return useBindless; } }; +#endif diff --git a/PBR/TextureManager/sycl_texture.cpp b/PBR/TextureManager/sycl_texture.cpp index 385c2a2..39e0e0f 100644 --- a/PBR/TextureManager/sycl_texture.cpp +++ b/PBR/TextureManager/sycl_texture.cpp @@ -65,6 +65,7 @@ int index = textures.size(); textures.push_back(std::move(texData)); pathToIndex[path] = index; + m_handlesDirty = true; stbi_image_free(data); @@ -120,4 +121,5 @@ textures.clear(); pathToIndex.clear(); + m_handlesDirty = true; } diff --git a/PBR/TextureManager/sycl_texture.hpp b/PBR/TextureManager/sycl_texture.hpp index bfb65aa..41d68d2 100644 --- a/PBR/TextureManager/sycl_texture.hpp +++ b/PBR/TextureManager/sycl_texture.hpp @@ -23,14 +23,17 @@ class SYCLTexture : public IDeviceTexture { std::vector textures; std::map pathToIndex; sycl::queue* m_queue; + bool m_handlesDirty = false; public: SYCLTexture(sycl::queue& queue); ~SYCLTexture(); int loadTexture(const std::string& path) override; - int getTextureCount() const override{}; + int getTextureCount() const override { return static_cast(textures.size()); } void cleanup() override; + bool handlesAreDirty() const { return m_handlesDirty; } + void markHandlesClean() { m_handlesDirty = false; } // MATCHING CUDA PATTERN - return host-side handles! std::vector getImageHandles() { @@ -43,4 +46,4 @@ class SYCLTexture : public IDeviceTexture { } }; -#endif // _SYCL_TEXTURE_HPP_ \ No newline at end of file +#endif // _SYCL_TEXTURE_HPP_ diff --git a/Samples/hdr_test/main.cpp b/Samples/hdr_test/main.cpp index a70dc88..14567ac 100644 --- a/Samples/hdr_test/main.cpp +++ b/Samples/hdr_test/main.cpp @@ -13,26 +13,23 @@ int main() { FunGTSceneManager scene_manager = myGame->getSceneManager(); - // FunGTCubeMap skybox = CubeMap::create(); - // skybox->setShaders(getAssetPath("resources/equirect_skybox.vs"), - // getAssetPath("resources/equirect_skybox.fs")); - // skybox->buildHDR(getAssetPath("assets_local/hdri/san_giuseppe_bridge_4k.hdr")); + FunGTCubeMap skybox = CubeMap::create(); + skybox->setShaders(getAssetPath("resources/equirect_skybox.vs"), + getAssetPath("resources/equirect_skybox.fs")); + skybox->buildHDR(getAssetPath("assets_local/hdri/small_empty_room_3_4k.hdr")); - // scene_manager->loadEnvironment(getAssetPath("assets_local/hdri/san_giuseppe_bridge_4k.hdr")); - // scene_manager->setIBLIntensity(1.3f); + scene_manager->loadEnvironment(getAssetPath("assets_local/hdri/small_empty_room_3_4k.hdr")); + scene_manager->setIBLIntensity(1.3f); ModelPaths nightStreet; - nightStreet.path = getAssetPath("assets_local/Obj/street/source/Wagon/#Ulica.obj"); + nightStreet.path = getAssetPath("assets_local/scene_final/scene.gltf"); FunGTSModel street = SimpleModel::create(); street->load(nightStreet); - street->position(0.f, 0.f, -40.f); - street->rotation(0.f, 0.f, 0.f); - street->scale(1.0); myGame->set([&]() { - //scene_manager->addRenderableObj(skybox); + scene_manager->addRenderableObj(skybox); scene_manager->addRenderableObj(street); }); diff --git a/Samples/iwocl/CMakeLists.txt b/Samples/iwocl/CMakeLists.txt index c71d202..104b542 100644 --- a/Samples/iwocl/CMakeLists.txt +++ b/Samples/iwocl/CMakeLists.txt @@ -223,6 +223,8 @@ include_directories( ${FUNGT_BASE_DIR}/GraphicsRenderBackend ${FUNGT_BASE_DIR}/GraphicsRenderBackend/opengl ${FUNGT_BASE_DIR}/MeshGPU + ${FUNGT_BASE_DIR}/TextureGPU + ${FUNGT_BASE_DIR}/PrimitiveGPU ) # ════════════════════════════════════════════════════════════════════════════ @@ -246,6 +248,8 @@ set(SOURCE_FILES ${FUNGT_BASE_DIR}/Material/material.cpp ${FUNGT_BASE_DIR}/Mesh/mesh.cpp ${FUNGT_BASE_DIR}/MeshGPU/mesh_gpu.cpp + ${FUNGT_BASE_DIR}/TextureGPU/texture_gpu.cpp + ${FUNGT_BASE_DIR}/PrimitiveGPU/primitive_gpu.cpp ${FUNGT_BASE_DIR}/Camera/camera.cpp ${FUNGT_BASE_DIR}/Geometries/cube.cpp ${FUNGT_BASE_DIR}/Geometries/plane.cpp From bf8910fc7b56392afd23f6a0441823fe5efa9c60 Mon Sep 17 00:00:00 2001 From: juanchuletas Date: Tue, 18 Aug 2026 22:44:54 -0600 Subject: [PATCH 3/7] Adding OpenCL backend : defining textures objects for bindless images support --- PBR/Render/include/compute_backends.hpp | 1 + PBR/Render/include/icompute_renderer.hpp | 8 ++ PBR/Render/include/opencl_renderer.hpp | 68 +++++++++ PBR/Render/src/compute_backends.cpp | 1 + PBR/Render/src/opencl_renderer.cpp | 171 +++++++++++++++++++++++ PBR/Space/space.cpp | 25 ++++ PBR/TextureManager/ocl_texture.hpp | 49 ------- PBR/TextureManager/opencl_texture.cpp | 124 ++++++++++++++++ PBR/TextureManager/opencl_texture.hpp | 62 ++++++++ PBR/main/CMakeLists.txt | 132 +++++++++-------- PBR/main/main.cpp | 59 -------- PBR/standalone/CMakeLists.txt | 167 ++++++++++++++++++++++ PBR/standalone/main.cpp | 56 ++++++++ 13 files changed, 759 insertions(+), 164 deletions(-) create mode 100644 PBR/Render/include/opencl_renderer.hpp create mode 100644 PBR/Render/src/opencl_renderer.cpp delete mode 100644 PBR/TextureManager/ocl_texture.hpp create mode 100644 PBR/TextureManager/opencl_texture.cpp create mode 100644 PBR/TextureManager/opencl_texture.hpp delete mode 100644 PBR/main/main.cpp create mode 100644 PBR/standalone/CMakeLists.txt create mode 100644 PBR/standalone/main.cpp diff --git a/PBR/Render/include/compute_backends.hpp b/PBR/Render/include/compute_backends.hpp index 742e5a7..239dd11 100644 --- a/PBR/Render/include/compute_backends.hpp +++ b/PBR/Render/include/compute_backends.hpp @@ -8,6 +8,7 @@ namespace Compute{ CUDA, SYCL, SYCL_CUDA, // SYCL targeting CUDA devices + OPENCL, }; } class ComputeRender { diff --git a/PBR/Render/include/icompute_renderer.hpp b/PBR/Render/include/icompute_renderer.hpp index 9f396f9..87ab267 100644 --- a/PBR/Render/include/icompute_renderer.hpp +++ b/PBR/Render/include/icompute_renderer.hpp @@ -11,9 +11,17 @@ class PBRCamera; class IComputeRenderer{ + protected: + IComputeRenderer() = default; public: virtual ~IComputeRenderer() = default; + //Avoid copy and move semantics for this interface + IComputeRenderer(const IComputeRenderer&) = delete; + IComputeRenderer& operator=(const IComputeRenderer&) = delete; + IComputeRenderer(IComputeRenderer&&) = delete; + IComputeRenderer& operator=(IComputeRenderer&&) = delete; + // virtual IDeviceTexture& textures() = 0; virtual std::vector RenderScene( int width, int height, diff --git a/PBR/Render/include/opencl_renderer.hpp b/PBR/Render/include/opencl_renderer.hpp new file mode 100644 index 0000000..98ed1c6 --- /dev/null +++ b/PBR/Render/include/opencl_renderer.hpp @@ -0,0 +1,68 @@ +#if !defined(_OPENCL_RENDERER_H_) +#define _OPENCL_RENDERER_H_ + +#include // MUST be first! +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include "icompute_renderer.hpp" +#include "PBR/TextureManager/opencl_texture.hpp" + +class OpenCL_Renderer : public IComputeRenderer { + + cl_platform_id m_oclplatform = nullptr; + cl_device_id m_ocldevice = nullptr; + cl_context m_oclcontext = nullptr; + cl_command_queue m_oclqueue = nullptr; + cl_program m_oclprogram = nullptr; + cl_kernel m_oclrenderKernel = nullptr; + cl_mem m_texturesObj = nullptr; + std::size_t m_numTextures = 0; + std::unique_ptr m_textureManager; + + void prepareTextures(); + + public: + OpenCL_Renderer()= default; + ~OpenCL_Renderer() override; + + void initialize(); + + IDeviceTexture& textures() override { + if (!m_textureManager) { + throw std::runtime_error( + "OpenCL texture manager requested before OpenCL initialization."); + } + return *m_textureManager; + } + + std::vector RenderScene( + int width, + int height, + const std::vector& triangles, + const std::vector& nodes, + const std::vector& lights, + const std::vector& emissiveTriIndices, + const PBRCamera& camera, + int samplesPerPixel, + int sampleOffset + ) override; + + void setOpenCLTextures(const std::vector& handles); + + +}; + + +#endif // _OPENCL_RENDERER_H_ diff --git a/PBR/Render/src/compute_backends.cpp b/PBR/Render/src/compute_backends.cpp index 441f98d..53fbe6a 100644 --- a/PBR/Render/src/compute_backends.cpp +++ b/PBR/Render/src/compute_backends.cpp @@ -8,6 +8,7 @@ const std::string ComputeRender::GetBackendName() { case Compute::Backend::CUDA: return "CUDA"; case Compute::Backend::SYCL: return "SYCL"; case Compute::Backend::SYCL_CUDA: return "SYCL_CUDA"; + case Compute::Backend::OPENCL: return "OPENCL"; case Compute::Backend::CPU: return "CPU"; default: return "Unknown"; } diff --git a/PBR/Render/src/opencl_renderer.cpp b/PBR/Render/src/opencl_renderer.cpp new file mode 100644 index 0000000..f6c9d08 --- /dev/null +++ b/PBR/Render/src/opencl_renderer.cpp @@ -0,0 +1,171 @@ +#include "opencl_renderer.hpp" + +void OpenCL_Renderer::initialize() +{ + if (m_oclcontext && m_oclqueue && m_textureManager) { + return; + } + + cl_uint platformCount = 0; + cl_int err = clGetPlatformIDs(0, nullptr, &platformCount); + if (err != CL_SUCCESS || platformCount == 0) { + throw std::runtime_error( + "OpenCL initialization failed: no OpenCL platforms available, error " + + std::to_string(err)); + } + + std::vector platforms(platformCount); + err = clGetPlatformIDs(platformCount, platforms.data(), nullptr); + if (err != CL_SUCCESS) { + throw std::runtime_error( + "OpenCL initialization failed while enumerating platforms, error " + + std::to_string(err)); + } + + for (cl_platform_id platform : platforms) { + if (!clGetExtensionFunctionAddressForPlatform( + platform, "clCreateImageWithPropertiesINTEL")) { + continue; + } + + cl_device_id device = nullptr; + err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, nullptr); + if (err == CL_SUCCESS && device) { + m_oclplatform = platform; + m_ocldevice = device; + break; + } + } + + if (!m_oclplatform || !m_ocldevice) { + throw std::runtime_error( + "OpenCL initialization failed: no GPU platform exposes the bindless image API."); + } + + char platformName[256] = {}; + char deviceName[256] = {}; + clGetPlatformInfo( + m_oclplatform, CL_PLATFORM_NAME, + sizeof(platformName), platformName, nullptr); + clGetDeviceInfo( + m_ocldevice, CL_DEVICE_NAME, + sizeof(deviceName), deviceName, nullptr); + std::cout << "OpenCL platform: " << platformName << std::endl; + std::cout << "OpenCL device: " << deviceName << std::endl; + + m_oclcontext = clCreateContext( + nullptr, 1, &m_ocldevice, nullptr, nullptr, &err); + if (err != CL_SUCCESS || !m_oclcontext) { + m_oclcontext = nullptr; + throw std::runtime_error( + "OpenCL initialization failed while creating the context, error " + + std::to_string(err)); + } + + m_oclqueue = clCreateCommandQueueWithProperties( + m_oclcontext, m_ocldevice, nullptr, &err); + if (err != CL_SUCCESS || !m_oclqueue) { + m_oclqueue = nullptr; + clReleaseContext(m_oclcontext); + m_oclcontext = nullptr; + throw std::runtime_error( + "OpenCL initialization failed while creating the command queue, error " + + std::to_string(err)); + } + + m_textureManager = std::make_unique( + m_oclcontext, m_oclplatform, true); + + std::cout << "OpenCL context and command queue initialized." << std::endl; +} + +OpenCL_Renderer::~OpenCL_Renderer() +{ + if (m_oclqueue) { + clFinish(m_oclqueue); + } + + m_textureManager.reset(); + + if (m_texturesObj) { + clReleaseMemObject(m_texturesObj); + m_texturesObj = nullptr; + } + m_numTextures = 0; + + if (m_oclrenderKernel) { + clReleaseKernel(m_oclrenderKernel); + m_oclrenderKernel = nullptr; + } + + if (m_oclprogram) { + clReleaseProgram(m_oclprogram); + m_oclprogram = nullptr; + } + + if (m_oclqueue) { + clReleaseCommandQueue(m_oclqueue); + m_oclqueue = nullptr; + } + + if (m_oclcontext) { + clReleaseContext(m_oclcontext); + m_oclcontext = nullptr; + } + + m_ocldevice = nullptr; + m_oclplatform = nullptr; +} + +std::vector OpenCL_Renderer::RenderScene(int width, int height, const std::vector& triangles, const std::vector& nodes, const std::vector& lights, const std::vector& emissiveTriIndices, const PBRCamera& camera, int samplesPerPixel, int sampleOffset) +{ + prepareTextures(); + return std::vector(); +} + +void OpenCL_Renderer::prepareTextures() +{ + if (!m_textureManager) { + throw std::runtime_error( + "OpenCL textures requested before OpenCL initialization."); + } + + if (m_textureManager->handlesAreDirty()) { + setOpenCLTextures(m_textureManager->getBindlessHandles()); + m_textureManager->markHandlesClean(); + } +} + +void OpenCL_Renderer::setOpenCLTextures(const std::vector& handles) +{ + if (!m_oclcontext) { + throw std::runtime_error( + "setOpenCLTextures: OpenCL context is not initialized."); + } + + std::cout << "*** SETTING OPENCL TEXTURE OBJECTS*** " << std::endl; + if (m_texturesObj) { + clReleaseMemObject(m_texturesObj); + m_texturesObj = nullptr; + } + + m_numTextures = handles.size(); + std::cout << "*** NUM OPENCL TEXTURE OBJECTS *** " << m_numTextures << std::endl; + + if (m_numTextures == 0) { + return; + } + + cl_int err = CL_SUCCESS; + m_texturesObj = clCreateBuffer( + m_oclcontext, + CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + m_numTextures * sizeof(uint64_t), + const_cast(handles.data()), + &err); + if (err != CL_SUCCESS || !m_texturesObj) { + throw std::runtime_error("setOpenCLTextures: clCreateBuffer failed for bindless handles, error " + std::to_string(err)); + } + + std::cout << " Uploaded " << m_numTextures << " bindless handles to GPU" << std::endl; +} diff --git a/PBR/Space/space.cpp b/PBR/Space/space.cpp index 4c24bdc..a906cc6 100644 --- a/PBR/Space/space.cpp +++ b/PBR/Space/space.cpp @@ -7,6 +7,9 @@ #ifdef FUNGT_USE_SYCL #include "PBR/Render/include/sycl_renderer.hpp" #endif +#ifdef FUNGT_USE_OPENCL +#include "PBR/Render/include/opencl_renderer.hpp" +#endif #define STB_IMAGE_WRITE_IMPLEMENTATION #include "../../vendor/stb_image/stb_image_write.h" Space::Space(){ @@ -19,6 +22,7 @@ Space::Space(){ std::cout << " CUDA = " << static_cast(Compute::Backend::CUDA) << std::endl; std::cout << " SYCL = " << static_cast(Compute::Backend::SYCL) << std::endl; std::cout << " CPU = " << static_cast(Compute::Backend::CPU) << std::endl; + std::cout << " OPENCL = " << static_cast(Compute::Backend::OPENCL) << std::endl; switch (ComputeRender::GetBackend()) { case Compute::Backend::CPU: @@ -48,6 +52,14 @@ Space::Space(){ break; } #endif +#ifdef FUNGT_USE_OPENCL + case Compute::Backend::OPENCL: + { + std::cout << "Using OpenCL to render scene" << std::endl; + m_computeRenderer = std::make_unique(); + break; + } +#endif default: throw std::runtime_error("Unknown Compute API!"); } @@ -136,6 +148,19 @@ void Space::InitComputeRenderBackend() } #endif +#ifdef FUNGT_USE_OPENCL + case Compute::Backend::OPENCL: + { + std::cout << "Initializing OpenCL backend" << std::endl; + auto* openclRenderer = dynamic_cast(m_computeRenderer.get()); + if (!openclRenderer) { + throw std::runtime_error("OpenCL renderer was not created."); + } + openclRenderer->initialize(); + break; + } +#endif + default: throw std::runtime_error("Unknown Compute API!"); } diff --git a/PBR/TextureManager/ocl_texture.hpp b/PBR/TextureManager/ocl_texture.hpp deleted file mode 100644 index d6d9f9e..0000000 --- a/PBR/TextureManager/ocl_texture.hpp +++ /dev/null @@ -1,49 +0,0 @@ -#if !defined(_OPENCL_TEXTURE_HPP_) -#define _OPENCL_TEXTURE_HPP_ -#include "idevice_texture.hpp" -#include -#include -#include -#include - -#define CL_MEM_BINDLESS_IMAGE_INTEL 0x10060 -#define CL_IMAGE_BINDLESS_HANDLE_INTEL 0x10061 - -struct OpenCLTextureData { - cl_mem image; - uint64_t bindlessHandle; // only meaningful if isBindless - bool isBindless; - int width, height; - std::string path; -}; - -class OpenCLTexture : public IDeviceTexture { -private: - std::vector textures; - std::map pathToIndex; - - cl_context context; - cl_platform_id platform; - bool useBindless; // decided once at construction, e.g. from driver capability check - - typedef cl_mem(CL_API_CALL* clCreateImageWithPropertiesINTEL_fn)( - cl_context, const cl_bitfield*, cl_mem_flags, - const cl_image_format*, const cl_image_desc*, void*, cl_int*); - clCreateImageWithPropertiesINTEL_fn createBindlessImage = nullptr; - -public: - OpenCLTexture(cl_context context, cl_platform_id platform, bool preferBindless); - ~OpenCLTexture(); - - int loadTexture(const std::string& path) override; - int getTextureCount() const override; - void cleanup() override; - std::vector getTextureObjects(); // analogous to CUDA's getTextureObjects - bool isBindlessMode() const { return useBindless; } -}; - -#endif - - - -#endif // _OPENCL_TEXTURE_HPP_ diff --git a/PBR/TextureManager/opencl_texture.cpp b/PBR/TextureManager/opencl_texture.cpp new file mode 100644 index 0000000..88c2fda --- /dev/null +++ b/PBR/TextureManager/opencl_texture.cpp @@ -0,0 +1,124 @@ +#include "opencl_texture.hpp" +#include "stb_image.h" + +OpenCLTexture::OpenCLTexture(cl_context context, cl_platform_id platform, bool useBindless) +: m_context{context}, m_useBindless{useBindless} +{ + if (m_useBindless) { + createBindlessImage = reinterpret_cast( + clGetExtensionFunctionAddressForPlatform(platform, "clCreateImageWithPropertiesINTEL")); + if (!createBindlessImage) { + throw std::runtime_error( + "OpenCLTexture: bindless image API is not available on the selected platform."); + } + } + std::cout << "OpenCLTexture initialized (" << (m_useBindless ? "bindless" : "bound") << ")" << std::endl; +} + +OpenCLTexture::~OpenCLTexture() +{ + cleanup(); +} + +int OpenCLTexture::loadTexture(const std::string& path) +{ + auto it = m_pathToIndex.find(path); + if (it != m_pathToIndex.end()) { + std::cout << " [OpenCL] Texture cached: " << path << " (index " << it->second << ")" << std::endl; + return it->second; + } + + int width, height, channels; + unsigned char* data = stbi_load(path.c_str(), &width, &height, &channels, 4); + if (!data) { + std::cerr << " [OpenCL] Failed to load: " << path << std::endl; + std::cerr << " [OpenCL] Error: " << stbi_failure_reason() << std::endl; + return -1; + } + + std::cout << " [OpenCL] Loaded: " << path << " (" << width << "x" << height + << ", " << channels << " channels)" << std::endl; + + OpenCLTextureData ocltex{}; + ocltex.width = width; + ocltex.height = height; + ocltex.path = path; + ocltex.isBindless = m_useBindless; + + cl_image_format format{}; + format.image_channel_order = CL_sRGBA; + format.image_channel_data_type = CL_UNORM_INT8; + + cl_image_desc desc{}; + desc.image_type = CL_MEM_OBJECT_IMAGE2D; + desc.image_width = static_cast(width); + desc.image_height = static_cast(height); + + cl_int err = CL_SUCCESS; + + if (m_useBindless) { + const cl_bitfield props[] = { + CL_MEM_BINDLESS_IMAGE_INTEL, 1, 0 + }; + ocltex.image = createBindlessImage(m_context, props, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + &format, &desc, data, &err); + } + else { + ocltex.image = clCreateImage(m_context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + &format, &desc, data, &err); + } + + stbi_image_free(data); + + if (err != CL_SUCCESS || !ocltex.image) { + std::cerr << " [OpenCL] Failed to create image: " << path << " (err " << err << ")" << std::endl; + return -1; + } + + if (m_useBindless) { + err = clGetImageInfo(ocltex.image, CL_IMAGE_BINDLESS_HANDLE_INTEL, + sizeof(ocltex.bindlessHandle), &ocltex.bindlessHandle, nullptr); + if (err != CL_SUCCESS) { + std::cerr << " [OpenCL] Failed to get bindless image handle: " + << path << " (err " << err << ")" << std::endl; + clReleaseMemObject(ocltex.image); + return -1; + } + } + + int idx = static_cast(m_textures.size()); + m_textures.push_back(ocltex); + if (m_useBindless) { + m_bindlessHandles.push_back(ocltex.bindlessHandle); + m_handlesDirty = true; + } + m_pathToIndex[path] = idx; + std::cout << " textures size : " << m_textures.size() << std::endl; + std::cout << " [OpenCL] Texture index: " << idx << std::endl; + return idx; +} + +std::vector OpenCLTexture::getTextureObjects() const { + std::vector objs; + objs.reserve(m_textures.size()); + for (const auto& tex : m_textures) { + objs.push_back(tex.image); + } + return objs; +} + +void OpenCLTexture::cleanup() +{ + for (auto& texture : m_textures) { + if (texture.image) { + clReleaseMemObject(texture.image); + texture.image = nullptr; + } + texture.bindlessHandle = 0; + } + + m_textures.clear(); + m_bindlessHandles.clear(); + m_pathToIndex.clear(); + m_handlesDirty = true; +} diff --git a/PBR/TextureManager/opencl_texture.hpp b/PBR/TextureManager/opencl_texture.hpp new file mode 100644 index 0000000..e52ca55 --- /dev/null +++ b/PBR/TextureManager/opencl_texture.hpp @@ -0,0 +1,62 @@ +#if !defined(_OPENCL_TEXTURE_HPP_) +#define _OPENCL_TEXTURE_HPP_ +#include "idevice_texture.hpp" +#include +#include +#include +#include +#include +#include +#include +#include + +#ifndef CL_MEM_BINDLESS_IMAGE_INTEL +#define CL_MEM_BINDLESS_IMAGE_INTEL 0x10060 +#endif +#ifndef CL_IMAGE_BINDLESS_HANDLE_INTEL +#define CL_IMAGE_BINDLESS_HANDLE_INTEL 0x10061 +#endif + +struct OpenCLTextureData { + cl_mem image; + uint64_t bindlessHandle; // only meaningful if uses bindless images + bool isBindless; + int width, height; + std::string path; +}; + +class OpenCLTexture : public IDeviceTexture { +private: + std::vector m_textures; + std::vector m_bindlessHandles; + std::map m_pathToIndex; + + cl_context m_context = nullptr; + bool m_useBindless = false; + bool m_handlesDirty = false; + typedef cl_mem(CL_API_CALL* clCreateImageWithPropertiesINTEL_fn)( + cl_context, const cl_bitfield*, cl_mem_flags, + const cl_image_format*, const cl_image_desc*, void*, cl_int*); + clCreateImageWithPropertiesINTEL_fn createBindlessImage = nullptr; + +public: + OpenCLTexture(cl_context context, cl_platform_id platform, bool useBindless = true); + ~OpenCLTexture() override; + + OpenCLTexture(const OpenCLTexture&) = delete; + OpenCLTexture& operator=(const OpenCLTexture&) = delete; + OpenCLTexture(OpenCLTexture&&) = delete; + OpenCLTexture& operator=(OpenCLTexture&&) = delete; + + int loadTexture(const std::string& path) override; + int getTextureCount() const override { return static_cast(m_textures.size() ); }; + void cleanup() override; + std::vector getTextureObjects() const; + const std::vector& getBindlessHandles() const { return m_bindlessHandles; } + bool isBindlessMode() const { return m_useBindless; } + bool handlesAreDirty() const { return m_handlesDirty; } + void markHandlesClean() { m_handlesDirty = false; } +}; + + +#endif // _OPENCL_TEXTURE_HPP_ diff --git a/PBR/main/CMakeLists.txt b/PBR/main/CMakeLists.txt index 19e2aab..e356a75 100644 --- a/PBR/main/CMakeLists.txt +++ b/PBR/main/CMakeLists.txt @@ -1,8 +1,9 @@ cmake_minimum_required(VERSION 3.15) -project(pbr_render) +project(pbr_libraries) option(FUNGT_USE_CUDA "Enable CUDA backend" ON) option(FUNGT_USE_SYCL "Enable SYCL backend" ON) +option(FUNGT_USE_OPENCL "Enable OpenCL backend" ON) set(SYCL_CUDA_ARCH sm_75 CACHE STRING "NVIDIA GPU architecture for SYCL CUDA backend") set_property(CACHE SYCL_CUDA_ARCH PROPERTY STRINGS sm_50 sm_52 sm_53 sm_60 sm_61 sm_62 sm_70 sm_72 sm_75 sm_80 sm_86 sm_87 sm_89 sm_90) @@ -34,15 +35,13 @@ endif() # Build configuration # ──────────────────────────────────────────────────────── set(CMAKE_BUILD_TYPE Release) -set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/release/linux) - set(CMAKE_CXX_STANDARD 17) set(CMAKE_CXX_STANDARD_REQUIRED ON) # ════════════════════════════════════════════════════════════════════════════ # SYCL TOOLCHAIN DETECTION # ════════════════════════════════════════════════════════════════════════════ -if(UNIX) +if(FUNGT_USE_SYCL AND UNIX) if(DEFINED SYCL_COMPILER_PATH) if(NOT EXISTS "${SYCL_COMPILER_PATH}") message(FATAL_ERROR "SYCL_COMPILER_PATH does not exist: ${SYCL_COMPILER_PATH}") @@ -64,7 +63,7 @@ if(UNIX) " 2. Provide your own: cmake -DSYCL_COMPILER_PATH=/path/to/clang++ ..\n") endif() endif() -elseif(WIN32) +elseif(FUNGT_USE_SYCL AND WIN32) if(DEFINED SYCL_COMPILER_PATH) if(NOT EXISTS "${SYCL_COMPILER_PATH}") message(FATAL_ERROR "SYCL_COMPILER_PATH does not exist: ${SYCL_COMPILER_PATH}") @@ -142,6 +141,7 @@ set(FUNGT_INCLUDES # CORE PBR LIBRARY (no GPU-specific code) # ──────────────────────────────────────────────────────── set(PBR_CORE_FILES + ${FUNGT_BASE_DIR}/PBR/Space/space.cpp ${FUNGT_BASE_DIR}/PBR/Intersection/intersection.cpp ${FUNGT_BASE_DIR}/PBR/BVH/bvh_builder.cpp ${FUNGT_BASE_DIR}/PBR/Render/src/compute_backends.cpp @@ -213,6 +213,8 @@ if(FUNGT_USE_CUDA) ) target_compile_definitions(pbr_cuda PUBLIC FUNGT_USE_CUDA) + target_compile_definitions(pbr_core PUBLIC FUNGT_USE_CUDA) + target_include_directories(pbr_core PRIVATE ${CUDAToolkit_INCLUDE_DIRS}) message(STATUS "CUDA library configured") message(STATUS " Files: cuda_renderer.cu, cuda_texture.cpp") @@ -257,6 +259,8 @@ if(FUNGT_USE_SYCL) target_link_options(pbr_sycl PRIVATE -fsycl -fsycl-targets=${SYCL_TARGETS} ${SYCL_CUDA_ARCH_ARGS}) target_compile_definitions(pbr_sycl PUBLIC FUNGT_USE_SYCL) + target_compile_definitions(pbr_core PUBLIC FUNGT_USE_SYCL) + target_compile_options(pbr_core PRIVATE -fsycl) message(STATUS "SYCL library configured") message(STATUS " Files: sycl_renderer.cpp, sycl_texture.cpp") @@ -269,7 +273,29 @@ endif() # ──────────────────────────────────────────────────────── # Find required packages # ──────────────────────────────────────────────────────── -if(FUNGT_USE_SYCL) +if(FUNGT_USE_OPENCL) + set(FUNGT_OPENCL_HEADERS + "$ENV{HOME}/Documents/Development/neo-dev/compute-runtime/third_party/opencl_headers" + CACHE PATH "Path containing the custom CL headers") + set(FUNGT_OPENCL_LOADER + "/usr/lib/x86_64-linux-gnu/libOpenCL.so.1" + CACHE FILEPATH "OpenCL ICD loader used by the standalone PBR test") + + if(NOT EXISTS "${FUNGT_OPENCL_HEADERS}/CL/cl.h") + message(FATAL_ERROR + "Custom OpenCL headers not found at: ${FUNGT_OPENCL_HEADERS}") + endif() + if(NOT EXISTS "${FUNGT_OPENCL_LOADER}") + message(FATAL_ERROR + "OpenCL ICD loader not found at: ${FUNGT_OPENCL_LOADER}") + endif() + + set(OpenCL_INCLUDE_DIRS "${FUNGT_OPENCL_HEADERS}") + set(OpenCL_LIBRARIES "${FUNGT_OPENCL_LOADER}") + message(STATUS "OpenCL backend uses custom compute-runtime headers:") + message(STATUS " Include: ${OpenCL_INCLUDE_DIRS}") + message(STATUS " Loader: ${OpenCL_LIBRARIES}") +elseif(FUNGT_USE_SYCL) set(BUNDLED_OPENCL_INCLUDE "${BUNDLED_SYCL_DIR}/include") find_library(BUNDLED_OPENCL_LIBRARY NAMES OpenCL @@ -292,13 +318,41 @@ if(FUNGT_USE_SYCL) message(FATAL_ERROR "OpenCL library not found!") endif() endif() -else() - find_package(OpenCL REQUIRED) - if(OpenCL_FOUND) - message(STATUS "OpenCL found: ${OpenCL_INCLUDE_DIRS}") - else() - message(FATAL_ERROR "OpenCL library not found!") - endif() +endif() + +# ──────────────────────────────────────────────────────── +# OPENCL BACKEND LIBRARY +# ──────────────────────────────────────────────────────── +if(FUNGT_USE_OPENCL) + set(OPENCL_FILES + ${FUNGT_BASE_DIR}/PBR/Render/src/opencl_renderer.cpp + ${FUNGT_BASE_DIR}/PBR/TextureManager/opencl_texture.cpp + ) + + add_library(pbr_opencl STATIC ${OPENCL_FILES}) + + target_include_directories(pbr_opencl BEFORE PUBLIC + ${FUNGT_OPENCL_HEADERS} + ${FUNGT_INCLUDES} + ${FUNLIB_DIR}/include + ) + + target_compile_definitions(pbr_opencl PUBLIC + FUNGT_USE_OPENCL + CL_TARGET_OPENCL_VERSION=300 + ) + target_compile_definitions(pbr_core PUBLIC + FUNGT_USE_OPENCL + CL_TARGET_OPENCL_VERSION=300 + ) + target_include_directories(pbr_core BEFORE PRIVATE + ${FUNGT_OPENCL_HEADERS} + ) + + target_link_libraries(pbr_opencl PUBLIC ${FUNGT_OPENCL_LOADER}) + + message(STATUS "OpenCL library configured") + message(STATUS " Files: opencl_renderer.cpp, opencl_texture.cpp") endif() find_package(assimp REQUIRED) @@ -315,52 +369,12 @@ else() message(FATAL_ERROR "OpenGL library not found!") endif() -# ──────────────────────────────────────────────────────── -# FINAL EXECUTABLE -# ──────────────────────────────────────────────────────── -add_executable(pbr_render - ${CMAKE_CURRENT_SOURCE_DIR}/main.cpp - ${FUNGT_BASE_DIR}/PBR/Space/space.cpp -) - -if(FUNGT_USE_SYCL) - target_compile_options(pbr_render PRIVATE -fsycl) - target_link_options(pbr_render PRIVATE -fsycl) - target_compile_definitions(pbr_render PRIVATE FUNGT_USE_SYCL) -endif() - -if(FUNGT_USE_CUDA) - target_compile_definitions(pbr_render PRIVATE FUNGT_USE_CUDA) - # THE FIX: explicitly give bundled clang++ the CUDA toolkit include path - target_include_directories(pbr_render PRIVATE ${CUDAToolkit_INCLUDE_DIRS}) - set_target_properties(pbr_render PROPERTIES - CUDA_ARCHITECTURES "native" - ) - -endif() - -target_include_directories(pbr_render PRIVATE - ${FUNGT_INCLUDES} - ${OpenCL_INCLUDE_DIRS} -) - -target_link_libraries(pbr_render PRIVATE - pbr_core - $<$:pbr_cuda> - $<$:pbr_sycl> - OpenGL::GL - assimp::assimp - funlib - dl - ${OpenCL_LIBRARIES} -) - # ──────────────────────────────────────────────────────── # Build summary # ──────────────────────────────────────────────────────── message(STATUS "") message(STATUS "══════════════════════════════════════") -message(STATUS " RaySpace Render Build Configuration") +message(STATUS " PBR Library Build Configuration") message(STATUS "══════════════════════════════════════") message(STATUS "Build type: ${CMAKE_BUILD_TYPE}") message(STATUS "FunGT base: ${FUNGT_BASE_DIR}") @@ -377,6 +391,12 @@ if(FUNGT_USE_SYCL) else() message(STATUS "SYCL backend: DISABLED") endif() -message(STATUS "Output directory: ${CMAKE_RUNTIME_OUTPUT_DIRECTORY}") +if(FUNGT_USE_OPENCL) + message(STATUS "OpenCL backend: ENABLED") + message(STATUS "OpenCL headers: ${FUNGT_OPENCL_HEADERS}") + message(STATUS "OpenCL loader: ${FUNGT_OPENCL_LOADER}") +else() + message(STATUS "OpenCL backend: DISABLED") +endif() message(STATUS "══════════════════════════════════════") -message(STATUS "") \ No newline at end of file +message(STATUS "") diff --git a/PBR/main/main.cpp b/PBR/main/main.cpp deleted file mode 100644 index 86e7afb..0000000 --- a/PBR/main/main.cpp +++ /dev/null @@ -1,59 +0,0 @@ -#include "../Space/space.hpp" -const int IMAGE_WIDTH = 800; -const int IMAGE_HEIGHT = 400; - -int main(){ - - //Path to your shaders and models: - ModelPaths monkey; - DisplayGraphics::SetBackend(Backend::OpenGL); - monkey.path = getAssetPath("Obj/Woody/woody-head.obj"); - monkey.vs_path = getAssetPath("resources/luxoball_vs.glsl"); - monkey.fs_path = getAssetPath("resources/luxoball_fs.glsl"); - //ComputeRender::Init(); - - //CUDA backend: - ComputeRender::SetBackend(Compute::Backend::CUDA); - std::cout << "Backend in use: " << ComputeRender::GetBackendName() << std::endl; - //Model loading: - std::shared_ptr monkey_model = SimpleModel::create(); - monkey_model->LoadModelData(monkey); - - - - //std::shared_ptr> txtMgr = std::make_shared(); - // std::vector triangleList = monkey_model->getTriangleList(); - // std::cout<<"Total triangles : "<(totalEnd - totalStart).count(); - // ============ PRINT RESULTS ============ - std::cout << "\n========== TIMING RESULTS ==========" << std::endl; - std::cout << "Total time: " << totalTime << " ms" << std::endl; - std::cout << "====================================\n" << std::endl; - - return 0; -} \ No newline at end of file diff --git a/PBR/standalone/CMakeLists.txt b/PBR/standalone/CMakeLists.txt new file mode 100644 index 0000000..7e4bdb3 --- /dev/null +++ b/PBR/standalone/CMakeLists.txt @@ -0,0 +1,167 @@ +cmake_minimum_required(VERSION 3.15) + +get_filename_component( + FUNGT_DEFAULT_BASE_DIR + "${CMAKE_CURRENT_LIST_DIR}/../.." + ABSOLUTE +) + +set(FUNGT_BASE_DIR + "${FUNGT_DEFAULT_BASE_DIR}" + CACHE PATH "FunGT base directory" +) + +option(FUNGT_USE_CUDA "Link the CUDA PBR backend" ON) +option(FUNGT_USE_SYCL "Link the SYCL PBR backend" ON) +option(FUNGT_USE_OPENCL "Link the OpenCL PBR backend" ON) +set(SYCL_CUDA_ARCH sm_75 CACHE STRING "NVIDIA GPU architecture for SYCL") + +set(SYCL_COMPILER_PATH + "/home/juanchuletas/Documents/Development/sycl_workspace/llvm/build/bin/clang++" + CACHE FILEPATH "SYCL-capable C++ compiler") + +if(FUNGT_USE_SYCL) + if(NOT EXISTS "${SYCL_COMPILER_PATH}") + message(FATAL_ERROR + "SYCL compiler not found at ${SYCL_COMPILER_PATH}") + endif() + set(CMAKE_CXX_COMPILER + "${SYCL_COMPILER_PATH}" + CACHE FILEPATH "C++ compiler" FORCE) +endif() + +project(pbr_standalone LANGUAGES CXX) + +set(CMAKE_CXX_STANDARD 17) +set(CMAKE_CXX_STANDARD_REQUIRED ON) +set(CMAKE_RUNTIME_OUTPUT_DIRECTORY "${CMAKE_CURRENT_SOURCE_DIR}/release/linux") + +set(PBR_LIBRARY_DIR + "${FUNGT_BASE_DIR}/PBR/main/build" + CACHE PATH "Directory containing the prebuilt PBR static libraries" +) +set(FUNGT_OPENCL_HEADERS + "$ENV{HOME}/Documents/Development/neo-dev/compute-runtime/third_party/opencl_headers" + CACHE PATH "Path containing the custom CL headers" +) +set(FUNGT_OPENCL_LOADER + "/usr/lib/x86_64-linux-gnu/libOpenCL.so.1" + CACHE FILEPATH "OpenCL ICD loader" +) + +function(import_pbr_library target archive) + if(NOT EXISTS "${PBR_LIBRARY_DIR}/${archive}") + message(FATAL_ERROR + "Missing ${archive}. Build the PBR libraries first in ${PBR_LIBRARY_DIR}.") + endif() + + add_library(${target} STATIC IMPORTED) + set_target_properties(${target} PROPERTIES + IMPORTED_LOCATION "${PBR_LIBRARY_DIR}/${archive}" + ) +endfunction() + +import_pbr_library(pbr_core libpbr_core.a) + +if(FUNGT_USE_CUDA) + enable_language(CUDA) + import_pbr_library(pbr_cuda libpbr_cuda.a) + find_package(CUDAToolkit REQUIRED) +endif() + +if(FUNGT_USE_SYCL) + import_pbr_library(pbr_sycl libpbr_sycl.a) +endif() + +if(FUNGT_USE_OPENCL) + import_pbr_library(pbr_opencl libpbr_opencl.a) + if(NOT EXISTS "${FUNGT_OPENCL_HEADERS}/CL/cl.h") + message(FATAL_ERROR + "Custom OpenCL headers not found at ${FUNGT_OPENCL_HEADERS}") + endif() + if(NOT EXISTS "${FUNGT_OPENCL_LOADER}") + message(FATAL_ERROR + "OpenCL ICD loader not found at ${FUNGT_OPENCL_LOADER}") + endif() +endif() + +set(FUNLIB_DIR "${FUNGT_BASE_DIR}/vendor/funlib") +add_library(funlib STATIC IMPORTED) +set_target_properties(funlib PROPERTIES + IMPORTED_LOCATION "${FUNLIB_DIR}/lib/libfunlib.a" + INTERFACE_INCLUDE_DIRECTORIES "${FUNLIB_DIR}/include" +) + +find_package(assimp REQUIRED) +find_package(OpenGL REQUIRED) + +add_executable(pbr_standalone main.cpp) + +if(FUNGT_USE_CUDA) + set_target_properties(pbr_standalone PROPERTIES + CUDA_ARCHITECTURES "native" + CUDA_SEPARABLE_COMPILATION ON + CUDA_RESOLVE_DEVICE_SYMBOLS ON + ) +endif() + +target_include_directories(pbr_standalone BEFORE PRIVATE + ${FUNGT_OPENCL_HEADERS} + ${FUNGT_BASE_DIR} + ${FUNGT_BASE_DIR}/PBR/Render/include + ${FUNGT_BASE_DIR}/PBR/TextureManager + ${FUNGT_BASE_DIR}/vendor/stb_image + ${FUNGT_BASE_DIR}/vendor/glm/include + ${FUNLIB_DIR}/include +) + +target_compile_definitions(pbr_standalone PRIVATE + $<$:FUNGT_USE_CUDA> + $<$:FUNGT_USE_SYCL> + $<$:FUNGT_USE_OPENCL> + $<$:CL_TARGET_OPENCL_VERSION=300> +) + +if(FUNGT_USE_CUDA) + target_include_directories(pbr_standalone PRIVATE + ${CUDAToolkit_INCLUDE_DIRS} + ) +endif() + +if(FUNGT_USE_SYCL) + target_compile_options(pbr_standalone PRIVATE -fsycl) + if(FUNGT_USE_CUDA) + set(SYCL_TARGETS nvptx64-nvidia-cuda,spir64) + set(SYCL_CUDA_ARCH_ARGS + "-Xsycl-target-backend=nvptx64-nvidia-cuda" + "--offload-arch=${SYCL_CUDA_ARCH}" + ) + else() + set(SYCL_TARGETS spir64) + set(SYCL_CUDA_ARCH_ARGS "") + endif() + target_link_options(pbr_standalone PRIVATE + -fsycl + -fsycl-targets=${SYCL_TARGETS} + ${SYCL_CUDA_ARCH_ARGS} + ) +endif() + +target_link_libraries(pbr_standalone PRIVATE + pbr_core + $<$:pbr_cuda> + $<$:pbr_sycl> + $<$:pbr_opencl> + OpenGL::GL + assimp::assimp + funlib + dl + $<$:${FUNGT_OPENCL_LOADER}> + $<$:CUDA::cudart_static> + $<$:CUDA::cuda_driver> +) + +message(STATUS "PBR libraries: ${PBR_LIBRARY_DIR}") +message(STATUS "CUDA backend: ${FUNGT_USE_CUDA}") +message(STATUS "SYCL backend: ${FUNGT_USE_SYCL}") +message(STATUS "OpenCL backend: ${FUNGT_USE_OPENCL}") diff --git a/PBR/standalone/main.cpp b/PBR/standalone/main.cpp new file mode 100644 index 0000000..31986f8 --- /dev/null +++ b/PBR/standalone/main.cpp @@ -0,0 +1,56 @@ +#include "PBR/Space/space.hpp" +const int IMAGE_WIDTH = 800; +const int IMAGE_HEIGHT = 400; + +int main() +{ + // ComputeRender::SetBackend(Compute::Backend::OPENCL); + // std::cout << "Initializing standalone PBR backend: " + // << ComputeRender::GetBackendName() << std::endl; + + // Space space; + // space.InitComputeRenderBackend(); + + // std::cout << "OpenCL backend initialized successfully." << std::endl; + + // return 0; + ModelPaths monkey; + DisplayGraphics::SetBackend(Backend::OpenGL); + monkey.path = getAssetPath("assets_local/Obj/Woody/woody-head.obj"); + //monkey.vs_path = getAssetPath("resources/luxoball_vs.glsl"); + //monkey.fs_path = getAssetPath("resources/luxoball_fs.glsl"); + + ComputeRender::SetBackend(Compute::Backend::CUDA); + std::cout << "Backend in use: " << ComputeRender::GetBackendName() << std::endl; + + std::shared_ptr monkey_model = SimpleModel::create(); + monkey_model->LoadModelData(monkey); + + PBRCamera camera( + fungt::Vec3(0, 2.5, 30), + fungt::Vec3(0, 1.8, 0), + fungt::Vec3(0, 1, 0), + 50.0f, + float(IMAGE_WIDTH) / float(IMAGE_HEIGHT) + ); + + Space space(camera); + space.InitComputeRenderBackend(); + space.LoadModelToRender(*monkey_model); + space.setSamples(100); + space.BuildBVH(); + + auto totalStart = std::chrono::high_resolution_clock::now(); + auto framebuffer = space.Render(IMAGE_WIDTH, IMAGE_HEIGHT); + auto totalEnd = std::chrono::high_resolution_clock::now(); + + Space::SaveFrameBufferAsPNG(framebuffer, IMAGE_WIDTH, IMAGE_HEIGHT); + + auto totalTime = std::chrono::duration_cast( + totalEnd - totalStart).count(); + std::cout << "\n========== TIMING RESULTS ==========" << std::endl; + std::cout << "Total time: " << totalTime << " ms" << std::endl; + std::cout << "====================================\n" << std::endl; + + return 0; +} From d276bbce6de9050386f6d6de8531f726bd915a8c Mon Sep 17 00:00:00 2001 From: juanchuletas Date: Thu, 20 Aug 2026 19:37:32 -0600 Subject: [PATCH 4/7] feature: OpenCL backend with bindless images --- PBR/PBRCamera/pbr_camera.hpp | 18 +- PBR/Render/include/opencl_renderer.hpp | 50 +- PBR/Render/ocl_kernels/bvh.cl | 121 ++++ .../ocl_kernels/evaluate_cook_torrance.cl | 121 ++++ PBR/Render/ocl_kernels/intersection.cl | 89 +++ PBR/Render/ocl_kernels/path_sampling.cl | 89 +++ PBR/Render/ocl_kernels/path_tracer.cl | 272 +++++++++ PBR/Render/ocl_kernels/rayspace_camera.cl | 26 + PBR/Render/ocl_kernels/rng.cl | 34 ++ PBR/Render/ocl_kernels/sample_framebuffer.cl | 11 + PBR/Render/ocl_kernels/surface_hit.cl | 77 +++ PBR/Render/ocl_kernels/texture_sampling.cl | 20 + PBR/Render/shared/opencl/fgt_opencl_data.h | 140 +++++ .../opencl/fgt_opencl_data_translator.hpp | 207 +++++++ PBR/Render/src/opencl_renderer.cpp | 564 +++++++++++++++++- PBR/TextureManager/opencl_texture.cpp | 14 +- PBR/TextureManager/opencl_texture.hpp | 2 + PBR/standalone/main.cpp | 101 +++- 18 files changed, 1919 insertions(+), 37 deletions(-) create mode 100644 PBR/Render/ocl_kernels/bvh.cl create mode 100644 PBR/Render/ocl_kernels/evaluate_cook_torrance.cl create mode 100644 PBR/Render/ocl_kernels/intersection.cl create mode 100644 PBR/Render/ocl_kernels/path_sampling.cl create mode 100644 PBR/Render/ocl_kernels/path_tracer.cl create mode 100644 PBR/Render/ocl_kernels/rayspace_camera.cl create mode 100644 PBR/Render/ocl_kernels/rng.cl create mode 100644 PBR/Render/ocl_kernels/sample_framebuffer.cl create mode 100644 PBR/Render/ocl_kernels/surface_hit.cl create mode 100644 PBR/Render/ocl_kernels/texture_sampling.cl create mode 100644 PBR/Render/shared/opencl/fgt_opencl_data.h create mode 100644 PBR/Render/shared/opencl/fgt_opencl_data_translator.hpp diff --git a/PBR/PBRCamera/pbr_camera.hpp b/PBR/PBRCamera/pbr_camera.hpp index 323842f..80e99b2 100644 --- a/PBR/PBRCamera/pbr_camera.hpp +++ b/PBR/PBRCamera/pbr_camera.hpp @@ -20,6 +20,9 @@ class PBRCamera { lowerLeftCorner = fungt::Vec3(-2, -1.5, -1); horizontal = fungt::Vec3(4, 0, 0); vertical = fungt::Vec3(0, 3, 0); + u = fungt::Vec3(1, 0, 0); + v = fungt::Vec3(0, 1, 0); + w = fungt::Vec3(0, 0, 1); lensRadius = 0; } @@ -84,8 +87,21 @@ class PBRCamera { horizontal = fungt::Vec3(viewportWidth, 0, 0); vertical = fungt::Vec3(0, viewportHeight, 0); lowerLeftCorner = origin - horizontal / 2 - vertical / 2 - fungt::Vec3(0, 0, focalLength); + u = fungt::Vec3(1, 0, 0); + v = fungt::Vec3(0, 1, 0); + w = fungt::Vec3(0, 0, 1); + lensRadius = 0.0f; } + fgt_device const fungt::Vec3& getOrigin() const { return origin; } + fgt_device const fungt::Vec3& getLowerLeftCorner() const { return lowerLeftCorner; } + fgt_device const fungt::Vec3& getHorizontal() const { return horizontal; } + fgt_device const fungt::Vec3& getVertical() const { return vertical; } + fgt_device const fungt::Vec3& getBasisU() const { return u; } + fgt_device const fungt::Vec3& getBasisV() const { return v; } + fgt_device const fungt::Vec3& getBasisW() const { return w; } + fgt_device float getLensRadius() const { return lensRadius; } + // Get ray (with optional depth of field) fgt_device fungt::Ray getRay(float u, float v) const { fungt::Vec3 rd = randomInUnitDisk() * lensRadius; @@ -114,4 +130,4 @@ namespace sycl { } #endif -#endif // _PBR_CAMERA_H_ \ No newline at end of file +#endif // _PBR_CAMERA_H_ diff --git a/PBR/Render/include/opencl_renderer.hpp b/PBR/Render/include/opencl_renderer.hpp index 98ed1c6..766d13d 100644 --- a/PBR/Render/include/opencl_renderer.hpp +++ b/PBR/Render/include/opencl_renderer.hpp @@ -15,10 +15,42 @@ #include #include #include - +#include +#include +#include +#include "Path_Manager/path_manager.hpp" #include "icompute_renderer.hpp" #include "PBR/TextureManager/opencl_texture.hpp" +const std::vector OpenCLRendererSourceFiles = { + "PBR/Render/shared/opencl/fgt_opencl_data.h", + "PBR/Render/ocl_kernels/rng.cl", + "PBR/Render/ocl_kernels/rayspace_camera.cl", + "PBR/Render/ocl_kernels/intersection.cl", + "PBR/Render/ocl_kernels/bvh.cl", + "PBR/Render/ocl_kernels/surface_hit.cl", + "PBR/Render/ocl_kernels/texture_sampling.cl", + "PBR/Render/ocl_kernels/evaluate_cook_torrance.cl", + "PBR/Render/ocl_kernels/path_sampling.cl", + "PBR/Render/ocl_kernels/path_tracer.cl", + "PBR/Render/ocl_kernels/sample_framebuffer.cl" +}; + +struct OpenCLRaySpaceBuffer { + cl_mem triangleGeometry = nullptr; + cl_mem triangleShading = nullptr; + cl_mem materials = nullptr; + cl_mem bvhNodes = nullptr; + cl_mem lights = nullptr; + cl_mem emissiveTriangles = nullptr; + + std::size_t numTriangles = 0; + std::size_t numMaterials = 0; + std::size_t numBVHNodes = 0; + std::size_t numLights = 0; + std::size_t numEmissiveTriangles = 0; +}; + class OpenCL_Renderer : public IComputeRenderer { cl_platform_id m_oclplatform = nullptr; @@ -27,11 +59,24 @@ class OpenCL_Renderer : public IComputeRenderer { cl_command_queue m_oclqueue = nullptr; cl_program m_oclprogram = nullptr; cl_kernel m_oclrenderKernel = nullptr; + cl_kernel m_oclsampleKernel = nullptr; cl_mem m_texturesObj = nullptr; + cl_mem m_textureDimensions = nullptr; std::size_t m_numTextures = 0; std::unique_ptr m_textureManager; + OpenCLRaySpaceBuffer m_raySpaceBuffer; + bool m_sceneUploaded = false; void prepareTextures(); + void setOpenCLTextures( + const std::vector& handles, + const std::vector>& dimensions); + void uploadScene( + const std::vector& triangles, + const std::vector& nodes, + const std::vector& lights, + const std::vector& emissiveTriIndices); + void releaseRaySpaceBuffer(OpenCLRaySpaceBuffer& buffers) noexcept; public: OpenCL_Renderer()= default; @@ -59,8 +104,7 @@ class OpenCL_Renderer : public IComputeRenderer { int sampleOffset ) override; - void setOpenCLTextures(const std::vector& handles); - + void buildOCLPrograms(); }; diff --git a/PBR/Render/ocl_kernels/bvh.cl b/PBR/Render/ocl_kernels/bvh.cl new file mode 100644 index 0000000..6f86200 --- /dev/null +++ b/PBR/Render/ocl_kernels/bvh.cl @@ -0,0 +1,121 @@ +#define FGT_BVH_STACK_SIZE 32 +#define FGT_FLOAT_MAX 3.402823466e+38f + +inline bool fgt_trace_closest_bvh( + fgt_ray ray, + __global const fgt_triangle_geom* triangle_geometry, + __global const fgt_bvh_node* bvh_nodes, + int num_bvh_nodes, + __private fgt_geometry_hit* closest_hit) +{ + if (num_bvh_nodes <= 0) { + return false; + } + + int stack[FGT_BVH_STACK_SIZE]; + int stack_size = 0; + stack[stack_size++] = 0; + + bool found_hit = false; + float closest_distance = FGT_FLOAT_MAX; + + while (stack_size > 0) { + const int node_index = stack[--stack_size]; + if (node_index < 0 || node_index >= num_bvh_nodes) { + continue; + } + + const fgt_bvh_node node = bvh_nodes[node_index]; + if (!fgt_intersect_aabb( + ray, + node.bounds, + 0.001f, + closest_distance)) { + continue; + } + + if (node.triangle_count > 0) { + for (int offset = 0; offset < node.triangle_count; ++offset) { + const int triangle_index = + node.first_triangle_index + offset; + fgt_geometry_hit candidate; + + if (fgt_intersect_triangle( + ray, + &triangle_geometry[triangle_index], + 0.001f, + closest_distance, + &candidate)) { + found_hit = true; + closest_distance = candidate.distance; + candidate.triangle_index = triangle_index; + *closest_hit = candidate; + } + } + } + else { + stack[stack_size++] = node.left_child; + stack[stack_size++] = node.right_child; + } + } + + return found_hit; +} + +inline bool fgt_trace_shadow_bvh( + fgt_ray ray, + __global const fgt_triangle_geom* triangle_geometry, + __global const fgt_bvh_node* bvh_nodes, + int num_bvh_nodes, + float maximum_distance) +{ + if (num_bvh_nodes <= 0) { + return false; + } + + int stack[FGT_BVH_STACK_SIZE]; + int stack_size = 0; + stack[stack_size++] = 0; + + while (stack_size > 0) { + const int node_index = stack[--stack_size]; + if (node_index < 0 || node_index >= num_bvh_nodes) { + continue; + } + + const fgt_bvh_node node = bvh_nodes[node_index]; + if (!fgt_intersect_aabb( + ray, + node.bounds, + 0.001f, + maximum_distance)) { + continue; + } + + if (node.triangle_count > 0) { + for (int offset = 0; offset < node.triangle_count; ++offset) { + const int triangle_index = + node.first_triangle_index + offset; + fgt_geometry_hit candidate; + + if (fgt_intersect_triangle( + ray, + &triangle_geometry[triangle_index], + 0.001f, + maximum_distance, + &candidate)) { + return true; + } + } + } + else { + stack[stack_size++] = node.left_child; + stack[stack_size++] = node.right_child; + } + } + + return false; +} + +#undef FGT_BVH_STACK_SIZE +#undef FGT_FLOAT_MAX diff --git a/PBR/Render/ocl_kernels/evaluate_cook_torrance.cl b/PBR/Render/ocl_kernels/evaluate_cook_torrance.cl new file mode 100644 index 0000000..d3a8308 --- /dev/null +++ b/PBR/Render/ocl_kernels/evaluate_cook_torrance.cl @@ -0,0 +1,121 @@ +#define FGT_PI_F 3.14159265358979323846f + +inline float fgt_distribution_ggx(float normal_dot_half, float roughness) +{ + const float alpha = roughness * roughness; + const float alpha_squared = alpha * alpha; + const float normal_dot_half_squared = + normal_dot_half * normal_dot_half; + float denominator = + normal_dot_half_squared * (alpha_squared - 1.0f) + 1.0f; + denominator = fmax(denominator, 1.0e-6f); + + return alpha_squared / + (FGT_PI_F * denominator * denominator + 1.0e-8f); +} + +inline float fgt_geometry_schlick_ggx( + float normal_dot_direction, + float k) +{ + return normal_dot_direction / + (normal_dot_direction * (1.0f - k) + k + 1.0e-8f); +} + +inline float fgt_geometry_smith( + float normal_dot_view, + float normal_dot_light, + float roughness) +{ + const float remapped_roughness = roughness + 1.0f; + const float k = + (remapped_roughness * remapped_roughness) / 8.0f; + + return + fgt_geometry_schlick_ggx(normal_dot_view, k) * + fgt_geometry_schlick_ggx(normal_dot_light, k); +} + +inline float3 fgt_fresnel_schlick(float3 f0, float view_dot_half) +{ + const float value = 1.0f - view_dot_half; + const float value_squared = value * value; + const float value_fifth = + value_squared * value_squared * value; + + return f0 + ((float3)(1.0f) - f0) * value_fifth; +} + +inline void fgt_material_to_pbr( + fgt_material_data material, + __private float3* base_color, + __private float* metallic, + __private float* roughness, + __private float3* f0) +{ + *base_color = (float3)( + material.base_color[0], + material.base_color[1], + material.base_color[2]); + *metallic = clamp(material.metallic, 0.0f, 1.0f); + *roughness = clamp(material.roughness, 0.05f, 1.0f); + + const float3 dielectric_f0 = (float3)(material.reflectance); + *f0 = mix(dielectric_f0, *base_color, *metallic); +} + +inline float3 fgt_evaluate_cook_torrance( + float3 normal, + float3 view_direction, + float3 light_direction, + fgt_material_data material, + float3 light_radiance) +{ + float3 base_color; + float metallic; + float roughness; + float3 f0; + fgt_material_to_pbr( + material, + &base_color, + &metallic, + &roughness, + &f0); + + const float3 half_direction = + normalize(view_direction + light_direction); + const float normal_dot_light = + fmax(dot(normal, light_direction), 0.0f); + const float normal_dot_view = + fmax(dot(normal, view_direction), 0.0f); + const float normal_dot_half = + fmax(dot(normal, half_direction), 0.0f); + const float view_dot_half = + fmax(dot(view_direction, half_direction), 0.0f); + + if (normal_dot_light <= 0.0f || normal_dot_view <= 0.0f) { + return (float3)(0.0f); + } + + const float distribution = fmin( + fgt_distribution_ggx(normal_dot_half, roughness), + 100.0f); + const float geometry = fgt_geometry_smith( + normal_dot_view, + normal_dot_light, + roughness); + const float3 fresnel = fgt_fresnel_schlick(f0, view_dot_half); + + const float3 numerator = fresnel * (distribution * geometry); + const float denominator = + 4.0f * normal_dot_view * normal_dot_light + 1.0e-6f; + const float3 specular = numerator / denominator; + + const float3 diffuse_weight = + ((float3)(1.0f) - fresnel) * (1.0f - metallic); + const float3 diffuse = base_color / FGT_PI_F; + + return + (diffuse_weight * diffuse + specular) * + light_radiance * normal_dot_light; +} diff --git a/PBR/Render/ocl_kernels/intersection.cl b/PBR/Render/ocl_kernels/intersection.cl new file mode 100644 index 0000000..85ceae9 --- /dev/null +++ b/PBR/Render/ocl_kernels/intersection.cl @@ -0,0 +1,89 @@ +typedef struct fgt_geometry_hit { + float distance; + float3 point; + float3 barycentric; + float3 geometric_normal; + int triangle_index; +} fgt_geometry_hit; + +inline bool fgt_intersect_triangle( + fgt_ray ray, + __global const fgt_triangle_geom* triangle, + float minimum_distance, + float maximum_distance, + __private fgt_geometry_hit* hit) +{ + const float epsilon = 1.0e-8f; + const float3 vertex0 = fgt_vec4_xyz(triangle->v0); + const float3 edge1 = fgt_vec4_xyz(triangle->edge1); + const float3 edge2 = fgt_vec4_xyz(triangle->edge2); + + const float3 direction_cross_edge2 = cross(ray.direction, edge2); + const float determinant = dot(edge1, direction_cross_edge2); + if (fabs(determinant) < epsilon) { + return false; + } + + const float inverse_determinant = 1.0f / determinant; + const float3 origin_from_vertex = ray.origin - vertex0; + const float barycentric_u = + inverse_determinant * dot(origin_from_vertex, direction_cross_edge2); + if (barycentric_u < 0.0f || barycentric_u > 1.0f) { + return false; + } + + const float3 origin_cross_edge1 = cross(origin_from_vertex, edge1); + const float barycentric_v = + inverse_determinant * dot(ray.direction, origin_cross_edge1); + if (barycentric_v < 0.0f || + barycentric_u + barycentric_v > 1.0f) { + return false; + } + + const float distance = + inverse_determinant * dot(edge2, origin_cross_edge1); + if (distance < minimum_distance || distance > maximum_distance) { + return false; + } + + hit->distance = distance; + hit->point = ray.origin + ray.direction * distance; + hit->barycentric = (float3)( + 1.0f - barycentric_u - barycentric_v, + barycentric_u, + barycentric_v); + hit->geometric_normal = normalize(cross(edge1, edge2)); + return true; +} + +inline bool fgt_intersect_aabb( + fgt_ray ray, + fgt_aabb bounds, + float minimum_distance, + float maximum_distance) +{ + const float3 bounds_min = fgt_vec4_xyz(bounds.bounds_min); + const float3 bounds_max = fgt_vec4_xyz(bounds.bounds_max); + + for (int axis = 0; axis < 3; ++axis) { + const float inverse_direction = 1.0f / ray.direction[axis]; + float near_distance = + (bounds_min[axis] - ray.origin[axis]) * inverse_direction; + float far_distance = + (bounds_max[axis] - ray.origin[axis]) * inverse_direction; + + if (inverse_direction < 0.0f) { + const float temporary = near_distance; + near_distance = far_distance; + far_distance = temporary; + } + + minimum_distance = fmax(minimum_distance, near_distance); + maximum_distance = fmin(maximum_distance, far_distance); + if (maximum_distance < minimum_distance) { + return false; + } + } + + return true; +} diff --git a/PBR/Render/ocl_kernels/path_sampling.cl b/PBR/Render/ocl_kernels/path_sampling.cl new file mode 100644 index 0000000..ce74202 --- /dev/null +++ b/PBR/Render/ocl_kernels/path_sampling.cl @@ -0,0 +1,89 @@ +inline float3 fgt_sample_hemisphere( + float3 normal, + __private fgt_rng* rng) +{ + const float u = fgt_rng_next_float(rng); + const float v = fgt_rng_next_float(rng); + const float z = sqrt(1.0f - u); + const float radius = sqrt(u); + const float phi = 2.0f * FGT_PI_F * v; + + const float x = radius * cos(phi); + const float y = radius * sin(phi); + const float3 helper = fabs(normal.x) > 0.1f + ? (float3)(0.0f, 1.0f, 0.0f) + : (float3)(1.0f, 0.0f, 0.0f); + const float3 tangent = normalize(cross(helper, normal)); + const float3 bitangent = cross(normal, tangent); + + return normalize(tangent * x + bitangent * y + normal * z); +} + +inline bool fgt_sample_emissive_triangle( + __global const fgt_triangle_geom* triangle_geometry, + __global const fgt_triangle_shading* triangle_shading, + __global const fgt_material_data* materials, + __global const fgt_int32* emissive_triangles, + int num_triangles, + int num_materials, + int num_emissive_triangles, + __private fgt_rng* rng, + __private float3* light_position, + __private float3* light_normal, + __private float3* light_emission, + __private float* light_pdf) +{ + if (num_emissive_triangles <= 0) { + return false; + } + + const uint random_index = + fgt_rng_next_u32(rng) % (uint)num_emissive_triangles; + const int triangle_index = emissive_triangles[random_index]; + if (triangle_index < 0 || triangle_index >= num_triangles) { + return false; + } + + const fgt_triangle_geom geometry = triangle_geometry[triangle_index]; + const fgt_triangle_shading shading = triangle_shading[triangle_index]; + if (shading.material_index < 0 || + shading.material_index >= num_materials) { + return false; + } + + float barycentric_y = fgt_rng_next_float(rng); + float barycentric_z = fgt_rng_next_float(rng); + if (barycentric_y + barycentric_z > 1.0f) { + barycentric_y = 1.0f - barycentric_y; + barycentric_z = 1.0f - barycentric_z; + } + const float barycentric_x = + 1.0f - barycentric_y - barycentric_z; + + const float3 v0 = (float3)( + geometry.v0.x, geometry.v0.y, geometry.v0.z); + const float3 edge1 = (float3)( + geometry.edge1.x, geometry.edge1.y, geometry.edge1.z); + const float3 edge2 = (float3)( + geometry.edge2.x, geometry.edge2.y, geometry.edge2.z); + const float3 normal0 = fgt_storage_vec3_to_float3(shading.n0); + const float3 normal1 = fgt_storage_vec3_to_float3(shading.n1); + const float3 normal2 = fgt_storage_vec3_to_float3(shading.n2); + const float area = 0.5f * length(cross(edge1, edge2)); + if (area < 1.0e-8f) { + return false; + } + + const fgt_material_data material = materials[shading.material_index]; + *light_position = v0 + edge1 * barycentric_y + edge2 * barycentric_z; + *light_normal = normalize( + normal0 * barycentric_x + + normal1 * barycentric_y + + normal2 * barycentric_z); + *light_emission = (float3)( + material.base_color[0], + material.base_color[1], + material.base_color[2]) * material.emission; + *light_pdf = 1.0f / ((float)num_emissive_triangles * area); + return true; +} diff --git a/PBR/Render/ocl_kernels/path_tracer.cl b/PBR/Render/ocl_kernels/path_tracer.cl new file mode 100644 index 0000000..e02faed --- /dev/null +++ b/PBR/Render/ocl_kernels/path_tracer.cl @@ -0,0 +1,272 @@ +inline float3 fgt_path_trace_cook_torrance( + fgt_ray ray, + __global const fgt_triangle_geom* triangle_geometry, + __global const fgt_triangle_shading* triangle_shading, + __global const fgt_material_data* materials, + __global const fgt_bvh_node* bvh_nodes, + __global const fgt_light* lights, + __global const fgt_int32* emissive_triangles, + __global const ulong* texture_handles, + __global const int2* texture_dimensions, + int num_triangles, + int num_materials, + int num_bvh_nodes, + int num_lights, + int num_emissive_triangles, + int num_textures, + __private fgt_rng* rng) +{ + float3 throughput = (float3)(1.0f); + float3 radiance = (float3)(0.0f); + + for (int bounce = 0; bounce < 6; ++bounce) { + fgt_geometry_hit geometry_hit; + if (!fgt_trace_closest_bvh( + ray, + triangle_geometry, + bvh_nodes, + num_bvh_nodes, + &geometry_hit)) { + radiance += throughput * (float3)(0.3f); + break; + } + + fgt_surface_hit hit; + if (!fgt_resolve_surface_hit( + geometry_hit, + triangle_shading, + materials, + num_triangles, + num_materials, + &hit)) { + break; + } + + const int texture_index = + hit.material.base_color_texture_index; + if (texture_index >= 0 && texture_index < num_textures) { + const float3 texture_color = fgt_sample_bindless_texture( + texture_handles, + texture_dimensions, + texture_index, + hit.uv); + hit.material.base_color[0] = texture_color.x; + hit.material.base_color[1] = texture_color.y; + hit.material.base_color[2] = texture_color.z; + } + + const float3 normal = normalize(hit.shading_normal); + const float3 view_direction = normalize(-ray.direction); + float3 base_color; + float metallic; + float roughness; + float3 f0; + fgt_material_to_pbr( + hit.material, + &base_color, + &metallic, + &roughness, + &f0); + + if (hit.material.emission > 0.0f) { + radiance += + throughput * base_color * hit.material.emission; + } + + float3 direct_light = (float3)(0.0f); + for (int light_index = 0; + light_index < num_lights; + ++light_index) { + const fgt_light light = lights[light_index]; + const float3 light_position = fgt_storage_vec3_to_float3( + light.position); + const float3 light_intensity = fgt_storage_vec3_to_float3( + light.intensity); + const float3 to_light = light_position - hit.point; + const float distance_squared = dot(to_light, to_light); + if (distance_squared <= 1.0e-12f) { + continue; + } + + const float distance = sqrt(distance_squared); + const float3 light_direction = to_light / distance; + fgt_ray shadow_ray; + shadow_ray.origin = + hit.point + hit.geometric_normal * 0.001f; + shadow_ray.direction = light_direction; + + if (fgt_trace_shadow_bvh( + shadow_ray, + triangle_geometry, + bvh_nodes, + num_bvh_nodes, + distance - 0.001f)) { + continue; + } + + const float3 light_radiance = + light_intensity / (distance_squared + 1.0e-6f); + direct_light += fgt_evaluate_cook_torrance( + normal, + view_direction, + light_direction, + hit.material, + light_radiance); + } + radiance += throughput * direct_light; + + // Sample emissive triangles only on early bounces. + if (num_emissive_triangles > 0 && bounce < 3) { + float3 light_position; + float3 light_normal; + float3 light_emission; + float light_pdf; + if (fgt_sample_emissive_triangle( + triangle_geometry, + triangle_shading, + materials, + emissive_triangles, + num_triangles, + num_materials, + num_emissive_triangles, + rng, + &light_position, + &light_normal, + &light_emission, + &light_pdf)) { + const float3 to_light = light_position - hit.point; + const float distance_squared = dot(to_light, to_light); + if (distance_squared > 1.0e-12f) { + const float distance = sqrt(distance_squared); + const float3 light_direction = to_light / distance; + const float surface_cosine = + dot(normal, light_direction); + const float light_cosine = + dot(light_normal, -light_direction); + + if (surface_cosine > 0.0f && light_cosine > 0.0f) { + fgt_ray shadow_ray; + shadow_ray.origin = + hit.point + hit.geometric_normal * 0.001f; + shadow_ray.direction = light_direction; + const bool blocked = fgt_trace_shadow_bvh( + shadow_ray, + triangle_geometry, + bvh_nodes, + num_bvh_nodes, + distance - 0.001f); + + if (!blocked) { + const float3 nee = + fgt_evaluate_cook_torrance( + normal, + view_direction, + light_direction, + hit.material, + light_emission); + const float geometry_term = + light_cosine / distance_squared; + radiance += throughput * nee * + geometry_term / light_pdf; + } + } + } + } + } + + const float3 new_direction = + fgt_sample_hemisphere(normal, rng); + const float3 average_fresnel = fgt_fresnel_schlick( + f0, + fmax(dot(view_direction, normal), 0.0f)); + const float3 diffuse_weight = + ((float3)(1.0f) - average_fresnel) * (1.0f - metallic); + throughput *= diffuse_weight * base_color; + + ray.origin = hit.point + normal * 0.001f; + ray.direction = new_direction; + + if (bounce > 2) { + const float probability = fmin( + 0.95f, + fmax(throughput.x, + fmax(throughput.y, throughput.z))); + if (probability <= 0.0f || + fgt_rng_next_float(rng) > probability) { + break; + } + throughput /= probability; + } + } + + return radiance; +} + +__kernel void path_tracer( + __global fgt_vec4* framebuffer, + __global const fgt_triangle_geom* triangle_geometry, + __global const fgt_triangle_shading* triangle_shading, + __global const fgt_material_data* materials, + __global const fgt_bvh_node* bvh_nodes, + __global const fgt_light* lights, + __global const fgt_int32* emissive_triangles, + __global const ulong* texture_handles, + __global const int2* texture_dimensions, + int num_triangles, + int num_materials, + int num_bvh_nodes, + int num_lights, + int num_emissive_triangles, + int num_textures, + __constant const fgt_rayspace_camera* camera, + int width, + int height, + int sample_index) +{ + const int x = (int)get_global_id(0); + const int y = (int)get_global_id(1); + + if (x >= width || y >= height) { + return; + } + + const int pixel_index = y * width + x; + const ulong seed = + (ulong)pixel_index * 1337UL + + (ulong)sample_index * 7919UL + + 123UL; + fgt_rng rng = fgt_rng_create(seed, 1UL); + + const float width_denominator = (float)max(width - 1, 1); + const float height_denominator = (float)max(height - 1, 1); + const float image_u = + ((float)x + fgt_rng_next_float(&rng)) / width_denominator; + const float image_v = + ((float)y + fgt_rng_next_float(&rng)) / height_denominator; + const fgt_ray ray = fgt_camera_get_ray(camera, image_u, image_v); + + const float3 contribution = fgt_path_trace_cook_torrance( + ray, + triangle_geometry, + triangle_shading, + materials, + bvh_nodes, + lights, + emissive_triangles, + texture_handles, + texture_dimensions, + num_triangles, + num_materials, + num_bvh_nodes, + num_lights, + num_emissive_triangles, + num_textures, + &rng); + + fgt_vec4 accumulated = framebuffer[pixel_index]; + accumulated.x += contribution.x; + accumulated.y += contribution.y; + accumulated.z += contribution.z; + accumulated.w = 1.0f; + framebuffer[pixel_index] = accumulated; +} diff --git a/PBR/Render/ocl_kernels/rayspace_camera.cl b/PBR/Render/ocl_kernels/rayspace_camera.cl new file mode 100644 index 0000000..3fd557a --- /dev/null +++ b/PBR/Render/ocl_kernels/rayspace_camera.cl @@ -0,0 +1,26 @@ +typedef struct fgt_ray { + float3 origin; + float3 direction; +} fgt_ray; + +inline float3 fgt_vec4_xyz(fgt_vec4 value) +{ + return (float3)(value.x, value.y, value.z); +} + +inline fgt_ray fgt_camera_get_ray( + __constant const fgt_rayspace_camera* camera, + float image_u, + float image_v) +{ + const float3 origin = fgt_vec4_xyz(camera->origin_lens_radius); + const float3 lower_left = fgt_vec4_xyz(camera->lower_left_corner); + const float3 horizontal = fgt_vec4_xyz(camera->horizontal); + const float3 vertical = fgt_vec4_xyz(camera->vertical); + + fgt_ray ray; + ray.origin = origin; + ray.direction = normalize( + lower_left + image_u * horizontal + image_v * vertical - origin); + return ray; +} diff --git a/PBR/Render/ocl_kernels/rng.cl b/PBR/Render/ocl_kernels/rng.cl new file mode 100644 index 0000000..a0de03a --- /dev/null +++ b/PBR/Render/ocl_kernels/rng.cl @@ -0,0 +1,34 @@ +typedef struct fgt_rng { + ulong state; + ulong increment; +} fgt_rng; + +inline uint fgt_rng_next_u32(__private fgt_rng* rng) +{ + const ulong old_state = rng->state; + rng->state = old_state * 6364136223846793005UL + rng->increment; + + const uint xorshifted = + (uint)(((old_state >> 18U) ^ old_state) >> 27U); + const uint rotation = (uint)(old_state >> 59U); + + return (xorshifted >> rotation) | + (xorshifted << ((-rotation) & 31U)); +} + +inline fgt_rng fgt_rng_create(ulong seed, ulong sequence) +{ + fgt_rng rng; + rng.state = 0UL; + rng.increment = (sequence << 1U) | 1UL; + fgt_rng_next_u32(&rng); + rng.state += seed; + fgt_rng_next_u32(&rng); + return rng; +} + +inline float fgt_rng_next_float(__private fgt_rng* rng) +{ + return (float)(fgt_rng_next_u32(rng) >> 8U) * + (1.0f / 16777216.0f); +} diff --git a/PBR/Render/ocl_kernels/sample_framebuffer.cl b/PBR/Render/ocl_kernels/sample_framebuffer.cl new file mode 100644 index 0000000..367ab54 --- /dev/null +++ b/PBR/Render/ocl_kernels/sample_framebuffer.cl @@ -0,0 +1,11 @@ +__kernel void sample_framebuffer( + __global fgt_vec4* framebuffer, + int samples_per_pixel){ + int idx = get_global_id(0); + const float inverse_samples = 1.0f / (float)samples_per_pixel; + framebuffer[idx].x *= inverse_samples; + framebuffer[idx].y *= inverse_samples; + framebuffer[idx].z *= inverse_samples; + framebuffer[idx].w = 1.0f; + +} diff --git a/PBR/Render/ocl_kernels/surface_hit.cl b/PBR/Render/ocl_kernels/surface_hit.cl new file mode 100644 index 0000000..c3459cf --- /dev/null +++ b/PBR/Render/ocl_kernels/surface_hit.cl @@ -0,0 +1,77 @@ +typedef struct fgt_surface_hit { + float distance; + float3 point; + float3 geometric_normal; + float3 shading_normal; + float2 uv; + fgt_material_data material; + int triangle_index; +} fgt_surface_hit; + +inline float3 fgt_storage_vec3_to_float3(fgt_vec3 value) +{ + return (float3)(value.x, value.y, value.z); +} + +inline float2 fgt_storage_vec2_to_float2(fgt_vec2 value) +{ + return (float2)(value.x, value.y); +} + +inline bool fgt_resolve_surface_hit( + fgt_geometry_hit geometry_hit, + __global const fgt_triangle_shading* triangle_shading, + __global const fgt_material_data* materials, + int num_triangles, + int num_materials, + __private fgt_surface_hit* surface_hit) +{ + const int triangle_index = geometry_hit.triangle_index; + if (triangle_index < 0 || triangle_index >= num_triangles) { + return false; + } + + const fgt_triangle_shading shading = + triangle_shading[triangle_index]; + if (shading.material_index < 0 || + shading.material_index >= num_materials) { + return false; + } + + const float barycentric_x = geometry_hit.barycentric.x; + const float barycentric_y = geometry_hit.barycentric.y; + const float barycentric_z = geometry_hit.barycentric.z; + + const float3 normal0 = fgt_storage_vec3_to_float3(shading.n0); + const float3 normal1 = fgt_storage_vec3_to_float3(shading.n1); + const float3 normal2 = fgt_storage_vec3_to_float3(shading.n2); + const float3 interpolated_normal = + normal0 * barycentric_x + + normal1 * barycentric_y + + normal2 * barycentric_z; + + const float3 geometric_normal = geometry_hit.geometric_normal; + float3 shading_normal = geometric_normal; + if (dot(interpolated_normal, interpolated_normal) > 1.0e-12f) { + shading_normal = normalize(interpolated_normal); + } + if (dot(shading_normal, geometric_normal) < 0.0f) { + shading_normal = -shading_normal; + } + + const float2 uv0 = fgt_storage_vec2_to_float2(shading.uv0); + const float2 uv1 = fgt_storage_vec2_to_float2(shading.uv1); + const float2 uv2 = fgt_storage_vec2_to_float2(shading.uv2); + + surface_hit->distance = geometry_hit.distance; + surface_hit->point = geometry_hit.point; + surface_hit->geometric_normal = geometric_normal; + surface_hit->shading_normal = shading_normal; + surface_hit->uv = + uv0 * barycentric_x + + uv1 * barycentric_y + + uv2 * barycentric_z; + surface_hit->material = materials[shading.material_index]; + surface_hit->triangle_index = triangle_index; + return true; +} diff --git a/PBR/Render/ocl_kernels/texture_sampling.cl b/PBR/Render/ocl_kernels/texture_sampling.cl new file mode 100644 index 0000000..dbb809b --- /dev/null +++ b/PBR/Render/ocl_kernels/texture_sampling.cl @@ -0,0 +1,20 @@ +float4 __builtin_IB_OCL_2d_ld_ro(long image_id, int2 coord, int lod); + +inline float3 fgt_sample_bindless_texture( + __global const ulong* texture_handles, + __global const int2* texture_dimensions, + int texture_index, + float2 uv) +{ + const long handle = (long)texture_handles[texture_index]; + const int2 size = texture_dimensions[texture_index]; + const float2 repeated_uv = uv - floor(uv); + const int2 coordinate = min( + convert_int2(repeated_uv * convert_float2(size)), + size - 1); + + return __builtin_IB_OCL_2d_ld_ro( + handle, + coordinate, + 0).xyz; +} diff --git a/PBR/Render/shared/opencl/fgt_opencl_data.h b/PBR/Render/shared/opencl/fgt_opencl_data.h new file mode 100644 index 0000000..5f331cc --- /dev/null +++ b/PBR/Render/shared/opencl/fgt_opencl_data.h @@ -0,0 +1,140 @@ +#ifndef FGT_OPENCL_DATA_H +#define FGT_OPENCL_DATA_H + +/* + * Shared scene-transfer ABI for the C++ OpenCL host and OpenCL C kernels. + * Keep this header free of C++ classes, constructors, namespaces and methods. + */ +#if defined(__OPENCL_C_VERSION__) || defined(__OPENCL_VERSION__) +typedef int fgt_int32; +#define FGT_DATA_ALIGN_16 __attribute__((aligned(16))) +#else +#include +#include +#include +typedef std::int32_t fgt_int32; +#define FGT_DATA_ALIGN_16 alignas(16) +#endif + +typedef struct fgt_vec2 { + float x; + float y; +} fgt_vec2; + +/* Compact storage type. Do not replace this with OpenCL float3. */ +typedef struct fgt_vec3 { + float x; + float y; + float z; +} fgt_vec3; + +/* Explicitly aligned type for vector-friendly geometry loads. */ +typedef struct FGT_DATA_ALIGN_16 fgt_vec4 { + float x; + float y; + float z; + float w; +} fgt_vec4; + +typedef struct FGT_DATA_ALIGN_16 fgt_material_data { + float base_color[3]; + float metallic; + + float roughness; + float reflectance; + float emission; + fgt_int32 base_color_texture_index; +} fgt_material_data; + +/* Hot intersection data: v0 plus precomputed triangle edges. */ +typedef struct FGT_DATA_ALIGN_16 fgt_triangle_geom { + fgt_vec4 v0; + fgt_vec4 edge1; + fgt_vec4 edge2; +} fgt_triangle_geom; + +/* Cold data, fetched only after an intersection is accepted. */ +typedef struct FGT_DATA_ALIGN_16 fgt_triangle_shading { + fgt_vec3 n0; + fgt_vec3 n1; + fgt_vec3 n2; + + fgt_vec2 uv0; + fgt_vec2 uv1; + fgt_vec2 uv2; + + fgt_int32 material_index; +} fgt_triangle_shading; + +typedef struct FGT_DATA_ALIGN_16 fgt_aabb { + fgt_vec4 bounds_min; + fgt_vec4 bounds_max; +} fgt_aabb; + +typedef struct FGT_DATA_ALIGN_16 fgt_bvh_node { + fgt_aabb bounds; + fgt_int32 left_child; + fgt_int32 right_child; + fgt_int32 first_triangle_index; + fgt_int32 triangle_count; +} fgt_bvh_node; + +#define FGT_LIGHT_POINT 0 +#define FGT_LIGHT_DIRECTIONAL 1 +#define FGT_LIGHT_AREA 2 + +typedef struct FGT_DATA_ALIGN_16 fgt_light { + fgt_vec3 position; + fgt_int32 type; + + fgt_vec3 intensity; + fgt_int32 padding; +} fgt_light; + +/* Camera fields use 16-byte rows to keep the by-value/buffer ABI unambiguous. */ +typedef struct FGT_DATA_ALIGN_16 fgt_rayspace_camera { + fgt_vec4 origin_lens_radius; + fgt_vec4 lower_left_corner; + fgt_vec4 horizontal; + fgt_vec4 vertical; + fgt_vec4 basis_u; + fgt_vec4 basis_v; + fgt_vec4 basis_w; +} fgt_rayspace_camera; + +#if !defined(__OPENCL_C_VERSION__) && !defined(__OPENCL_VERSION__) +static_assert(sizeof(fgt_int32) == 4); + +static_assert(sizeof(fgt_vec2) == 8); +static_assert(sizeof(fgt_vec3) == 12); +static_assert(sizeof(fgt_vec4) == 16); +static_assert(alignof(fgt_vec4) == 16); + +static_assert(sizeof(fgt_material_data) == 32); +static_assert(alignof(fgt_material_data) == 16); +static_assert(offsetof(fgt_material_data, metallic) == 12); +static_assert(offsetof(fgt_material_data, roughness) == 16); +static_assert(offsetof(fgt_material_data, base_color_texture_index) == 28); + +static_assert(sizeof(fgt_triangle_geom) == 48); +static_assert(alignof(fgt_triangle_geom) == 16); +static_assert(sizeof(fgt_triangle_shading) == 64); +static_assert(alignof(fgt_triangle_shading) == 16); +static_assert(offsetof(fgt_triangle_shading, material_index) == 60); + +static_assert(sizeof(fgt_aabb) == 32); +static_assert(sizeof(fgt_bvh_node) == 48); +static_assert(sizeof(fgt_light) == 32); +static_assert(sizeof(fgt_rayspace_camera) == 112); + +static_assert(std::is_standard_layout_v); +static_assert(std::is_trivially_copyable_v); +static_assert(std::is_standard_layout_v); +static_assert(std::is_trivially_copyable_v); +static_assert(std::is_standard_layout_v); +static_assert(std::is_trivially_copyable_v); +#endif + +#undef FGT_DATA_ALIGN_16 + +#endif diff --git a/PBR/Render/shared/opencl/fgt_opencl_data_translator.hpp b/PBR/Render/shared/opencl/fgt_opencl_data_translator.hpp new file mode 100644 index 0000000..5a1eb08 --- /dev/null +++ b/PBR/Render/shared/opencl/fgt_opencl_data_translator.hpp @@ -0,0 +1,207 @@ +#ifndef FGT_OPENCL_DATA_TRANSLATOR_HPP +#define FGT_OPENCL_DATA_TRANSLATOR_HPP + +#include "PBR/Render/shared/opencl/fgt_opencl_data.h" + +#include "PBR/BVH/bvh_node.hpp" +#include "PBR/Light/light.hpp" +#include "PBR/PBRCamera/pbr_camera.hpp" +#include "Triangle/triangle.hpp" +#include "Vector/vector3.hpp" + +#include +#include +#include +#include + +namespace fgt::opencl { + +struct fgt_opencl_scene_data { + std::vector triangle_geometry; + std::vector triangle_shading; + std::vector materials; + std::vector bvh_nodes; + std::vector lights; +}; + +inline fgt_vec3 translate_vec3(const fungt::Vec3& value) +{ + return {value.x, value.y, value.z}; +} + +inline fgt_vec4 translate_vec4(const fungt::Vec3& value, float w = 0.0f) +{ + return {value.x, value.y, value.z, w}; +} + +inline fgt_material_data translate_material(const MaterialData& material) +{ + return { + { + material.baseColor[0], + material.baseColor[1], + material.baseColor[2] + }, + material.metallic, + material.roughness, + material.reflectance, + material.emission, + static_cast(material.baseColorTexIdx) + }; +} + +inline bool materials_equal( + const fgt_material_data& lhs, + const fgt_material_data& rhs) +{ + return lhs.base_color[0] == rhs.base_color[0] && + lhs.base_color[1] == rhs.base_color[1] && + lhs.base_color[2] == rhs.base_color[2] && + lhs.metallic == rhs.metallic && + lhs.roughness == rhs.roughness && + lhs.reflectance == rhs.reflectance && + lhs.emission == rhs.emission && + lhs.base_color_texture_index == rhs.base_color_texture_index; +} + +inline fgt_int32 find_or_add_material( + const MaterialData& source, + std::vector& materials) +{ + const fgt_material_data translated = translate_material(source); + + for (std::size_t index = 0; index < materials.size(); ++index) { + if (materials_equal(materials[index], translated)) { + return static_cast(index); + } + } + + if (materials.size() >= + static_cast(std::numeric_limits::max())) { + throw std::overflow_error( + "OpenCL material table exceeds the signed 32-bit index range."); + } + + materials.push_back(translated); + return static_cast(materials.size() - 1); +} + +inline fgt_triangle_geom translate_triangle_geometry(const Triangle& triangle) +{ + const fungt::Vec3 edge1 = triangle.v1 - triangle.v0; + const fungt::Vec3 edge2 = triangle.v2 - triangle.v0; + + return { + translate_vec4(triangle.v0), + translate_vec4(edge1), + translate_vec4(edge2) + }; +} + +inline fgt_triangle_shading translate_triangle_shading( + const Triangle& triangle, + fgt_int32 material_index) +{ + return { + translate_vec3(triangle.n0), + translate_vec3(triangle.n1), + translate_vec3(triangle.n2), + {triangle.uvs[0][0], triangle.uvs[0][1]}, + {triangle.uvs[1][0], triangle.uvs[1][1]}, + {triangle.uvs[2][0], triangle.uvs[2][1]}, + material_index + }; +} + +inline fgt_aabb translate_aabb(const AABB& bounds) +{ + return { + translate_vec4(bounds.m_min), + translate_vec4(bounds.m_max) + }; +} + +inline fgt_bvh_node translate_bvh_node(const BVHNode& node) +{ + return { + translate_aabb(node.m_boundingBox), + static_cast(node.leftChild), + static_cast(node.rightChild), + static_cast(node.firstTriIdx), + static_cast(node.triCount) + }; +} + +inline fgt_int32 translate_light_type(LightType type) +{ + switch (type) { + case LightType::Point: + return FGT_LIGHT_POINT; + case LightType::Directional: + return FGT_LIGHT_DIRECTIONAL; + case LightType::Area: + return FGT_LIGHT_AREA; + } + + throw std::invalid_argument("Unknown light type for OpenCL translation."); +} + +inline fgt_light translate_light(const Light& light) +{ + return { + translate_vec3(light.m_pos), + translate_light_type(light.m_type), + translate_vec3(light.m_intensity), + 0 + }; +} + +inline fgt_rayspace_camera translate_rayspace_camera(const PBRCamera& camera) +{ + return { + translate_vec4(camera.getOrigin(), camera.getLensRadius()), + translate_vec4(camera.getLowerLeftCorner()), + translate_vec4(camera.getHorizontal()), + translate_vec4(camera.getVertical()), + translate_vec4(camera.getBasisU()), + translate_vec4(camera.getBasisV()), + translate_vec4(camera.getBasisW()) + }; +} + +inline fgt_opencl_scene_data translate_scene_data( + const std::vector& triangles, + const std::vector& nodes, + const std::vector& lights) +{ + fgt_opencl_scene_data result; + + result.triangle_geometry.reserve(triangles.size()); + result.triangle_shading.reserve(triangles.size()); + + for (const Triangle& triangle : triangles) { + const fgt_int32 material_index = + find_or_add_material(triangle.material, result.materials); + + result.triangle_geometry.push_back( + translate_triangle_geometry(triangle)); + result.triangle_shading.push_back( + translate_triangle_shading(triangle, material_index)); + } + + result.bvh_nodes.reserve(nodes.size()); + for (const BVHNode& node : nodes) { + result.bvh_nodes.push_back(translate_bvh_node(node)); + } + + result.lights.reserve(lights.size()); + for (const Light& light : lights) { + result.lights.push_back(translate_light(light)); + } + + return result; +} + +} // namespace fgt::opencl + +#endif diff --git a/PBR/Render/src/opencl_renderer.cpp b/PBR/Render/src/opencl_renderer.cpp index f6c9d08..6465c9a 100644 --- a/PBR/Render/src/opencl_renderer.cpp +++ b/PBR/Render/src/opencl_renderer.cpp @@ -1,4 +1,7 @@ #include "opencl_renderer.hpp" +#include "PBR/Render/shared/opencl/fgt_opencl_data_translator.hpp" +#include +#include void OpenCL_Renderer::initialize() { @@ -62,6 +65,8 @@ void OpenCL_Renderer::initialize() std::to_string(err)); } + // This is an in order queuee + // Commands run in the order we add them. m_oclqueue = clCreateCommandQueueWithProperties( m_oclcontext, m_ocldevice, nullptr, &err); if (err != CL_SUCCESS || !m_oclqueue) { @@ -76,7 +81,8 @@ void OpenCL_Renderer::initialize() m_textureManager = std::make_unique( m_oclcontext, m_oclplatform, true); - std::cout << "OpenCL context and command queue initialized." << std::endl; + // Defer kernel build until textures are loaded for bindless to work + std::cout << "OpenCL context and command queue initialized (kernel build deferred)." << std::endl; } OpenCL_Renderer::~OpenCL_Renderer() @@ -87,16 +93,27 @@ OpenCL_Renderer::~OpenCL_Renderer() m_textureManager.reset(); + if (m_textureDimensions) { + clReleaseMemObject(m_textureDimensions); + m_textureDimensions = nullptr; + } if (m_texturesObj) { clReleaseMemObject(m_texturesObj); m_texturesObj = nullptr; } m_numTextures = 0; + releaseRaySpaceBuffer(m_raySpaceBuffer); + m_sceneUploaded = false; + if (m_oclrenderKernel) { clReleaseKernel(m_oclrenderKernel); m_oclrenderKernel = nullptr; } + if (m_oclsampleKernel) { + clReleaseKernel(m_oclsampleKernel); + m_oclsampleKernel = nullptr; + } if (m_oclprogram) { clReleaseProgram(m_oclprogram); @@ -117,10 +134,397 @@ OpenCL_Renderer::~OpenCL_Renderer() m_oclplatform = nullptr; } -std::vector OpenCL_Renderer::RenderScene(int width, int height, const std::vector& triangles, const std::vector& nodes, const std::vector& lights, const std::vector& emissiveTriIndices, const PBRCamera& camera, int samplesPerPixel, int sampleOffset) +std::vector OpenCL_Renderer::RenderScene(int width, + int height, + const std::vector& triangles, + const std::vector& nodes, + const std::vector& lights, + const std::vector& emissiveTriIndices, + const PBRCamera& camera, + int samplesPerPixel, + int sampleOffset) { + if (width <= 0 || height <= 0) { + throw std::invalid_argument( + "OpenCL RenderScene requires positive image dimensions."); + } + if (samplesPerPixel <= 0) { + throw std::invalid_argument( + "OpenCL RenderScene requires at least one sample per pixel."); + } + if (!m_oclcontext || !m_oclqueue) { + throw std::runtime_error( + "OpenCL RenderScene called before renderer initialization."); + } + prepareTextures(); - return std::vector(); + + if (!m_oclrenderKernel) { + throw std::runtime_error( + "OpenCL RenderScene: kernel not built (no textures loaded?)."); + } + if (!m_sceneUploaded) { + uploadScene(triangles, nodes, lights, emissiveTriIndices); + } + + const fgt_rayspace_camera cameraData = + fgt::opencl::translate_rayspace_camera(camera); + const std::size_t imageSize = + static_cast(width) * static_cast(height); + + cl_int err = CL_SUCCESS; + cl_mem framebufferBuffer = clCreateBuffer( + m_oclcontext, + CL_MEM_READ_WRITE, + imageSize * sizeof(fgt_vec4), + nullptr, + &err); + if (err != CL_SUCCESS || !framebufferBuffer) { + throw std::runtime_error( + "OpenCL RenderScene failed to create framebuffer, error " + + std::to_string(err)); + } + + cl_mem cameraBuffer = clCreateBuffer( + m_oclcontext, + CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + sizeof(fgt_rayspace_camera), + const_cast(&cameraData), + &err); + if (err != CL_SUCCESS || !cameraBuffer) { + clReleaseMemObject(framebufferBuffer); + throw std::runtime_error( + "OpenCL RenderScene failed to create camera buffer, error " + + std::to_string(err)); + } + + const fgt_vec4 zeroPixel{0.0f, 0.0f, 0.0f, 0.0f}; + // This clear runs before all sample kernels. + err = clEnqueueFillBuffer( + m_oclqueue, + framebufferBuffer, + &zeroPixel, + sizeof(zeroPixel), + 0, + imageSize * sizeof(fgt_vec4), + 0, + nullptr, + nullptr); + if (err != CL_SUCCESS) { + clReleaseMemObject(cameraBuffer); + clReleaseMemObject(framebufferBuffer); + throw std::runtime_error( + "OpenCL RenderScene failed to clear framebuffer, error " + + std::to_string(err)); + } + + const cl_int numTriangles = static_cast(m_raySpaceBuffer.numTriangles); + const cl_int numMaterials = static_cast(m_raySpaceBuffer.numMaterials); + const cl_int numBVHNodes = static_cast(m_raySpaceBuffer.numBVHNodes); + const cl_int numLights = static_cast(m_raySpaceBuffer.numLights); + const cl_int numEmissiveTriangles = static_cast(m_raySpaceBuffer.numEmissiveTriangles); + const cl_int numTextures = static_cast(m_numTextures); + + cl_uint argument = 0; + //Lambda to setkernel arguments + const auto setKernelArg = [&](cl_kernel kernel, std::size_t size, const void* value) { + const cl_int argumentError = + clSetKernelArg(kernel, argument++, size, value); + if (argumentError != CL_SUCCESS) { + throw std::runtime_error( + "OpenCL RenderScene failed to set kernel argument " + + std::to_string(argument - 1) + ", error " + + std::to_string(argumentError)); + } + }; + + try { + setKernelArg(m_oclrenderKernel, sizeof(cl_mem), &framebufferBuffer); + setKernelArg(m_oclrenderKernel, sizeof(cl_mem), &m_raySpaceBuffer.triangleGeometry); + setKernelArg(m_oclrenderKernel, sizeof(cl_mem), &m_raySpaceBuffer.triangleShading); + setKernelArg(m_oclrenderKernel, sizeof(cl_mem), &m_raySpaceBuffer.materials); + setKernelArg(m_oclrenderKernel, sizeof(cl_mem), &m_raySpaceBuffer.bvhNodes); + setKernelArg(m_oclrenderKernel, sizeof(cl_mem), &m_raySpaceBuffer.lights); + setKernelArg(m_oclrenderKernel, sizeof(cl_mem), &m_raySpaceBuffer.emissiveTriangles); + setKernelArg(m_oclrenderKernel, sizeof(cl_mem), &m_texturesObj); + setKernelArg(m_oclrenderKernel, sizeof(cl_mem), &m_textureDimensions); + setKernelArg(m_oclrenderKernel, sizeof(cl_int), &numTriangles); + setKernelArg(m_oclrenderKernel, sizeof(cl_int), &numMaterials); + setKernelArg(m_oclrenderKernel, sizeof(cl_int), &numBVHNodes); + setKernelArg(m_oclrenderKernel, sizeof(cl_int), &numLights); + setKernelArg(m_oclrenderKernel, sizeof(cl_int), &numEmissiveTriangles); + setKernelArg(m_oclrenderKernel, sizeof(cl_int), &numTextures); + setKernelArg(m_oclrenderKernel, sizeof(cl_mem), &cameraBuffer); + setKernelArg(m_oclrenderKernel, sizeof(cl_int), &width); + setKernelArg(m_oclrenderKernel, sizeof(cl_int), &height); + + const cl_int initialSampleIndex = sampleOffset; + setKernelArg(m_oclrenderKernel, sizeof(cl_int), &initialSampleIndex); + + const std::size_t localSize[2] = {16, 16}; + const std::size_t globalSize[2] = { + (static_cast(width) + localSize[0] - 1) / + localSize[0] * localSize[0], + (static_cast(height) + localSize[1] - 1) / + localSize[1] * localSize[1] + }; + + constexpr cl_uint sampleIndexArgument = 18; + for (int sample = 0; sample < samplesPerPixel; ++sample) { + const cl_int sampleIndex = sampleOffset + sample; + err = clSetKernelArg( + m_oclrenderKernel, + sampleIndexArgument, + sizeof(sampleIndex), + &sampleIndex); + if (err != CL_SUCCESS) { + throw std::runtime_error( + "OpenCL RenderScene failed to set sample index, error " + + std::to_string(err)); + } + + // No wait is needed + // The next sample runs after this sample. + err = clEnqueueNDRangeKernel( + m_oclqueue, + m_oclrenderKernel, + 2, + nullptr, + globalSize, + localSize, + 0, + nullptr, + nullptr); + if (err != CL_SUCCESS) { + throw std::runtime_error( + "OpenCL RenderScene failed to enqueue path_tracer sample " + + std::to_string(sampleIndex) + ", error " + + std::to_string(err)); + } + } // When this ends framebuffer is filled with samplesPerPixel samples per pixel. + + // Now add args for the second kernel. + argument = 0; + + setKernelArg(m_oclsampleKernel, sizeof(cl_mem), &framebufferBuffer); + setKernelArg(m_oclsampleKernel, sizeof(cl_int), &samplesPerPixel); + + // Divide each pixel by samplesPerPixel. + const std::size_t sampleGlobalSize[1] = {imageSize}; + + err = clEnqueueNDRangeKernel( + m_oclqueue, + m_oclsampleKernel, + 1, + nullptr, + sampleGlobalSize, + nullptr, + 0, + nullptr, + nullptr); + if (err != CL_SUCCESS) { + throw std::runtime_error( + "OpenCL RenderScene failed to enqueue sample_framebuffer kernel, error " + + std::to_string(err)); + } + // Get the framebuffer back to host memory + std::vector finalframebuffer(imageSize); + // CL_TRUE waits for this read and all earlier kernels. + err = clEnqueueReadBuffer( + m_oclqueue, + framebufferBuffer, + CL_TRUE, + 0, + imageSize * sizeof(fgt_vec4), + finalframebuffer.data(), + 0, + nullptr, + nullptr); + if (err != CL_SUCCESS) { + throw std::runtime_error( + "OpenCL RenderScene failed to read framebuffer, error " + + std::to_string(err)); + } + + std::vector framebuffer(imageSize); + for (std::size_t index = 0; index < imageSize; ++index) { + framebuffer[index] = fungt::Vec3( + finalframebuffer[index].x, + finalframebuffer[index].y, + finalframebuffer[index].z); + } + + clReleaseMemObject(cameraBuffer); + clReleaseMemObject(framebufferBuffer); + + return framebuffer; + } + catch (...) { + clReleaseMemObject(cameraBuffer); + clReleaseMemObject(framebufferBuffer); + throw; + } +} + +void OpenCL_Renderer::uploadScene( + const std::vector& triangles, + const std::vector& nodes, + const std::vector& lights, + const std::vector& emissiveTriIndices) +{ + if (!m_oclcontext) { + throw std::runtime_error( + "uploadScene: OpenCL context is not initialized."); + } + + const fgt::opencl::fgt_opencl_scene_data sceneData = + fgt::opencl::translate_scene_data(triangles, nodes, lights); + + std::cout << "[uploadScene] Materials texture indices: "; + for (size_t i = 0; i < sceneData.materials.size(); ++i) { + std::cout << "mat" << i << "=" << sceneData.materials[i].base_color_texture_index << " "; + } + std::cout << std::endl; + + std::vector emissiveIndices; + emissiveIndices.reserve(emissiveTriIndices.size()); + for (int index : emissiveTriIndices) { + emissiveIndices.push_back(static_cast(index)); + } + + OpenCLRaySpaceBuffer newBuffers; + newBuffers.numTriangles = sceneData.triangle_geometry.size(); + newBuffers.numMaterials = sceneData.materials.size(); + newBuffers.numBVHNodes = sceneData.bvh_nodes.size(); + newBuffers.numLights = sceneData.lights.size(); + newBuffers.numEmissiveTriangles = emissiveIndices.size(); + + cl_int err = CL_SUCCESS; + + if (!sceneData.triangle_geometry.empty()) { + newBuffers.triangleGeometry = clCreateBuffer( + m_oclcontext, + CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + sceneData.triangle_geometry.size() * sizeof(fgt_triangle_geom), + const_cast(sceneData.triangle_geometry.data()), + &err); + if (err != CL_SUCCESS || !newBuffers.triangleGeometry) { + releaseRaySpaceBuffer(newBuffers); + throw std::runtime_error( + "uploadScene: failed to create triangle geometry buffer, error " + + std::to_string(err)); + } + } + + if (!sceneData.triangle_shading.empty()) { + newBuffers.triangleShading = clCreateBuffer( + m_oclcontext, + CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + sceneData.triangle_shading.size() * sizeof(fgt_triangle_shading), + const_cast(sceneData.triangle_shading.data()), + &err); + if (err != CL_SUCCESS || !newBuffers.triangleShading) { + releaseRaySpaceBuffer(newBuffers); + throw std::runtime_error( + "uploadScene: failed to create triangle shading buffer, error " + + std::to_string(err)); + } + } + + if (!sceneData.materials.empty()) { + newBuffers.materials = clCreateBuffer( + m_oclcontext, + CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + sceneData.materials.size() * sizeof(fgt_material_data), + const_cast(sceneData.materials.data()), + &err); + if (err != CL_SUCCESS || !newBuffers.materials) { + releaseRaySpaceBuffer(newBuffers); + throw std::runtime_error( + "uploadScene: failed to create material buffer, error " + + std::to_string(err)); + } + } + + if (!sceneData.bvh_nodes.empty()) { + newBuffers.bvhNodes = clCreateBuffer( + m_oclcontext, + CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + sceneData.bvh_nodes.size() * sizeof(fgt_bvh_node), + const_cast(sceneData.bvh_nodes.data()), + &err); + if (err != CL_SUCCESS || !newBuffers.bvhNodes) { + releaseRaySpaceBuffer(newBuffers); + throw std::runtime_error( + "uploadScene: failed to create BVH buffer, error " + + std::to_string(err)); + } + } + + if (!sceneData.lights.empty()) { + newBuffers.lights = clCreateBuffer( + m_oclcontext, + CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + sceneData.lights.size() * sizeof(fgt_light), + const_cast(sceneData.lights.data()), + &err); + if (err != CL_SUCCESS || !newBuffers.lights) { + releaseRaySpaceBuffer(newBuffers); + throw std::runtime_error( + "uploadScene: failed to create light buffer, error " + + std::to_string(err)); + } + } + + if (!emissiveIndices.empty()) { + newBuffers.emissiveTriangles = clCreateBuffer( + m_oclcontext, + CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + emissiveIndices.size() * sizeof(fgt_int32), + emissiveIndices.data(), + &err); + if (err != CL_SUCCESS || !newBuffers.emissiveTriangles) { + releaseRaySpaceBuffer(newBuffers); + throw std::runtime_error( + "uploadScene: failed to create emissive triangle buffer, error " + + std::to_string(err)); + } + } + + releaseRaySpaceBuffer(m_raySpaceBuffer); + m_raySpaceBuffer = newBuffers; + m_sceneUploaded = true; + + std::cout << "OpenCL RaySpace scene uploaded: " + << m_raySpaceBuffer.numTriangles << " triangles, " + << m_raySpaceBuffer.numMaterials << " materials, " + << m_raySpaceBuffer.numBVHNodes << " BVH nodes, " + << m_raySpaceBuffer.numLights << " lights." << std::endl; +} + +void OpenCL_Renderer::releaseRaySpaceBuffer( + OpenCLRaySpaceBuffer& buffers) noexcept +{ + if (buffers.emissiveTriangles) { + clReleaseMemObject(buffers.emissiveTriangles); + } + if (buffers.lights) { + clReleaseMemObject(buffers.lights); + } + if (buffers.bvhNodes) { + clReleaseMemObject(buffers.bvhNodes); + } + if (buffers.materials) { + clReleaseMemObject(buffers.materials); + } + if (buffers.triangleShading) { + clReleaseMemObject(buffers.triangleShading); + } + if (buffers.triangleGeometry) { + clReleaseMemObject(buffers.triangleGeometry); + } + + buffers = {}; } void OpenCL_Renderer::prepareTextures() @@ -130,42 +534,156 @@ void OpenCL_Renderer::prepareTextures() "OpenCL textures requested before OpenCL initialization."); } - if (m_textureManager->handlesAreDirty()) { - setOpenCLTextures(m_textureManager->getBindlessHandles()); + static bool kernelBuilt = false; + if (m_textureManager->handlesAreDirty() || + !m_texturesObj || !m_textureDimensions) { + setOpenCLTextures( + m_textureManager->getBindlessHandles(), + m_textureManager->getTextureDimensions()); m_textureManager->markHandlesClean(); } + if (!kernelBuilt && m_textureManager->getTextureCount() > 0) { + std::cout << "[prepareTextures] Building kernel AFTER textures loaded" << std::endl; + buildOCLPrograms(); + kernelBuilt = true; + } } -void OpenCL_Renderer::setOpenCLTextures(const std::vector& handles) +void OpenCL_Renderer::setOpenCLTextures( + const std::vector& handles, + const std::vector>& dimensions) { - if (!m_oclcontext) { - throw std::runtime_error( - "setOpenCLTextures: OpenCL context is not initialized."); + static_assert(sizeof(std::array) == sizeof(cl_int) * 2); + if (handles.size() != dimensions.size()) { + throw std::invalid_argument( + "setOpenCLTextures: handle and dimension counts do not match."); + } + if (handles.empty()) { + return; + } + if (m_textureDimensions) { + clReleaseMemObject(m_textureDimensions); } - - std::cout << "*** SETTING OPENCL TEXTURE OBJECTS*** " << std::endl; if (m_texturesObj) { clReleaseMemObject(m_texturesObj); - m_texturesObj = nullptr; } - - m_numTextures = handles.size(); - std::cout << "*** NUM OPENCL TEXTURE OBJECTS *** " << m_numTextures << std::endl; - - if (m_numTextures == 0) { - return; - } - cl_int err = CL_SUCCESS; m_texturesObj = clCreateBuffer( m_oclcontext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, - m_numTextures * sizeof(uint64_t), + handles.size() * sizeof(uint64_t), const_cast(handles.data()), &err); if (err != CL_SUCCESS || !m_texturesObj) { - throw std::runtime_error("setOpenCLTextures: clCreateBuffer failed for bindless handles, error " + std::to_string(err)); + throw std::runtime_error( + "setOpenCLTextures: failed to create handle buffer, error " + + std::to_string(err)); + } + + m_textureDimensions = clCreateBuffer( + m_oclcontext, + CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + dimensions.size() * sizeof(std::array), + const_cast*>(dimensions.data()), + &err); + if (err != CL_SUCCESS || !m_textureDimensions) { + clReleaseMemObject(m_texturesObj); + throw std::runtime_error( + "setOpenCLTextures: failed to create dimension buffer, error " + + std::to_string(err)); + } + m_numTextures = handles.size(); +} + +void OpenCL_Renderer::buildOCLPrograms() +{ + if (!m_oclcontext) { + throw std::runtime_error( + "buildOCLPrograms: OpenCL context is not initialized."); + } + std::vector source_pointers; + std::vector source_lengths; + std::vector file_contents; + file_contents.reserve(OpenCLRendererSourceFiles.size()); + for (size_t i = 0; i < OpenCLRendererSourceFiles.size(); ++i) { + const std::string& sourceFile = getAssetPath(OpenCLRendererSourceFiles[i]); + std::ifstream file(sourceFile); + if (!file.is_open()) { + throw std::runtime_error("Failed to open OpenCL source file: " + sourceFile); + } + std::string buffer{ + std::istreambuf_iterator(file), + std::istreambuf_iterator() }; + buffer.push_back('\n'); + file_contents.push_back(std::move(buffer)); + } + for (const auto& content : file_contents) { + source_pointers.push_back(content.c_str()); + source_lengths.push_back(content.size()); + } + cl_int err = CL_SUCCESS; + m_oclprogram = clCreateProgramWithSource( + m_oclcontext, + static_cast(source_pointers.size()), + source_pointers.data(), + source_lengths.data(), + &err); + if (err != CL_SUCCESS || !m_oclprogram) { + throw std::runtime_error("Failed to create OpenCL program, error " + std::to_string(err)); + } + //The flags are set to enable bindless images and advanced bindless mode + err = clBuildProgram( + m_oclprogram, + 1, + &m_ocldevice, + "-cl-intel-use-bindless-images " + "-cl-intel-use-bindless-advanced-mode", + nullptr, + nullptr); + if (err != CL_SUCCESS) { + std::size_t logSize = 0; + clGetProgramBuildInfo( + m_oclprogram, + m_ocldevice, + CL_PROGRAM_BUILD_LOG, + 0, + nullptr, + &logSize); + + std::string buildLog(logSize, '\0'); + if (logSize > 0) { + clGetProgramBuildInfo( + m_oclprogram, + m_ocldevice, + CL_PROGRAM_BUILD_LOG, + logSize, + buildLog.data(), + nullptr); + } + + clReleaseProgram(m_oclprogram); + m_oclprogram = nullptr; + throw std::runtime_error( + "Failed to build OpenCL program, error " + + std::to_string(err) + "\n" + buildLog); + } + + m_oclrenderKernel = clCreateKernel(m_oclprogram,"path_tracer", &err); + + if (err != CL_SUCCESS || !m_oclrenderKernel) { + m_oclrenderKernel = nullptr; + clReleaseProgram(m_oclprogram); + m_oclprogram = nullptr; + throw std::runtime_error("Failed to create OpenCL path_tracer kernel, error " + std::to_string(err)); + } + m_oclsampleKernel = clCreateKernel(m_oclprogram, "sample_framebuffer", &err); + if (err != CL_SUCCESS || !m_oclsampleKernel) { + m_oclsampleKernel = nullptr; + clReleaseKernel(m_oclrenderKernel); + m_oclrenderKernel = nullptr; + clReleaseProgram(m_oclprogram); + m_oclprogram = nullptr; + throw std::runtime_error("Failed to create OpenCL sample_framebuffer, error " + std::to_string(err)); } - std::cout << " Uploaded " << m_numTextures << " bindless handles to GPU" << std::endl; } diff --git a/PBR/TextureManager/opencl_texture.cpp b/PBR/TextureManager/opencl_texture.cpp index 88c2fda..be337cd 100644 --- a/PBR/TextureManager/opencl_texture.cpp +++ b/PBR/TextureManager/opencl_texture.cpp @@ -11,6 +11,7 @@ OpenCLTexture::OpenCLTexture(cl_context context, cl_platform_id platform, bool u throw std::runtime_error( "OpenCLTexture: bindless image API is not available on the selected platform."); } + // No early dummy image - it interferes with bindless handle resolution } std::cout << "OpenCLTexture initialized (" << (m_useBindless ? "bindless" : "bound") << ")" << std::endl; } @@ -94,7 +95,8 @@ int OpenCLTexture::loadTexture(const std::string& path) } m_pathToIndex[path] = idx; std::cout << " textures size : " << m_textures.size() << std::endl; - std::cout << " [OpenCL] Texture index: " << idx << std::endl; + std::cout << " [OpenCL] Texture index: " << idx + << ", handle: " << ocltex.bindlessHandle << std::endl; return idx; } @@ -107,6 +109,16 @@ std::vector OpenCLTexture::getTextureObjects() const { return objs; } +std::vector> OpenCLTexture::getTextureDimensions() const +{ + std::vector> dimensions; + dimensions.reserve(m_textures.size()); + for (const auto& texture : m_textures) { + dimensions.push_back({texture.width, texture.height}); + } + return dimensions; +} + void OpenCLTexture::cleanup() { for (auto& texture : m_textures) { diff --git a/PBR/TextureManager/opencl_texture.hpp b/PBR/TextureManager/opencl_texture.hpp index e52ca55..87d410e 100644 --- a/PBR/TextureManager/opencl_texture.hpp +++ b/PBR/TextureManager/opencl_texture.hpp @@ -9,6 +9,7 @@ #include #include #include +#include #ifndef CL_MEM_BINDLESS_IMAGE_INTEL #define CL_MEM_BINDLESS_IMAGE_INTEL 0x10060 @@ -52,6 +53,7 @@ class OpenCLTexture : public IDeviceTexture { int getTextureCount() const override { return static_cast(m_textures.size() ); }; void cleanup() override; std::vector getTextureObjects() const; + std::vector> getTextureDimensions() const; const std::vector& getBindlessHandles() const { return m_bindlessHandles; } bool isBindlessMode() const { return m_useBindless; } bool handlesAreDirty() const { return m_handlesDirty; } diff --git a/PBR/standalone/main.cpp b/PBR/standalone/main.cpp index 31986f8..810df8b 100644 --- a/PBR/standalone/main.cpp +++ b/PBR/standalone/main.cpp @@ -1,9 +1,77 @@ #include "PBR/Space/space.hpp" +#include "PBR/Render/shared/opencl/fgt_opencl_data.h" +#include +#include +#include + const int IMAGE_WIDTH = 800; const int IMAGE_HEIGHT = 400; +PBRCamera frameScene( + const SimpleModel& model, + float aspectRatio, + float verticalFov = 50.0f, + float padding = 1.2f) +{ + glm::vec3 boundsMin; + glm::vec3 boundsMax; + if (!model.getWorldBounds(boundsMin, boundsMax)) { + throw std::runtime_error("Cannot frame a model with no vertices."); + } + + const glm::vec3 center = (boundsMin + boundsMax) * 0.5f; + const float radius = std::max( + glm::length(boundsMax - boundsMin) * 0.5f, + 0.001f); + const float verticalFovRadians = glm::radians(verticalFov); + const float horizontalFovRadians = 2.0f * std::atan( + std::tan(verticalFovRadians * 0.5f) * aspectRatio); + const float limitingFov = std::min( + verticalFovRadians, + horizontalFovRadians); + const float distance = + padding * radius / std::sin(limitingFov * 0.5f); + + return PBRCamera( + fungt::Vec3(center.x, center.y, center.z + distance), + fungt::Vec3(center.x, center.y, center.z), + fungt::Vec3(0.0f, 1.0f, 0.0f), + verticalFov, + aspectRatio); +} + +SceneLight frameSceneLight(const SimpleModel& model) +{ + glm::vec3 boundsMin; + glm::vec3 boundsMax; + if (!model.getWorldBounds(boundsMin, boundsMax)) { + throw std::runtime_error("Cannot light a model with no vertices."); + } + + const glm::vec3 center = (boundsMin + boundsMax) * 0.5f; + const float radius = std::max( + glm::length(boundsMax - boundsMin) * 0.5f, + 0.001f); + + SceneLight light; + light.type = SceneLightType::Point; + light.name = "Framed model light"; + light.position = center + glm::vec3(-0.8f, 1.2f, 0.8f) * radius; + light.color = glm::vec3(1.0f); + light.power = 10.0f * radius * radius; + return light; +} + int main() { + std::cout << std::boolalpha + << "sizeof(Vec3) == sizeof(fgt_vec3): " + << (sizeof(fungt::Vec3) == sizeof(fgt_vec3)) << '\n' + << "alignof(Vec3) == alignof(fgt_vec3): " + << (alignof(fungt::Vec3) == alignof(fgt_vec3)) << '\n' + << "Vec3 is trivially copyable: " + << std::is_trivially_copyable_v << std::endl; + // ComputeRender::SetBackend(Compute::Backend::OPENCL); // std::cout << "Initializing standalone PBR backend: " // << ComputeRender::GetBackendName() << std::endl; @@ -16,26 +84,25 @@ int main() // return 0; ModelPaths monkey; DisplayGraphics::SetBackend(Backend::OpenGL); - monkey.path = getAssetPath("assets_local/Obj/Woody/woody-head.obj"); + monkey.path = getAssetPath("assets_local/opencl_logo2/scene.gltf"); //monkey.vs_path = getAssetPath("resources/luxoball_vs.glsl"); //monkey.fs_path = getAssetPath("resources/luxoball_fs.glsl"); - ComputeRender::SetBackend(Compute::Backend::CUDA); + ComputeRender::SetBackend(Compute::Backend::OPENCL); std::cout << "Backend in use: " << ComputeRender::GetBackendName() << std::endl; std::shared_ptr monkey_model = SimpleModel::create(); monkey_model->LoadModelData(monkey); + monkey_model->position(0.0f, -5.0f, 2.0f); + monkey_model->rotation(0.0f, 25.0f, 0.0f); - PBRCamera camera( - fungt::Vec3(0, 2.5, 30), - fungt::Vec3(0, 1.8, 0), - fungt::Vec3(0, 1, 0), - 50.0f, - float(IMAGE_WIDTH) / float(IMAGE_HEIGHT) - ); + const float aspectRatio = + float(IMAGE_WIDTH) / float(IMAGE_HEIGHT); + PBRCamera camera = frameScene(*monkey_model, aspectRatio); Space space(camera); space.InitComputeRenderBackend(); + space.loadLightsFromScene({frameSceneLight(*monkey_model)}); space.LoadModelToRender(*monkey_model); space.setSamples(100); space.BuildBVH(); @@ -44,6 +111,22 @@ int main() auto framebuffer = space.Render(IMAGE_WIDTH, IMAGE_HEIGHT); auto totalEnd = std::chrono::high_resolution_clock::now(); + const std::size_t expectedPixels = + static_cast(IMAGE_WIDTH) * IMAGE_HEIGHT; + if (framebuffer.size() != expectedPixels) { + std::cerr << "OpenCL framebuffer size mismatch: expected " + << expectedPixels << ", received " + << framebuffer.size() << std::endl; + return 1; + } + + std::cout << "OpenCL skeleton returned " << framebuffer.size() + << " pixels." << std::endl; + std::cout << "First pixel: (" + << framebuffer.front().x << ", " + << framebuffer.front().y << ", " + << framebuffer.front().z << ")" << std::endl; + Space::SaveFrameBufferAsPNG(framebuffer, IMAGE_WIDTH, IMAGE_HEIGHT); auto totalTime = std::chrono::duration_cast( From 255f96bc96e1109729ede6f8a5c884ae741b0377 Mon Sep 17 00:00:00 2001 From: juanchuletas Date: Fri, 21 Aug 2026 08:23:30 -0600 Subject: [PATCH 5/7] feature: adding the opencl backend to the progressive path tracer viewport --- CMakeLists.txt | 31 +- GUI/render_action_window.cpp | 70 ++- GUI/render_action_window.hpp | 8 +- GUI/render_settings_window.hpp | 4 +- InfoDevice/gpu_device_info.cpp | 44 +- InfoDevice/gpu_device_info.hpp | 4 +- PBR/Render/include/opencl_renderer.hpp | 32 +- PBR/Render/ocl_kernels/path_tracer.cl | 4 +- PBR/Render/ocl_kernels/sample_framebuffer.cl | 24 +- PBR/Render/src/opencl_renderer.cpp | 461 ++++++++++++++++-- PBR/Space/space.cpp | 69 ++- PBR/Space/space.hpp | 11 +- PBR/main/CMakeLists.txt | 9 +- PBR/standalone/CMakeLists.txt | 6 +- Samples/iwocl/CMakeLists.txt | 31 +- Samples/iwocl/main.cpp | 4 +- .../opengl/opengl_progressive_path_tracer.cpp | 15 + .../opengl/opengl_progressive_path_tracer.hpp | 1 + ViewPort/opengl/opengl_viewport.cpp | 80 ++- ViewPort/opengl/opengl_viewport.hpp | 3 + ViewPort/progressive_path_tracer_viewport.cpp | 42 +- ViewPort/progressive_path_tracer_viewport.hpp | 6 +- ViewPort/viewport.hpp | 22 +- funGT/fungt.cpp | 43 +- 24 files changed, 866 insertions(+), 158 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 104b542..92c5d9d 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -6,8 +6,10 @@ project(FunGT) # ════════════════════════════════════════════════════════════════════════════ option(FUNGT_USE_CUDA "Enable CUDA backend for NVIDIA GPUs" ON) option(FUNGT_USE_SYCL "Enable SYCL backend for Intel GPUs" ON) -set(SYCL_CUDA_ARCH sm_75 CACHE STRING "NVIDIA GPU architecture for SYCL CUDA backend") -set_property(CACHE SYCL_CUDA_ARCH PROPERTY STRINGS sm_50 sm_52 sm_53 sm_60 sm_61 sm_62 sm_70 sm_72 sm_75 sm_80 sm_86 sm_87 sm_89 sm_90) +option(FUNGT_USE_OPENCL "Enable OpenCL backend" ON) +set(FUNGT_CUDA_ARCH 89 CACHE STRING "NVIDIA GPU compute capability") +set_property(CACHE FUNGT_CUDA_ARCH PROPERTY STRINGS 50 52 53 60 61 62 70 72 75 80 86 87 89 90) +set(SYCL_CUDA_ARCH "sm_${FUNGT_CUDA_ARCH}") # Allow FUNGT_BASE_DIR to be set externally (command line, environment, or cache) # Priority: 1. Cache variable, 2. Environment variable, 3. Default fallback @@ -145,6 +147,15 @@ if(UNIX) set(PBR_SYCL_LIB pbr_sycl) message(STATUS "PBR SYCL renderer library: ${PBR_LIB_DIR}/libpbr_sycl.a") endif() + + if(FUNGT_USE_OPENCL) + add_library(pbr_opencl STATIC IMPORTED) + set_target_properties(pbr_opencl PROPERTIES + IMPORTED_LOCATION "${PBR_LIB_DIR}/libpbr_opencl.a" + ) + set(PBR_OPENCL_LIB pbr_opencl) + message(STATUS "PBR OpenCL renderer library: ${PBR_LIB_DIR}/libpbr_opencl.a") + endif() endif() if (WIN32) @@ -372,7 +383,7 @@ if(FUNGT_USE_CUDA) set_target_properties(FunGT PROPERTIES CUDA_RESOLVE_DEVICE_SYMBOLS ON CUDA_SEPARABLE_COMPILATION ON - CUDA_ARCHITECTURES "native" + CUDA_ARCHITECTURES ${FUNGT_CUDA_ARCH} ) endif() # ════════════════════════════════════════════════════════════════════════════ @@ -491,11 +502,12 @@ elseif (UNIX) dl funlib imgui - ${OpenCL_LIBRARIES} ${CUDA_BACKEND_LIB} pbr_core ${PBR_CUDA_LIB} ${PBR_SYCL_LIB} + ${PBR_OPENCL_LIB} + ${OpenCL_LIBRARIES} $<$:CUDA::cudart_static> $<$:CUDA::cuda_driver> ) @@ -533,6 +545,9 @@ elseif (UNIX) if(FUNGT_USE_SYCL) target_compile_definitions(FunGT PRIVATE FUNGT_USE_SYCL) endif() + if(FUNGT_USE_OPENCL) + target_compile_definitions(FunGT PRIVATE FUNGT_USE_OPENCL) + endif() endif() @@ -543,12 +558,10 @@ message(STATUS "") message(STATUS "════════════════════════════════════════") message(STATUS " FunGT Build Configuration") message(STATUS "════════════════════════════════════════") -message(STATUS " Main Compiler: ${CMAKE_CXX_COMPILER}") +message(STATUS " Main Compiler: Intel DPC++") message(STATUS " CUDA Backend: ${FUNGT_USE_CUDA}") message(STATUS " SYCL Backend: ${FUNGT_USE_SYCL}") +message(STATUS " OpenCL Backend: ${FUNGT_USE_OPENCL}") message(STATUS " Build Type: ${CMAKE_BUILD_TYPE}") -if(FUNGT_USE_CUDA) - message(STATUS " CUDA Compiled: Separately with g++") -endif() message(STATUS "════════════════════════════════════════") -message(STATUS "") \ No newline at end of file +message(STATUS "") diff --git a/GUI/render_action_window.cpp b/GUI/render_action_window.cpp index e55c2c7..779dbdb 100644 --- a/GUI/render_action_window.cpp +++ b/GUI/render_action_window.cpp @@ -41,8 +41,27 @@ void RenderWindow::onImGuiRender() { // VIEWPORT PREVIEW ImGui::SeparatorText("Viewport Preview"); +#ifdef FUNGT_USE_OPENCL + if (ImGui::Checkbox("Use OpenCL Interop", &m_useOpenCLPreview)) { + if (m_viewportPathTrace && m_viewport) { + if (m_useOpenCLPreview) { + ComputeRender::SetBackend(Compute::Backend::OPENCL); + } else { + selectComputeBackend(); + } + m_viewport->enablePathTracing(true); + m_viewport->resetAccumulation(); + } + } +#endif + if (ImGui::Checkbox("Enable RaySpace Preview", &m_viewportPathTrace)) { if (m_viewport) { + if (m_useOpenCLPreview) { + ComputeRender::SetBackend(Compute::Backend::OPENCL); + } else { + selectComputeBackend(); + } m_viewport->enablePathTracing(m_viewportPathTrace); if (m_viewportPathTrace) { m_viewport->resetAccumulation(); @@ -158,29 +177,7 @@ void RenderWindow::triggerRender() { m_isRendering = true; try { - // SET COMPUTE BACKEND FROM SELECTED DEVICE - { - const auto& devices = m_gpuManager->getDevices(); - int activeIdx = m_gpuManager->getActiveDeviceIndex(); - if (activeIdx >= 0 && activeIdx < static_cast(devices.size())) { - const auto& dev = devices[activeIdx]; - if (dev.backend == fungt::GPUBackend::CUDA) { - ComputeRender::SetBackend(Compute::Backend::CUDA); - } - else if (dev.backend == fungt::GPUBackend::SYCL) { - bool isNvidia = dev.vendor.find("NVIDIA") != std::string::npos || - dev.vendor.find("nvidia") != std::string::npos; - ComputeRender::SetBackend(isNvidia ? Compute::Backend::SYCL_CUDA - : Compute::Backend::SYCL); - } - else { - ComputeRender::SetBackend(Compute::Backend::CPU); - } - } - else { - ComputeRender::SetBackend(Compute::Backend::CUDA); - } - } + selectComputeBackend(); std::cout << "Backend: " << ComputeRender::GetBackendName() << std::endl; // SYNC CAMERA FROM VIEWPORT @@ -268,4 +265,29 @@ void RenderWindow::triggerRender() { } m_isRendering = false; -} \ No newline at end of file +} + +void RenderWindow::selectComputeBackend() +{ + const auto& devices = m_gpuManager->getDevices(); + const int activeIdx = m_gpuManager->getActiveDeviceIndex(); + if (activeIdx < 0 || activeIdx >= static_cast(devices.size())) { + ComputeRender::SetBackend(Compute::Backend::CPU); + return; + } + + const auto& device = devices[activeIdx]; + if (device.backend == fungt::GPUBackend::CUDA) { + ComputeRender::SetBackend(Compute::Backend::CUDA); + } else if (device.backend == fungt::GPUBackend::SYCL) { + const bool isNvidia = + device.vendor.find("NVIDIA") != std::string::npos || + device.vendor.find("nvidia") != std::string::npos; + ComputeRender::SetBackend( + isNvidia ? Compute::Backend::SYCL_CUDA : Compute::Backend::SYCL); + } else if (device.backend == fungt::GPUBackend::OPENCL) { + ComputeRender::SetBackend(Compute::Backend::OPENCL); + } else { + ComputeRender::SetBackend(Compute::Backend::CPU); + } +} diff --git a/GUI/render_action_window.hpp b/GUI/render_action_window.hpp index 74ae34a..b3159ca 100644 --- a/GUI/render_action_window.hpp +++ b/GUI/render_action_window.hpp @@ -23,6 +23,11 @@ class RenderWindow : public ImGuiWindow { bool m_useViewportSize = false; bool m_isRendering = false; bool m_viewportPathTrace = false; +#ifdef FUNGT_USE_OPENCL + bool m_useOpenCLPreview = true; +#else + bool m_useOpenCLPreview = false; +#endif int m_previewSamples = 32; ViewPort* m_viewport = nullptr; @@ -39,6 +44,7 @@ class RenderWindow : public ImGuiWindow { int m_viewportHeight = 1080; void triggerRender(); + void selectComputeBackend(); public: RenderWindow(std::shared_ptr sceneManager, Camera* camera, ViewPort* viewport, @@ -58,4 +64,4 @@ class RenderWindow : public ImGuiWindow { void onImGuiRender() override; }; -#endif // _RENDER_WINDOW_H_ \ No newline at end of file +#endif // _RENDER_WINDOW_H_ diff --git a/GUI/render_settings_window.hpp b/GUI/render_settings_window.hpp index 90a6138..5c13c35 100644 --- a/GUI/render_settings_window.hpp +++ b/GUI/render_settings_window.hpp @@ -45,6 +45,8 @@ class RenderSettingsWindow : public ImGuiWindow { ImGui::Spacing(); renderBackendSection("SYCL Devices", fungt::GPUBackend::SYCL); ImGui::Spacing(); + renderBackendSection("OpenCL Devices", fungt::GPUBackend::OPENCL); + ImGui::Spacing(); renderBackendSection("OpenGL Fallback", fungt::GPUBackend::OPENGL); } } @@ -116,4 +118,4 @@ class RenderSettingsWindow : public ImGuiWindow { } }; -#endif // _RENDER_SETTINGS_WINDOW_H_ \ No newline at end of file +#endif // _RENDER_SETTINGS_WINDOW_H_ diff --git a/InfoDevice/gpu_device_info.cpp b/InfoDevice/gpu_device_info.cpp index e00f1ae..bcbb24d 100644 --- a/InfoDevice/gpu_device_info.cpp +++ b/InfoDevice/gpu_device_info.cpp @@ -1,5 +1,8 @@ #include "gpu_device_info.hpp" #include +#ifdef FUNGT_USE_OPENCL +#include +#endif // ════════════════════════════════════════════════════════════════════════════ // GPU Device Manager Implementation @@ -59,6 +62,45 @@ void GPUDeviceManager::initialize() { std::cout << " SYCL backend disabled (compile with -DFUNGT_USE_SYCL=ON)" << std::endl; #endif + // Find OpenCL GPU devices. +#ifdef FUNGT_USE_OPENCL + cl_uint platformCount = 0; + if (clGetPlatformIDs(0, nullptr, &platformCount) == CL_SUCCESS) { + std::vector platforms(platformCount); + clGetPlatformIDs(platformCount, platforms.data(), nullptr); + + for (cl_platform_id platform : platforms) { + cl_uint deviceCount = 0; + if (clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 0, nullptr, &deviceCount) != CL_SUCCESS) { + continue; + } + + std::vector devices(deviceCount); + clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, deviceCount, devices.data(), nullptr); + for (cl_uint i = 0; i < deviceCount; ++i) { + char name[256] = {}; + char vendor[256] = {}; + cl_ulong memory = 0; + cl_uint computeUnits = 0; + clGetDeviceInfo(devices[i], CL_DEVICE_NAME, sizeof(name), name, nullptr); + clGetDeviceInfo(devices[i], CL_DEVICE_VENDOR, sizeof(vendor), vendor, nullptr); + clGetDeviceInfo(devices[i], CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(memory), &memory, nullptr); + clGetDeviceInfo(devices[i], CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(computeUnits), &computeUnits, nullptr); + + fungt::GPUDeviceInfo info; + info.id = static_cast(i); + info.name = name; + info.vendor = vendor; + info.memory_bytes = static_cast(memory); + info.compute_units = static_cast(computeUnits); + info.backend = fungt::GPUBackend::OPENCL; + info.isActive = false; + all_devices_.push_back(std::move(info)); + } + } + } +#endif + // Add OpenGL fallback device (always available) fungt::GPUDeviceInfo opengl_device; opengl_device.id = static_cast(all_devices_.size()); @@ -106,4 +148,4 @@ void GPUDeviceManager::setActiveDevice(int deviceIndex) { int GPUDeviceManager::getActiveDeviceIndex() const { return active_device_index_; -} \ No newline at end of file +} diff --git a/InfoDevice/gpu_device_info.hpp b/InfoDevice/gpu_device_info.hpp index 13dab23..3762a6f 100644 --- a/InfoDevice/gpu_device_info.hpp +++ b/InfoDevice/gpu_device_info.hpp @@ -13,6 +13,7 @@ namespace fungt { enum class GPUBackend { CUDA, SYCL, + OPENCL, OPENGL // Fallback }; @@ -40,6 +41,7 @@ namespace fungt { switch (backend) { case GPUBackend::CUDA: return "CUDA"; case GPUBackend::SYCL: return "SYCL"; + case GPUBackend::OPENCL: return "OpenCL"; case GPUBackend::OPENGL: return "OpenGL"; default: return "Unknown"; } @@ -88,4 +90,4 @@ class GPUDeviceManager { int active_device_index_; }; -#endif // _GPU_DEVICE_INFO_HPP_ \ No newline at end of file +#endif // _GPU_DEVICE_INFO_HPP_ diff --git a/PBR/Render/include/opencl_renderer.hpp b/PBR/Render/include/opencl_renderer.hpp index 766d13d..54f93e4 100644 --- a/PBR/Render/include/opencl_renderer.hpp +++ b/PBR/Render/include/opencl_renderer.hpp @@ -1,8 +1,8 @@ #if !defined(_OPENCL_RENDERER_H_) #define _OPENCL_RENDERER_H_ - #include // MUST be first! #include +#include #include #include #include @@ -60,6 +60,7 @@ class OpenCL_Renderer : public IComputeRenderer { cl_program m_oclprogram = nullptr; cl_kernel m_oclrenderKernel = nullptr; cl_kernel m_oclsampleKernel = nullptr; + cl_kernel m_ocl_ogldisplayKernel = nullptr; cl_mem m_texturesObj = nullptr; cl_mem m_textureDimensions = nullptr; std::size_t m_numTextures = 0; @@ -67,6 +68,14 @@ class OpenCL_Renderer : public IComputeRenderer { OpenCLRaySpaceBuffer m_raySpaceBuffer; bool m_sceneUploaded = false; + //For OpenGL interop / progressive path tracer (PBO approach) + cl_mem m_oglSharedBuffer = nullptr; + GLuint m_oglBufferID = 0; + int m_oglBufferWidth = 0; + int m_oglBufferHeight = 0; + cl_mem m_oclAccumulationBuffer = nullptr; + int m_accumulationWidth = 0; + int m_accumulationHeight = 0; void prepareTextures(); void setOpenCLTextures( const std::vector& handles, @@ -82,7 +91,7 @@ class OpenCL_Renderer : public IComputeRenderer { OpenCL_Renderer()= default; ~OpenCL_Renderer() override; - void initialize(); + void initialize(bool useOpenGLInterop = false); IDeviceTexture& textures() override { if (!m_textureManager) { @@ -104,7 +113,26 @@ class OpenCL_Renderer : public IComputeRenderer { int sampleOffset ) override; + //Maybe we need another render scene function + //that works with OpenGL interop + + void RenderSceneOpenGLInterop( + int width, + int height, + const std::vector& triangles, + const std::vector& nodes, + const std::vector& lights, + const std::vector& emissiveTriIndices, + const PBRCamera& camera, + int samplesPerPixel, + int sampleOffset, + GLuint glBufferID + ); + void releaseOpenGLInteropResources() noexcept; + void invalidateScene() { m_sceneUploaded = false; } + private: void buildOCLPrograms(); + void setKernelArg(cl_kernel kernel, cl_uint &argIndex,size_t argSize,const void* argValue); }; diff --git a/PBR/Render/ocl_kernels/path_tracer.cl b/PBR/Render/ocl_kernels/path_tracer.cl index e02faed..8b37dba 100644 --- a/PBR/Render/ocl_kernels/path_tracer.cl +++ b/PBR/Render/ocl_kernels/path_tracer.cl @@ -203,7 +203,7 @@ inline float3 fgt_path_trace_cook_torrance( } __kernel void path_tracer( - __global fgt_vec4* framebuffer, + __global float4* framebuffer, __global const fgt_triangle_geom* triangle_geometry, __global const fgt_triangle_shading* triangle_shading, __global const fgt_material_data* materials, @@ -263,7 +263,7 @@ __kernel void path_tracer( num_textures, &rng); - fgt_vec4 accumulated = framebuffer[pixel_index]; + float4 accumulated = framebuffer[pixel_index]; accumulated.x += contribution.x; accumulated.y += contribution.y; accumulated.z += contribution.z; diff --git a/PBR/Render/ocl_kernels/sample_framebuffer.cl b/PBR/Render/ocl_kernels/sample_framebuffer.cl index 367ab54..d2c1d8f 100644 --- a/PBR/Render/ocl_kernels/sample_framebuffer.cl +++ b/PBR/Render/ocl_kernels/sample_framebuffer.cl @@ -1,11 +1,25 @@ __kernel void sample_framebuffer( - __global fgt_vec4* framebuffer, + __global float4* framebuffer, int samples_per_pixel){ int idx = get_global_id(0); const float inverse_samples = 1.0f / (float)samples_per_pixel; - framebuffer[idx].x *= inverse_samples; - framebuffer[idx].y *= inverse_samples; - framebuffer[idx].z *= inverse_samples; - framebuffer[idx].w = 1.0f; + float4 pixel = framebuffer[idx]; + pixel.xyz *= inverse_samples; + pixel.w = 1.0f; + framebuffer[idx] = pixel; +} +__kernel void present_framebuffer( + __global const float4* src_buffer, + __global float4* dst_buffer, + int accumulated_samples) +{ + const int idx = get_global_id(0); + const float inv = 1.0f / (float)accumulated_samples; + const float4 pixel = src_buffer[idx]; + const float gamma = 1.0f / 2.2f; + float r = clamp(pixel.x * inv, 0.0f, 1.0f); + float g = clamp(pixel.y * inv, 0.0f, 1.0f); + float b = clamp(pixel.z * inv, 0.0f, 1.0f); + dst_buffer[idx] = (float4)(pow(r, gamma), pow(g, gamma), pow(b, gamma), 1.0f); } diff --git a/PBR/Render/src/opencl_renderer.cpp b/PBR/Render/src/opencl_renderer.cpp index 6465c9a..a467d64 100644 --- a/PBR/Render/src/opencl_renderer.cpp +++ b/PBR/Render/src/opencl_renderer.cpp @@ -3,12 +3,11 @@ #include #include -void OpenCL_Renderer::initialize() +void OpenCL_Renderer::initialize(bool useOpenGLInterop) { if (m_oclcontext && m_oclqueue && m_textureManager) { return; } - cl_uint platformCount = 0; cl_int err = clGetPlatformIDs(0, nullptr, &platformCount); if (err != CL_SUCCESS || platformCount == 0) { @@ -25,18 +24,76 @@ void OpenCL_Renderer::initialize() std::to_string(err)); } - for (cl_platform_id platform : platforms) { - if (!clGetExtensionFunctionAddressForPlatform( - platform, "clCreateImageWithPropertiesINTEL")) { - continue; + std::array contextProperties{}; + const cl_context_properties* properties = nullptr; + + if (useOpenGLInterop) { + const auto glContext = glXGetCurrentContext(); + const auto glDisplay = glXGetCurrentDisplay(); + if (!glContext || !glDisplay) { + throw std::runtime_error( + "OpenGL context is not current in this thread."); + } + + using GetGLContextInfo = cl_int (CL_API_CALL *)( + const cl_context_properties*, + cl_gl_context_info, + std::size_t, + void*, + std::size_t*); + + for (cl_platform_id platform : platforms) { + if (!clGetExtensionFunctionAddressForPlatform( + platform, "clCreateImageWithPropertiesINTEL")) { + continue; + } + + contextProperties = { + CL_GL_CONTEXT_KHR, + reinterpret_cast(glContext), + CL_GLX_DISPLAY_KHR, + reinterpret_cast(glDisplay), + CL_CONTEXT_PLATFORM, + reinterpret_cast(platform), + 0 + }; + + const auto getGLContextInfo = reinterpret_cast( + clGetExtensionFunctionAddressForPlatform( + platform, "clGetGLContextInfoKHR")); + if (!getGLContextInfo) { + continue; + } + + cl_device_id glDevice = nullptr; + err = getGLContextInfo( + contextProperties.data(), + CL_CURRENT_DEVICE_FOR_GL_CONTEXT_KHR, + sizeof(glDevice), + &glDevice, + nullptr); + if (err == CL_SUCCESS && glDevice) { + m_oclplatform = platform; + m_ocldevice = glDevice; + properties = contextProperties.data(); + break; + } } + } else { + for (cl_platform_id platform : platforms) { + if (!clGetExtensionFunctionAddressForPlatform( + platform, "clCreateImageWithPropertiesINTEL")) { + continue; + } - cl_device_id device = nullptr; - err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, nullptr); - if (err == CL_SUCCESS && device) { - m_oclplatform = platform; - m_ocldevice = device; - break; + cl_device_id device = nullptr; + err = clGetDeviceIDs( + platform, CL_DEVICE_TYPE_GPU, 1, &device, nullptr); + if (err == CL_SUCCESS && device) { + m_oclplatform = platform; + m_ocldevice = device; + break; + } } } @@ -55,9 +112,15 @@ void OpenCL_Renderer::initialize() sizeof(deviceName), deviceName, nullptr); std::cout << "OpenCL platform: " << platformName << std::endl; std::cout << "OpenCL device: " << deviceName << std::endl; - + if (useOpenGLInterop) { + char extensions[2048] = {}; + clGetDeviceInfo(m_ocldevice, CL_DEVICE_EXTENSIONS, sizeof(extensions), extensions, nullptr); + if (std::string(extensions).find("cl_khr_gl_sharing") == std::string::npos) { + throw std::runtime_error("OpenCL device does not support OpenGL interoperability."); + } + } m_oclcontext = clCreateContext( - nullptr, 1, &m_ocldevice, nullptr, nullptr, &err); + properties, 1, &m_ocldevice, nullptr, nullptr, &err); if (err != CL_SUCCESS || !m_oclcontext) { m_oclcontext = nullptr; throw std::runtime_error( @@ -83,6 +146,7 @@ void OpenCL_Renderer::initialize() // Defer kernel build until textures are loaded for bindless to work std::cout << "OpenCL context and command queue initialized (kernel build deferred)." << std::endl; + } OpenCL_Renderer::~OpenCL_Renderer() @@ -114,17 +178,34 @@ OpenCL_Renderer::~OpenCL_Renderer() clReleaseKernel(m_oclsampleKernel); m_oclsampleKernel = nullptr; } + if (m_ocl_ogldisplayKernel) { + clReleaseKernel(m_ocl_ogldisplayKernel); + m_ocl_ogldisplayKernel = nullptr; + } if (m_oclprogram) { clReleaseProgram(m_oclprogram); m_oclprogram = nullptr; } + if (m_oclAccumulationBuffer) { + clReleaseMemObject(m_oclAccumulationBuffer); + m_oclAccumulationBuffer = nullptr; + m_accumulationWidth = 0; + m_accumulationHeight = 0; + } + if (m_oglSharedBuffer) { + clReleaseMemObject(m_oglSharedBuffer); + m_oglSharedBuffer = nullptr; + m_oglBufferID = 0; + m_oglBufferWidth = 0; + m_oglBufferHeight = 0; + } + if (m_oclqueue) { clReleaseCommandQueue(m_oclqueue); m_oclqueue = nullptr; } - if (m_oclcontext) { clReleaseContext(m_oclcontext); m_oclcontext = nullptr; @@ -134,6 +215,28 @@ OpenCL_Renderer::~OpenCL_Renderer() m_oclplatform = nullptr; } +void OpenCL_Renderer::releaseOpenGLInteropResources() noexcept +{ + if (m_oclqueue) { + clFinish(m_oclqueue); + } + + if (m_oglSharedBuffer) { + clReleaseMemObject(m_oglSharedBuffer); + m_oglSharedBuffer = nullptr; + } + m_oglBufferID = 0; + m_oglBufferWidth = 0; + m_oglBufferHeight = 0; + + if (m_oclAccumulationBuffer) { + clReleaseMemObject(m_oclAccumulationBuffer); + m_oclAccumulationBuffer = nullptr; + } + m_accumulationWidth = 0; + m_accumulationHeight = 0; +} + std::vector OpenCL_Renderer::RenderScene(int width, int height, const std::vector& triangles, @@ -226,40 +329,28 @@ std::vector OpenCL_Renderer::RenderScene(int width, const cl_int numTextures = static_cast(m_numTextures); cl_uint argument = 0; - //Lambda to setkernel arguments - const auto setKernelArg = [&](cl_kernel kernel, std::size_t size, const void* value) { - const cl_int argumentError = - clSetKernelArg(kernel, argument++, size, value); - if (argumentError != CL_SUCCESS) { - throw std::runtime_error( - "OpenCL RenderScene failed to set kernel argument " + - std::to_string(argument - 1) + ", error " + - std::to_string(argumentError)); - } - }; - try { - setKernelArg(m_oclrenderKernel, sizeof(cl_mem), &framebufferBuffer); - setKernelArg(m_oclrenderKernel, sizeof(cl_mem), &m_raySpaceBuffer.triangleGeometry); - setKernelArg(m_oclrenderKernel, sizeof(cl_mem), &m_raySpaceBuffer.triangleShading); - setKernelArg(m_oclrenderKernel, sizeof(cl_mem), &m_raySpaceBuffer.materials); - setKernelArg(m_oclrenderKernel, sizeof(cl_mem), &m_raySpaceBuffer.bvhNodes); - setKernelArg(m_oclrenderKernel, sizeof(cl_mem), &m_raySpaceBuffer.lights); - setKernelArg(m_oclrenderKernel, sizeof(cl_mem), &m_raySpaceBuffer.emissiveTriangles); - setKernelArg(m_oclrenderKernel, sizeof(cl_mem), &m_texturesObj); - setKernelArg(m_oclrenderKernel, sizeof(cl_mem), &m_textureDimensions); - setKernelArg(m_oclrenderKernel, sizeof(cl_int), &numTriangles); - setKernelArg(m_oclrenderKernel, sizeof(cl_int), &numMaterials); - setKernelArg(m_oclrenderKernel, sizeof(cl_int), &numBVHNodes); - setKernelArg(m_oclrenderKernel, sizeof(cl_int), &numLights); - setKernelArg(m_oclrenderKernel, sizeof(cl_int), &numEmissiveTriangles); - setKernelArg(m_oclrenderKernel, sizeof(cl_int), &numTextures); - setKernelArg(m_oclrenderKernel, sizeof(cl_mem), &cameraBuffer); - setKernelArg(m_oclrenderKernel, sizeof(cl_int), &width); - setKernelArg(m_oclrenderKernel, sizeof(cl_int), &height); + setKernelArg(m_oclrenderKernel,argument, sizeof(cl_mem), &framebufferBuffer); + setKernelArg(m_oclrenderKernel,argument, sizeof(cl_mem), &m_raySpaceBuffer.triangleGeometry); + setKernelArg(m_oclrenderKernel,argument, sizeof(cl_mem), &m_raySpaceBuffer.triangleShading); + setKernelArg(m_oclrenderKernel,argument, sizeof(cl_mem), &m_raySpaceBuffer.materials); + setKernelArg(m_oclrenderKernel,argument, sizeof(cl_mem), &m_raySpaceBuffer.bvhNodes); + setKernelArg(m_oclrenderKernel,argument, sizeof(cl_mem), &m_raySpaceBuffer.lights); + setKernelArg(m_oclrenderKernel,argument, sizeof(cl_mem), &m_raySpaceBuffer.emissiveTriangles); + setKernelArg(m_oclrenderKernel,argument, sizeof(cl_mem), &m_texturesObj); + setKernelArg(m_oclrenderKernel,argument, sizeof(cl_mem), &m_textureDimensions); + setKernelArg(m_oclrenderKernel,argument, sizeof(cl_int), &numTriangles); + setKernelArg(m_oclrenderKernel,argument, sizeof(cl_int), &numMaterials); + setKernelArg(m_oclrenderKernel,argument, sizeof(cl_int), &numBVHNodes); + setKernelArg(m_oclrenderKernel,argument, sizeof(cl_int), &numLights); + setKernelArg(m_oclrenderKernel,argument, sizeof(cl_int), &numEmissiveTriangles); + setKernelArg(m_oclrenderKernel,argument, sizeof(cl_int), &numTextures); + setKernelArg(m_oclrenderKernel,argument, sizeof(cl_mem), &cameraBuffer); + setKernelArg(m_oclrenderKernel,argument, sizeof(cl_int), &width); + setKernelArg(m_oclrenderKernel,argument, sizeof(cl_int), &height); const cl_int initialSampleIndex = sampleOffset; - setKernelArg(m_oclrenderKernel, sizeof(cl_int), &initialSampleIndex); + setKernelArg(m_oclrenderKernel,argument,sizeof(cl_int), &initialSampleIndex); const std::size_t localSize[2] = {16, 16}; const std::size_t globalSize[2] = { @@ -283,8 +374,6 @@ std::vector OpenCL_Renderer::RenderScene(int width, std::to_string(err)); } - // No wait is needed - // The next sample runs after this sample. err = clEnqueueNDRangeKernel( m_oclqueue, m_oclrenderKernel, @@ -306,8 +395,8 @@ std::vector OpenCL_Renderer::RenderScene(int width, // Now add args for the second kernel. argument = 0; - setKernelArg(m_oclsampleKernel, sizeof(cl_mem), &framebufferBuffer); - setKernelArg(m_oclsampleKernel, sizeof(cl_int), &samplesPerPixel); + setKernelArg(m_oclsampleKernel, argument,sizeof(cl_mem), &framebufferBuffer); + setKernelArg(m_oclsampleKernel, argument, sizeof(cl_int), &samplesPerPixel); // Divide each pixel by samplesPerPixel. const std::size_t sampleGlobalSize[1] = {imageSize}; @@ -534,7 +623,6 @@ void OpenCL_Renderer::prepareTextures() "OpenCL textures requested before OpenCL initialization."); } - static bool kernelBuilt = false; if (m_textureManager->handlesAreDirty() || !m_texturesObj || !m_textureDimensions) { setOpenCLTextures( @@ -542,10 +630,9 @@ void OpenCL_Renderer::prepareTextures() m_textureManager->getTextureDimensions()); m_textureManager->markHandlesClean(); } - if (!kernelBuilt && m_textureManager->getTextureCount() > 0) { + if (!m_oclrenderKernel && m_textureManager->getTextureCount() > 0) { std::cout << "[prepareTextures] Building kernel AFTER textures loaded" << std::endl; buildOCLPrograms(); - kernelBuilt = true; } } @@ -595,6 +682,246 @@ void OpenCL_Renderer::setOpenCLTextures( m_numTextures = handles.size(); } +void OpenCL_Renderer::RenderSceneOpenGLInterop(int width, + int height, + const std::vector& triangles, + const std::vector& nodes, + const std::vector& lights, + const std::vector& emissiveTriIndices, + const PBRCamera& camera, int samplesPerPixel, + int sampleOffset, + GLuint glBufferID) +{ + if (width <= 0 || height <= 0) { + throw std::invalid_argument( + "OpenCL RenderSceneOpenGLInterop requires positive image dimensions."); + } + if (samplesPerPixel <= 0) { + throw std::invalid_argument( + "OpenCL RenderSceneOpenGLInterop requires at least one sample."); + } + if (!m_oclcontext || !m_oclqueue) { + throw std::runtime_error( + "OpenCL RenderSceneOpenGLInterop called before renderer initialization."); + } + if (glBufferID == 0) { + throw std::invalid_argument( + "OpenCL RenderSceneOpenGLInterop requires a valid OpenGL buffer."); + } + + prepareTextures(); + + if (!m_sceneUploaded) { + uploadScene(triangles, nodes, lights, emissiveTriIndices); + clFinish(m_oclqueue); + } + + if (!m_oglSharedBuffer || + m_oglBufferID != glBufferID || + m_oglBufferWidth != width || + m_oglBufferHeight != height) { + if (m_oglSharedBuffer) { + clReleaseMemObject(m_oglSharedBuffer); + m_oglSharedBuffer = nullptr; + m_oglBufferID = 0; + m_oglBufferWidth = 0; + m_oglBufferHeight = 0; + } + cl_int err = CL_SUCCESS; + m_oglSharedBuffer = clCreateFromGLBuffer( + m_oclcontext, + CL_MEM_WRITE_ONLY, + glBufferID, + &err); + if (err != CL_SUCCESS || !m_oglSharedBuffer) { + m_oglSharedBuffer = nullptr; + throw std::runtime_error( + "Failed to register OpenGL buffer with OpenCL, error " + + std::to_string(err)); + } + m_oglBufferID = glBufferID; + m_oglBufferWidth = width; + m_oglBufferHeight = height; + } + //Get camera data + const fgt_rayspace_camera cameraData = + fgt::opencl::translate_rayspace_camera(camera); + + const std::size_t imageSize = + static_cast(width) * static_cast(height); + + cl_int err = CL_SUCCESS; + bool accumulationBufferCreated = false; + if (!m_oclAccumulationBuffer || + m_accumulationWidth != width || + m_accumulationHeight != height) { + if (m_oclAccumulationBuffer) { + clReleaseMemObject(m_oclAccumulationBuffer); + m_oclAccumulationBuffer = nullptr; + } + + m_oclAccumulationBuffer = clCreateBuffer( + m_oclcontext, + CL_MEM_READ_WRITE, + imageSize * sizeof(fgt_vec4), + nullptr, + &err); + if (err != CL_SUCCESS || !m_oclAccumulationBuffer) { + throw std::runtime_error( + "OpenCL RenderSceneOpenGLInterop failed to create accumulation buffer, error " + + std::to_string(err)); + } + m_accumulationWidth = width; + m_accumulationHeight = height; + accumulationBufferCreated = true; + } + + // Clear the sum when a new accumulation starts. + if (accumulationBufferCreated || sampleOffset == 0) { + const fgt_vec4 zeroPixel{0.0f, 0.0f, 0.0f, 0.0f}; + err = clEnqueueFillBuffer( + m_oclqueue, + m_oclAccumulationBuffer, + &zeroPixel, + sizeof(zeroPixel), + 0, + imageSize * sizeof(fgt_vec4), + 0, + nullptr, + nullptr); + if (err != CL_SUCCESS) { + throw std::runtime_error( + "OpenCL RenderSceneOpenGLInterop failed to clear accumulation buffer, error " + + std::to_string(err)); + } + } + + cl_mem cameraBuffer = clCreateBuffer( + m_oclcontext, + CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + sizeof(fgt_rayspace_camera), + const_cast(&cameraData), + &err); + if (err != CL_SUCCESS || !cameraBuffer) { + throw std::runtime_error( + "OpenCL RenderSceneOpenGLInterop failed to create camera buffer, error " + + std::to_string(err)); + } + + const cl_int numTriangles = static_cast(m_raySpaceBuffer.numTriangles); + const cl_int numMaterials = static_cast(m_raySpaceBuffer.numMaterials); + const cl_int numBVHNodes = static_cast(m_raySpaceBuffer.numBVHNodes); + const cl_int numLights = static_cast(m_raySpaceBuffer.numLights); + const cl_int numEmissiveTriangles = static_cast(m_raySpaceBuffer.numEmissiveTriangles); + const cl_int numTextures = static_cast(m_numTextures); + cl_uint argument = 0; + setKernelArg(m_oclrenderKernel, argument, sizeof(cl_mem), &m_oclAccumulationBuffer); + setKernelArg(m_oclrenderKernel, argument, sizeof(cl_mem), &m_raySpaceBuffer.triangleGeometry); + setKernelArg(m_oclrenderKernel, argument, sizeof(cl_mem), &m_raySpaceBuffer.triangleShading); + setKernelArg(m_oclrenderKernel, argument, sizeof(cl_mem), &m_raySpaceBuffer.materials); + setKernelArg(m_oclrenderKernel, argument, sizeof(cl_mem), &m_raySpaceBuffer.bvhNodes); + setKernelArg(m_oclrenderKernel, argument, sizeof(cl_mem), &m_raySpaceBuffer.lights); + setKernelArg(m_oclrenderKernel, argument, sizeof(cl_mem), &m_raySpaceBuffer.emissiveTriangles); + setKernelArg(m_oclrenderKernel, argument, sizeof(cl_mem), &m_texturesObj); + setKernelArg(m_oclrenderKernel, argument, sizeof(cl_mem), &m_textureDimensions); + setKernelArg(m_oclrenderKernel, argument, sizeof(cl_int), &numTriangles); + setKernelArg(m_oclrenderKernel, argument, sizeof(cl_int), &numMaterials); + setKernelArg(m_oclrenderKernel, argument, sizeof(cl_int), &numBVHNodes); + setKernelArg(m_oclrenderKernel, argument, sizeof(cl_int), &numLights); + setKernelArg(m_oclrenderKernel, argument, sizeof(cl_int), &numEmissiveTriangles); + setKernelArg(m_oclrenderKernel, argument, sizeof(cl_int), &numTextures); + setKernelArg(m_oclrenderKernel, argument, sizeof(cl_mem), &cameraBuffer); + setKernelArg(m_oclrenderKernel, argument, sizeof(cl_int), &width); + setKernelArg(m_oclrenderKernel, argument, sizeof(cl_int), &height); + + const std::size_t localSize[2] = {16, 16}; + const std::size_t globalSize[2] = { + (static_cast(width) + localSize[0] - 1) / + localSize[0] * localSize[0], + (static_cast(height) + localSize[1] - 1) / + localSize[1] * localSize[1] + }; + + // Add each new sample to the accumulation buffer. + for (int sample = 0; sample < samplesPerPixel; ++sample) { + const cl_int sampleIndex = sampleOffset + sample; + err = clSetKernelArg( + m_oclrenderKernel, + argument, + sizeof(sampleIndex), + &sampleIndex); + if (err != CL_SUCCESS) { + clReleaseMemObject(cameraBuffer); + throw std::runtime_error( + "OpenCL RenderSceneOpenGLInterop failed to set sample index, error " + + std::to_string(err)); + } + + err = clEnqueueNDRangeKernel( + m_oclqueue, + m_oclrenderKernel, + 2, + nullptr, + globalSize, + localSize, + 0, + nullptr, + nullptr); + if (err != CL_SUCCESS) { + clReleaseMemObject(cameraBuffer); + throw std::runtime_error( + "OpenCL RenderSceneOpenGLInterop failed to enqueue path tracer, error " + + std::to_string(err)); + } + } + + clReleaseMemObject(cameraBuffer); + + // Acquire the shared GL buffer + glFinish(); + err = clEnqueueAcquireGLObjects( + m_oclqueue, 1, &m_oglSharedBuffer, 0, nullptr, nullptr); + if (err != CL_SUCCESS) { + throw std::runtime_error( + "OpenCL failed to acquire the OpenGL buffer, error " + + std::to_string(err)); + } + + // Copy and normalize from accumulation buffer to shared buffer + const cl_int accumulatedSamples = sampleOffset + samplesPerPixel; + const size_t pixelCount = static_cast(width) * static_cast(height); + cl_uint arg = 0; + setKernelArg(m_ocl_ogldisplayKernel, arg, sizeof(cl_mem), &m_oclAccumulationBuffer); + setKernelArg(m_ocl_ogldisplayKernel, arg, sizeof(cl_mem), &m_oglSharedBuffer); + setKernelArg(m_ocl_ogldisplayKernel, arg, sizeof(cl_int), &accumulatedSamples); + + err = clEnqueueNDRangeKernel( + m_oclqueue, m_ocl_ogldisplayKernel, 1, nullptr, + &pixelCount, nullptr, 0, nullptr, nullptr); + if (err != CL_SUCCESS) { + clEnqueueReleaseGLObjects( + m_oclqueue, 1, &m_oglSharedBuffer, 0, nullptr, nullptr); + clFinish(m_oclqueue); + throw std::runtime_error( + "OpenCL failed to run display kernel, error " + std::to_string(err)); + } + + err = clEnqueueReleaseGLObjects( + m_oclqueue, 1, &m_oglSharedBuffer, 0, nullptr, nullptr); + if (err != CL_SUCCESS) { + throw std::runtime_error( + "OpenCL failed to release the OpenGL buffer, error " + + std::to_string(err)); + } + err = clFinish(m_oclqueue); + if (err != CL_SUCCESS) { + throw std::runtime_error( + "OpenCL failed while finishing interop commands, error " + + std::to_string(err)); + } + +} + void OpenCL_Renderer::buildOCLPrograms() { if (!m_oclcontext) { @@ -686,4 +1013,32 @@ void OpenCL_Renderer::buildOCLPrograms() throw std::runtime_error("Failed to create OpenCL sample_framebuffer, error " + std::to_string(err)); } + m_ocl_ogldisplayKernel = clCreateKernel(m_oclprogram, "present_framebuffer", &err); + if (err != CL_SUCCESS || !m_ocl_ogldisplayKernel) { + m_ocl_ogldisplayKernel = nullptr; + clReleaseKernel(m_oclsampleKernel); + m_oclsampleKernel = nullptr; + clReleaseKernel(m_oclrenderKernel); + m_oclrenderKernel = nullptr; + clReleaseProgram(m_oclprogram); + m_oclprogram = nullptr; + throw std::runtime_error("Failed to create OpenCL present_framebuffer, error " + std::to_string(err)); + } + +} + +void OpenCL_Renderer::setKernelArg( + cl_kernel kernel, + cl_uint& argIndex, + size_t argSize, + const void* argValue) +{ + const cl_int argumentError = + clSetKernelArg(kernel, argIndex++, argSize, argValue); + if (argumentError != CL_SUCCESS) { + throw std::runtime_error( + "OpenCL RenderScene failed to set kernel argument " + + std::to_string(argIndex - 1) + ", error " + + std::to_string(argumentError)); + } } diff --git a/PBR/Space/space.cpp b/PBR/Space/space.cpp index a906cc6..169a6e6 100644 --- a/PBR/Space/space.cpp +++ b/PBR/Space/space.cpp @@ -106,7 +106,7 @@ std::vector Space::Render(const int width, const int height,int sam return frameBuffer; } -void Space::InitComputeRenderBackend() +void Space::InitComputeRenderBackend(bool oglInterop) { if (!m_computeRenderer) { std::cout << "Error: Renderer not created!" << std::endl; @@ -156,7 +156,7 @@ void Space::InitComputeRenderBackend() if (!openclRenderer) { throw std::runtime_error("OpenCL renderer was not created."); } - openclRenderer->initialize(); + openclRenderer->initialize(oglInterop); break; } #endif @@ -166,6 +166,46 @@ void Space::InitComputeRenderBackend() } } +void Space::RenderOpenGLInterop( + int width, + int height, + int sampleOffset, + uint32_t glBufferID) +{ +#ifdef FUNGT_USE_OPENCL + auto* openclRenderer = dynamic_cast(m_computeRenderer.get()); + if (!openclRenderer) { + throw std::runtime_error( + "OpenGL interop requires the OpenCL renderer."); + } + + openclRenderer->RenderSceneOpenGLInterop( + width, + height, + m_triangles, + m_bvh_nodes, + m_lights, + m_emissiveTriIndices, + m_camera, + m_samplesPerPixel, + sampleOffset, + glBufferID); +#else + throw std::runtime_error( + "OpenCL support is not enabled."); +#endif +} + +void Space::ReleaseOpenGLInteropResources() +{ +#ifdef FUNGT_USE_OPENCL + auto* openclRenderer = dynamic_cast(m_computeRenderer.get()); + if (openclRenderer) { + openclRenderer->releaseOpenGLInteropResources(); + } +#endif +} + void Space::LoadModelToRender(const SimpleModel& Simplemodel) { Model& model = Simplemodel.getModel(); @@ -358,6 +398,19 @@ void Space::loadLightsFromScene(const std::vector& sceneLights) m_lights.push_back(Light(pos, intensity, type)); } } + +void Space::invalidateScene() +{ +#ifdef FUNGT_USE_OPENCL + if (ComputeRender::GetBackend() == Compute::Backend::OPENCL) { + auto* openclRenderer = dynamic_cast(m_computeRenderer.get()); + if (openclRenderer) { + openclRenderer->invalidateScene(); + } + } +#endif +} + void Space::SaveFrameBufferAsPNG(const std::vector& framebuffer, int width, int height) { std::vector pixels(width * height * 3); @@ -370,10 +423,14 @@ void Space::SaveFrameBufferAsPNG(const std::vector& framebuffer, in fungt::Vec3 color = framebuffer[srcIdx]; - // Clamp and gamma correct - color.x = std::pow(std::clamp(color.x, 0.0f, 1.0f), 1.0f / 2.2f); - color.y = std::pow(std::clamp(color.y, 0.0f, 1.0f), 1.0f / 2.2f); - color.z = std::pow(std::clamp(color.z, 0.0f, 1.0f), 1.0f / 2.2f); + //if (Compute::Backend::OPENCL != ComputeRender::GetBackend()) { + + // Clamp and gamma correct + color.x = std::pow(std::clamp(color.x, 0.0f, 1.0f), 1.0f / 2.2f); + color.y = std::pow(std::clamp(color.y, 0.0f, 1.0f), 1.0f / 2.2f); + color.z = std::pow(std::clamp(color.z, 0.0f, 1.0f), 1.0f / 2.2f); + //} + pixels[dstIdx * 3 + 0] = static_cast(255.99f * color.x); pixels[dstIdx * 3 + 1] = static_cast(255.99f * color.y); diff --git a/PBR/Space/space.hpp b/PBR/Space/space.hpp index 0410a84..4a460b8 100644 --- a/PBR/Space/space.hpp +++ b/PBR/Space/space.hpp @@ -1,5 +1,6 @@ #if !defined(_SPACE_H_) #define _SPACE_H_ +#include #include #include "Triangle/triangle.hpp" @@ -38,16 +39,24 @@ class Space { std::vector Render(const int width, const int height,int sampleOffset = 0); - void InitComputeRenderBackend(); + void InitComputeRenderBackend(bool oglInterop = false); + void RenderOpenGLInterop( + int width, + int height, + int sampleOffset, + uint32_t glBufferID); + void ReleaseOpenGLInteropResources(); void LoadModelToRender(const SimpleModel& model); void LoadGeometryToRender(const SimpleGeometry& geometry); void loadLightsFromScene(const std::vector& sceneLights); + void invalidateScene(); void static SaveFrameBufferAsPNG(const std::vector& framebuffer, int width, int height); static void SaveFrameBufferAsPNG(const std::vector& framebuffer, int width, int height, const std::string& filename); void BuildBVH(); void setSamples(int numOfSamples); + void setCamera(const PBRCamera& camera) { m_camera = camera; } void ClearSpace() { m_triangles.clear(); m_bvh_nodes.clear(); diff --git a/PBR/main/CMakeLists.txt b/PBR/main/CMakeLists.txt index e356a75..adbaabb 100644 --- a/PBR/main/CMakeLists.txt +++ b/PBR/main/CMakeLists.txt @@ -4,8 +4,9 @@ project(pbr_libraries) option(FUNGT_USE_CUDA "Enable CUDA backend" ON) option(FUNGT_USE_SYCL "Enable SYCL backend" ON) option(FUNGT_USE_OPENCL "Enable OpenCL backend" ON) -set(SYCL_CUDA_ARCH sm_75 CACHE STRING "NVIDIA GPU architecture for SYCL CUDA backend") -set_property(CACHE SYCL_CUDA_ARCH PROPERTY STRINGS sm_50 sm_52 sm_53 sm_60 sm_61 sm_62 sm_70 sm_72 sm_75 sm_80 sm_86 sm_87 sm_89 sm_90) +set(FUNGT_CUDA_ARCH 89 CACHE STRING "NVIDIA GPU compute capability") +set_property(CACHE FUNGT_CUDA_ARCH PROPERTY STRINGS 50 52 53 60 61 62 70 72 75 80 86 87 89 90) +set(SYCL_CUDA_ARCH "sm_${FUNGT_CUDA_ARCH}") # ──────────────────────────────────────────────────────── # FUNGT_BASE_DIR validation @@ -209,7 +210,7 @@ if(FUNGT_USE_CUDA) set_target_properties(pbr_cuda PROPERTIES CUDA_SEPARABLE_COMPILATION ON - CUDA_ARCHITECTURES "native" + CUDA_ARCHITECTURES ${FUNGT_CUDA_ARCH} ) target_compile_definitions(pbr_cuda PUBLIC FUNGT_USE_CUDA) @@ -382,7 +383,7 @@ message(STATUS "Default compiler: Bundled Intel clang (SYCL-capable)") if(FUNGT_USE_CUDA) message(STATUS "CUDA backend: ENABLED (nvcc)") message(STATUS "CUDA includes: ${CUDAToolkit_INCLUDE_DIRS}") - message(STATUS "CUDA architecture: ${SYCL_CUDA_ARCH}") + message(STATUS "CUDA architecture: ${FUNGT_CUDA_ARCH}") else() message(STATUS "CUDA backend: DISABLED") endif() diff --git a/PBR/standalone/CMakeLists.txt b/PBR/standalone/CMakeLists.txt index 7e4bdb3..ef63b10 100644 --- a/PBR/standalone/CMakeLists.txt +++ b/PBR/standalone/CMakeLists.txt @@ -14,7 +14,9 @@ set(FUNGT_BASE_DIR option(FUNGT_USE_CUDA "Link the CUDA PBR backend" ON) option(FUNGT_USE_SYCL "Link the SYCL PBR backend" ON) option(FUNGT_USE_OPENCL "Link the OpenCL PBR backend" ON) -set(SYCL_CUDA_ARCH sm_75 CACHE STRING "NVIDIA GPU architecture for SYCL") +set(FUNGT_CUDA_ARCH 89 CACHE STRING "NVIDIA GPU compute capability") +set_property(CACHE FUNGT_CUDA_ARCH PROPERTY STRINGS 50 52 53 60 61 62 70 72 75 80 86 87 89 90) +set(SYCL_CUDA_ARCH "sm_${FUNGT_CUDA_ARCH}") set(SYCL_COMPILER_PATH "/home/juanchuletas/Documents/Development/sycl_workspace/llvm/build/bin/clang++" @@ -99,7 +101,7 @@ add_executable(pbr_standalone main.cpp) if(FUNGT_USE_CUDA) set_target_properties(pbr_standalone PROPERTIES - CUDA_ARCHITECTURES "native" + CUDA_ARCHITECTURES ${FUNGT_CUDA_ARCH} CUDA_SEPARABLE_COMPILATION ON CUDA_RESOLVE_DEVICE_SYMBOLS ON ) diff --git a/Samples/iwocl/CMakeLists.txt b/Samples/iwocl/CMakeLists.txt index 104b542..92c5d9d 100644 --- a/Samples/iwocl/CMakeLists.txt +++ b/Samples/iwocl/CMakeLists.txt @@ -6,8 +6,10 @@ project(FunGT) # ════════════════════════════════════════════════════════════════════════════ option(FUNGT_USE_CUDA "Enable CUDA backend for NVIDIA GPUs" ON) option(FUNGT_USE_SYCL "Enable SYCL backend for Intel GPUs" ON) -set(SYCL_CUDA_ARCH sm_75 CACHE STRING "NVIDIA GPU architecture for SYCL CUDA backend") -set_property(CACHE SYCL_CUDA_ARCH PROPERTY STRINGS sm_50 sm_52 sm_53 sm_60 sm_61 sm_62 sm_70 sm_72 sm_75 sm_80 sm_86 sm_87 sm_89 sm_90) +option(FUNGT_USE_OPENCL "Enable OpenCL backend" ON) +set(FUNGT_CUDA_ARCH 89 CACHE STRING "NVIDIA GPU compute capability") +set_property(CACHE FUNGT_CUDA_ARCH PROPERTY STRINGS 50 52 53 60 61 62 70 72 75 80 86 87 89 90) +set(SYCL_CUDA_ARCH "sm_${FUNGT_CUDA_ARCH}") # Allow FUNGT_BASE_DIR to be set externally (command line, environment, or cache) # Priority: 1. Cache variable, 2. Environment variable, 3. Default fallback @@ -145,6 +147,15 @@ if(UNIX) set(PBR_SYCL_LIB pbr_sycl) message(STATUS "PBR SYCL renderer library: ${PBR_LIB_DIR}/libpbr_sycl.a") endif() + + if(FUNGT_USE_OPENCL) + add_library(pbr_opencl STATIC IMPORTED) + set_target_properties(pbr_opencl PROPERTIES + IMPORTED_LOCATION "${PBR_LIB_DIR}/libpbr_opencl.a" + ) + set(PBR_OPENCL_LIB pbr_opencl) + message(STATUS "PBR OpenCL renderer library: ${PBR_LIB_DIR}/libpbr_opencl.a") + endif() endif() if (WIN32) @@ -372,7 +383,7 @@ if(FUNGT_USE_CUDA) set_target_properties(FunGT PROPERTIES CUDA_RESOLVE_DEVICE_SYMBOLS ON CUDA_SEPARABLE_COMPILATION ON - CUDA_ARCHITECTURES "native" + CUDA_ARCHITECTURES ${FUNGT_CUDA_ARCH} ) endif() # ════════════════════════════════════════════════════════════════════════════ @@ -491,11 +502,12 @@ elseif (UNIX) dl funlib imgui - ${OpenCL_LIBRARIES} ${CUDA_BACKEND_LIB} pbr_core ${PBR_CUDA_LIB} ${PBR_SYCL_LIB} + ${PBR_OPENCL_LIB} + ${OpenCL_LIBRARIES} $<$:CUDA::cudart_static> $<$:CUDA::cuda_driver> ) @@ -533,6 +545,9 @@ elseif (UNIX) if(FUNGT_USE_SYCL) target_compile_definitions(FunGT PRIVATE FUNGT_USE_SYCL) endif() + if(FUNGT_USE_OPENCL) + target_compile_definitions(FunGT PRIVATE FUNGT_USE_OPENCL) + endif() endif() @@ -543,12 +558,10 @@ message(STATUS "") message(STATUS "════════════════════════════════════════") message(STATUS " FunGT Build Configuration") message(STATUS "════════════════════════════════════════") -message(STATUS " Main Compiler: ${CMAKE_CXX_COMPILER}") +message(STATUS " Main Compiler: Intel DPC++") message(STATUS " CUDA Backend: ${FUNGT_USE_CUDA}") message(STATUS " SYCL Backend: ${FUNGT_USE_SYCL}") +message(STATUS " OpenCL Backend: ${FUNGT_USE_OPENCL}") message(STATUS " Build Type: ${CMAKE_BUILD_TYPE}") -if(FUNGT_USE_CUDA) - message(STATUS " CUDA Compiled: Separately with g++") -endif() message(STATUS "════════════════════════════════════════") -message(STATUS "") \ No newline at end of file +message(STATUS "") diff --git a/Samples/iwocl/main.cpp b/Samples/iwocl/main.cpp index c7f7ce3..55a6ac6 100644 --- a/Samples/iwocl/main.cpp +++ b/Samples/iwocl/main.cpp @@ -9,7 +9,7 @@ int main() { //Path to your models: ModelPaths iwocl; - iwocl.path = getAssetPath("demo_assets/iwocl/iwocl.obj"); + iwocl.path = getAssetPath("assets_local/opencl_logo2/scene.gltf"); //Backend: DisplayGraphics::SetBackend(Backend::OpenGL); @@ -25,7 +25,7 @@ int main() { // Creates an animation object FunGTSModel iwocl_model = SimpleModel::create(); iwocl_model->load(iwocl); - iwocl_model->position(0.f, 0.f, 0.f); + iwocl_model->position(0.f, 5.f, 0.f); iwocl_model->rotation(0.f, 0.f, 0.f); iwocl_model->scale(5.0); diff --git a/ViewPort/opengl/opengl_progressive_path_tracer.cpp b/ViewPort/opengl/opengl_progressive_path_tracer.cpp index 25131eb..02083ea 100644 --- a/ViewPort/opengl/opengl_progressive_path_tracer.cpp +++ b/ViewPort/opengl/opengl_progressive_path_tracer.cpp @@ -40,3 +40,18 @@ void OpenGLProgressivePathTracer::renderSample(int sample, uint32_t targetTextur GL_RGBA, GL_FLOAT, gammaCorrected.data()); glBindTexture(GL_TEXTURE_2D, 0); } + +void OpenGLProgressivePathTracer::renderSampleInterop(int sample, uint32_t targetPBO) +{ + if (!m_initialized || !m_space) { + std::cerr << "Path tracer not initialized!" << std::endl; + return; + } + + m_space->setSamples(1); + m_space->RenderOpenGLInterop( + m_width, + m_height, + sample, + targetPBO); +} diff --git a/ViewPort/opengl/opengl_progressive_path_tracer.hpp b/ViewPort/opengl/opengl_progressive_path_tracer.hpp index 0a18227..38a00cf 100644 --- a/ViewPort/opengl/opengl_progressive_path_tracer.hpp +++ b/ViewPort/opengl/opengl_progressive_path_tracer.hpp @@ -10,6 +10,7 @@ class OpenGLProgressivePathTracer : public ProgressivePathTracer { ~OpenGLProgressivePathTracer() = default; void renderSample(int sample, uint32_t targetTexture) override; + void renderSampleInterop(int sample, uint32_t targetPBO) override; }; #endif // OPENGL_PROGRESSIVE_PATH_TRACER_HPP diff --git a/ViewPort/opengl/opengl_viewport.cpp b/ViewPort/opengl/opengl_viewport.cpp index 81224f0..9581a47 100644 --- a/ViewPort/opengl/opengl_viewport.cpp +++ b/ViewPort/opengl/opengl_viewport.cpp @@ -21,6 +21,11 @@ void OpenGLViewPort::onAttach() glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_EDGE); glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_EDGE); glBindTexture(GL_TEXTURE_2D, 0); + + glGenBuffers(1, &m_pathTracePBO); + glBindBuffer(GL_PIXEL_UNPACK_BUFFER, m_pathTracePBO); + glBufferData(GL_PIXEL_UNPACK_BUFFER, 1280 * 720 * 4 * sizeof(float), nullptr, GL_DYNAMIC_DRAW); + glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0); } void OpenGLViewPort::onDetach() @@ -29,6 +34,10 @@ void OpenGLViewPort::onDetach() glDeleteTextures(1, &m_pathTraceTexture); m_pathTraceTexture = 0; } + if (m_pathTracePBO != 0) { + glDeleteBuffers(1, &m_pathTracePBO); + m_pathTracePBO = 0; + } } void OpenGLViewPort::onUpdate() @@ -60,11 +69,11 @@ void OpenGLViewPort::onUpdate() m_frameBuffer->unbind(); - if (m_pathTraceMode && m_PathTraceFunc && m_currentSample < m_maxPreviewSamples) { + if (m_pathTraceMode && m_ActivePathTraceFunc && m_currentSample < m_maxPreviewSamples) { int width = static_cast(m_viewportSize.x); int height = static_cast(m_viewportSize.y); - m_PathTraceFunc(width, height, m_currentSample); + m_ActivePathTraceFunc(width, height, m_currentSample); m_currentSample++; if (m_currentSample <= 5 || m_currentSample == m_maxPreviewSamples) { @@ -92,28 +101,44 @@ void OpenGLViewPort::onImGuiRender() if (diffX > 1.0f || diffY > 1.0f) { if (viewportPanelSize.x > 32 && viewportPanelSize.y > 32) { - m_viewportSize = viewportPanelSize; - pendingSize = viewportPanelSize; - lastResizeRequest = currentTime; - pendingResize = true; + const float pendingDiffX = std::abs(viewportPanelSize.x - pendingSize.x); + const float pendingDiffY = std::abs(viewportPanelSize.y - pendingSize.y); + if (!pendingResize || pendingDiffX > 1.0f || pendingDiffY > 1.0f) { + pendingSize = viewportPanelSize; + lastResizeRequest = currentTime; + pendingResize = true; + } } } bool isResizing = ImGui::IsMouseDragging(ImGuiMouseButton_Left); if (pendingResize && !isResizing && (currentTime - lastResizeRequest) > 0.25) { + // Release the OpenCL wrapper before changing the GL texture storage. + if (m_PathTraceReleaseInteropFunc) { + m_PathTraceReleaseInteropFunc(); + } + FrameBuffSpec spec{ - static_cast(m_viewportSize.x), - static_cast(m_viewportSize.y), + static_cast(pendingSize.x), + static_cast(pendingSize.y), 1 }; m_resizeBuffer = FrameBuffer::create(spec); glBindTexture(GL_TEXTURE_2D, m_pathTraceTexture); glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA32F, - static_cast(m_viewportSize.x), - static_cast(m_viewportSize.y), + static_cast(pendingSize.x), + static_cast(pendingSize.y), 0, GL_RGBA, GL_FLOAT, nullptr); glBindTexture(GL_TEXTURE_2D, 0); + + glBindBuffer(GL_PIXEL_UNPACK_BUFFER, m_pathTracePBO); + glBufferData(GL_PIXEL_UNPACK_BUFFER, + static_cast(pendingSize.x) * static_cast(pendingSize.y) * 4 * sizeof(float), + nullptr, GL_DYNAMIC_DRAW); + glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0); + + m_viewportSize = pendingSize; m_currentSample = 0; pendingResize = false; } @@ -124,18 +149,37 @@ void OpenGLViewPort::onImGuiRender() else texID = m_frameBuffer->GetColorAttachmentRendererID(); - if (m_pathTraceMode) { - ImGui::SetCursorPos(ImVec2(10, 30)); - ImGui::TextColored(ImVec4(1.0f, 0.8f, 0.0f, 1.0f), - "RaySpace: %d/%d samples", m_currentSample, m_maxPreviewSamples); - ImGui::TextColored(ImVec4(1.0f, 0.8f, 0.0f, 1.0f), " Using SYCL on iGPU"); - } - + const ImVec2 imagePosition = ImGui::GetCursorScreenPos(); ImGui::Image((void*)(intptr_t)texID, - m_viewportSize, + viewportPanelSize, ImVec2{ 0, 1 }, ImVec2{ 1, 0 }); + if (m_pathTraceMode) { + const ImU32 textColor = ImGui::ColorConvertFloat4ToU32( + ImVec4(1.0f, 0.8f, 0.0f, 1.0f)); + auto* drawList = ImGui::GetWindowDrawList(); + drawList->AddText( + ImVec2(imagePosition.x + 10.0f, imagePosition.y + 10.0f), + textColor, + ("RaySpace: " + std::to_string(m_currentSample) + "/" + + std::to_string(m_maxPreviewSamples) + " samples").c_str()); + drawList->AddText( + ImVec2(imagePosition.x + 10.0f, imagePosition.y + 30.0f), + textColor, + ("Using " + ComputeRender::GetBackendName()).c_str()); + } + ImGui::End(); ImGui::PopStyleColor(); ImGui::PopStyleVar(); } + +void OpenGLViewPort::copyPBOToTexture(int width, int height) +{ + glBindBuffer(GL_PIXEL_UNPACK_BUFFER, m_pathTracePBO); + glBindTexture(GL_TEXTURE_2D, m_pathTraceTexture); + glTexSubImage2D(GL_TEXTURE_2D, 0, 0, 0, width, height, + GL_RGBA, GL_FLOAT, nullptr); + glBindTexture(GL_TEXTURE_2D, 0); + glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0); +} diff --git a/ViewPort/opengl/opengl_viewport.hpp b/ViewPort/opengl/opengl_viewport.hpp index 7fa4dd2..d2433fd 100644 --- a/ViewPort/opengl/opengl_viewport.hpp +++ b/ViewPort/opengl/opengl_viewport.hpp @@ -11,6 +11,7 @@ class OpenGLViewPort : public ViewPort { std::shared_ptr m_frameBuffer; std::shared_ptr m_resizeBuffer; GLuint m_pathTraceTexture = 0; + GLuint m_pathTracePBO = 0; public: OpenGLViewPort(); @@ -22,6 +23,8 @@ class OpenGLViewPort : public ViewPort { void onImGuiRender() override; uint32_t getPathTraceTexture() const override { return m_pathTraceTexture; } + uint32_t getPathTracePBO() const override { return m_pathTracePBO; } + void copyPBOToTexture(int width, int height) override; }; #endif // _OPENGL_VIEWPORT_H_ diff --git a/ViewPort/progressive_path_tracer_viewport.cpp b/ViewPort/progressive_path_tracer_viewport.cpp index c7448e3..78e9325 100644 --- a/ViewPort/progressive_path_tracer_viewport.cpp +++ b/ViewPort/progressive_path_tracer_viewport.cpp @@ -17,10 +17,9 @@ std::unique_ptr ProgressivePathTracer::create() void ProgressivePathTracer::initialize(Camera* viewportCam, std::shared_ptr sceneManager, - int width, int height) + int width, int height, bool ogl_interop) { std::cout << "Initializing progressive path tracer: " << width << "x" << height << std::endl; - ComputeRender::SetBackend(Compute::Backend::SYCL); std::cout << "Using backend: " << ComputeRender::GetBackendName() << std::endl; m_width = width; m_height = height; @@ -42,7 +41,7 @@ void ProgressivePathTracer::initialize(Camera* viewportCam, PBRCamera pbrCam(pbrPos, pbrLookAt, pbrUp, fov, aspect); m_space = std::make_unique(pbrCam); - m_space->InitComputeRenderBackend(); + m_space->InitComputeRenderBackend(ogl_interop); m_space->loadLightsFromScene(sceneManager->getLights()); const auto& objects = sceneManager->getRenderable(); @@ -82,3 +81,40 @@ void ProgressivePathTracer::reset() std::fill(m_accumBuffer.begin(), m_accumBuffer.end(), 0.0f); m_initialized = false; } + +void ProgressivePathTracer::updateCamera( + Camera* viewportCam, + int width, + int height) +{ + if (!m_space || !viewportCam || width <= 0 || height <= 0) { + return; + } + + const glm::vec3 position = viewportCam->getPosition(); + const glm::vec3 lookAt = position + viewportCam->getFront(); + const glm::vec3 up = viewportCam->getUp(); + + m_width = width; + m_height = height; + m_space->setCamera(PBRCamera( + fungt::Vec3(position.x, position.y, position.z), + fungt::Vec3(lookAt.x, lookAt.y, lookAt.z), + fungt::Vec3(up.x, up.y, up.z), + viewportCam->getFOV(), + static_cast(width) / static_cast(height))); +} + +void ProgressivePathTracer::reloadLights(std::shared_ptr sceneManager) +{ + if (m_space && sceneManager) { + m_space->loadLightsFromScene(sceneManager->getLights()); + } +} + +void ProgressivePathTracer::releaseOpenGLInteropResources() +{ + if (m_space) { + m_space->ReleaseOpenGLInteropResources(); + } +} diff --git a/ViewPort/progressive_path_tracer_viewport.hpp b/ViewPort/progressive_path_tracer_viewport.hpp index 3357d2f..e413f9a 100644 --- a/ViewPort/progressive_path_tracer_viewport.hpp +++ b/ViewPort/progressive_path_tracer_viewport.hpp @@ -27,9 +27,13 @@ class ProgressivePathTracer { void initialize(Camera* viewportCam, std::shared_ptr sceneManager, - int width, int height); + int width, int height, bool ogl_interop = false); + void updateCamera(Camera* viewportCam, int width, int height); + void reloadLights(std::shared_ptr sceneManager); virtual void renderSample(int sample, uint32_t targetTexture) = 0; + virtual void renderSampleInterop(int sample, uint32_t targetPBO) = 0; + void releaseOpenGLInteropResources(); void reset(); diff --git a/ViewPort/viewport.hpp b/ViewPort/viewport.hpp index 917fc43..9bd3a8f 100644 --- a/ViewPort/viewport.hpp +++ b/ViewPort/viewport.hpp @@ -2,6 +2,7 @@ #define _VIEWPORT_H_ #include "../Layer/layer.hpp" #include "../Renders/display_graphics.hpp" +#include "PBR/Render/include/compute_backends.hpp" #include "../include/imgui_headers.hpp" #include #include @@ -12,6 +13,9 @@ class ViewPort : public Layer { ImVec2 m_viewportSize = ImVec2(1280, 720); std::function m_RenderFunc; std::function m_PathTraceFunc; + std::function m_PathTraceInteropFunc; + std::function m_ActivePathTraceFunc; + std::function m_PathTraceReleaseInteropFunc; bool m_pathTraceMode = false; int m_currentSample = 0; int m_maxPreviewSamples = 32; @@ -28,11 +32,27 @@ class ViewPort : public Layer { void setPathTraceFunction(const std::function& func) { m_PathTraceFunc = func; } + void setPathTraceInteropFunction(const std::function& func) { + m_PathTraceInteropFunc = func; + } + void setPathTraceReleaseInteropFunction(const std::function& func) { + m_PathTraceReleaseInteropFunc = func; + } virtual uint32_t getPathTraceTexture() const { return 0; } + virtual uint32_t getPathTracePBO() const { return 0; } + virtual void copyPBOToTexture(int width, int height) {} ImVec2 getViewPortSize() { return m_viewportSize; } - void enablePathTracing(bool enable) { m_pathTraceMode = enable; } + void enablePathTracing(bool enable) { + m_pathTraceMode = enable; + if (enable) { + m_ActivePathTraceFunc = + ComputeRender::GetBackend() == Compute::Backend::OPENCL + ? m_PathTraceInteropFunc + : m_PathTraceFunc; + } + } bool isPathTracing() const { return m_pathTraceMode; } void setMaxPreviewSamples(int s) { m_maxPreviewSamples = s; } void resetAccumulation() { m_currentSample = 0; } diff --git a/funGT/fungt.cpp b/funGT/fungt.cpp index 74bdbf4..54c4a80 100644 --- a/funGT/fungt.cpp +++ b/funGT/fungt.cpp @@ -169,20 +169,39 @@ void FunGT::set(const std::function& renderLambda){ m_lightGizmoRenderer->render(ProjectionMatrix); } }); - m_ViewPortLayer->setPathTraceFunction([this](int width, int height, int sample) { - auto* viewport = m_layerStack.get(); - if (!viewport) return; - - uint32_t pathTraceTexture = viewport->getPathTraceTexture(); - - // Initialize on first sample or if not initialized - if (sample == 0 || !m_progressiveTracer->isInitialized()) { - m_progressiveTracer->initialize(&m_camera, m_sceneManager, width, height); - } + m_ViewPortLayer->setPathTraceInteropFunction([this](int width, int height, int sample) { + auto* viewport = m_layerStack.get(); + if (!viewport) return; + + if (!m_progressiveTracer->isInitialized()) { + m_progressiveTracer->initialize( + &m_camera, m_sceneManager, width, height, true); + } else if (sample == 0) { + m_progressiveTracer->updateCamera( + &m_camera, width, height); + } - // Render one sample - m_progressiveTracer->renderSample(sample, pathTraceTexture); + m_progressiveTracer->renderSampleInterop( + sample, viewport->getPathTracePBO()); + viewport->copyPBOToTexture(width, height); + }); + m_ViewPortLayer->setPathTraceFunction([this](int width, int height, int sample) { + auto* viewport = m_layerStack.get(); + if (!viewport) return; + + if (!m_progressiveTracer->isInitialized()) { + m_progressiveTracer->initialize( + &m_camera, m_sceneManager, width, height); + } else if (sample == 0) { + m_progressiveTracer->updateCamera( + &m_camera, width, height); + } + m_progressiveTracer->renderSample( + sample, viewport->getPathTraceTexture()); + }); + m_ViewPortLayer->setPathTraceReleaseInteropFunction([this]() { + m_progressiveTracer->releaseOpenGLInteropResources(); }); m_layerStack.PushLayer(std::move(m_ViewPortLayer)); } From e5f7d0ac04c5a8480e200ac16743bd345dbf4590 Mon Sep 17 00:00:00 2001 From: juanchuletas Date: Sat, 22 Aug 2026 18:09:13 -0600 Subject: [PATCH 6/7] fix : lights were not updated in every progressive path tracer preview --- PBR/Render/include/opencl_renderer.hpp | 2 +- .../opencl/fgt_opencl_data_translator.hpp | 14 +- PBR/Render/src/opencl_renderer.cpp | 65 +- Samples/opencl/ball_lamp/CMakeLists.txt | 567 ++++++++++++++++++ Samples/opencl/ball_lamp/main.cpp | 143 +++++ ViewPort/opengl/opengl_viewport.cpp | 22 +- funGT/fungt.cpp | 14 +- 7 files changed, 776 insertions(+), 51 deletions(-) create mode 100644 Samples/opencl/ball_lamp/CMakeLists.txt create mode 100644 Samples/opencl/ball_lamp/main.cpp diff --git a/PBR/Render/include/opencl_renderer.hpp b/PBR/Render/include/opencl_renderer.hpp index 54f93e4..418bd43 100644 --- a/PBR/Render/include/opencl_renderer.hpp +++ b/PBR/Render/include/opencl_renderer.hpp @@ -83,8 +83,8 @@ class OpenCL_Renderer : public IComputeRenderer { void uploadScene( const std::vector& triangles, const std::vector& nodes, - const std::vector& lights, const std::vector& emissiveTriIndices); + void loadSceneLigths(const std::vector &lights); void releaseRaySpaceBuffer(OpenCLRaySpaceBuffer& buffers) noexcept; public: diff --git a/PBR/Render/shared/opencl/fgt_opencl_data_translator.hpp b/PBR/Render/shared/opencl/fgt_opencl_data_translator.hpp index 5a1eb08..eda5d21 100644 --- a/PBR/Render/shared/opencl/fgt_opencl_data_translator.hpp +++ b/PBR/Render/shared/opencl/fgt_opencl_data_translator.hpp @@ -21,7 +21,6 @@ struct fgt_opencl_scene_data { std::vector triangle_shading; std::vector materials; std::vector bvh_nodes; - std::vector lights; }; inline fgt_vec3 translate_vec3(const fungt::Vec3& value) @@ -168,11 +167,9 @@ inline fgt_rayspace_camera translate_rayspace_camera(const PBRCamera& camera) translate_vec4(camera.getBasisW()) }; } - inline fgt_opencl_scene_data translate_scene_data( const std::vector& triangles, - const std::vector& nodes, - const std::vector& lights) + const std::vector& nodes) { fgt_opencl_scene_data result; @@ -194,9 +191,14 @@ inline fgt_opencl_scene_data translate_scene_data( result.bvh_nodes.push_back(translate_bvh_node(node)); } - result.lights.reserve(lights.size()); + return result; +} +inline std::vector load_scene_lights(const std::vector& lights) { + + std::vector result; + result.reserve(lights.size()); for (const Light& light : lights) { - result.lights.push_back(translate_light(light)); + result.push_back(translate_light(light)); } return result; diff --git a/PBR/Render/src/opencl_renderer.cpp b/PBR/Render/src/opencl_renderer.cpp index a467d64..2fc292c 100644 --- a/PBR/Render/src/opencl_renderer.cpp +++ b/PBR/Render/src/opencl_renderer.cpp @@ -267,7 +267,9 @@ std::vector OpenCL_Renderer::RenderScene(int width, "OpenCL RenderScene: kernel not built (no textures loaded?)."); } if (!m_sceneUploaded) { - uploadScene(triangles, nodes, lights, emissiveTriIndices); + uploadScene(triangles, nodes, emissiveTriIndices); + loadSceneLigths(lights); + } const fgt_rayspace_camera cameraData = @@ -458,7 +460,6 @@ std::vector OpenCL_Renderer::RenderScene(int width, void OpenCL_Renderer::uploadScene( const std::vector& triangles, const std::vector& nodes, - const std::vector& lights, const std::vector& emissiveTriIndices) { if (!m_oclcontext) { @@ -467,7 +468,7 @@ void OpenCL_Renderer::uploadScene( } const fgt::opencl::fgt_opencl_scene_data sceneData = - fgt::opencl::translate_scene_data(triangles, nodes, lights); + fgt::opencl::translate_scene_data(triangles, nodes); std::cout << "[uploadScene] Materials texture indices: "; for (size_t i = 0; i < sceneData.materials.size(); ++i) { @@ -485,7 +486,6 @@ void OpenCL_Renderer::uploadScene( newBuffers.numTriangles = sceneData.triangle_geometry.size(); newBuffers.numMaterials = sceneData.materials.size(); newBuffers.numBVHNodes = sceneData.bvh_nodes.size(); - newBuffers.numLights = sceneData.lights.size(); newBuffers.numEmissiveTriangles = emissiveIndices.size(); cl_int err = CL_SUCCESS; @@ -550,21 +550,6 @@ void OpenCL_Renderer::uploadScene( } } - if (!sceneData.lights.empty()) { - newBuffers.lights = clCreateBuffer( - m_oclcontext, - CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, - sceneData.lights.size() * sizeof(fgt_light), - const_cast(sceneData.lights.data()), - &err); - if (err != CL_SUCCESS || !newBuffers.lights) { - releaseRaySpaceBuffer(newBuffers); - throw std::runtime_error( - "uploadScene: failed to create light buffer, error " + - std::to_string(err)); - } - } - if (!emissiveIndices.empty()) { newBuffers.emissiveTriangles = clCreateBuffer( m_oclcontext, @@ -587,8 +572,37 @@ void OpenCL_Renderer::uploadScene( std::cout << "OpenCL RaySpace scene uploaded: " << m_raySpaceBuffer.numTriangles << " triangles, " << m_raySpaceBuffer.numMaterials << " materials, " - << m_raySpaceBuffer.numBVHNodes << " BVH nodes, " - << m_raySpaceBuffer.numLights << " lights." << std::endl; + << m_raySpaceBuffer.numBVHNodes << " BVH nodes, " << std::endl; +} + +void OpenCL_Renderer::loadSceneLigths(const std::vector &lights) +{ + cl_int err = CL_SUCCESS; + std::vector scene_lights = fgt::opencl::load_scene_lights(lights); + + if (m_raySpaceBuffer.lights) { + clReleaseMemObject(m_raySpaceBuffer.lights); + m_raySpaceBuffer.lights = nullptr; + } + + m_raySpaceBuffer.numLights = scene_lights.size(); + if (!scene_lights.empty()) { + m_raySpaceBuffer.lights = clCreateBuffer( + m_oclcontext, + CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + scene_lights.size() * sizeof(fgt_light), + const_cast(scene_lights.data()), + &err); + if (err != CL_SUCCESS || !m_raySpaceBuffer.lights) { + throw std::runtime_error( + "loadSceneLigths: failed to create light buffer, error " + + std::to_string(err)); + } + } + std::cout << " Lights Loaded to the Scene : " << m_raySpaceBuffer.numLights<< std::endl; + if (!scene_lights.empty()) { + std::cout << " [DEBUG] Color intensity : " << scene_lights[0].intensity.x << " , " << scene_lights[0].intensity.y << " ," << scene_lights[0].intensity.z << std::endl; + } } void OpenCL_Renderer::releaseRaySpaceBuffer( @@ -712,10 +726,17 @@ void OpenCL_Renderer::RenderSceneOpenGLInterop(int width, prepareTextures(); if (!m_sceneUploaded) { - uploadScene(triangles, nodes, lights, emissiveTriIndices); + uploadScene(triangles, nodes, emissiveTriIndices); + clFinish(m_oclqueue); + } + + // Update lights when a new accumulation starts. + if (sampleOffset == 0 || !m_raySpaceBuffer.lights) { + loadSceneLigths(lights); clFinish(m_oclqueue); } + if (!m_oglSharedBuffer || m_oglBufferID != glBufferID || m_oglBufferWidth != width || diff --git a/Samples/opencl/ball_lamp/CMakeLists.txt b/Samples/opencl/ball_lamp/CMakeLists.txt new file mode 100644 index 0000000..92c5d9d --- /dev/null +++ b/Samples/opencl/ball_lamp/CMakeLists.txt @@ -0,0 +1,567 @@ +cmake_minimum_required(VERSION 3.15) +project(FunGT) + +# ════════════════════════════════════════════════════════════════════════════ +# GPU BACKEND OPTIONS +# ════════════════════════════════════════════════════════════════════════════ +option(FUNGT_USE_CUDA "Enable CUDA backend for NVIDIA GPUs" ON) +option(FUNGT_USE_SYCL "Enable SYCL backend for Intel GPUs" ON) +option(FUNGT_USE_OPENCL "Enable OpenCL backend" ON) +set(FUNGT_CUDA_ARCH 89 CACHE STRING "NVIDIA GPU compute capability") +set_property(CACHE FUNGT_CUDA_ARCH PROPERTY STRINGS 50 52 53 60 61 62 70 72 75 80 86 87 89 90) +set(SYCL_CUDA_ARCH "sm_${FUNGT_CUDA_ARCH}") + +# Allow FUNGT_BASE_DIR to be set externally (command line, environment, or cache) +# Priority: 1. Cache variable, 2. Environment variable, 3. Default fallback +if(NOT DEFINED FUNGT_BASE_DIR) + # Check if it's set as an environment variable + if(DEFINED ENV{FUNGT_BASE_DIR}) + set(FUNGT_BASE_DIR $ENV{FUNGT_BASE_DIR} CACHE PATH "FunGT base directory") + message(STATUS "Using FUNGT_BASE_DIR from environment: ${FUNGT_BASE_DIR}") + else() + # Use a default fallback (you must specify this via -DFUNGT_BASE_DIR or environment) + message(FATAL_ERROR "FUNGT_BASE_DIR not specified. Please set it via:\n" + " cmake -DFUNGT_BASE_DIR=/path/to/FunGT ..\n" + " or export FUNGT_BASE_DIR=/path/to/FunGT") + endif() +else() + message(STATUS "Using pre-defined FUNGT_BASE_DIR: ${FUNGT_BASE_DIR}") +endif() + +# Validate that the base directory exists and contains expected subdirectories +if(NOT EXISTS "${FUNGT_BASE_DIR}") + message(FATAL_ERROR "FUNGT_BASE_DIR does not exist: ${FUNGT_BASE_DIR}") +endif() + +if(NOT EXISTS "${FUNGT_BASE_DIR}/VertexGL" OR NOT EXISTS "${FUNGT_BASE_DIR}/Shaders") + message(FATAL_ERROR "FUNGT_BASE_DIR does not appear to be a valid FunGT directory: ${FUNGT_BASE_DIR}") +endif() + +# Set the C++ standard +set(CMAKE_CXX_STANDARD 17) +set(CMAKE_CXX_STANDARD_REQUIRED ON) +# ════════════════════════════════════════════════════════════════════════════ +# SYCL TOOLCHAIN DETECTION +# ════════════════════════════════════════════════════════════════════════════ +# Priority: +# 1. User-provided SYCL_COMPILER_PATH (cmake -DSYCL_COMPILER_PATH=...) +# 2. Bundled SYCL toolchain in lib/linux_x64/dpcpp/ +# 3. Error - user must provide SYCL compiler + +if(UNIX) + if(DEFINED SYCL_COMPILER_PATH) + # User provided their own SYCL compiler + if(NOT EXISTS "${SYCL_COMPILER_PATH}") + message(FATAL_ERROR "SYCL_COMPILER_PATH does not exist: ${SYCL_COMPILER_PATH}") + endif() + set(CMAKE_CXX_COMPILER ${SYCL_COMPILER_PATH}) + message(STATUS "Using user-provided SYCL compiler: ${SYCL_COMPILER_PATH}") + else() + # Try bundled SYCL toolchain + set(BUNDLED_SYCL_DIR "${FUNGT_BASE_DIR}/toolchain/sycl/linux_x64/dpcpp") + set(BUNDLED_SYCL_COMPILER "${BUNDLED_SYCL_DIR}/bin/clang++") + + if(EXISTS "${BUNDLED_SYCL_COMPILER}") + set(CMAKE_CXX_COMPILER ${BUNDLED_SYCL_COMPILER}) + message(STATUS "Using bundled SYCL compiler: ${BUNDLED_SYCL_COMPILER}") + else() + message(FATAL_ERROR + "SYCL compiler not found!\n" + "Please either:\n" + " 1. Use bundled toolchain (should be at ${BUNDLED_SYCL_DIR})\n" + " 2. Provide your own: cmake -DSYCL_COMPILER_PATH=/path/to/clang++ ..\n") + endif() + endif() +elseif(WIN32) + if(DEFINED SYCL_COMPILER_PATH) + # User provided their own SYCL compiler + if(NOT EXISTS "${SYCL_COMPILER_PATH}") + message(FATAL_ERROR "SYCL_COMPILER_PATH does not exist: ${SYCL_COMPILER_PATH}") + endif() + set(CMAKE_CXX_COMPILER ${SYCL_COMPILER_PATH}) + message(STATUS "Using user-provided SYCL compiler: ${SYCL_COMPILER_PATH}") + else() + # Try bundled SYCL toolchain + set(BUNDLED_SYCL_DIR "${FUNGT_BASE_DIR}/lib/windows_x64/dpcpp") + set(BUNDLED_SYCL_COMPILER "${BUNDLED_SYCL_DIR}/bin/clang++.exe") + + if(EXISTS "${BUNDLED_SYCL_COMPILER}") + set(CMAKE_CXX_COMPILER ${BUNDLED_SYCL_COMPILER}) + message(STATUS "Using bundled SYCL compiler: ${BUNDLED_SYCL_COMPILER}") + else() + message(FATAL_ERROR + "SYCL compiler not found!\n" + "Please either:\n" + " 1. Use bundled toolchain (should be at ${BUNDLED_SYCL_DIR})\n" + " 2. Provide your own: cmake -DSYCL_COMPILER_PATH=C:/path/to/clang++.exe ..\n") + endif() + endif() +endif() + +# ════════════════════════════════════════════════════════════════════════════ +# PBR RENDERER LIBRARIES - PREBUILT FROM PBR/main/build +# ════════════════════════════════════════════════════════════════════════════ +set(PBR_LIB_DIR ${FUNGT_BASE_DIR}/PBR/main/build) + +if(UNIX) + # Set the build type to Release + set(CMAKE_BUILD_TYPE Release) + set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/release/linux) + + # Add the prebuilt ImGui library + add_library(imgui STATIC IMPORTED) + set_target_properties(imgui PROPERTIES + IMPORTED_LOCATION "${FUNGT_BASE_DIR}/vendor/imgui/lib/libimgui.a" + INTERFACE_INCLUDE_DIRECTORIES "${FUNGT_BASE_DIR}/vendor/imgui" + ) + + # Add funlib + set(FUNLIB_DIR ${FUNGT_BASE_DIR}/vendor/funlib) + add_library(funlib STATIC IMPORTED GLOBAL) + set_target_properties(funlib PROPERTIES + IMPORTED_LOCATION ${FUNLIB_DIR}/lib/libfunlib.a + INTERFACE_INCLUDE_DIRECTORIES ${FUNLIB_DIR}/include + ) + # PBR Core (always required) + add_library(pbr_core STATIC IMPORTED) + set_target_properties(pbr_core PROPERTIES + IMPORTED_LOCATION "${PBR_LIB_DIR}/libpbr_core.a" + ) + + # PBR CUDA renderer (if enabled) + if(FUNGT_USE_CUDA) + add_library(pbr_cuda STATIC IMPORTED) + set_target_properties(pbr_cuda PROPERTIES + IMPORTED_LOCATION "${PBR_LIB_DIR}/libpbr_cuda.a" + ) + set(PBR_CUDA_LIB pbr_cuda) + message(STATUS "PBR CUDA renderer library: ${PBR_LIB_DIR}/libpbr_cuda.a") + endif() + + # PBR SYCL renderer (if enabled) + if(FUNGT_USE_SYCL) + add_library(pbr_sycl STATIC IMPORTED) + set_target_properties(pbr_sycl PROPERTIES + IMPORTED_LOCATION "${PBR_LIB_DIR}/libpbr_sycl.a" + ) + set(PBR_SYCL_LIB pbr_sycl) + message(STATUS "PBR SYCL renderer library: ${PBR_LIB_DIR}/libpbr_sycl.a") + endif() + + if(FUNGT_USE_OPENCL) + add_library(pbr_opencl STATIC IMPORTED) + set_target_properties(pbr_opencl PROPERTIES + IMPORTED_LOCATION "${PBR_LIB_DIR}/libpbr_opencl.a" + ) + set(PBR_OPENCL_LIB pbr_opencl) + message(STATUS "PBR OpenCL renderer library: ${PBR_LIB_DIR}/libpbr_opencl.a") + endif() +endif() + +if (WIN32) + # Add the macro definition + add_definitions(-DGLM_ENABLE_EXPERIMENTAL) + # Set vcpkg toolchain file + set(CMAKE_TOOLCHAIN_FILE "C:/Users/juang/Documents/Development/vcpkg/scripts/buildsystems/vcpkg.cmake" CACHE STRING "Vcpkg toolchain file") + # Set CMake prefix path for vcpkg packages + set(CMAKE_PREFIX_PATH "C:/Users/juang/Documents/Development/vcpkg/installed/x64-windows") + set(GLFW_BUILD_STATIC ON) + message(STATUS "Working Dir: " ${CMAKE_CURRENT_SOURCE_DIR}) +endif() + +# Include directories +include_directories( + ${FUNGT_BASE_DIR} + ${FUNGT_BASE_DIR}/GT + ${FUNGT_BASE_DIR}/VertexGL + ${FUNGT_BASE_DIR}/Shaders + ${FUNGT_BASE_DIR}/funGT + ${FUNGT_BASE_DIR}/AnimatedModel + ${FUNGT_BASE_DIR}/Textures + ${FUNGT_BASE_DIR}/vendor/stb_image + ${FUNGT_BASE_DIR}/vendor/glm/include + ${FUNGT_BASE_DIR}/Imgui_Setup + ${FUNGT_BASE_DIR}/Material + ${FUNGT_BASE_DIR}/Mesh + ${FUNGT_BASE_DIR}/Camera + ${FUNGT_BASE_DIR}/Geometries + ${FUNGT_BASE_DIR}/Model + ${FUNGT_BASE_DIR}/Helpers + ${FUNGT_BASE_DIR}/Animation + ${FUNGT_BASE_DIR}/Bone + ${FUNGT_BASE_DIR}/Matrix + ${FUNGT_BASE_DIR}/SceneManager + ${FUNGT_BASE_DIR}/CubeMap + ${FUNGT_BASE_DIR}/ParticleSimulation + ${FUNGT_BASE_DIR}/Random + ${FUNGT_BASE_DIR}/Path_Manager + ${FUNGT_BASE_DIR}/InfoWindow + ${FUNGT_BASE_DIR}/GUI + ${FUNGT_BASE_DIR}/GUI/physics + ${FUNGT_BASE_DIR}/SimpleModel + ${FUNGT_BASE_DIR}/SimpleGeometry + ${FUNGT_BASE_DIR}/Physics/CollisionManager + ${FUNGT_BASE_DIR}/Physics/Contact + ${FUNGT_BASE_DIR}/Physics/RigidBody + ${FUNGT_BASE_DIR}/Physics/Shapes + ${FUNGT_BASE_DIR}/Physics/Integrators + ${FUNGT_BASE_DIR}/Physics/ContactHelpers + ${FUNGT_BASE_DIR}/Physics/PhysicsWorld + ${FUNGT_BASE_DIR}/Physics/AnimationCreator + ${FUNGT_BASE_DIR}/Physics/Clothing + ${FUNGT_BASE_DIR}/Quaternion + ${FUNGT_BASE_DIR}/Vector + ${FUNGT_BASE_DIR}/ViewPort + ${FUNGT_BASE_DIR}/Renders + ${FUNGT_BASE_DIR}/Platform/OpenGL + ${FUNGT_BASE_DIR}/InfoDevice + ${FUNGT_BASE_DIR}/PBR/PBRCamera + ${FUNGT_BASE_DIR}/PBR/Ray + ${FUNGT_BASE_DIR}/PBR/HitData + ${FUNGT_BASE_DIR}/PBR/Intersection + ${FUNGT_BASE_DIR}/PBR/Light + ${FUNGT_BASE_DIR}/PBR/Space + ${FUNGT_BASE_DIR}/PBR/Render/include + ${FUNGT_BASE_DIR}/PBR/Render/brdf + ${FUNGT_BASE_DIR}/PBR/Render/shared + ${FUNGT_BASE_DIR}/PBR/TextureManager + ${FUNGT_BASE_DIR}/PBR/TriangleExtractor + ${FUNGT_BASE_DIR}/PBR/BVH + ${FUNGT_BASE_DIR}/Triangle + ${FUNGT_BASE_DIR}/Gizmo/LightGizmo + ${FUNGT_BASE_DIR}/Lights + ${FUNGT_BASE_DIR}/IBL + ${FUNGT_BASE_DIR}/GraphicsRenderBackend + ${FUNGT_BASE_DIR}/GraphicsRenderBackend/opengl + ${FUNGT_BASE_DIR}/MeshGPU + ${FUNGT_BASE_DIR}/TextureGPU + ${FUNGT_BASE_DIR}/PrimitiveGPU +) + +# ════════════════════════════════════════════════════════════════════════════ +# COMMON SOURCE FILES (compiled with SYCL compiler) +# ════════════════════════════════════════════════════════════════════════════ +set(SOURCE_FILES + ${CMAKE_CURRENT_SOURCE_DIR}/main.cpp + ${FUNGT_BASE_DIR}/GT/graphicsTool.cpp + ${FUNGT_BASE_DIR}/VertexGL/vertexArrayObjects.cpp + ${FUNGT_BASE_DIR}/VertexGL/vertexBuffers.cpp + ${FUNGT_BASE_DIR}/VertexGL/vertexIndices.cpp + ${FUNGT_BASE_DIR}/VertexGL/uniform_buffer_object.cpp + ${FUNGT_BASE_DIR}/VertexGL/shaderStorageBufferObejct.cpp + ${FUNGT_BASE_DIR}/Shaders/shader.cpp + ${FUNGT_BASE_DIR}/Shaders/opengl/opengl_shader.cpp + ${FUNGT_BASE_DIR}/funGT/fungt.cpp + ${FUNGT_BASE_DIR}/AnimatedModel/animated_model.cpp + ${FUNGT_BASE_DIR}/Textures/textures.cpp + ${FUNGT_BASE_DIR}/vendor/stb_image/stb_image.cpp + ${FUNGT_BASE_DIR}/Imgui_Setup/imgui_setup.cpp + ${FUNGT_BASE_DIR}/Material/material.cpp + ${FUNGT_BASE_DIR}/Mesh/mesh.cpp + ${FUNGT_BASE_DIR}/MeshGPU/mesh_gpu.cpp + ${FUNGT_BASE_DIR}/TextureGPU/texture_gpu.cpp + ${FUNGT_BASE_DIR}/PrimitiveGPU/primitive_gpu.cpp + ${FUNGT_BASE_DIR}/Camera/camera.cpp + ${FUNGT_BASE_DIR}/Geometries/cube.cpp + ${FUNGT_BASE_DIR}/Geometries/plane.cpp + ${FUNGT_BASE_DIR}/Geometries/primitives.cpp + ${FUNGT_BASE_DIR}/Geometries/pyramid.cpp + ${FUNGT_BASE_DIR}/Geometries/shape.cpp + ${FUNGT_BASE_DIR}/Geometries/inf_grid.cpp + ${FUNGT_BASE_DIR}/Geometries/square.cpp + ${FUNGT_BASE_DIR}/Geometries/sphere.cpp + ${FUNGT_BASE_DIR}/Geometries/box.cpp + ${FUNGT_BASE_DIR}/Model/model.cpp + ${FUNGT_BASE_DIR}/Helpers/helpers.cpp + ${FUNGT_BASE_DIR}/Animation/animation.cpp + ${FUNGT_BASE_DIR}/Bone/bone.cpp + ${FUNGT_BASE_DIR}/Matrix/matrix4x4f.cpp + ${FUNGT_BASE_DIR}/Matrix/matrix3x3f.cpp + ${FUNGT_BASE_DIR}/SceneManager/scene_manager.cpp + ${FUNGT_BASE_DIR}/CubeMap/cube_map.cpp + ${FUNGT_BASE_DIR}/IBL/ibl_probe.cpp + ${FUNGT_BASE_DIR}/ParticleSimulation/particle_simulation.cpp + ${FUNGT_BASE_DIR}/ParticleSimulation/particle_simulation_rtc.cpp + ${FUNGT_BASE_DIR}/ParticleSimulation/particle_simulation_sycl_ops.cpp + ${FUNGT_BASE_DIR}/ParticleSimulation/particle_rtc_sycl_ops.cpp + ${FUNGT_BASE_DIR}/Random/random.cpp + ${FUNGT_BASE_DIR}/Path_Manager/path_manager.cpp + ${FUNGT_BASE_DIR}/InfoWindow/infowindow.cpp + ${FUNGT_BASE_DIR}/GUI/gui.cpp + ${FUNGT_BASE_DIR}/GUI/physics/simulation_controller_window.cpp + ${FUNGT_BASE_DIR}/GUI/demo_particles_window.cpp + ${FUNGT_BASE_DIR}/GUI/render_action_window.cpp + ${FUNGT_BASE_DIR}/GUI/physics/debug_rigidbody_renderer.cpp + ${FUNGT_BASE_DIR}/GUI/particle_rtc_window.cpp + ${FUNGT_BASE_DIR}/SimpleModel/simple_model.cpp + ${FUNGT_BASE_DIR}/SimpleGeometry/simple_geometry.cpp + ${FUNGT_BASE_DIR}/Physics/Collisions/simple_collision.cpp + ${FUNGT_BASE_DIR}/Physics/Collisions/manifold_collision.cpp + ${FUNGT_BASE_DIR}/Physics/CollisionManager/collision_manager.cpp + ${FUNGT_BASE_DIR}/Physics/Collider/collider.cpp + ${FUNGT_BASE_DIR}/Physics/AnimationCreator/key_frame_recorder.cpp + ${FUNGT_BASE_DIR}/Physics/AnimationCreator/animation_exporter.cpp + ${FUNGT_BASE_DIR}/ViewPort/viewport.cpp + ${FUNGT_BASE_DIR}/ViewPort/opengl/opengl_viewport.cpp + ${FUNGT_BASE_DIR}/ViewPort/opengl/opengl_progressive_path_tracer.cpp + ${FUNGT_BASE_DIR}/ViewPort/progressive_path_tracer_viewport.cpp + ${FUNGT_BASE_DIR}/Renders/framebuffer.cpp + ${FUNGT_BASE_DIR}/Renders/display_graphics.cpp + ${FUNGT_BASE_DIR}/Platform/OpenGL/OGLframeBuffer.cpp + ${FUNGT_BASE_DIR}/InfoDevice/gpu_device_info.cpp + ${FUNGT_BASE_DIR}/PBR/Intersection/intersection.cpp + ${FUNGT_BASE_DIR}/PBR/Space/space.cpp + ${FUNGT_BASE_DIR}/PBR/BVH/bvh_builder.cpp + ${FUNGT_BASE_DIR}/PBR/TriangleExtractor/triangle_extractor.cpp + ${FUNGT_BASE_DIR}/Gizmo/LightGizmo/light_gizmo_renderer.cpp + ${FUNGT_BASE_DIR}/GraphicsRenderBackend/graphics_render_device.cpp + ${FUNGT_BASE_DIR}/GraphicsRenderBackend/gpu_buffer.cpp + ${FUNGT_BASE_DIR}/GraphicsRenderBackend/gpu_texture.cpp + ${FUNGT_BASE_DIR}/GraphicsRenderBackend/opengl/opengl_device.cpp + ${FUNGT_BASE_DIR}/GraphicsRenderBackend/opengl/opengl_buffer.cpp + ${FUNGT_BASE_DIR}/GraphicsRenderBackend/opengl/opengl_texture.cpp +) + +# ════════════════════════════════════════════════════════════════════════════ +# GPU DEVICE DETECTION BACKENDS +# ════════════════════════════════════════════════════════════════════════════ + +# CUDA backend (compiled with g++, NOT SYCL compiler) +if(FUNGT_USE_CUDA) + find_package(CUDAToolkit REQUIRED) + + if(CUDAToolkit_FOUND) + message(STATUS "CUDA Toolkit found: ${CUDAToolkit_VERSION}") + + # Create separate library for CUDA backend + add_library(cuda_backend STATIC + ${FUNGT_BASE_DIR}/InfoDevice/cuda_backend.cpp + ) + + # CRITICAL: Compile with standard g++, NOT SYCL compiler + set_source_files_properties( + ${FUNGT_BASE_DIR}/InfoDevice/cuda_backend.cpp + PROPERTIES + COMPILE_FLAGS "-std=c++17" + ) + + target_include_directories(cuda_backend PRIVATE + ${CUDAToolkit_INCLUDE_DIRS} + ${FUNGT_BASE_DIR}/InfoDevice + ) + + target_link_libraries(cuda_backend PRIVATE + CUDA::cudart + ) + + set(CUDA_BACKEND_LIB cuda_backend) + add_compile_definitions(FUNGT_USE_CUDA) + message(STATUS "CUDA device detection enabled (compiled with g++)") + enable_language(CUDA) + else() + message(FATAL_ERROR "CUDA requested but CUDA Toolkit not found!") + endif() +else() + message(STATUS "CUDA backend disabled") +endif() + +# SYCL backend (compiled with SYCL compiler) +if(FUNGT_USE_SYCL) + list(APPEND SOURCE_FILES ${FUNGT_BASE_DIR}/InfoDevice/sycl_backend.cpp) + message(STATUS "SYCL device detection enabled (compiled with SYCL compiler)") +else() + message(STATUS "SYCL backend disabled") +endif() + +# ════════════════════════════════════════════════════════════════════════════ +# CREATE EXECUTABLE +# ════════════════════════════════════════════════════════════════════════════ +add_executable(FunGT ${SOURCE_FILES}) +if(FUNGT_USE_CUDA) + target_compile_definitions(FunGT PRIVATE FUNGT_USE_CUDA) + target_include_directories(FunGT PRIVATE ${CUDAToolkit_INCLUDE_DIRS}) + set_target_properties(FunGT PROPERTIES + CUDA_RESOLVE_DEVICE_SYMBOLS ON + CUDA_SEPARABLE_COMPILATION ON + CUDA_ARCHITECTURES ${FUNGT_CUDA_ARCH} + ) +endif() +# ════════════════════════════════════════════════════════════════════════════ +# PLATFORM-SPECIFIC CONFIGURATIONS +# ════════════════════════════════════════════════════════════════════════════ +if (WIN32) + # Use vcpkg packages on Windows + find_package(OpenGL REQUIRED) + find_package(GLEW REQUIRED) + find_package(glfw3 REQUIRED) + find_package(assimp REQUIRED) + + target_link_libraries(FunGT + PRIVATE + OpenGL::GL + glfw + GLEW::GLEW + assimp::assimp + ) + +elseif (UNIX) + # ════════════════════════════════════════════════════════════════════════ + # LINUX-SPECIFIC SETTINGS + # ════════════════════════════════════════════════════════════════════════ + message(STATUS "Configuring for Linux") + + # Find required packages + if(FUNGT_USE_SYCL) + + # Try bundled toolchain OpenCL first + set(BUNDLED_OPENCL_INCLUDE "${BUNDLED_SYCL_DIR}/include") + find_library(BUNDLED_OPENCL_LIBRARY + NAMES OpenCL + PATHS "${BUNDLED_SYCL_DIR}/lib" + NO_DEFAULT_PATH + ) + + if(EXISTS "${BUNDLED_OPENCL_INCLUDE}/CL/opencl.h" AND BUNDLED_OPENCL_LIBRARY) + + set(OpenCL_INCLUDE_DIRS "${BUNDLED_OPENCL_INCLUDE}") + set(OpenCL_LIBRARIES "${BUNDLED_OPENCL_LIBRARY}") + + message(STATUS "OpenCL (bundled) found:") + message(STATUS " Include: ${OpenCL_INCLUDE_DIRS}") + message(STATUS " Library: ${OpenCL_LIBRARIES}") + + else() + + message(STATUS "Bundled OpenCL not found, trying system OpenCL...") + + find_package(OpenCL REQUIRED) + + if(OpenCL_FOUND) + message(STATUS "OpenCL (system) found: ${OpenCL_INCLUDE_DIRS}") + else() + message(FATAL_ERROR "OpenCL library not found!") + endif() + + endif() + + else() + + # If SYCL disabled, just use system OpenCL + find_package(OpenCL REQUIRED) + if(OpenCL_FOUND) + message(STATUS "OpenCL found: ${OpenCL_INCLUDE_DIRS}") + else() + message(FATAL_ERROR "OpenCL library not found!") + endif() + + endif() + + find_package(CURL REQUIRED) + if (CURL_FOUND) + message(STATUS "CURL found") + else() + message(FATAL_ERROR "CURL library not found!") + endif() + find_package(OpenGL REQUIRED) + if (OpenGL_FOUND) + message(STATUS "OpenGL found") + else() + message(FATAL_ERROR "OpenGL library not found!") + endif() + + find_package(glfw3 REQUIRED) + if (glfw3_FOUND) + message(STATUS "GLFW found") + else() + message(FATAL_ERROR "GLFW library not found!") + endif() + + find_package(GLEW REQUIRED) + if (GLEW_FOUND) + message(STATUS "GLEW found") + else() + message(FATAL_ERROR "GLEW library not found!") + endif() + + find_package(assimp REQUIRED) + if (assimp_FOUND) + message(STATUS "Assimp found") + else() + message(FATAL_ERROR "Assimp library not found!") + endif() + + # Link libraries + target_include_directories(FunGT PRIVATE ${OpenCL_INCLUDE_DIRS}) + target_link_libraries(FunGT + PRIVATE + CURL::libcurl + OpenGL::GL + glfw + GLEW::GLEW + assimp::assimp + dl + funlib + imgui + ${CUDA_BACKEND_LIB} + pbr_core + ${PBR_CUDA_LIB} + ${PBR_SYCL_LIB} + ${PBR_OPENCL_LIB} + ${OpenCL_LIBRARIES} + $<$:CUDA::cudart_static> + $<$:CUDA::cuda_driver> + ) + + # ════════════════════════════════════════════════════════════════════════ + # SYCL compilation flags (per-file, not blanket) + # ════════════════════════════════════════════════════════════════════════ + if(CUDAToolkit_FOUND) + message(STATUS "CUDA SYCL backend enabled") + set(SYCL_TARGETS nvptx64-nvidia-cuda,spir64) + set(SYCL_CUDA_ARCH_ARGS + "-Xsycl-target-backend=nvptx64-nvidia-cuda" + "--offload-arch=${SYCL_CUDA_ARCH}" + ) + else() + set(SYCL_TARGETS spir64) + set(SYCL_CUDA_ARCH_ARGS "") + endif() + + set(SYCL_SOURCES + ${FUNGT_BASE_DIR}/InfoDevice/sycl_backend.cpp + ${FUNGT_BASE_DIR}/ParticleSimulation/particle_simulation_sycl_ops.cpp + ${FUNGT_BASE_DIR}/ParticleSimulation/particle_rtc_sycl_ops.cpp + ${FUNGT_BASE_DIR}/PBR/Space/space.cpp + ) + + list(JOIN SYCL_CUDA_ARCH_ARGS " " SYCL_CUDA_ARCH_STR) + + set_source_files_properties(${SYCL_SOURCES} + PROPERTIES COMPILE_FLAGS "-fsycl -fsycl-targets=${SYCL_TARGETS} ${SYCL_CUDA_ARCH_STR}" + ) + + target_link_options(FunGT PRIVATE -fsycl -fsycl-targets=${SYCL_TARGETS} ${SYCL_CUDA_ARCH_ARGS}) + + if(FUNGT_USE_SYCL) + target_compile_definitions(FunGT PRIVATE FUNGT_USE_SYCL) + endif() + if(FUNGT_USE_OPENCL) + target_compile_definitions(FunGT PRIVATE FUNGT_USE_OPENCL) + endif() + +endif() + +# ════════════════════════════════════════════════════════════════════════════ +# BUILD SUMMARY +# ════════════════════════════════════════════════════════════════════════════ +message(STATUS "") +message(STATUS "════════════════════════════════════════") +message(STATUS " FunGT Build Configuration") +message(STATUS "════════════════════════════════════════") +message(STATUS " Main Compiler: Intel DPC++") +message(STATUS " CUDA Backend: ${FUNGT_USE_CUDA}") +message(STATUS " SYCL Backend: ${FUNGT_USE_SYCL}") +message(STATUS " OpenCL Backend: ${FUNGT_USE_OPENCL}") +message(STATUS " Build Type: ${CMAKE_BUILD_TYPE}") +message(STATUS "════════════════════════════════════════") +message(STATUS "") diff --git a/Samples/opencl/ball_lamp/main.cpp b/Samples/opencl/ball_lamp/main.cpp new file mode 100644 index 0000000..f439b5d --- /dev/null +++ b/Samples/opencl/ball_lamp/main.cpp @@ -0,0 +1,143 @@ +#include "funGT/fungt.hpp" +#include +#include + +const unsigned int SCREEN_WIDTH = 2100; +const unsigned int SCREEN_HEIGHT = 1200; + +int main() { + std::string path = findProjectRoot(); + std::cout << path << std::endl; + srand(time(0)); + + ModelPaths model_ball, model_lamp; + model_lamp.path = getAssetPath("assets_local/Obj/LuxoLamp/Luxo.obj"); + model_ball.path = getAssetPath("assets_local/Obj/LuxoBall/luxoball.obj"); + + DisplayGraphics::SetBackend(Backend::OpenGL); + FunGTScene myGame = FunGT::createScene(SCREEN_WIDTH, SCREEN_HEIGHT); + myGame->setBackgroundColor(); + myGame->initGL(); + + myGame->createPhysicsWorld(); + spCollisionManager myCollision = myGame->getCollisionManager(); + myCollision->showCollidableBodies(false); + + auto ground = std::make_shared( + std::make_unique(80.0f, 1.0f, 80.0f), 0.0f + ); + ground->m_pos = fungt::Vec3(0, -0.5f, 0.f); + ground->m_restitution = 0.4f; + ground->m_friction = 0.6f; + myCollision->add(ground); + + auto leftWall = std::make_shared( + std::make_unique(7.0f, 16.0f, 7.0f), 0.0f + ); + leftWall->m_pos = fungt::Vec3(0.0f, 8.0f, 0.f); + leftWall->m_restitution = 0.7f; + leftWall->m_friction = 0.3f; + myCollision->add(leftWall); + + auto ball = std::make_shared( + std::make_unique(2.3), 1.0f + ); + ball->m_pos = fungt::Vec3(0.0f, 2.3f, 15.0f); + ball->m_vel = fungt::Vec3(0.0f, 0.0f, -15.0f); + + ball->m_angularVel = fungt::Vec3(-1.0, 0, 0); // ≈ (6.52, 0, 0) + + ball->m_restitution = 0.8f; + ball->m_friction = 0.3f; + myCollision->add(ball); + + FunGTSModel ballModel = SimpleModel::create(); + ballModel->load(model_ball); + ballModel->position(0.0f, 2.3f, 15.f); + ballModel->scale(2.3f); + ballModel->setAnimationID("ball"); + ballModel->addCollisionProperty(ball); + + FunGTSModel lamp = SimpleModel::create(); + lamp->load(model_lamp); + lamp->position(0.f, 0.f, 0.f); + lamp->setAnimationID("lamp"); + + FunGTSGeom groundPlane = SimpleGeometry::create(Geometry::Plane); + groundPlane->load(getAssetPath("assets_local/img/floor.png")); + groundPlane->position(0.0, 0.0, 0.0); + + //Back wall: + + FunGTSGeom backWall = SimpleGeometry::create(Geometry::Plane); + backWall->load(getAssetPath("assets_local/img/leftwall.png")); + backWall->rotation(90.0f, 0.0f, 0.0f); + backWall->position(0.0f, 0.0f, -40.0f); + FunGTSGeom left_Wall = SimpleGeometry::create(Geometry::Plane); + left_Wall->load(getAssetPath("assets_local/img/leftwall.png")); + left_Wall->rotation(90.0f, 0.0f, -90.0f); + left_Wall->position(-40.0f, 0.0f, 0.0f); + + + FunGTSceneManager scene_manager = myGame->getSceneManager(); + FunGTSimController simController = myGame->getSimulationController(); + + simController->captureSnapshot(ball); + + myGame->set([&]() { + scene_manager->addRenderableObj(groundPlane); + scene_manager->addRenderableObj(backWall); + scene_manager->addRenderableObj(left_Wall); + scene_manager->addRenderableObj(ballModel); + scene_manager->addRenderableObj(lamp); + }); + + auto animController = myGame->getAnimationController(); + + animController->setObjectBakingEnabled("ball", true); + + auto animWindow = myGame->getAnimationWindow(); + + float lastTime = glfwGetTime(); + FunGTPhysicsWorld physics = myGame->getPhysicsWorld(); + + myGame->render([&]() { + float currentTime = glfwGetTime(); + float dt = currentTime - lastTime; + lastTime = currentTime; + + int currentFrame = static_cast(simController->getCurrentTime() * 30.0f); + + if (simController->isPlaying()) { + float adjustedDt = dt * simController->getPlaybackSpeed(); + + if (animWindow && animWindow->isRecordingMode()) { + + + for (auto& obj : scene_manager->getRenderable()) { + if (auto model = std::dynamic_pointer_cast(obj)) { + model->setUsePhysicsRigidBody(true); // Ensure physics is enabled for all models (optional, based on your design) + } + } + + physics->runColliders(adjustedDt); + animController->recordPhysicsFrame(currentFrame); + } + else if (animWindow && animWindow->isPlaybackMode()) { + // Disable physics for all + for (auto& obj : scene_manager->getRenderable()) { + if (auto model = std::dynamic_pointer_cast(obj)) { + model->setUsePhysicsRigidBody(false); + } + } + animController->updateFrame(currentFrame); + } + + simController->updateTime(dt); + } + + scene_manager->renderScene(); + }); + + return 0; +} \ No newline at end of file diff --git a/ViewPort/opengl/opengl_viewport.cpp b/ViewPort/opengl/opengl_viewport.cpp index 9581a47..616bcb0 100644 --- a/ViewPort/opengl/opengl_viewport.cpp +++ b/ViewPort/opengl/opengl_viewport.cpp @@ -101,13 +101,10 @@ void OpenGLViewPort::onImGuiRender() if (diffX > 1.0f || diffY > 1.0f) { if (viewportPanelSize.x > 32 && viewportPanelSize.y > 32) { - const float pendingDiffX = std::abs(viewportPanelSize.x - pendingSize.x); - const float pendingDiffY = std::abs(viewportPanelSize.y - pendingSize.y); - if (!pendingResize || pendingDiffX > 1.0f || pendingDiffY > 1.0f) { - pendingSize = viewportPanelSize; - lastResizeRequest = currentTime; - pendingResize = true; - } + m_viewportSize = viewportPanelSize; + pendingSize = viewportPanelSize; + lastResizeRequest = currentTime; + pendingResize = true; } } @@ -120,25 +117,24 @@ void OpenGLViewPort::onImGuiRender() } FrameBuffSpec spec{ - static_cast(pendingSize.x), - static_cast(pendingSize.y), + static_cast(m_viewportSize.x), + static_cast(m_viewportSize.y), 1 }; m_resizeBuffer = FrameBuffer::create(spec); glBindTexture(GL_TEXTURE_2D, m_pathTraceTexture); glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA32F, - static_cast(pendingSize.x), - static_cast(pendingSize.y), + static_cast(m_viewportSize.x), + static_cast(m_viewportSize.y), 0, GL_RGBA, GL_FLOAT, nullptr); glBindTexture(GL_TEXTURE_2D, 0); glBindBuffer(GL_PIXEL_UNPACK_BUFFER, m_pathTracePBO); glBufferData(GL_PIXEL_UNPACK_BUFFER, - static_cast(pendingSize.x) * static_cast(pendingSize.y) * 4 * sizeof(float), + static_cast(m_viewportSize.x) * static_cast(m_viewportSize.y) * 4 * sizeof(float), nullptr, GL_DYNAMIC_DRAW); glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0); - m_viewportSize = pendingSize; m_currentSample = 0; pendingResize = false; } diff --git a/funGT/fungt.cpp b/funGT/fungt.cpp index 54c4a80..3e427a7 100644 --- a/funGT/fungt.cpp +++ b/funGT/fungt.cpp @@ -179,6 +179,7 @@ void FunGT::set(const std::function& renderLambda){ } else if (sample == 0) { m_progressiveTracer->updateCamera( &m_camera, width, height); + m_progressiveTracer->reloadLights(m_sceneManager); } m_progressiveTracer->renderSampleInterop( @@ -189,16 +190,11 @@ void FunGT::set(const std::function& renderLambda){ auto* viewport = m_layerStack.get(); if (!viewport) return; - if (!m_progressiveTracer->isInitialized()) { - m_progressiveTracer->initialize( - &m_camera, m_sceneManager, width, height); - } else if (sample == 0) { - m_progressiveTracer->updateCamera( - &m_camera, width, height); + if (sample == 0 || !m_progressiveTracer->isInitialized()) { + m_progressiveTracer->initialize(&m_camera, m_sceneManager, width, height); } - - m_progressiveTracer->renderSample( - sample, viewport->getPathTraceTexture()); + //Render one Sample + m_progressiveTracer->renderSample(sample, viewport->getPathTraceTexture()); }); m_ViewPortLayer->setPathTraceReleaseInteropFunction([this]() { m_progressiveTracer->releaseOpenGLInteropResources(); From 8da1e1e7b14bf68f2a73776d9ff04f6ae1d21301 Mon Sep 17 00:00:00 2001 From: juanchuletas Date: Sun, 23 Aug 2026 21:03:12 -0600 Subject: [PATCH 7/7] feat: adding dirty flags to avoid not required scene buffer uploads --- GUI/lights_editor_window.hpp | 29 +++--- GUI/material_editor_window.hpp | 17 +++- PBR/Render/include/opencl_renderer.hpp | 4 +- PBR/Render/ocl_kernels/path_tracer.cl | 9 +- .../opencl/fgt_opencl_data_translator.hpp | 30 ++++++ PBR/Render/src/opencl_renderer.cpp | 99 ++++++++++++++++++- PBR/Space/space.cpp | 99 ++++++++++++++++++- PBR/Space/space.hpp | 3 + SceneManager/scene_manager.hpp | 27 ++++- ViewPort/progressive_path_tracer_viewport.cpp | 7 ++ ViewPort/progressive_path_tracer_viewport.hpp | 1 + funGT/fungt.cpp | 12 ++- 12 files changed, 313 insertions(+), 24 deletions(-) diff --git a/GUI/lights_editor_window.hpp b/GUI/lights_editor_window.hpp index fb22c72..5975429 100644 --- a/GUI/lights_editor_window.hpp +++ b/GUI/lights_editor_window.hpp @@ -74,6 +74,7 @@ class LightEditorWindow : public ImGuiWindow { if (m_selectedIndex >= 0 && m_selectedIndex < (int)lights.size()) { if (ImGui::Button("Remove Selected")) { lights.erase(lights.begin() + m_selectedIndex); + m_sceneManager->markRaySpaceDirty(SceneManager::RaySpaceDirtyLights); m_selectedIndex = -1; ImGui::End(); return; @@ -106,15 +107,17 @@ class LightEditorWindow : public ImGuiWindow { ImGui::Spacing(); ImGui::Separator(); + bool lightChanged = false; + // Common properties ImGui::Text("Position"); - ImGui::DragFloat3("##Pos", &light.position.x, 0.1f); + lightChanged |= ImGui::DragFloat3("##Pos", &light.position.x, 0.1f); ImGui::Text("Color"); - ImGui::ColorEdit3("##Color", &light.color.x, ImGuiColorEditFlags_Float); + lightChanged |= ImGui::ColorEdit3("##Color", &light.color.x, ImGuiColorEditFlags_Float); ImGui::Text("Power"); - ImGui::DragFloat("##Power", &light.power, 0.1f, 0.0f, 10000.f); + lightChanged |= ImGui::DragFloat("##Power", &light.power, 0.1f, 0.0f, 10000.f); ImGui::Spacing(); ImGui::Separator(); @@ -124,33 +127,37 @@ class LightEditorWindow : public ImGuiWindow { { case SceneLightType::Point: ImGui::Text("Radius"); - ImGui::DragFloat("##Radius", &light.radius, 0.01f, 0.0f, 100.f); + lightChanged |= ImGui::DragFloat("##Radius", &light.radius, 0.01f, 0.0f, 100.f); break; case SceneLightType::Sun: ImGui::Text("Direction"); - ImGui::DragFloat3("##Dir", &light.direction.x, 0.01f, -1.f, 1.f); + lightChanged |= ImGui::DragFloat3("##Dir", &light.direction.x, 0.01f, -1.f, 1.f); break; case SceneLightType::Spot: ImGui::Text("Direction"); - ImGui::DragFloat3("##SpotDir", &light.direction.x, 0.01f, -1.f, 1.f); + lightChanged |= ImGui::DragFloat3("##SpotDir", &light.direction.x, 0.01f, -1.f, 1.f); ImGui::Text("Inner Angle"); - ImGui::DragFloat("##Inner", &light.innerAngle, 0.5f, 0.f, light.outerAngle); + lightChanged |= ImGui::DragFloat("##Inner", &light.innerAngle, 0.5f, 0.f, light.outerAngle); ImGui::Text("Outer Angle"); - ImGui::DragFloat("##Outer", &light.outerAngle, 0.5f, light.innerAngle, 90.f); + lightChanged |= ImGui::DragFloat("##Outer", &light.outerAngle, 0.5f, light.innerAngle, 90.f); break; case SceneLightType::Area: ImGui::Text("Normal"); - ImGui::DragFloat3("##Normal", &light.normal.x, 0.01f, -1.f, 1.f); + lightChanged |= ImGui::DragFloat3("##Normal", &light.normal.x, 0.01f, -1.f, 1.f); ImGui::Text("Size"); - ImGui::DragFloat2("##Size", &light.size.x, 0.1f, 0.1f, 100.f); + lightChanged |= ImGui::DragFloat2("##Size", &light.size.x, 0.1f, 0.1f, 100.f); break; } + if (lightChanged) { + m_sceneManager->markRaySpaceDirty(SceneManager::RaySpaceDirtyLights); + } + ImGui::End(); } }; -#endif // _LIGHT_EDITOR_WINDOW_H_ \ No newline at end of file +#endif // _LIGHT_EDITOR_WINDOW_H_ diff --git a/GUI/material_editor_window.hpp b/GUI/material_editor_window.hpp index ddecc55..92f356f 100644 --- a/GUI/material_editor_window.hpp +++ b/GUI/material_editor_window.hpp @@ -131,28 +131,30 @@ class MaterialEditorWindow : public ImGuiWindow { ImGui::Separator(); ImGui::Spacing(); + bool materialChanged = false; + ImGui::Text("Base Color"); - ImGui::ColorEdit3("##BaseColor", &mat.m_baseColor.x, ImGuiColorEditFlags_Float); + materialChanged |= ImGui::ColorEdit3("##BaseColor", &mat.m_baseColor.x, ImGuiColorEditFlags_Float); ImGui::Spacing(); ImGui::Separator(); ImGui::Spacing(); ImGui::Text("Metallic"); - ImGui::SliderFloat("##Metallic", &mat.m_metallic, 0.0f, 1.0f, "%.3f"); + materialChanged |= ImGui::SliderFloat("##Metallic", &mat.m_metallic, 0.0f, 1.0f, "%.3f"); ImGui::Text("Roughness"); - ImGui::SliderFloat("##Roughness", &mat.m_roughness, 0.05f, 1.0f, "%.3f"); + materialChanged |= ImGui::SliderFloat("##Roughness", &mat.m_roughness, 0.05f, 1.0f, "%.3f"); ImGui::Text("Reflectance"); - ImGui::SliderFloat("##Reflectance", &mat.m_reflectance, 0.0f, 1.0f, "%.3f"); + materialChanged |= ImGui::SliderFloat("##Reflectance", &mat.m_reflectance, 0.0f, 1.0f, "%.3f"); ImGui::Spacing(); ImGui::Separator(); ImGui::Spacing(); // Emission ImGui::Text("Emission"); - ImGui::SliderFloat("##Emission", &mat.m_emission, 0.0f, 50.0f, "%.1f"); + materialChanged |= ImGui::SliderFloat("##Emission", &mat.m_emission, 0.0f, 50.0f, "%.1f"); ImGui::Spacing(); ImGui::Separator(); @@ -160,6 +162,11 @@ class MaterialEditorWindow : public ImGuiWindow { // Reset button if (ImGui::Button("Reset to Default", ImVec2(-1, 0))) { mat = Material::createDefaultMaterial(); + materialChanged = true; + } + + if (materialChanged) { + m_sceneManager->markRaySpaceDirty(SceneManager::RaySpaceDirtyShading); } // Info section diff --git a/PBR/Render/include/opencl_renderer.hpp b/PBR/Render/include/opencl_renderer.hpp index 418bd43..d4fb29b 100644 --- a/PBR/Render/include/opencl_renderer.hpp +++ b/PBR/Render/include/opencl_renderer.hpp @@ -126,8 +126,10 @@ class OpenCL_Renderer : public IComputeRenderer { const PBRCamera& camera, int samplesPerPixel, int sampleOffset, - GLuint glBufferID + GLuint glBufferID, + bool sceneShadingDirty ); + void updateSceneShading(const std::vector& triangles); void releaseOpenGLInteropResources() noexcept; void invalidateScene() { m_sceneUploaded = false; } private: diff --git a/PBR/Render/ocl_kernels/path_tracer.cl b/PBR/Render/ocl_kernels/path_tracer.cl index 8b37dba..9a8d766 100644 --- a/PBR/Render/ocl_kernels/path_tracer.cl +++ b/PBR/Render/ocl_kernels/path_tracer.cl @@ -157,17 +157,20 @@ inline float3 fgt_path_trace_cook_torrance( distance - 0.001f); if (!blocked) { + const float3 emissive_light = + light_emission / + (distance_squared + 1.0e-6f); const float3 nee = fgt_evaluate_cook_torrance( normal, view_direction, light_direction, hit.material, - light_emission); + emissive_light); const float geometry_term = - light_cosine / distance_squared; + light_cosine / light_pdf; radiance += throughput * nee * - geometry_term / light_pdf; + geometry_term; } } } diff --git a/PBR/Render/shared/opencl/fgt_opencl_data_translator.hpp b/PBR/Render/shared/opencl/fgt_opencl_data_translator.hpp index eda5d21..3a4207c 100644 --- a/PBR/Render/shared/opencl/fgt_opencl_data_translator.hpp +++ b/PBR/Render/shared/opencl/fgt_opencl_data_translator.hpp @@ -23,6 +23,12 @@ struct fgt_opencl_scene_data { std::vector bvh_nodes; }; +struct fgt_opencl_shading_data { + std::vector triangle_shading; + std::vector materials; + std::vector emissive_triangles; +}; + inline fgt_vec3 translate_vec3(const fungt::Vec3& value) { return {value.x, value.y, value.z}; @@ -167,6 +173,30 @@ inline fgt_rayspace_camera translate_rayspace_camera(const PBRCamera& camera) translate_vec4(camera.getBasisW()) }; } + +inline fgt_opencl_shading_data translate_shading_data( + const std::vector& triangles) +{ + fgt_opencl_shading_data result; + result.triangle_shading.reserve(triangles.size()); + + for (std::size_t index = 0; index < triangles.size(); ++index) { + const Triangle& triangle = triangles[index]; + const fgt_int32 material_index = + find_or_add_material(triangle.material, result.materials); + + result.triangle_shading.push_back( + translate_triangle_shading(triangle, material_index)); + + if (triangle.material.emission > 0.0f) { + result.emissive_triangles.push_back( + static_cast(index)); + } + } + + return result; +} + inline fgt_opencl_scene_data translate_scene_data( const std::vector& triangles, const std::vector& nodes) diff --git a/PBR/Render/src/opencl_renderer.cpp b/PBR/Render/src/opencl_renderer.cpp index 2fc292c..94db2b1 100644 --- a/PBR/Render/src/opencl_renderer.cpp +++ b/PBR/Render/src/opencl_renderer.cpp @@ -605,6 +605,96 @@ void OpenCL_Renderer::loadSceneLigths(const std::vector &lights) } } +void OpenCL_Renderer::updateSceneShading( + const std::vector& triangles) +{ + if (!m_oclcontext || !m_sceneUploaded) { + throw std::runtime_error( + "updateSceneShading: OpenCL scene is not uploaded."); + } + if (triangles.size() != m_raySpaceBuffer.numTriangles) { + throw std::invalid_argument( + "updateSceneShading: triangle count changed."); + } + + const fgt::opencl::fgt_opencl_shading_data shadingData = + fgt::opencl::translate_shading_data(triangles); + + cl_mem newTriangleShading = nullptr; + cl_mem newMaterials = nullptr; + cl_mem newEmissiveTriangles = nullptr; + cl_int err = CL_SUCCESS; + + if (!shadingData.triangle_shading.empty()) { + newTriangleShading = clCreateBuffer( + m_oclcontext, + CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + shadingData.triangle_shading.size() * sizeof(fgt_triangle_shading), + const_cast( + shadingData.triangle_shading.data()), + &err); + if (err != CL_SUCCESS || !newTriangleShading) { + throw std::runtime_error( + "updateSceneShading: failed to create triangle shading buffer, error " + + std::to_string(err)); + } + } + + if (!shadingData.materials.empty()) { + newMaterials = clCreateBuffer( + m_oclcontext, + CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + shadingData.materials.size() * sizeof(fgt_material_data), + const_cast(shadingData.materials.data()), + &err); + if (err != CL_SUCCESS || !newMaterials) { + if (newTriangleShading) { + clReleaseMemObject(newTriangleShading); + } + throw std::runtime_error( + "updateSceneShading: failed to create material buffer, error " + + std::to_string(err)); + } + } + + if (!shadingData.emissive_triangles.empty()) { + newEmissiveTriangles = clCreateBuffer( + m_oclcontext, + CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + shadingData.emissive_triangles.size() * sizeof(fgt_int32), + const_cast(shadingData.emissive_triangles.data()), + &err); + if (err != CL_SUCCESS || !newEmissiveTriangles) { + if (newMaterials) { + clReleaseMemObject(newMaterials); + } + if (newTriangleShading) { + clReleaseMemObject(newTriangleShading); + } + throw std::runtime_error( + "updateSceneShading: failed to create emissive triangle buffer, error " + + std::to_string(err)); + } + } + + if (m_raySpaceBuffer.emissiveTriangles) { + clReleaseMemObject(m_raySpaceBuffer.emissiveTriangles); + } + if (m_raySpaceBuffer.materials) { + clReleaseMemObject(m_raySpaceBuffer.materials); + } + if (m_raySpaceBuffer.triangleShading) { + clReleaseMemObject(m_raySpaceBuffer.triangleShading); + } + + m_raySpaceBuffer.triangleShading = newTriangleShading; + m_raySpaceBuffer.materials = newMaterials; + m_raySpaceBuffer.emissiveTriangles = newEmissiveTriangles; + m_raySpaceBuffer.numMaterials = shadingData.materials.size(); + m_raySpaceBuffer.numEmissiveTriangles = + shadingData.emissive_triangles.size(); +} + void OpenCL_Renderer::releaseRaySpaceBuffer( OpenCLRaySpaceBuffer& buffers) noexcept { @@ -704,7 +794,8 @@ void OpenCL_Renderer::RenderSceneOpenGLInterop(int width, const std::vector& emissiveTriIndices, const PBRCamera& camera, int samplesPerPixel, int sampleOffset, - GLuint glBufferID) + GLuint glBufferID, + bool sceneShadingDirty) { if (width <= 0 || height <= 0) { throw std::invalid_argument( @@ -730,6 +821,12 @@ void OpenCL_Renderer::RenderSceneOpenGLInterop(int width, clFinish(m_oclqueue); } + // Update material/emission buffers only when the scene changed. + if (sceneShadingDirty) { + updateSceneShading(triangles); + clFinish(m_oclqueue); + } + // Update lights when a new accumulation starts. if (sampleOffset == 0 || !m_raySpaceBuffer.lights) { loadSceneLigths(lights); diff --git a/PBR/Space/space.cpp b/PBR/Space/space.cpp index 169a6e6..d79c836 100644 --- a/PBR/Space/space.cpp +++ b/PBR/Space/space.cpp @@ -189,7 +189,9 @@ void Space::RenderOpenGLInterop( m_camera, m_samplesPerPixel, sampleOffset, - glBufferID); + glBufferID, + m_sceneShadingDirty); + m_sceneShadingDirty = false; #else throw std::runtime_error( "OpenCL support is not enabled."); @@ -399,6 +401,101 @@ void Space::loadLightsFromScene(const std::vector& sceneLights) } } +void Space::loadShadingFromScene(const std::vector>& renderables) +{ + std::vector sceneMaterials; + sceneMaterials.reserve(m_triangles.size()); + + for (const auto& obj : renderables) { + auto simpleModel = std::dynamic_pointer_cast(obj); + if (simpleModel) { + const auto& meshes = simpleModel->getModel().getMeshes(); + for (const auto& meshPtr : meshes) { + MaterialData material; + material.baseColor[0] = 0.922f; + material.baseColor[1] = 0.467f; + material.baseColor[2] = 0.882f; + material.metallic = 0.0f; + material.roughness = 0.5f; + material.reflectance = 0.05f; + material.emission = 0.0f; + material.baseColorTexIdx = -1; + + if (!meshPtr->m_texture.empty()) { + material.baseColorTexIdx = + m_computeRenderer->textures().loadTexture(meshPtr->m_texture[0].getPath()); + } + + if (!meshPtr->m_material.empty()) { + const auto& source = meshPtr->m_material[0]; + material.baseColor[0] = source.m_baseColor.x; + material.baseColor[1] = source.m_baseColor.y; + material.baseColor[2] = source.m_baseColor.z; + material.metallic = source.m_metallic; + material.roughness = source.m_roughness; + material.reflectance = source.m_reflectance; + material.emission = source.m_emission; + } + + const std::size_t triangleCount = meshPtr->m_index.size() / 3; + sceneMaterials.insert(sceneMaterials.end(), triangleCount, material); + } + continue; + } + + auto simpleGeometry = std::dynamic_pointer_cast(obj); + if (simpleGeometry) { + auto primitive = simpleGeometry->getPrimitive(); + if (!primitive) { + continue; + } + const Primitive* constPrimitive = primitive.get(); + const std::vector& indices = constPrimitive->getIndices(); + + MaterialData material; + const auto& source = simpleGeometry->getMaterial(); + material.baseColor[0] = source.baseColor.x; + material.baseColor[1] = source.baseColor.y; + material.baseColor[2] = source.baseColor.z; + material.metallic = source.metallic; + material.roughness = source.roughness; + material.reflectance = 0.05f; + material.emission = 0.0f; + material.baseColorTexIdx = -1; + + if (simpleGeometry->isTexturized()) { + material.baseColorTexIdx = + m_computeRenderer->textures().loadTexture(constPrimitive->texture.getPath()); + } + + const std::size_t triangleCount = indices.size() / 3; + sceneMaterials.insert(sceneMaterials.end(), triangleCount, material); + } + } + + if (sceneMaterials.size() != m_triangles.size()) { + std::cerr << "Space::loadShadingFromScene() triangle count mismatch: " + << sceneMaterials.size() << " scene materials for " + << m_triangles.size() << " triangles" << std::endl; + return; + } + + for (std::size_t i = 0; i < m_triangles.size(); ++i) { + const std::size_t sourceIndex = + m_bvh_indices.empty() ? i : static_cast(m_bvh_indices[i]); + m_triangles[i].material = sceneMaterials[sourceIndex]; + } + + m_emissiveTriIndices.clear(); + for (std::size_t i = 0; i < m_triangles.size(); ++i) { + if (m_triangles[i].material.emission > 0.0f) { + m_emissiveTriIndices.push_back(static_cast(i)); + } + } + + m_sceneShadingDirty = true; +} + void Space::invalidateScene() { #ifdef FUNGT_USE_OPENCL diff --git a/PBR/Space/space.hpp b/PBR/Space/space.hpp index 4a460b8..3c523d5 100644 --- a/PBR/Space/space.hpp +++ b/PBR/Space/space.hpp @@ -15,6 +15,7 @@ #include "PBR/Render/include/compute_backends.hpp" #include "PBR/Render/include/icompute_renderer.hpp" #include "PBR/Render/include/cpu_renderer.hpp" +#include "Renderable/renderable.hpp" #include "PBR/BVH/bvh_builder.hpp" #include "SimpleGeometry/simple_geometry.hpp" @@ -30,6 +31,7 @@ class Space { std::vector m_bvh_nodes; std::vector m_bvh_indices; std::vector m_emissiveTriIndices; + bool m_sceneShadingDirty = false; public: Space(); @@ -49,6 +51,7 @@ class Space { void LoadModelToRender(const SimpleModel& model); void LoadGeometryToRender(const SimpleGeometry& geometry); void loadLightsFromScene(const std::vector& sceneLights); + void loadShadingFromScene(const std::vector>& renderables); void invalidateScene(); void static SaveFrameBufferAsPNG(const std::vector& framebuffer, int width, int height); static void SaveFrameBufferAsPNG(const std::vector& framebuffer, diff --git a/SceneManager/scene_manager.hpp b/SceneManager/scene_manager.hpp index 59b01f7..78484ce 100644 --- a/SceneManager/scene_manager.hpp +++ b/SceneManager/scene_manager.hpp @@ -11,6 +11,7 @@ #include "GraphicsRenderBackend/gpu_buffer.hpp" #include "VertexGL/shaderStorageBufferObejct.hpp" #include +#include #include struct ShaderBucket { @@ -20,6 +21,15 @@ struct ShaderBucket { class SceneManager{ + public: + enum RaySpaceDirty : uint32_t { + RaySpaceDirtyNone = 0, + RaySpaceDirtyLights = 1u << 0, + RaySpaceDirtyShading = 1u << 1, + RaySpaceDirtyGeometry = 1u << 2, + RaySpaceDirtyTextures = 1u << 3 + }; + private: std::unique_ptr m_shader; std::vector> m_VectorOfRenderNodes; @@ -46,6 +56,11 @@ class SceneManager{ bool m_modelMatrixSSBOInitialized = false; std::unordered_map m_modelMatrixIndex; glm::vec3 m_ambientColor = glm::vec3(0.3f, 0.3f, 0.3f); + uint32_t m_raySpaceDirty = + RaySpaceDirtyLights | + RaySpaceDirtyShading | + RaySpaceDirtyGeometry | + RaySpaceDirtyTextures; static constexpr int MAX_LIGHTS = 32; std::unique_ptr m_matricesUBO; @@ -83,10 +98,20 @@ class SceneManager{ glm::vec3& getLightSpecular() { return m_lightSpecular; } //Light - void addLight(const SceneLight& light) { m_lights.push_back(light); } + void addLight(const SceneLight& light) { + m_lights.push_back(light); + markRaySpaceDirty(RaySpaceDirtyLights); + } const std::vector& getLights() const { return m_lights; } std::vector& getLights() { return m_lights; } size_t getLightCount() const { return m_lights.size(); } + void markRaySpaceDirty(uint32_t flags) { m_raySpaceDirty |= flags; } + bool isRaySpaceDirty(uint32_t flags) const { + return (m_raySpaceDirty & flags) != 0; + } + void clearRaySpaceDirty(uint32_t flags) { + m_raySpaceDirty &= ~flags; + } }; diff --git a/ViewPort/progressive_path_tracer_viewport.cpp b/ViewPort/progressive_path_tracer_viewport.cpp index 78e9325..1c75b4b 100644 --- a/ViewPort/progressive_path_tracer_viewport.cpp +++ b/ViewPort/progressive_path_tracer_viewport.cpp @@ -112,6 +112,13 @@ void ProgressivePathTracer::reloadLights(std::shared_ptr sceneMana } } +void ProgressivePathTracer::reloadSceneShading(std::shared_ptr sceneManager) +{ + if (m_space && sceneManager) { + m_space->loadShadingFromScene(sceneManager->getRenderable()); + } +} + void ProgressivePathTracer::releaseOpenGLInteropResources() { if (m_space) { diff --git a/ViewPort/progressive_path_tracer_viewport.hpp b/ViewPort/progressive_path_tracer_viewport.hpp index e413f9a..bc2c926 100644 --- a/ViewPort/progressive_path_tracer_viewport.hpp +++ b/ViewPort/progressive_path_tracer_viewport.hpp @@ -30,6 +30,7 @@ class ProgressivePathTracer { int width, int height, bool ogl_interop = false); void updateCamera(Camera* viewportCam, int width, int height); void reloadLights(std::shared_ptr sceneManager); + void reloadSceneShading(std::shared_ptr sceneManager); virtual void renderSample(int sample, uint32_t targetTexture) = 0; virtual void renderSampleInterop(int sample, uint32_t targetPBO) = 0; diff --git a/funGT/fungt.cpp b/funGT/fungt.cpp index 3e427a7..02a04e9 100644 --- a/funGT/fungt.cpp +++ b/funGT/fungt.cpp @@ -176,10 +176,20 @@ void FunGT::set(const std::function& renderLambda){ if (!m_progressiveTracer->isInitialized()) { m_progressiveTracer->initialize( &m_camera, m_sceneManager, width, height, true); + m_sceneManager->clearRaySpaceDirty( + SceneManager::RaySpaceDirtyLights | + SceneManager::RaySpaceDirtyShading); } else if (sample == 0) { m_progressiveTracer->updateCamera( &m_camera, width, height); - m_progressiveTracer->reloadLights(m_sceneManager); + if (m_sceneManager->isRaySpaceDirty(SceneManager::RaySpaceDirtyLights)) { + m_progressiveTracer->reloadLights(m_sceneManager); + m_sceneManager->clearRaySpaceDirty(SceneManager::RaySpaceDirtyLights); + } + if (m_sceneManager->isRaySpaceDirty(SceneManager::RaySpaceDirtyShading)) { + m_progressiveTracer->reloadSceneShading(m_sceneManager); + m_sceneManager->clearRaySpaceDirty(SceneManager::RaySpaceDirtyShading); + } } m_progressiveTracer->renderSampleInterop(