Merge

.vscode/imgui.ini

@@ -0,0 +1,4 @@
+[Window][Debug##Default]
+Pos=60,60
+Size=400,400
+

.vscode/launch.json

@@ -0,0 +1,21 @@
+{
+    // Use IntelliSense to learn about possible attributes.
+    // Hover to view descriptions of existing attributes.
+    // For more information, visit: https://go.microsoft.com/fwlink/?linkid=830387
+    "version": "0.2.0",
+    "configurations": [
+        {
+            "name": "(lldb) Launch",
+            "type": "cppdbg",
+            "request": "launch",
+            "program": "${workspaceFolder}/build/RayTracer",
+            "args": [],
+            "stopAtEntry": false,
+            "cwd": "${workspaceFolder}/build",
+            "environment": [],
+            "externalConsole": false,
+            "MIMode": "lldb"
+        }
+
+    ]
+}

CMakeLists.txt

@@ -12,26 +12,29 @@ set(CMAKE_TOOLCHAIN_FILE ${CMAKE_CURRENT_SOURCE_DIR}/external/vcpkg/scripts/buil
 project(RayTracer)
 
 find_package(Vulkan REQUIRED)
+find_package(VulkanMemoryAllocator CONFIG REQUIRED)
 find_package(glew CONFIG REQUIRED)
 find_package(assimp CONFIG REQUIRED)
 find_package(glfw3 CONFIG REQUIRED)
 find_package(glm CONFIG REQUIRED)
 find_package(Ktx CONFIG REQUIRED)
 find_package(imgui CONFIG REQUIRED)
-find_package(OpenMP REQUIRED)
 
 add_executable(RayTracer "")
 target_include_directories(RayTracer PUBLIC src/)
-target_include_directories(RayTracer PUBLIC ${VCPKG_INSTALLED_DIR}/x64-windows/include)
 target_link_libraries(RayTracer PUBLIC Vulkan::Vulkan)
 target_link_libraries(RayTracer PUBLIC Vulkan::Headers)
+target_link_libraries(RayTracer PUBLIC GPUOpen::VulkanMemoryAllocator)
 target_link_libraries(RayTracer PUBLIC assimp::assimp)
 target_link_libraries(RayTracer PUBLIC glfw)
 target_link_libraries(RayTracer PUBLIC imgui::imgui)
 target_link_libraries(RayTracer PUBLIC GLEW::GLEW)
 target_link_libraries(RayTracer PUBLIC glm::glm)
 target_link_libraries(RayTracer PUBLIC KTX::ktx)
+if(WIN32)
+target_include_directories(RayTracer PUBLIC ${VCPKG_INSTALLED_DIR}/x64-windows/include)
 target_link_libraries(RayTracer PUBLIC ${VCPKG_INSTALLED_DIR}/x64-windows/lib/slang.lib)
-target_link_libraries(RayTracer PUBLIC OpenMP::OpenMP_CXX)
-
+elseif(APPLE)
+target_include_directories(RayTracer PUBLIC ${VCPKG_INSTALLED_DIR}/arm64-osx/include)
+endif()
 add_subdirectory(src/)

res/shaders/ClosestHit.slang

@@ -0,0 +1,125 @@
+import Common;
+
+
+[shader("closesthit")]
+void closestHit(inout RayPayload hitValue, in BuiltInTriangleIntersectionAttributes attr)
+{
+    hitValue.hit = true;
+    // todo: replace with anyhit shader
+    if(hitValue.anyHit)
+        return;
+    const float3 barycentricCoords = float3(1.0f - attr.barycentrics.x - attr.barycentrics.y, attr.barycentrics.x, attr.barycentrics.y);
+
+    ModelReference m = pParams.modelData[InstanceID()];
+
+    // offset into the index buffer
+    uint indexOffset = m.indicesOffset;
+    // added to indices to reference correct part of global mesh pool
+    uint vertexOffset = m.positionOffset;
+
+    uint vertexIndex0 = vertexOffset + pParams.indexBuffer[indexOffset + 3 * PrimitiveIndex() + 0];
+    uint vertexIndex1 = vertexOffset + pParams.indexBuffer[indexOffset + 3 * PrimitiveIndex() + 1];
+    uint vertexIndex2 = vertexOffset + pParams.indexBuffer[indexOffset + 3 * PrimitiveIndex() + 2];
+
+    Vertex attr0 = loadVertex(vertexIndex0);
+    Vertex attr1 = loadVertex(vertexIndex1);
+    Vertex attr2 = loadVertex(vertexIndex2);
+
+    Vertex vert = Vertex.interpolate(attr0, attr1, attr2, barycentricCoords);
+
+    float3 normalLight = dot(vert.normal, WorldRayDirection()) < 0 ? vert.normal : -vert.normal;
+
+    MaterialParameter mat; // TOOD:
+    
+    hitValue.depth++;
+    float3 localAccRad = float3(0);
+    float3 rnd = rand01(uint3(vertexIndex0, vertexIndex1, vertexIndex2));
+    //float kt = ka + ks;
+    //float s = -log(rnd.z) / kt;
+    //float3 xs = r.o + s * r.d;
+    //if (s < t) {
+    //    float p = kt * rnd.z;
+    //    if (depth > 5) {
+    //        if (rnd.z >= p) break;
+    //        else accmat /= p;
+    //    }
+    //    float3 ldirect = nextEventEstimation(accmat, r.d, xs, -r.d, kt, true, rnd);
+    //    accrad += (fogEmm + ks * ldirect) / kt;
+    //    accmat *= ks / kt;
+    //    rayDesc.Origin = xs;
+    //    rayDesc.Direction = float3(
+    //        cos(2*PI*rnd.x)*sqrt(1-rnd.y*rnd.y),
+    //        sin(2*PI*rnd.x)*sqrt(1-rnd.y*rnd.y),
+    //        rnd.y
+    //    );
+    //    continue;
+    //}
+
+    //float p = max(max(mat.albedo.x, mat.albedo.y), mat.albedo.z);
+    //if(hitValue.depth > 5) {
+    //    if (rnd.z >= p) return;
+    //    else hitValue.accmat /= p;
+    //}
+
+
+    //-- Ideal DIFFUSE reflection
+    //if(bool(useNEE)) {
+    //    accrad += nextEventEstimation(accmat, r.d, params.x, params.nl, kt, false, rnd);
+    //}
+    for(uint i = 0; i < pSamps.numDirectionalLights; ++i) {
+        float3 x = vert.position;
+        float3 l = -pParams.directionalLights[i].direction.xyz;
+        RayDesc rayDesc;
+        rayDesc.TMax = 10000.0f;
+        rayDesc.TMin = 0.001f;
+        rayDesc.Origin = x;
+        rayDesc.Direction = l;
+        RayPayload payload;
+        payload.depth = hitValue.depth;
+        payload.emissive = 1;
+        payload.anyHit = true;
+        TraceRay(pParams.scene, 0, 0xff, 0, 0, 0, rayDesc, payload);
+
+        // we have missed all geometry, so directional light is affecting us
+        if(!payload.hit) {
+            localAccRad += mat.shade(vert.normal, -WorldRayDirection(), -pParams.directionalLights[i].direction, pParams.directionalLights[i].color);
+        }
+    }
+    for(uint i = 0; i < pSamps.numPointLights; ++i) {
+        RayPayload payload;
+        float3 x = vert.position;
+        float3 l = pParams.pointLights[i].position - vert.position;
+        // todo: cancel if light too far away to affect
+
+        RayDesc rayDesc;
+        rayDesc.TMax = 1.0f;
+        rayDesc.TMin = 0.001f;
+        rayDesc.Origin = x;
+        rayDesc.Direction = l;
+        TraceRay(pParams.scene, 0, 0xff, 0, 0, 0, rayDesc, payload);
+    
+        // hitting only after the light
+        if(!payload.hit) {
+            localAccRad += mat.shade(vert.normal, -WorldRayDirection(), normalize(l), pParams.pointLights[i].color);
+        }
+    }
+    hitValue.light += localAccRad;
+    // Indirect Illumination: cosine-weighted importance sampling
+    if(hitValue.depth < 12) {
+    	float r1 = 2 * PI * rnd.x, r2 = rnd.y, r2s = sqrt(r2);
+        float3 w = normalLight;
+        float3 u = normalize((cross(abs(w.x)>0.1 ? float3(0,1,0) : float3(1,0,0), w)));
+        float3 v = cross(w,u);
+        RayDesc rayDesc;
+        rayDesc.TMax = 10000.0f;
+        rayDesc.TMin = 0.001f;
+        rayDesc.Origin = vert.position;
+        rayDesc.Direction = normalize(u*cos(r1)*r2s + v * sin(r1)*r2s + w * sqrt(1 - r2));
+        RayPayload payload;
+        payload.light = float3(0);
+        payload.emissive = 0;   // in the next bounce, consider reflective part only!
+        payload.depth = hitValue.depth+1;
+        payload.anyHit = false;
+        TraceRay(pParams.scene, 0, 0xff, 0, 0, 0, rayDesc, payload);
+    }
+}

res/shaders/Common.slang

@@ -0,0 +1,121 @@
+const static float PI = 3.1415926535897932f;
+
+struct Camera
+{
+    float3 cameraPosition;
+    float f;
+    float3 cameraForward;
+    float S_O;
+    float3 fogEmm;
+    float ks;
+    float A;
+    float ka;
+};
+
+struct MaterialParameter
+{
+    float3 albedo = float3(1, 1, 1);
+    float alpha = 1;
+    float3 specularColor = float3(1, 1, 1);
+    float shininess = 0.04;
+    float3 emissive = float3(0, 0, 0);
+    float3 shade(float3 normal, float3 viewDir, float3 lightDir, float3 lightColor)
+    {
+        float diffuse = max(dot(normal, lightDir), 0);
+        float3 h = normalize(lightDir + viewDir);
+        float specular = pow(clamp(dot(normal, h), 0, 1), shininess);
+
+        return (albedo * diffuse * lightColor);
+    }
+};
+
+struct ModelReference
+{
+  uint32_t positionOffset = 0;
+  uint32_t indicesOffset = 0;
+  uint32_t numIndices = 0;
+};
+
+struct PointLight
+{
+  float3 position = float3(0, 0, 0);
+  float3 color = float3(1, 1, 1);
+  float attenuation = 1;
+};
+
+struct DirectionalLight
+{
+  float3 direction = float3(0, 1, 0);
+  float3 color = float3(1, 1, 1);
+};
+
+struct RaytracingParams
+{
+    Camera cam;
+    RaytracingAccelerationStructure scene;
+    RWTexture2D<float4> radianceAccumulator;
+    RWTexture2D<float4> image;
+    StructuredBuffer<ModelReference> modelData;
+    StructuredBuffer<MaterialParameter> materialData;
+    StructuredBuffer<float> positions;
+    StructuredBuffer<float> texCoords;
+    StructuredBuffer<float> normals;
+    StructuredBuffer<DirectionalLight> directionalLights;
+    StructuredBuffer<PointLight> pointLights;
+    StructuredBuffer<uint32_t> indexBuffer;
+};
+ParameterBlock<RaytracingParams> pParams;
+
+struct Vertex
+{
+    float3 position;
+    float2 texCoords;
+    float3 normal;
+    static Vertex interpolate(Vertex f0, Vertex f1, Vertex f2, float3 barycentricCoords)
+    {
+        Vertex vert;
+        vert.position = f0.position * barycentricCoords.x + f1.position * barycentricCoords.y + f2.position * barycentricCoords.z;
+        vert.texCoords = f0.texCoords * barycentricCoords.x + f1.texCoords * barycentricCoords.y + f2.texCoords * barycentricCoords.z;
+        vert.normal = f0.normal * barycentricCoords.x + f1.normal * barycentricCoords.y + f2.normal * barycentricCoords.z;
+        return vert;
+    }
+};
+
+Vertex loadVertex(uint32_t vertexIndex)
+{
+    Vertex vert;
+    vert.position = float3(pParams.positions[vertexIndex * 3 + 0], pParams.positions[vertexIndex * 3 + 1], pParams.positions[vertexIndex * 3 + 2]);
+    vert.texCoords = float2(pParams.texCoords[vertexIndex * 2 + 0], pParams.texCoords[vertexIndex * 2 + 1]);
+    vert.normal = float3(pParams.normals[vertexIndex * 3 + 0], pParams.normals[vertexIndex * 3 + 1], pParams.normals[vertexIndex * 3 + 2]);
+    return vert;
+}
+
+struct SampleParams
+{
+    uint pass;
+    uint samplesPerPixel;
+    uint numDirectionalLights;
+    uint numPointLights;
+};
+layout(push_constant)
+ConstantBuffer<SampleParams> pSamps;
+
+struct Ray
+{
+    float3 o;
+    float3 d;
+};
+
+struct RayPayload
+{
+    float3 light;
+    float emissive;
+    uint depth;
+    bool hit;
+    bool anyHit;
+};
+
+float3 rand01(uint3 x){                   // pseudo-random number generator
+    for (int i=3; i-->0;) x = ((x>>8U)^x.yzx)*1103515245U;
+    return float3(x)*(1.0/float(0xffffffffU));
+}

res/shaders/Miss.slang

@@ -0,0 +1,8 @@
+import Common;
+
+[shader("miss")]
+void miss(inout RayPayload p)
+{
+    p.light = float3(0, 0, 0);
+    p.hit = false;
+}

res/shaders/RayGen.slang

@@ -0,0 +1,57 @@
+import Common;
+
+[shader("raygeneration")]
+void raygen()
+{
+    if(pSamps.pass == pSamps.samplesPerPixel) return;
+    uint2 pix = DispatchRaysIndex().xy;
+    uint2 imgdim = DispatchRaysDimensions().xy;
+
+    //-- define cam
+    Ray cam = Ray(pParams.cam.cameraPosition, pParams.cam.cameraForward);
+    float3 cx = -normalize(cross(cam.d, abs(cam.d.y) < 0.9 ? float3(0, 1, 0) : float3(0, 0, 1))), cy = cross(cam.d, cx);
+    const float2 sdim = float2(0.036, 0.024);
+
+    float S_I = (pParams.cam.S_O * pParams.cam.f) / (pParams.cam.S_O - pParams.cam.f);
+
+    //-- sample sensor
+    float2 rnd2 = 2*rand01(uint3(pix, pSamps.pass)).xy;   // vvv tent filter sample  
+    float2 tent = float2(rnd2.x<1 ? sqrt(rnd2.x)-1 : 1-sqrt(2-rnd2.x), rnd2.y<1 ? sqrt(rnd2.y)-1 : 1-sqrt(2-rnd2.y));
+    float2 s = ((pix + 0.5 * (0.5 + float2((pSamps.pass/2)%2, pSamps.pass%2) + tent)) / float2(imgdim) - 0.5) * sdim;
+    float3 spos = cam.o + cx*s.x + cy*s.y, lc = cam.o + cam.d * 0.035;           // sample on 3d sensor plane
+    Ray r = Ray(lc, normalize(lc - spos));      // construct ray
+
+
+    //-- setup lens
+    float3 lensP = lc;
+    float3 lensN = -cam.d;
+    float3 lensX = cross(lensN, float3(0, 1, 0)); // the exact vector doesnt matter
+    float3 lensY = cross(lensN, lensX);
+    uint3 rndSeed = uint3(pix, pSamps.pass);
+    float2 rnd01 = rand01(rndSeed).xy;
+
+    float3 lensSample = lensP + rnd01.x * pParams.cam.A * lensX + rnd01.y * pParams.cam.A * lensY;
+
+    float3 focalPoint = cam.o + (pParams.cam.S_O + S_I) * cam.d;
+    float t = dot(focalPoint - r.o, lensN) / dot(r.d, lensN);
+    float3 focus = r.o + t * r.d;
+
+    RayDesc rayDesc;
+    rayDesc.Origin = lensSample;
+    rayDesc.Direction = normalize(focus - lensSample);
+    rayDesc.TMin = 0.001;
+    rayDesc.TMax = 10000.0;
+	
+    const uint maxDepth = 12;
+    RayPayload payload;
+    // initialize accumulated radiance and bxdf
+    payload.light=float3(0);
+    payload.emissive = 1;
+    payload.depth = 1;
+    payload.anyHit = false;
+    TraceRay(pParams.scene, 0, 0xff, 0, 0, 0, rayDesc, payload);
+
+    if(pSamps.pass == 0) pParams.radianceAccumulator[pix] = float4(0);
+    pParams.radianceAccumulator[pix] += float4(payload.light / pSamps.samplesPerPixel, 0);
+    pParams.image[pix] = float4(clamp(pParams.radianceAccumulator[pix].xyz, 0, 1), 1);
+}

src/ThreadPool.h

@@ -25,8 +25,8 @@ private:
     std::mutex queueLock;
     std::condition_variable queueCV;
     std::condition_variable completedCV;
-    uint32_t numRemaining;
-    uint32_t numRunning;
+    uint32_t numRemaining = 0;
+    uint32_t numRunning = 0;
     std::list<Batch> taskQueue;
     std::vector<std::thread> workers;
 };

src/gpu/CMakeLists.txt

@@ -1,4 +1,7 @@
 target_sources(RayTracer
 	PRIVATE
 		GPURenderer.h
-		GPURenderer.cpp)
+		GPURenderer.cpp
+		GPUScene.h
+		GPUScene.cpp
+)

src/gpu/GPURenderer.cpp

@@ -1,13 +1,20 @@
 #include "GPURenderer.h"
+#include "util/ModelLoader.h"
+#include "vulkan/vulkan_enums.hpp"
+#include "vulkan/vulkan_handles.hpp"
+#include "vulkan/vulkan_raii.hpp"
+#include "vulkan/vulkan_structs.hpp"
 #include <slang-com-ptr.h>
 #include <slang.h>
+#define VMA_IMPLEMENTATION
+#include "vk_mem_alloc.h"
 
 GPURenderer::GPURenderer()
-    : instance(nullptr), physicalDevice(nullptr), device(nullptr), queue(nullptr), cmdPool(nullptr), cmdBuffers(nullptr),
-      descriptorLayout(nullptr), descriptorSet(nullptr), descriptorPool(nullptr), pipelineLayout(nullptr), rayGen(nullptr),
-      closestHit(nullptr), miss(nullptr), pipeline(nullptr)
-
 {
+  createDevice();
+  createCommands();
+  createDescriptors();
+  createPipeline();
 }
 
 GPURenderer::~GPURenderer() {}
@@ -29,6 +36,12 @@ void GPURenderer::createDevice()
       }
     }
   }
+  auto properties = physicalDevice.getProperties2<vk::PhysicalDeviceProperties2, vk::PhysicalDeviceAccelerationStructurePropertiesKHR,
+                                                  vk::PhysicalDeviceRayTracingPipelinePropertiesKHR>();
+
+  accelerationProperties = properties.get<vk::PhysicalDeviceAccelerationStructurePropertiesKHR>();
+  rayTracingProperties = properties.get<vk::PhysicalDeviceRayTracingPipelinePropertiesKHR>();
+
   uint32_t computeQueueFamily = 0;
   auto queueProps = physicalDevice.getQueueFamilyProperties();
   for (uint32_t i = 0; i < queueProps.size(); ++i)
@@ -39,73 +52,512 @@ void GPURenderer::createDevice()
       break;
     }
   }
-  float queuePriority = 0.0f;
-  vk::DeviceQueueCreateInfo deviceQueueCreateInfo({}, computeQueueFamily, 1, &queuePriority);
-  vk::DeviceCreateInfo deviceCreateInfo({}, deviceQueueCreateInfo);
+  std::vector<float> queuePriority = {1.0f};
+  auto featureChain = physicalDevice.getFeatures2<vk::PhysicalDeviceFeatures2, vk::PhysicalDeviceRayTracingPipelineFeaturesKHR,
+                                                  vk::PhysicalDeviceAccelerationStructureFeaturesKHR>();
+  auto features = featureChain.get<vk::PhysicalDeviceFeatures2>();
+
+  vk::DeviceQueueCreateInfo deviceQueueCreateInfo({}, computeQueueFamily, queuePriority);
+  const char* extensions[] = {vk::KHRAccelerationStructureExtensionName, vk::KHRRayTracingPipelineExtensionName, vk::KHRDeferredHostOperationsExtensionName};
+  vk::DeviceCreateInfo deviceCreateInfo({}, deviceQueueCreateInfo, {}, extensions, nullptr, &features);
   device = Device(physicalDevice, deviceCreateInfo);
+
+  queue = Queue(device, computeQueueFamily, 0);
+
+  VmaVulkanFunctions vulkanFunctions = {
+      .vkGetInstanceProcAddr = &vkGetInstanceProcAddr,
+      .vkGetDeviceProcAddr = &vkGetDeviceProcAddr,
+  };
+
+  VmaAllocatorCreateInfo allocatorCreateInfo = {
+      .flags = VMA_ALLOCATOR_CREATE_EXT_MEMORY_BUDGET_BIT | VMA_ALLOCATOR_CREATE_BUFFER_DEVICE_ADDRESS_BIT,
+      .physicalDevice = *physicalDevice,
+      .device = *device,
+      .pVulkanFunctions = &vulkanFunctions,
+      .instance = *instance,
+      .vulkanApiVersion = VK_API_VERSION_1_3,
+  };
+
+  vmaCreateAllocator(&allocatorCreateInfo, &allocator);
 }
 
 void GPURenderer::createCommands()
 {
   vk::CommandPoolCreateInfo commandPoolCreateInfo({}, computeQueueFamily);
   cmdPool = CommandPool(device, commandPoolCreateInfo);
-
-  // allocate a CommandBuffer from the CommandPool
-  vk::CommandBufferAllocateInfo commandBufferAllocateInfo(cmdPool, vk::CommandBufferLevel::ePrimary, 10);
-  cmdBuffers = vk::raii::CommandBuffers(device, commandBufferAllocateInfo);
 }
 
 void GPURenderer::createDescriptors()
 {
-  vk::DescriptorSetLayoutBinding descriptorSetLayoutBinding(0, vk::DescriptorType::eUniformBuffer, 1, vk::ShaderStageFlagBits::eVertex);
-  vk::DescriptorSetLayoutCreateInfo descriptorSetLayoutCreateInfo({}, descriptorSetLayoutBinding);
+  vk::DescriptorSetLayoutBinding bindings[] = {
+      // camera
+      vk::DescriptorSetLayoutBinding(0, vk::DescriptorType::eUniformBuffer, 1,
+                                     vk::ShaderStageFlagBits::eRaygenKHR | vk::ShaderStageFlagBits::eClosestHitKHR),
+      // scene acceleration structure
+      vk::DescriptorSetLayoutBinding(1, vk::DescriptorType::eAccelerationStructureKHR, 1,
+                                     vk::ShaderStageFlagBits::eRaygenKHR | vk::ShaderStageFlagBits::eClosestHitKHR),
+      // radiance accumulator
+      vk::DescriptorSetLayoutBinding(2, vk::DescriptorType::eStorageImage, 1,
+                                     vk::ShaderStageFlagBits::eRaygenKHR | vk::ShaderStageFlagBits::eClosestHitKHR),
+      // image
+      vk::DescriptorSetLayoutBinding(3, vk::DescriptorType::eStorageImage, 1,
+                                     vk::ShaderStageFlagBits::eRaygenKHR | vk::ShaderStageFlagBits::eClosestHitKHR),
+      // model data
+      vk::DescriptorSetLayoutBinding(4, vk::DescriptorType::eStorageBuffer, 1,
+                                     vk::ShaderStageFlagBits::eRaygenKHR | vk::ShaderStageFlagBits::eClosestHitKHR),
+      // material data
+      vk::DescriptorSetLayoutBinding(5, vk::DescriptorType::eStorageBuffer, 1,
+                                     vk::ShaderStageFlagBits::eRaygenKHR | vk::ShaderStageFlagBits::eClosestHitKHR),
+      // positions
+      vk::DescriptorSetLayoutBinding(6, vk::DescriptorType::eStorageBuffer, 1,
+                                     vk::ShaderStageFlagBits::eRaygenKHR | vk::ShaderStageFlagBits::eClosestHitKHR),
+      // texcoords
+      vk::DescriptorSetLayoutBinding(7, vk::DescriptorType::eStorageBuffer, 1,
+                                     vk::ShaderStageFlagBits::eRaygenKHR | vk::ShaderStageFlagBits::eClosestHitKHR),
+      // normals
+      vk::DescriptorSetLayoutBinding(8, vk::DescriptorType::eStorageBuffer, 1,
+                                     vk::ShaderStageFlagBits::eRaygenKHR | vk::ShaderStageFlagBits::eClosestHitKHR),
+      // directional lights
+      vk::DescriptorSetLayoutBinding(9, vk::DescriptorType::eStorageBuffer, 1,
+                                     vk::ShaderStageFlagBits::eRaygenKHR | vk::ShaderStageFlagBits::eClosestHitKHR),
+      // point lights
+      vk::DescriptorSetLayoutBinding(10, vk::DescriptorType::eStorageBuffer, 1,
+                                     vk::ShaderStageFlagBits::eRaygenKHR | vk::ShaderStageFlagBits::eClosestHitKHR),
+      // index buffer
+      vk::DescriptorSetLayoutBinding(11, vk::DescriptorType::eStorageBuffer, 1,
+                                     vk::ShaderStageFlagBits::eRaygenKHR | vk::ShaderStageFlagBits::eClosestHitKHR),
+  };
+
+  vk::DescriptorSetLayoutCreateInfo descriptorSetLayoutCreateInfo({}, bindings);
   descriptorLayout = DescriptorSetLayout(device, descriptorSetLayoutCreateInfo);
 
+  auto descriptorPoolSizes = {
+      vk::DescriptorPoolSize(vk::DescriptorType::eUniformBuffer, 1),
+      vk::DescriptorPoolSize(vk::DescriptorType::eAccelerationStructureKHR, 1),
+      vk::DescriptorPoolSize(vk::DescriptorType::eStorageImage, 2),
+      vk::DescriptorPoolSize(vk::DescriptorType::eStorageBuffer, 8),
+  };
+  descriptorPool =
+      DescriptorPool(device, vk::DescriptorPoolCreateInfo({vk::DescriptorPoolCreateFlagBits::eFreeDescriptorSet}, 4, descriptorPoolSizes));
+
   // create a PipelineLayout using that DescriptorSetLayout
-  vk::PipelineLayoutCreateInfo pipelineLayoutCreateInfo({}, *descriptorLayout);
+  vk::PushConstantRange range =
+      vk::PushConstantRange(vk::ShaderStageFlagBits::eRaygenKHR | vk::ShaderStageFlagBits::eClosestHitKHR, 0, sizeof(SampleParams));
+  vk::PipelineLayoutCreateInfo pipelineLayoutCreateInfo({}, *descriptorLayout, range);
   pipelineLayout = PipelineLayout(device, pipelineLayoutCreateInfo);
 }
 
 using namespace slang;
 
-void GPURenderer::createShaders()
+template <typename T> constexpr T align(T size, T alignment) { return (size + alignment - 1) & ~(alignment - 1); }
+
+void GPURenderer::createPipeline()
 {
-  /*
   Slang::ComPtr<IGlobalSession> globalSession;
-  SlangGlobalSessionDesc desc = {};
-  createGlobalSession(&desc, globalSession.writeRef());
-  SessionDesc sessionDesc;
-  TargetDesc targetDesc;
-  targetDesc.format = SLANG_SPIRV;
-  targetDesc.profile = globalSession->findProfile("glsl_450");
-  sessionDesc.targets = &targetDesc;
-  sessionDesc.targetCount = 1;
-  const char* searchPaths[] = {"res/shaders/"};
-  sessionDesc.searchPaths = searchPaths;
-  sessionDesc.searchPathCount = 1;
-  /* ... fill in `sessionDesc` ...
+  createGlobalSession(globalSession.writeRef());
+  TargetDesc targetDesc = {
+      .format = SLANG_SPIRV,
+      .profile = globalSession->findProfile("glsl_450"),
+  };
+  const char* searchPaths[] = {"../res/shaders/"};
+  SessionDesc sessionDesc = {
+      .targets = &targetDesc,
+      .targetCount = 1,
+      .searchPaths = searchPaths,
+      .searchPathCount = 1,
+  };
   Slang::ComPtr<ISession> session;
   globalSession->createSession(sessionDesc, session.writeRef());
 
   Slang::ComPtr<IBlob> diagnostics;
-  IModule* module = session->loadModule("MyShaders", diagnostics.writeRef());
+  IModule* raygenModule = session->loadModule("RayGen", diagnostics.writeRef());
   if (diagnostics)
   {
     std::cout << (const char*)diagnostics->getBufferPointer() << std::endl;
   }
-  Slang::ComPtr<IEntryPoint> computeEntryPoint;
-  module->findEntryPointByName("myComputeMain", computeEntryPoint.writeRef());
-  IComponentType* components[] = {module, computeEntryPoint};
+  Slang::ComPtr<IEntryPoint> rayGenEntry;
+  raygenModule->findEntryPointByName("raygen", rayGenEntry.writeRef());
+
+  IModule* closestHitModule = session->loadModule("ClosestHit", diagnostics.writeRef());
+  if (diagnostics)
+  {
+    std::cout << (const char*)diagnostics->getBufferPointer() << std::endl;
+  }
+  Slang::ComPtr<IEntryPoint> closestHitEntry;
+  closestHitModule->findEntryPointByName("closestHit", closestHitEntry.writeRef());
+
+  IModule* missModule = session->loadModule("Miss", diagnostics.writeRef());
+  if (diagnostics)
+  {
+    std::cout << (const char*)diagnostics->getBufferPointer() << std::endl;
+  }
+  Slang::ComPtr<IEntryPoint> missEntry;
+  missModule->findEntryPointByName("miss", missEntry.writeRef());
+
+  IComponentType* components[] = {raygenModule, rayGenEntry, closestHitModule, closestHitEntry, missModule, missEntry};
   Slang::ComPtr<IComponentType> program;
-  session->createCompositeComponentType(components, 2, program.writeRef());
+  session->createCompositeComponentType(components, 6, program.writeRef());
+
   Slang::ComPtr<IComponentType> linkedProgram;
-  Slang::ComPtr<ISlangBlob> diagnosticBlob;
-  program->link(linkedProgram.writeRef(), diagnosticBlob.writeRef());
-  int entryPointIndex = 0; // only one entry point
-  int targetIndex = 0;     // only one target
-  Slang::ComPtr<IBlob> kernelBlob;
-  linkedProgram->getEntryPointCode(entryPointIndex, targetIndex, kernelBlob.writeRef(), diagnostics.writeRef());
-  */
+  program->link(linkedProgram.writeRef(), diagnostics.writeRef());
+
+  Slang::ComPtr<IBlob> rayGenCode;
+  linkedProgram->getEntryPointCode(0, 0, rayGenCode.writeRef(), diagnostics.writeRef());
+
+  Slang::ComPtr<IBlob> closestHitCode;
+  linkedProgram->getEntryPointCode(1, 0, closestHitCode.writeRef(), diagnostics.writeRef());
+
+  Slang::ComPtr<IBlob> missCode;
+  linkedProgram->getEntryPointCode(2, 0, missCode.writeRef(), diagnostics.writeRef());
+
+  rayGen =
+      ShaderModule(device, vk::ShaderModuleCreateInfo({}, rayGenCode->getBufferSize(), (const uint32_t*)rayGenCode->getBufferPointer()));
+
+  closestHit = ShaderModule(
+      device, vk::ShaderModuleCreateInfo({}, closestHitCode->getBufferSize(), (const uint32_t*)closestHitCode->getBufferPointer()));
+
+  miss = ShaderModule(device, vk::ShaderModuleCreateInfo({}, missCode->getBufferSize(), (const uint32_t*)missCode->getBufferPointer()));
+
+  std::vector<vk::PipelineShaderStageCreateInfo> shaderStages;
+  std::vector<vk::RayTracingShaderGroupCreateInfoKHR> shaderGroups;
+
+  {
+    shaderStages.push_back(vk::PipelineShaderStageCreateInfo({}, vk::ShaderStageFlagBits::eRaygenKHR, rayGen, "main"));
+    shaderGroups.push_back(vk::RayTracingShaderGroupCreateInfoKHR(vk::RayTracingShaderGroupTypeKHR::eGeneral, shaderStages.size() - 1,
+                                                                  vk::ShaderUnusedKHR, vk::ShaderUnusedKHR, vk::ShaderUnusedKHR));
+  }
+  {
+    shaderStages.push_back(vk::PipelineShaderStageCreateInfo({}, vk::ShaderStageFlagBits::eClosestHitKHR, closestHit, "main"));
+
+    uint32_t hitIndex = static_cast<uint32_t>(shaderStages.size() - 1);
+    uint32_t anyHitIndex = VK_SHADER_UNUSED_KHR;
+    uint32_t intersectionIndex = VK_SHADER_UNUSED_KHR;
+    // if (hitgroup.anyHitShader != nullptr)
+    //{
+    //   auto anyHit = hitgroup.anyHitShader.cast<AnyHitShader>();
+    //   shaderStages.add(VkPipelineShaderStageCreateInfo{
+    //       .sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO,
+    //       .pNext = nullptr,
+    //       .flags = 0,
+    //       .stage = VK_SHADER_STAGE_ANY_HIT_BIT_KHR,
+    //       .module = anyHit->getModuleHandle(),
+    //       .pName = anyHit->getEntryPointName(),
+    //       .pSpecializationInfo = nullptr,
+    //   });
+    // }
+    // if (hitgroup.intersectionShader != nullptr)
+    //{
+    //   auto intersect = hitgroup.intersectionShader.cast<IntersectionShader>();
+    //   shaderStages.add(VkPipelineShaderStageCreateInfo{
+    //       .sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO,
+    //       .pNext = nullptr,
+    //       .flags = 0,
+    //       .stage = VK_SHADER_STAGE_INTERSECTION_BIT_KHR,
+    //       .module = intersect->getModuleHandle(),
+    //       .pName = intersect->getEntryPointName(),
+    //       .pSpecializationInfo = nullptr,
+    //   });
+    // }
+    shaderGroups.push_back(vk::RayTracingShaderGroupCreateInfoKHR(vk::RayTracingShaderGroupTypeKHR::eTrianglesHitGroup, vk::ShaderUnusedKHR,
+                                                                  hitIndex, anyHitIndex, intersectionIndex));
+  }
+  {
+    shaderStages.push_back(vk::PipelineShaderStageCreateInfo({}, vk::ShaderStageFlagBits::eMissKHR, miss, "main"));
+    shaderGroups.push_back(vk::RayTracingShaderGroupCreateInfoKHR(vk::RayTracingShaderGroupTypeKHR::eGeneral, shaderStages.size() - 1,
+                                                                  vk::ShaderUnusedKHR, vk::ShaderUnusedKHR, vk::ShaderUnusedKHR));
+  }
+  pipeline = device.createRayTracingPipelineKHR(
+      nullptr, nullptr, vk::RayTracingPipelineCreateInfoKHR({}, shaderStages, shaderGroups, 12, nullptr, nullptr, nullptr, pipelineLayout));
+
+  const uint32_t handleSize = rayTracingProperties.shaderGroupHandleSize;
+  const uint32_t handleSizeAligned = align(rayTracingProperties.shaderGroupHandleSize, rayTracingProperties.shaderGroupHandleAlignment);
+  const uint32_t handleAlignment = rayTracingProperties.shaderGroupHandleAlignment;
+  const uint32_t sbtAlignment = rayTracingProperties.shaderGroupBaseAlignment;
+  const uint32_t groupCount = static_cast<uint32_t>(shaderGroups.size());
+  const uint32_t sbtSize = groupCount * handleSizeAligned;
+  const VkBufferUsageFlags sbtBufferUsage =
+      VK_BUFFER_USAGE_SHADER_BINDING_TABLE_BIT_KHR | VK_BUFFER_USAGE_TRANSFER_DST_BIT | VK_BUFFER_USAGE_SHADER_DEVICE_ADDRESS_BIT;
+  const VmaMemoryUsage sbtMemoryUsage = VMA_MEMORY_USAGE_AUTO;
+
+  uint64_t rayGenStride = handleSize;
+  uint64_t hitStride = handleSize;
+  uint64_t missStride = handleSize;
+  auto rayGenSBTInfo = VkBufferCreateInfo{
+      .sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO,
+      .pNext = nullptr,
+      .flags = 0,
+      .size = rayGenStride,
+      .usage = sbtBufferUsage,
+  };
+  auto rayGenSBTAllocInfo = VmaAllocationCreateInfo{
+      .usage = sbtMemoryUsage,
+  };
+  VkBuffer rayGenSBTBuf;
+  vmaCreateBufferWithAlignment(allocator, &rayGenSBTInfo, &rayGenSBTAllocInfo, sbtAlignment, &rayGenSBTBuf, &rayGenAlloc, nullptr);
+  rayGenSBT = Buffer(device, rayGenSBTBuf);
+  rayGenAddr =
+      vk::StridedDeviceAddressRegionKHR(device.getBufferAddress(vk::BufferDeviceAddressInfo(*rayGenSBT)), rayGenStride, rayGenStride);
+
+  auto closestHitSBTInfo = VkBufferCreateInfo{
+      .sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO,
+      .pNext = nullptr,
+      .flags = 0,
+      .size = hitStride,
+      .usage = sbtBufferUsage,
+  };
+  auto closestHitSBTAllocInfo = VmaAllocationCreateInfo{
+      .usage = sbtMemoryUsage,
+  };
+  VkBuffer closestHitSBTBuf;
+  vmaCreateBufferWithAlignment(allocator, &closestHitSBTInfo, &closestHitSBTAllocInfo, sbtAlignment, &closestHitSBTBuf, &closestHitAlloc,
+                               nullptr);
+  closestHitSBT = Buffer(device, closestHitSBTBuf);
+  closestHitAddr =
+      vk::StridedDeviceAddressRegionKHR(device.getBufferAddress(vk::BufferDeviceAddressInfo(*closestHitSBT)), hitStride, hitStride);
+
+  auto missSBTInfo = VkBufferCreateInfo{
+      .sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO,
+      .pNext = nullptr,
+      .flags = 0,
+      .size = missStride,
+      .usage = sbtBufferUsage,
+  };
+  auto missSBTAllocInfo = VmaAllocationCreateInfo{
+      .usage = sbtMemoryUsage,
+  };
+  VkBuffer missSBTBuf;
+  vmaCreateBufferWithAlignment(allocator, &missSBTInfo, &missSBTAllocInfo, sbtAlignment, &missSBTBuf, &missAlloc, nullptr);
+  missSBT = Buffer(device, missSBTBuf);
+  missAddr = vk::StridedDeviceAddressRegionKHR(device.getBufferAddress(vk::BufferDeviceAddressInfo(*missSBT)), missStride, missStride);
+
+  std::vector<unsigned char> sbt = pipeline.getRayTracingShaderGroupHandlesKHR<unsigned char>(0, shaderGroups.size(), sbtSize);
+
+  uploadToGPU(rayGenSBT, sbt.data(), rayGenStride);
+  uploadToGPU(closestHitSBT, sbt.data() + handleSize, handleSize);
+  uploadToGPU(missSBT, sbt.data() + handleSize * 2, handleSize);
 }
 
-void GPURenderer::render(Camera cam, RenderParameter param) {}
+void GPURenderer::uploadToGPU(Buffer& buffer, void* data, size_t size)
+{
+  VkBufferCreateInfo stagingBufInfo = {
+      .sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO,
+      .size = size,
+      .usage = VK_BUFFER_USAGE_TRANSFER_SRC_BIT,
+  };
+  VmaAllocationCreateInfo stagingAllocInfo = {
+      .flags = VMA_ALLOCATION_CREATE_HOST_ACCESS_SEQUENTIAL_WRITE_BIT | VMA_ALLOCATION_CREATE_MAPPED_BIT,
+      .usage = VMA_MEMORY_USAGE_AUTO,
+  };
+  VkBuffer stagingBuf;
+  VmaAllocation stagingAllocation;
+  vmaCreateBuffer(allocator, &stagingBufInfo, &stagingAllocInfo, &stagingBuf, &stagingAllocation, nullptr);
+  Buffer stagingBuffer = Buffer(device, stagingBuf);
+
+  vmaCopyMemoryToAllocation(allocator, data, stagingAllocation, 0, size);
+  CommandBuffer copyCmd =
+      std::move(device.allocateCommandBuffers(vk::CommandBufferAllocateInfo(cmdPool, vk::CommandBufferLevel::ePrimary, 1)).front());
+  copyCmd.begin(vk::CommandBufferBeginInfo(vk::CommandBufferUsageFlagBits::eOneTimeSubmit));
+  copyCmd.copyBuffer(stagingBuffer, buffer, vk::BufferCopy(0, 0, size));
+  copyCmd.end();
+  queue.submit(vk::SubmitInfo({}, {}, *copyCmd, {}));
+  device.waitIdle();
+}
+
+void GPURenderer::render(Camera cam, RenderParameter param)
+{
+  for (uint32_t samp = 0; samp < param.numSamples; ++samp)
+  {
+    semaphores.push_back(device.createSemaphore(vk::SemaphoreCreateInfo()));
+    fences.push_back(device.createFence(vk::FenceCreateInfo()));
+  }
+  // camera
+  {
+    VkBufferCreateInfo bufferInfo = {
+        .sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO,
+        .size = sizeof(GPUCamera),
+        .usage = VK_BUFFER_USAGE_UNIFORM_BUFFER_BIT | VK_BUFFER_USAGE_TRANSFER_DST_BIT,
+    };
+    VmaAllocationCreateInfo allocInfo = {
+        .flags = VMA_ALLOCATION_CREATE_DEDICATED_MEMORY_BIT,
+        .usage = VMA_MEMORY_USAGE_AUTO,
+    };
+
+    VkBuffer camBuf;
+    vmaCreateBuffer(allocator, &bufferInfo, &allocInfo, &camBuf, &cameraAllocation, nullptr);
+    cameraBuffer = Buffer(device, camBuf);
+    GPUCamera gpuCam = {
+        .cameraPosition = cam.position,
+        .f = cam.f,
+        .cameraForward = glm::normalize(cam.target - cam.position),
+        .S_O = cam.S_O,
+        .fogEmm = glm::vec3(0, 0, 0),
+        .ks = 0,
+        .A = cam.A,
+        .ka = 0,
+    };
+    uploadToGPU(cameraBuffer, &gpuCam, sizeof(GPUCamera));
+  }
+  // radiance accumulator
+  {
+    VkImageCreateInfo imageInfo = {
+        .sType = VK_STRUCTURE_TYPE_IMAGE_CREATE_INFO,
+        .imageType = VK_IMAGE_TYPE_2D,
+        .format = VK_FORMAT_R32G32B32A32_SFLOAT,
+        .extent =
+            {
+                .width = (uint32_t)param.width,
+                .height = (uint32_t)param.height,
+                .depth = 1,
+            },
+        .mipLevels = 1,
+        .arrayLayers = 1,
+        .samples = VK_SAMPLE_COUNT_1_BIT,
+        .tiling = VK_IMAGE_TILING_OPTIMAL,
+        .usage = VK_IMAGE_USAGE_STORAGE_BIT,
+        .initialLayout = VK_IMAGE_LAYOUT_GENERAL,
+    };
+
+    VmaAllocationCreateInfo allocCreateInfo = {
+        .flags = VMA_ALLOCATION_CREATE_DEDICATED_MEMORY_BIT,
+        .usage = VMA_MEMORY_USAGE_AUTO,
+    };
+    VkImage radianceImg;
+    vmaCreateImage(allocator, &imageInfo, &allocCreateInfo, &radianceImg, &radianceAllocation, nullptr);
+    radianceAccumulator = Image(device, radianceImg);
+    radianceView =
+        device.createImageView(vk::ImageViewCreateInfo({}, *radianceAccumulator, vk::ImageViewType::e2D, vk::Format::eR32G32B32A32Sfloat));
+  }
+  // image
+  {
+    VkImageCreateInfo imageInfo = {
+        .sType = VK_STRUCTURE_TYPE_IMAGE_CREATE_INFO,
+        .imageType = VK_IMAGE_TYPE_2D,
+        .format = VK_FORMAT_R32G32B32A32_SFLOAT,
+        .extent =
+            {
+                .width = (uint32_t)param.width,
+                .height = (uint32_t)param.height,
+                .depth = 1,
+            },
+        .mipLevels = 1,
+        .arrayLayers = 1,
+        .samples = VK_SAMPLE_COUNT_1_BIT,
+        .tiling = VK_IMAGE_TILING_OPTIMAL,
+        .usage = VK_IMAGE_USAGE_STORAGE_BIT | VK_IMAGE_USAGE_TRANSFER_SRC_BIT,
+        .initialLayout = VK_IMAGE_LAYOUT_GENERAL,
+    };
+
+    VmaAllocationCreateInfo allocCreateInfo = {
+        .flags = VMA_ALLOCATION_CREATE_DEDICATED_MEMORY_BIT,
+        .usage = VMA_MEMORY_USAGE_AUTO,
+    };
+    VkImage img;
+    vmaCreateImage(allocator, &imageInfo, &allocCreateInfo, &img, &imageAllocation, nullptr);
+    image = Image(device, img);
+    radianceView = device.createImageView(vk::ImageViewCreateInfo({}, *image, vk::ImageViewType::e2D, vk::Format::eR32G32B32A32Sfloat));
+  }
+
+  DescriptorSet descriptorSet =
+      std::move(device.allocateDescriptorSets(vk::DescriptorSetAllocateInfo(*descriptorPool, *descriptorLayout)).front());
+  std::vector<vk::WriteDescriptorSet> writes;
+  // have to use lists so the pointers arent invalidated by push
+  std::list<vk::DescriptorBufferInfo> buffers;
+  std::list<vk::WriteDescriptorSetAccelerationStructureKHR> accel;
+  std::list<vk::DescriptorImageInfo> images;
+  uint32_t bindingCounter = 0;
+  {
+    buffers.push_back(vk::DescriptorBufferInfo(cameraBuffer));
+    writes.push_back(vk::WriteDescriptorSet(*descriptorSet, bindingCounter++, 0, 1, vk::DescriptorType::eUniformBuffer, nullptr,
+                                            &buffers.back(), nullptr));
+  }
+  {
+    accel.push_back(vk::WriteDescriptorSetAccelerationStructureKHR(*((GPUScene*)scene.get())->accelerationStructure));
+    writes.push_back(vk::WriteDescriptorSet(*descriptorSet, bindingCounter++, 0, 1, vk::DescriptorType::eAccelerationStructureKHR, nullptr,
+                                            nullptr, nullptr, &accel.back()));
+  }
+  {
+    images.push_back(vk::DescriptorImageInfo({}, radianceView, vk::ImageLayout::eGeneral));
+    writes.push_back(vk::WriteDescriptorSet(*descriptorSet, bindingCounter++, 0, 1, vk::DescriptorType::eStorageImage, &images.back(),
+                                            nullptr, nullptr));
+  }
+  {
+    images.push_back(vk::DescriptorImageInfo({}, imageView, vk::ImageLayout::eGeneral));
+    writes.push_back(vk::WriteDescriptorSet(*descriptorSet, bindingCounter++, 0, 1, vk::DescriptorType::eStorageImage, &images.back(),
+                                            nullptr, nullptr));
+  }
+  {
+    buffers.push_back(vk::DescriptorBufferInfo(((GPUScene*)scene.get())->modelBuffer));
+    writes.push_back(vk::WriteDescriptorSet(*descriptorSet, bindingCounter++, 0, 1, vk::DescriptorType::eStorageBuffer, nullptr,
+                                            &buffers.back(), nullptr));
+  }
+  {
+    buffers.push_back(vk::DescriptorBufferInfo(((GPUScene*)scene.get())->materialBuffer));
+    writes.push_back(vk::WriteDescriptorSet(*descriptorSet, bindingCounter++, 0, 1, vk::DescriptorType::eStorageBuffer, nullptr,
+                                            &buffers.back(), nullptr));
+  }
+  {
+    buffers.push_back(vk::DescriptorBufferInfo(((GPUScene*)scene.get())->positionBuffer));
+    writes.push_back(vk::WriteDescriptorSet(*descriptorSet, bindingCounter++, 0, 1, vk::DescriptorType::eStorageBuffer, nullptr,
+                                            &buffers.back(), nullptr));
+  }
+  {
+    buffers.push_back(vk::DescriptorBufferInfo(((GPUScene*)scene.get())->texCoordsBuffer));
+    writes.push_back(vk::WriteDescriptorSet(*descriptorSet, bindingCounter++, 0, 1, vk::DescriptorType::eStorageBuffer, nullptr,
+                                            &buffers.back(), nullptr));
+  }
+  {
+    buffers.push_back(vk::DescriptorBufferInfo(((GPUScene*)scene.get())->normalsBuffer));
+    writes.push_back(vk::WriteDescriptorSet(*descriptorSet, bindingCounter++, 0, 1, vk::DescriptorType::eStorageBuffer, nullptr,
+                                            &buffers.back(), nullptr));
+  }
+  {
+    buffers.push_back(vk::DescriptorBufferInfo(((GPUScene*)scene.get())->directionalLightBuffer));
+    writes.push_back(vk::WriteDescriptorSet(*descriptorSet, bindingCounter++, 0, 1, vk::DescriptorType::eStorageBuffer, nullptr,
+                                            &buffers.back(), nullptr));
+  }
+  {
+    buffers.push_back(vk::DescriptorBufferInfo(((GPUScene*)scene.get())->pointLightBuffer));
+    writes.push_back(vk::WriteDescriptorSet(*descriptorSet, bindingCounter++, 0, 1, vk::DescriptorType::eStorageBuffer, nullptr,
+                                            &buffers.back(), nullptr));
+  }
+  {
+    buffers.push_back(vk::DescriptorBufferInfo(((GPUScene*)scene.get())->indexBuffer));
+    writes.push_back(vk::WriteDescriptorSet(*descriptorSet, bindingCounter++, 0, 1, vk::DescriptorType::eStorageBuffer, nullptr,
+                                            &buffers.back(), nullptr));
+  }
+  device.updateDescriptorSets(writes, {});
+
+  // allocate a CommandBuffer from the CommandPool
+  vk::CommandBufferAllocateInfo commandBufferAllocateInfo(*cmdPool, vk::CommandBufferLevel::ePrimary, param.numSamples);
+  cmdBuffers = CommandBuffers(device, commandBufferAllocateInfo);
+  for (uint32_t samp = 0; samp < param.numSamples; ++samp)
+  {
+    auto& cmd = cmdBuffers[samp];
+    cmd.begin(vk::CommandBufferBeginInfo(vk::CommandBufferUsageFlagBits::eOneTimeSubmit));
+    cmd.bindPipeline(vk::PipelineBindPoint::eRayTracingKHR, *pipeline);
+    cmd.bindDescriptorSets(vk::PipelineBindPoint::eRayTracingKHR, pipelineLayout, 0, *descriptorSet, {});
+    std::vector<SampleParams> sampleParams = {SampleParams{
+        .pass = samp,
+        .samplesPerPixel = param.numSamples,
+        .numDirectionalLights = (uint32_t)scene->directionalLights.size(),
+        .numPointLights = (uint32_t)scene->pointLights.size(),
+    }};
+    cmd.pushConstants<SampleParams>(pipelineLayout, vk::ShaderStageFlagBits::eRaygenKHR | vk::ShaderStageFlagBits::eClosestHitKHR, 0,
+                                    sampleParams);
+    cmd.traceRaysKHR(rayGenAddr, closestHitAddr, missAddr, {}, param.width, param.height, 1);
+    cmd.end();
+    if (samp == 0)
+    {
+      queue.submit(vk::SubmitInfo({}, {}, *cmd, *semaphores[samp]), *fences[samp]);
+    }
+    else
+    {
+      vk::PipelineStageFlags dstWaitMask = vk::PipelineStageFlagBits::eRayTracingShaderKHR;
+      queue.submit(vk::SubmitInfo(*semaphores[samp - 1], dstWaitMask, *cmd, *semaphores[samp]), *fences[samp]);
+    }
+  }
+  for (uint32_t samp = 0; samp < param.numSamples; ++samp)
+  {
+    assert(device.waitForFences(*fences[samp], true, 1000000) == vk::Result::eSuccess);
+  }
+}

src/gpu/GPURenderer.h

@@ -1,7 +1,9 @@
 #pragma once
+#include "gpu/GPUScene.h"
 #include "scene/Renderer.h"
 #include <vulkan/vulkan.hpp>
 #include <vulkan/vulkan_raii.hpp>
+#include <vma/vk_mem_alloc.h>
 
 using namespace vk::raii;
 
@@ -10,33 +12,82 @@ struct GPURenderer : public Renderer
 public:
   GPURenderer();
   virtual ~GPURenderer();
+  virtual void render(Camera cam, RenderParameter param) override;
 
 private:
+  struct GPUCamera
+  {
+    glm::vec3 cameraPosition;
+    float f;
+    glm::vec3 cameraForward;
+    float S_O;
+    glm::vec3 fogEmm;
+    float ks;
+    float A;
+    float ka;
+  };
+
+  struct SampleParams
+  {
+    uint32_t pass;
+    uint32_t samplesPerPixel;
+    uint32_t numDirectionalLights;
+    uint32_t numPointLights;
+  };
   void createDevice();
   void createCommands();
   void createDescriptors();
-  void createShaders();
+  void createPipeline();
 
   Context context;
-  Instance instance;
-  PhysicalDevice physicalDevice;
-  Device device;
-  Queue queue;
+  Instance instance = nullptr;
+  PhysicalDevice physicalDevice = nullptr;
+  Device device = nullptr;
+  Queue queue = nullptr;
+  VmaAllocator allocator = nullptr;
 
-  uint32_t computeQueueFamily;
-  CommandPool cmdPool;
-  CommandBuffers cmdBuffers;
+  vk::PhysicalDeviceAccelerationStructurePropertiesKHR accelerationProperties = {};
+  vk::PhysicalDeviceRayTracingPipelinePropertiesKHR rayTracingProperties = {};
 
-  DescriptorSetLayout descriptorLayout;
-  DescriptorSet descriptorSet;
-  DescriptorPool descriptorPool;
-  PipelineLayout pipelineLayout;
+  uint32_t computeQueueFamily = 0;
+  CommandPool cmdPool = nullptr;
+  CommandBuffers cmdBuffers = nullptr;
+  std::vector<Semaphore> semaphores;
+  std::vector<Fence> fences;
 
-  ShaderModule rayGen;
-  ShaderModule closestHit;
-  ShaderModule miss;
+  DescriptorSetLayout descriptorLayout = nullptr;
+  DescriptorSet descriptorSet = nullptr;
+  DescriptorPool descriptorPool = nullptr;
+  PipelineLayout pipelineLayout = nullptr;
 
-  Pipeline pipeline;
+  ShaderModule rayGen = nullptr;
+  ShaderModule closestHit = nullptr;
+  ShaderModule miss = nullptr;
 
-  virtual void render(Camera cam, RenderParameter param);
+  Pipeline pipeline = nullptr;
+
+  Buffer rayGenSBT = nullptr;
+  vk::StridedDeviceAddressRegionKHR rayGenAddr;
+  VmaAllocation rayGenAlloc;
+
+  Buffer closestHitSBT = nullptr;
+  vk::StridedDeviceAddressRegionKHR closestHitAddr;
+  VmaAllocation closestHitAlloc;
+
+  Buffer missSBT = nullptr;
+  vk::StridedDeviceAddressRegionKHR missAddr;
+  VmaAllocation missAlloc;
+
+  Buffer cameraBuffer = nullptr;
+  VmaAllocation cameraAllocation;
+
+  Image radianceAccumulator = nullptr;
+  ImageView radianceView = nullptr;
+  VmaAllocation radianceAllocation;
+
+  Image image = nullptr;
+  ImageView imageView = nullptr;
+  VmaAllocation imageAllocation;
+
+  void uploadToGPU(Buffer& buffer, void* data, size_t size);
 };

src/gpu/GPUScene.cpp

@@ -0,0 +1,212 @@
+#include "GPUScene.h"
+
+GPUScene::GPUScene(Device& device, VmaAllocator& allocator, CommandPool& cmdPool, Queue& queue)
+    : device(device), allocator(allocator), cmdPool(cmdPool), queue(queue)
+{
+}
+
+GPUScene::~GPUScene() {}
+
+void GPUScene::createRayTracingHierarchy()
+{
+  // upload geometry to gpu
+  createStorageBuffer(modelBuffer, modelAllocation, refs.data(), refs.size() * sizeof(ModelReference));
+  // createStorageBuffer(materialBuffer, materialAllocation, refs.data(), refs.size() * sizeof(ModelReference));
+  createStorageBuffer(positionBuffer, positionAllocation, positionPool.data(), positionPool.size() * sizeof(glm::vec3));
+  createStorageBuffer(texCoordsBuffer, texCoordsAllocation, texCoordsPool.data(), texCoordsPool.size() * sizeof(glm::vec2));
+  createStorageBuffer(normalsBuffer, normalsAllocation, normalsPool.data(), normalsPool.size() * sizeof(glm::vec3));
+  createStorageBuffer(directionalLightBuffer, directionalLightAllocation, directionalLights.data(),
+                      directionalLights.size() * sizeof(DirectionalLight));
+  createStorageBuffer(pointLightBuffer, pointLightAllocation, pointLights.data(), pointLights.size() * sizeof(PointLight));
+  createStorageBuffer(indexBuffer, indexAllocation, indicesPool.data(), indicesPool.size() * sizeof(glm::uvec3));
+
+  vk::DeviceAddress vertexBufferAddr = device.getBufferAddress(vk::BufferDeviceAddressInfo(positionBuffer));
+  vk::DeviceAddress indexBufferAddr = device.getBufferAddress(vk::BufferDeviceAddressInfo(indexBuffer));
+
+  std::vector<vk::AccelerationStructureInstanceKHR> instances(models.size());
+  {
+    std::vector<vk::AccelerationStructureGeometryKHR> geometries(models.size());
+    std::vector<vk::AccelerationStructureBuildGeometryInfoKHR> buildGeometries(models.size());
+    std::vector<vk::AccelerationStructureBuildSizesInfoKHR> buildSizes(models.size());
+    std::vector<VkBuffer> scratchBuffers(models.size());
+    std::vector<VmaAllocation> scratchAllocations(models.size());
+    std::vector<vk::AccelerationStructureBuildRangeInfoKHR> buildRanges(models.size());
+    std::vector<const vk::AccelerationStructureBuildRangeInfoKHR*> buildRangePointers(models.size());
+    blas.resize(models.size());
+    for (uint32_t i = 0; i < models.size(); ++i)
+    {
+      vk::DeviceOrHostAddressConstKHR vertexDataAddress = (vertexBufferAddr + refs[i].positionOffset * sizeof(glm::vec3));
+      vk::DeviceOrHostAddressConstKHR indexDataAddress = (indexBufferAddr + refs[i].indicesOffset + sizeof(glm::uvec3));
+
+      geometries[i] = vk::AccelerationStructureGeometryKHR(
+          vk::GeometryTypeKHR::eTriangles,
+          vk::AccelerationStructureGeometryTrianglesDataKHR(vk::Format::eR32G32B32Sfloat, vertexDataAddress, sizeof(glm::vec3),
+                                                            (uint32_t)refs[i].numPositions, vk::IndexType::eUint32, indexDataAddress),
+          vk::GeometryFlagBitsKHR::eOpaque);
+
+      buildGeometries[i] = vk::AccelerationStructureBuildGeometryInfoKHR(
+          vk::AccelerationStructureTypeKHR::eBottomLevel, vk::BuildAccelerationStructureFlagBitsKHR::ePreferFastTrace,
+          vk::BuildAccelerationStructureModeKHR::eBuild, {}, {}, 1, &geometries[i], nullptr);
+
+      buildSizes[i] = {
+          .sType = VK_STRUCTURE_TYPE_ACCELERATION_STRUCTURE_BUILD_SIZES_INFO_KHR,
+          .pNext = nullptr,
+      };
+      const uint32_t primitiveCount = refs[i].numIndices / 3;
+      buildSizes[i] =
+          device.getAccelerationStructureBuildSizesKHR(vk::AccelerationStructureBuildTypeKHR::eDevice, buildGeometries[i], primitiveCount);
+
+      VkBufferCreateInfo bufferInfo = {
+          .sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO,
+          .pNext = nullptr,
+          .flags = 0,
+          .size = buildSizes[i].accelerationStructureSize,
+          .usage = VK_BUFFER_USAGE_ACCELERATION_STRUCTURE_STORAGE_BIT_KHR | VK_BUFFER_USAGE_SHADER_DEVICE_ADDRESS_BIT,
+      };
+      VmaAllocationCreateInfo bufferAllocInfo = {
+          .usage = VMA_MEMORY_USAGE_AUTO,
+      };
+      VkBuffer buf;
+      vmaCreateBuffer(allocator, &bufferInfo, &bufferAllocInfo, &buf, &blas[i].alloc, nullptr);
+      blas[i].buffer = Buffer(device, buf);
+
+      vk::AccelerationStructureCreateInfoKHR blasInfo({}, blas[i].buffer, 0, buildSizes[i].accelerationStructureSize,
+                                                      vk::AccelerationStructureTypeKHR::eBottomLevel);
+      blas[i].handle = device.createAccelerationStructureKHR(blasInfo);
+
+      VkBufferCreateInfo scratchInfo = {
+          .sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO,
+          .pNext = nullptr,
+          .flags = 0,
+          .size = buildSizes[i].buildScratchSize,
+          .usage = VK_BUFFER_USAGE_STORAGE_BUFFER_BIT | VK_BUFFER_USAGE_SHADER_DEVICE_ADDRESS_BIT,
+      };
+      VmaAllocationCreateInfo scratchAllocInfo = {
+          .usage = VMA_MEMORY_USAGE_AUTO,
+      };
+      vmaCreateBufferWithAlignment(allocator, &scratchInfo, &scratchAllocInfo, 16, &scratchBuffers[i], &scratchAllocations[i], nullptr);
+      vk::DeviceAddress scratchAddr = device.getBufferAddress(vk::BufferDeviceAddressInfo(scratchBuffers[i]));
+      buildGeometries[i].dstAccelerationStructure = blas[i].handle;
+      buildGeometries[i].scratchData.deviceAddress = scratchAddr;
+
+      buildRanges[i] = VkAccelerationStructureBuildRangeInfoKHR{
+          .primitiveCount = primitiveCount,
+          .primitiveOffset = 0,
+          .firstVertex = 0,
+          .transformOffset = 0,
+      };
+      buildRangePointers[i] = &buildRanges[i];
+      vk::DeviceAddress blasAddr = device.getBufferAddress(vk::BufferDeviceAddressInfo(blas[i].buffer));
+      instances[i] = vk::AccelerationStructureInstanceKHR({}, i, 0xff, 0, {}, blasAddr);
+    }
+    vk::CommandBufferAllocateInfo commandBufferAllocateInfo(*cmdPool, vk::CommandBufferLevel::ePrimary, 10);
+    CommandBuffer cmdBuffer = std::move(CommandBuffers(device, commandBufferAllocateInfo).front());
+    cmdBuffer.begin(vk::CommandBufferBeginInfo(vk::CommandBufferUsageFlagBits::eOneTimeSubmit));
+    cmdBuffer.buildAccelerationStructuresKHR(buildGeometries, buildRangePointers);
+    cmdBuffer.end();
+    vk::SubmitInfo submitInfo;
+    queue.submit(submitInfo);
+    device.waitIdle();
+  }
+  createStorageBuffer(instanceBuffer, instanceAllocation, instances.data(), instances.size());
+  vk::DeviceAddress instancesAddress = device.getBufferAddress(vk::BufferDeviceAddressInfo(instanceBuffer));
+  vk::AccelerationStructureGeometryKHR geometry(vk::GeometryTypeKHR::eInstances,
+                                                vk::AccelerationStructureGeometryInstancesDataKHR(false, {instancesAddress}),
+                                                vk::GeometryFlagBitsKHR::eOpaque);
+  vk::AccelerationStructureBuildGeometryInfoKHR structureBuildGeometry(vk::AccelerationStructureTypeKHR::eTopLevel,
+                                                                       vk::BuildAccelerationStructureFlagBitsKHR::ePreferFastTrace,
+                                                                       vk::BuildAccelerationStructureModeKHR::eBuild, {}, {}, geometry);
+
+  const uint32_t primitiveCount = instances.size();
+  auto buildSizes =
+      device.getAccelerationStructureBuildSizesKHR(vk::AccelerationStructureBuildTypeKHR::eDevice, structureBuildGeometry, primitiveCount);
+
+  VkBuffer buffer;
+  auto tlasInfo = VkBufferCreateInfo{
+      .sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO,
+      .pNext = nullptr,
+      .flags = 0,
+      .size = buildSizes.accelerationStructureSize,
+      .usage = VK_BUFFER_USAGE_ACCELERATION_STRUCTURE_STORAGE_BIT_KHR | VK_BUFFER_USAGE_SHADER_DEVICE_ADDRESS_BIT,
+  };
+  auto tlasAlloc = VmaAllocationCreateInfo{
+      .usage = VMA_MEMORY_USAGE_AUTO,
+  };
+  vmaCreateBuffer(allocator, &tlasInfo, &tlasAlloc, &buffer, &accelerationAllocation, nullptr);
+  accelerationBuffer = Buffer(device, buffer);
+
+  accelerationStructure = device.createAccelerationStructureKHR(vk::AccelerationStructureCreateInfoKHR(
+      {}, accelerationBuffer, 0, buildSizes.accelerationStructureSize, vk::AccelerationStructureTypeKHR::eTopLevel));
+
+  auto scratchInfo = VkBufferCreateInfo{
+      .sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO,
+      .pNext = nullptr,
+      .flags = 0,
+      .size = buildSizes.buildScratchSize,
+      .usage = VK_BUFFER_USAGE_STORAGE_BUFFER_BIT | VK_BUFFER_USAGE_SHADER_DEVICE_ADDRESS_BIT,
+  };
+  auto scratchAllocInfo = VmaAllocationCreateInfo{
+      .usage = VMA_MEMORY_USAGE_AUTO,
+  };
+  VkBuffer scratchBuf;
+  VmaAllocation scratchAlloc;
+  vmaCreateBufferWithAlignment(allocator, &scratchInfo, &scratchAllocInfo, 64, &scratchBuf, &scratchAlloc, nullptr);
+  Buffer scratchBuffer = Buffer(device, scratchBuf);
+
+  vk::DeviceAddress scratchAddr = device.getBufferAddress(vk::BufferDeviceAddressInfo(scratchBuffer));
+  vk::AccelerationStructureBuildGeometryInfoKHR buildGeometry(
+      vk::AccelerationStructureTypeKHR::eTopLevel, vk::BuildAccelerationStructureFlagBitsKHR::ePreferFastTrace,
+      vk::BuildAccelerationStructureModeKHR::eBuild, {}, accelerationStructure, geometry, {}, {scratchAddr});
+
+  vk::AccelerationStructureBuildRangeInfoKHR buildRange(primitiveCount, 0, 0, 0);
+
+  vk::CommandBufferAllocateInfo commandBufferAllocateInfo(*cmdPool, vk::CommandBufferLevel::ePrimary, 10);
+  CommandBuffer cmdBuffer = std::move(CommandBuffers(device, commandBufferAllocateInfo).front());
+  cmdBuffer.begin(vk::CommandBufferBeginInfo(vk::CommandBufferUsageFlagBits::eOneTimeSubmit));
+  cmdBuffer.buildAccelerationStructuresKHR(buildGeometry, {&buildRange});
+  cmdBuffer.end();
+  vk::SubmitInfo submitInfo;
+  queue.submit(submitInfo);
+  device.waitIdle();
+}
+
+void GPUScene::createStorageBuffer(Buffer& buffer, VmaAllocation& alloc, void* data, size_t size)
+{
+  if (size == 0)
+    return;
+  VkBufferCreateInfo bufferCreateInfo = {VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO};
+  bufferCreateInfo.size = size;
+  bufferCreateInfo.usage = VK_BUFFER_USAGE_ACCELERATION_STRUCTURE_BUILD_INPUT_READ_ONLY_BIT_KHR | VK_BUFFER_USAGE_STORAGE_BUFFER_BIT |
+                           VK_BUFFER_USAGE_TRANSFER_DST_BIT | VK_BUFFER_USAGE_SHADER_DEVICE_ADDRESS_BIT;
+
+  VmaAllocationCreateInfo allocCreateInfo = {};
+  allocCreateInfo.usage = VMA_MEMORY_USAGE_AUTO;
+  allocCreateInfo.flags = VMA_ALLOCATION_CREATE_DEDICATED_MEMORY_BIT;
+
+  VkBuffer buf;
+  vmaCreateBuffer(allocator, &bufferCreateInfo, &allocCreateInfo, &buf, &alloc, nullptr);
+  buffer = Buffer(device, buf);
+
+  VkBufferCreateInfo stagingBufInfo = {
+      .sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO,
+      .size = size,
+      .usage = VK_BUFFER_USAGE_TRANSFER_SRC_BIT,
+  };
+  VmaAllocationCreateInfo stagingAllocInfo = {
+      .flags = VMA_ALLOCATION_CREATE_HOST_ACCESS_SEQUENTIAL_WRITE_BIT | VMA_ALLOCATION_CREATE_MAPPED_BIT,
+      .usage = VMA_MEMORY_USAGE_AUTO,
+  };
+  VkBuffer stagingBuf;
+  VmaAllocation stagingAllocation;
+  vmaCreateBuffer(allocator, &stagingBufInfo, &stagingAllocInfo, &stagingBuf, &stagingAllocation, nullptr);
+  Buffer stagingBuffer = Buffer(device, stagingBuf);
+
+  vmaCopyMemoryToAllocation(allocator, data, stagingAllocation, 0, size);
+  CommandBuffer copyCmd =
+      std::move(device.allocateCommandBuffers(vk::CommandBufferAllocateInfo(cmdPool, vk::CommandBufferLevel::ePrimary, 1)).front());
+  copyCmd.begin(vk::CommandBufferBeginInfo(vk::CommandBufferUsageFlagBits::eOneTimeSubmit));
+  copyCmd.copyBuffer(stagingBuffer, buffer, vk::BufferCopy(0, 0, size));
+  copyCmd.end();
+  queue.submit(vk::SubmitInfo({}, {}, *copyCmd, {}));
+  device.waitIdle();
+}

src/gpu/GPUScene.h

@@ -0,0 +1,64 @@
+#pragma once
+#include <vulkan/vulkan.hpp>
+#include <vulkan/vulkan_raii.hpp>
+#include <vma/vk_mem_alloc.h>
+#include "scene/Scene.h"
+
+using namespace vk::raii;
+
+class GPUScene : public Scene
+{
+public:
+  GPUScene(Device& device, VmaAllocator& allocator, CommandPool& cmdPool, Queue& queue);
+  virtual ~GPUScene();
+  virtual void createRayTracingHierarchy() override;
+
+private:
+  void createStorageBuffer(Buffer& buffer, VmaAllocation& alloc, void* data, size_t size);
+
+  Device& device;
+  VmaAllocator& allocator;
+  CommandPool& cmdPool;
+  Queue& queue;
+
+  // bottom level acceleration structure
+  struct BLAS
+  {
+    vk::AccelerationStructureKHR handle = nullptr;
+    vk::Buffer buffer = nullptr;
+    VmaAllocation alloc = nullptr;
+  };
+
+  AccelerationStructureKHR accelerationStructure = nullptr;
+  Buffer accelerationBuffer = nullptr;
+  VmaAllocation accelerationAllocation = nullptr;
+  Buffer instanceBuffer = nullptr;
+  VmaAllocation instanceAllocation = nullptr;
+
+  std::vector<BLAS> blas;
+
+  Buffer modelBuffer = nullptr;
+  VmaAllocation modelAllocation;
+
+  Buffer materialBuffer = nullptr;
+  VmaAllocation materialAllocation;
+
+  Buffer positionBuffer = nullptr;
+  VmaAllocation positionAllocation;
+
+  Buffer texCoordsBuffer = nullptr;
+  VmaAllocation texCoordsAllocation;
+
+  Buffer normalsBuffer = nullptr;
+  VmaAllocation normalsAllocation;
+
+  Buffer directionalLightBuffer = nullptr;
+  VmaAllocation directionalLightAllocation;
+
+  Buffer pointLightBuffer = nullptr;
+  VmaAllocation pointLightAllocation;
+
+  Buffer indexBuffer = nullptr;
+  VmaAllocation indexAllocation;
+  friend class GPURenderer;
+};

src/main.cpp

@@ -1,40 +1,47 @@
+#include "gpu/GPURenderer.h"
 #include "scene/Renderer.h"
 #include "util/ModelLoader.h"
 #include "window/Window.h"
-#include <iostream>
 #include <imgui.h>
+#include <iostream>
 
 int main()
 {
-    Renderer scene;
-    Window window(1920, 1080);
-    scene.startRender(
-        Camera{
-            .position = glm::vec3(-2, 5, 5),
-            .target = glm::vec3(0, 0, 0),
-        },
-        RenderParameter{
-            .width = 1920,
-            .height = 1080,
-            .numSamples = 10000,
-        });
-    while (true)
+  std::unique_ptr<Renderer> scene = std::make_unique<Renderer>();
+  Window window(1920, 1080);
+  Camera camera = Camera{
+      .position = glm::vec3(5, 1, 2),
+      .target = glm::vec3(0, 0, 0),
+      .S_O = 6,
+  };
+  RenderParameter render = RenderParameter{
+      .width = 1920,
+      .height = 1080,
+      .numSamples = 10000,
+  };
+  scene->startRender(camera, render);
+
+  while (true)
+  {
+    window.beginFrame();
+    ImGui::Text("Camera Parameters");
+    ImGui::InputFloat3("Position", &camera.position.x);
+    ImGui::InputFloat3("Target", &camera.target.x);
+    ImGui::InputFloat("Focal Length", &camera.f);
+    ImGui::InputFloat("Aperture", &camera.A);
+    ImGui::InputFloat("S_O", &camera.S_O);
+    ImGui::Text("Render Parameters");
+    ImGui::InputInt2("Dimensions", (int*)&render.width);
+    ImGui::InputInt("Samples", (int*)&render.numSamples);
+    if (ImGui::Button("Render"))
     {
-        window.beginFrame();
-        if (ImGui::Button("Render"))
-        {
-          scene.startRender(
-              Camera{
-                  .position = glm::vec3(5, 5, 5),
-                  .target = glm::vec3(0, 0, 0),
-              },
-              RenderParameter{
-                  .width = 1920,
-                  .height = 1080,
-                  .numSamples = 10000,
-              });
-        }
-        window.update(scene.getImage());
+      scene->startRender(camera, render);
     }
-    return 0;
+    ImGui::Text("Render Stats");
+    ImGui::Text("Last Sample Time:    %.3f ms", scene->getLastSampleTime());
+    ImGui::Text("Average Sample Time: %.3f ms", scene->getAverageSampleTime());
+    ImGui::PlotLines("Sample Times", scene->getSampleTimes().data(), scene->getSampleTimes().size(), 0, 0, FLT_MAX, FLT_MAX, ImVec2(0, 40));
+    window.update(scene->getImage());
+  }
+  return 0;
 }

src/scene/Renderer.cpp

@@ -1,4 +1,5 @@
 #include "Renderer.h"
+#include "gpu/GPUScene.h"
 #include "util/ModelLoader.h"
 #include <chrono>
 #include <iostream>
@@ -6,29 +7,34 @@
 
 Renderer::Renderer()
 {
-  bvh.addDirectionalLight(DirectionalLight{
+  scene = std::make_unique<Scene>();
+
+  scene->addDirectionalLight(DirectionalLight{
       .direction = glm::normalize(glm::vec3(-0.4f, -0.3f, -0.2f)),
       .color = glm::vec3(1, 1, 1),
   });
-  bvh.addModels(ModelLoader::loadModel("../res/models/cube.fbx"),
-                glm::mat4(glm::vec4(1.0f, 0.0f, 0.0f, 0.0f), glm::vec4(0.0f, 1.0f, 0.0f, 0.0f), glm::vec4(0.0f, 0.0f, 1.0f, 0.0f),
-                          glm::vec4(0.0f, 0.0f, 0.0f, 1.0f)));
-  bvh.generate();
+  scene->addModels(ModelLoader::loadModel("../res/models/cube.fbx"),
+                   glm::mat4(glm::vec4(1.0f, 0.0f, 0.0f, 0.0f), glm::vec4(0.0f, 1.0f, 0.0f, 0.0f), glm::vec4(0.0f, 0.0f, 1.0f, 0.0f),
+                             glm::vec4(0.0f, 0.0f, 0.0f, 1.0f)));
+  scene->generate();
 }
 
 Renderer::~Renderer() {}
 
 void Renderer::startRender(Camera cam, RenderParameter params)
 {
-  threadPool.cancel();
-  pendingCancel = true;
-  if (worker.joinable())
+  //threadPool.cancel();
+  if (running)
+  {
+    running = false;
     worker.join();
-  pendingCancel = false;
+  }
+  sampleTimes.clear();
   image.clear();
   accumulator.clear();
   image.resize(params.width * params.height);
   accumulator.resize(params.width * params.height);
+  running = true;
   worker = std::thread(&Renderer::render, this, cam, params);
 }
 
@@ -39,13 +45,11 @@ glm::vec3 rand01(glm::uvec3 x)
   return glm::vec3(x) * (1.0f / float(0xffffffffU));
 }
 
-thread_local glm::vec3 rnd01;
-
 void Renderer::render(Camera camera, RenderParameter params)
 {
   for (int samp = 0; samp < params.numSamples; ++samp)
   {
-    if (pendingCancel)
+    if (!running)
       return;
     auto start = std::chrono::high_resolution_clock::now();
     Batch batch;
@@ -57,6 +61,7 @@ void Renderer::render(Camera camera, RenderParameter params)
             // #pragma omp parallel for
             for (int h = 0; h < params.height; ++h)
             {
+              Payload payload;
               Ray cam = Ray(camera.position, glm::normalize(camera.target - camera.position));
               glm::vec3 cx =
                             glm::normalize(glm::cross(cam.direction, abs(cam.direction.y) < 0.9 ? glm::vec3(0, 1, 0) : glm::vec3(0, 0, 1))),
@@ -67,8 +72,9 @@ void Renderer::render(Camera camera, RenderParameter params)
 
               //-- sample sensor
               glm::uvec2 pix = glm::uvec2(w, h);
-              rnd01 = rand01(glm::uvec3(pix, samp));
-              glm::vec2 rnd2 = 2.0f * glm::vec2(rnd01); // vvv tent filter sample
+
+              payload.rnd01 = rand01(glm::uvec3(pix, samp));
+              glm::vec2 rnd2 = 2.0f * glm::vec2(payload.rnd01); // vvv tent filter sample
               glm::vec2 tent =
                   glm::vec2(rnd2.x < 1 ? sqrt(rnd2.x) - 1 : 1 - sqrt(2 - rnd2.x), rnd2.y < 1 ? sqrt(rnd2.y) - 1 : 1 - sqrt(2 - rnd2.y));
               glm::vec2 s =
@@ -76,7 +82,7 @@ void Renderer::render(Camera camera, RenderParameter params)
                    0.5f) *
                   sdim;
               glm::vec3 spos = cam.origin + cx * s.x + cy * s.y, lc = cam.origin + cam.direction * 0.035f; // sample on 3d sensor plane
-              Ray r = Ray(lc, normalize(lc - spos));                  // construct ray
+              Ray r = Ray(lc, normalize(lc - spos));                                                       // construct ray
 
               //-- setup lens
               glm::vec3 lensP = lc;
@@ -84,27 +90,27 @@ void Renderer::render(Camera camera, RenderParameter params)
               glm::vec3 lensX = glm::cross(lensN, glm::vec3(0, 1, 0)); // the exact vector doesnt matter
               glm::vec3 lensY = glm::cross(lensN, lensX);
 
-              glm::vec3 lensSample = lensP + rnd01.x * camera.A * lensX + rnd01.y * camera.A * lensY;
+              glm::vec3 lensSample = lensP + payload.rnd01.x * camera.A * lensX + payload.rnd01.y * camera.A * lensY;
 
               glm::vec3 focalPoint = cam.origin + (camera.S_O + S_I) * cam.direction;
               float t = glm::dot(focalPoint - r.origin, lensN) / glm::dot(r.direction, lensN);
               glm::vec3 focus = r.origin + t * r.direction;
-              //r = Ray(lensSample, normalize(focus - lensSample)); // TODO: Fix lens
+              // r = Ray(lensSample, normalize(focus - lensSample)); // TODO: Fix lens
 
-              Payload payload;
-              bvh.traceRay(r, payload, 1e-4, 1e20);
+              scene->traceRay(r, payload, 1e-4, 1e20);
 
-              accumulator[w + h * params.width] += payload.accumulatedRadiance;
+              accumulator[w + h * params.width] += payload.accumulatedRadiance / float(params.numSamples);
             }
             co_return;
           }(w, samp));
     }
     threadPool.runBatch(std::move(batch));
     auto end = std::chrono::high_resolution_clock::now();
-    std::cout << std::chrono::duration_cast<std::chrono::milliseconds>(end - start).count() << std::endl;
+    sampleTimes.push_back(std::chrono::duration_cast<std::chrono::microseconds>(end - start).count() / 1000.0f);
+    float resolver = float(params.numSamples) / float(samp+1);
     for (uint32_t i = 0; i < accumulator.size(); ++i)
     {
-      image[i] = glm::pow(glm::max((accumulator[i] / float(samp+1)), 0.0f), glm::vec3(0.45f));
+      image[i] = glm::pow(glm::max(accumulator[i] * resolver, 0.0f), glm::vec3(0.45f));
     }
   }
 }

src/scene/Renderer.h

@@ -1,33 +1,44 @@
 #pragma once
 #include "Scene.h"
-#include "window/Window.h"
-#include "util/Camera.h"
 #include "ThreadPool.h"
+#include "util/Camera.h"
+#include "window/Window.h"
+#include <numeric>
 
 struct RenderParameter
 {
-    int width;
-    int height;
-    int numSamples;
+  uint32_t width;
+  uint32_t height;
+  uint32_t numSamples;
 };
 
 class Renderer
 {
-  public:
-    Renderer();
-    virtual ~Renderer();
-    void startRender(Camera cam, RenderParameter params);
-    constexpr const std::vector<glm::vec3>& getImage() const { return image; }
-  private:
-    virtual void render(Camera cam, RenderParameter params);
-    ThreadPool threadPool;
-    std::thread worker;
-    std::atomic_bool pendingCancel = false;
-    // the thing being displayed
-    std::vector<glm::vec3> image;
-    // radiance accumulator
-    std::vector<glm::vec3> accumulator;
-    std::vector<PointLight> pointLights;
-    std::vector<DirectionalLight> directionalLights;
-    Scene bvh;
+public:
+  Renderer();
+  virtual ~Renderer();
+  void startRender(Camera cam, RenderParameter params);
+  constexpr const std::vector<glm::vec3>& getImage() const { return image; }
+  constexpr const std::vector<float>& getSampleTimes() const { return sampleTimes; }
+  constexpr const float getLastSampleTime() const { return sampleTimes.empty() ? 0 : sampleTimes.back(); }
+  constexpr const float getAverageSampleTime() const
+  {
+    return std::accumulate(sampleTimes.begin(), sampleTimes.end(), 0.0f) / sampleTimes.size();
+  }
+
+protected:
+  virtual void render(Camera cam, RenderParameter params);
+  ThreadPool threadPool;
+  std::thread worker;
+  std::atomic_bool running = false;
+  std::vector<float> sampleTimes;
+  float lastSampleTime;
+  float averageSampleTime;
+  // the thing being displayed
+  std::vector<glm::vec3> image;
+  // radiance accumulator
+  std::vector<glm::vec3> accumulator;
+  std::vector<PointLight> pointLights;
+  std::vector<DirectionalLight> directionalLights;
+  std::unique_ptr<Scene> scene;
 };

src/scene/Scene.cpp

@@ -20,12 +20,13 @@ void Scene::addModels(std::vector<PModel> _models, glm::mat4 transform)
 
 void Scene::generate()
 {
-  std::vector<PNode> pendingNodes;
-  while (!models.empty())
+  // todo: clear everything
+  for (uint32_t i = 0; i < models.size(); ++i)
   {
-    auto& model = models.back();
+    auto& model = models[i];
     ModelReference ref = {
         .positionOffset = (uint32_t)positionPool.size(),
+        .numPositions = (uint32_t)model->positions.size(),
         .indicesOffset = (uint32_t)indicesPool.size(),
         .numIndices = (uint32_t)model->indices.size(),
     };
@@ -33,6 +34,7 @@ void Scene::generate()
     {
       positionPool.push_back(model->positions[i]);
       texCoordsPool.push_back(model->texCoords[i]);
+      normalsPool.push_back(model->normals[i]);
     }
     for (uint32_t i = 0; i < model->indices.size(); ++i)
     {
@@ -41,8 +43,78 @@ void Scene::generate()
       edgesPool.push_back(model->edges[i * 2 + 1]);
       faceNormalsPool.push_back(glm::normalize(model->faceNormals[i]));
     }
+    refs.push_back(ref);
+  }
+  createRayTracingHierarchy();
+}
+
+void Scene::traceRay(Ray ray, Payload& payload, const float tmin, const float tmax) const noexcept
+{
+  IntersectionInfo info = generateIntersections(hierarchy, ray, tmin, tmax);
+
+  if (info.hitInfo.t < std::numeric_limits<float>::max())
+  {
+    // russian roulette ray termination
+    float p = std::max(std::max(info.brdf.albedo.x, info.brdf.albedo.y), info.brdf.albedo.z);
+
+    if (payload.depth >= 12)
+    {
+      return;
+    }
+    else if (payload.depth > 5)
+    {
+      if (payload.rnd01.z >= p)
+        return;
+      else
+        payload.accumulatedMaterial /= p;
+    }
+    // emissive
+    payload.accumulatedRadiance += payload.accumulatedMaterial * info.brdf.emissive * payload.emissive;
+    payload.accumulatedMaterial *= info.brdf.albedo;
+
+    // direct lighting
+    for (const auto& d : directionalLights)
+    {
+      // if there is an intersection, the light is occluded so no lighting
+      if (!testIntersection(hierarchy, Ray(info.hitInfo.position, -d.direction), 1e-4, 1e20))
+      {
+        payload.accumulatedRadiance += info.brdf.evaluate(info.hitInfo, -ray.direction, -d.direction, d.color);
+      }
+    }
+    for (const auto& p : pointLights)
+    {
+      glm::vec3 lightDir = p.position - info.hitInfo.position;
+      // if (!testIntersection(hierarchy, Ray(info.hitInfo.position, -lightDir), 1e-4, 1))
+      {
+        float d = glm::length(lightDir);
+        float illuminance = std::max(1 - d / p.attenuation, 0.0f);
+
+        payload.accumulatedRadiance += illuminance * info.brdf.evaluate(info.hitInfo, -ray.direction, lightDir, p.color);
+      }
+    }
+
+    // TODO: Next Event Estimation for mesh lights
+
+    // indirect lighting
+    float r1 = 2 * std::numbers::pi * payload.rnd01.x;
+    float r2 = payload.rnd01.y;
+    float r2s = sqrt(r2);
+    glm::vec3 w = info.hitInfo.normalLight;
+    glm::vec3 u = glm::normalize(glm::cross(std::abs(w.x) > 0.1 ? glm::vec3(0, 1, 0) : glm::vec3(1, 0, 0), w));
+    glm::vec3 v = glm::cross(w, u);
+    ray = Ray(info.hitInfo.position, glm::normalize(u * cos(r1) * r2s + v * sin(r1) * r2s + w * sqrt(1 - r2)));
+    payload.emissive = 0;
+    payload.depth++;
+    traceRay(ray, payload, tmin, tmax);
+  }
+}
+
+void Scene::createRayTracingHierarchy()
+{
+  std::vector<PNode> pendingNodes;
+  for (const auto& [model, ref] : std::views::zip(models, refs))
+  {
     pendingNodes.push_back(std::make_unique<Node>(model->boundingBox, ref));
-    models.pop_back();
   }
   while (pendingNodes.size() > 1)
   {
@@ -77,69 +149,6 @@ void Scene::generate()
   hierarchy = std::move(pendingNodes[0]);
 }
 
-extern glm::vec3 rnd01;
-
-void Scene::traceRay(Ray ray, Payload& payload, const float tmin, const float tmax) const noexcept
-{
-  IntersectionInfo info = generateIntersections(hierarchy, ray, tmin, tmax);
-  
-  if (info.hitInfo.t < std::numeric_limits<float>::max())
-  {
-    // russian roulette ray termination
-    float p = std::max(std::max(info.brdf.albedo.x, info.brdf.albedo.y), info.brdf.albedo.z);
-
-    if (payload.depth >= 12)
-    {
-      return;
-    }
-    else if (payload.depth > 5)
-    {
-      if (rnd01.z >= p)
-        return;
-      else
-        payload.accumulatedMaterial /= p;
-    }
-    // emissive
-    payload.accumulatedRadiance += payload.accumulatedMaterial * info.brdf.emissive * payload.emissive;
-    payload.accumulatedMaterial *= info.brdf.albedo;
-
-    // direct lighting
-    for (const auto& d : directionalLights)
-    {
-      // if there is an intersection, the light is occluded so no lighting
-      if (!testIntersection(hierarchy, Ray(info.hitInfo.position, -d.direction), 1e-4, 1e20))
-      {
-        payload.accumulatedRadiance += info.brdf.evaluate(info.hitInfo, -ray.direction, -d.direction, d.color);
-      }
-    }
-    for (const auto& p : pointLights)
-    {
-      glm::vec3 lightDir = p.position - info.hitInfo.position;
-      // if (!testIntersection(hierarchy, Ray(info.hitInfo.position, -lightDir), 1e-4, 1))
-      {
-        float d = glm::length(lightDir);
-        float illuminance = std::max(1 - d / p.attenuation, 0.0f);
-
-        payload.accumulatedRadiance += illuminance * info.brdf.evaluate(info.hitInfo, -ray.direction, lightDir, p.color);
-      }
-    }
-
-    // TODO: Next Event Estimation for mesh lights
-
-    // indirect lighting
-    float r1 = 2 * std::numbers::pi * rnd01.x;
-    float r2 = rnd01.y;
-    float r2s = sqrt(r2);
-    glm::vec3 w = info.hitInfo.normalLight;
-    glm::vec3 u = glm::normalize(glm::cross(std::abs(w.x) > 0.1 ? glm::vec3(0, 1, 0) : glm::vec3(1, 0, 0), w));
-    glm::vec3 v = glm::cross(w, u);
-    ray = Ray(info.hitInfo.position, glm::normalize(u * cos(r1) * r2s + v * sin(r1) * r2s + w * sqrt(1 - r2)));
-    payload.emissive = 0;
-    payload.depth++;
-    traceRay(ray, payload, tmin, tmax);
-  }
-}
-
 bool Scene::testIntersection(const PNode& currentNode, const Ray ray, const float tmin, float tmax) const noexcept
 {
   if (!currentNode->aabb.intersects(ray, tmin, tmax))
@@ -156,8 +165,7 @@ bool Scene::testIntersection(const PNode& currentNode, const Ray ray, const floa
   return leftResults || rightResults;
 }
 
-IntersectionInfo Scene::generateIntersections(const PNode& currentNode, const Ray ray, const float tmin,
-                                                           float tmax) const noexcept
+IntersectionInfo Scene::generateIntersections(const PNode& currentNode, const Ray ray, const float tmin, float tmax) const noexcept
 {
   if (!currentNode->aabb.intersects(ray, tmin, tmax))
   {
@@ -221,8 +229,7 @@ bool Scene::testModel(const ModelReference& reference, const Ray ray, const floa
 
   return false;
 }
-IntersectionInfo Scene::intersectModel(const ModelReference& reference, const Ray ray, const float tmin,
-                                                    float tmax) const noexcept
+IntersectionInfo Scene::intersectModel(const ModelReference& reference, const Ray ray, const float tmin, float tmax) const noexcept
 {
   IntersectionInfo intersection = {};
 

src/scene/Scene.h

@@ -9,6 +9,7 @@
 struct ModelReference
 {
   uint32_t positionOffset = 0;
+  uint32_t numPositions = 0;
   uint32_t indicesOffset = 0;
   uint32_t numIndices = 0;
 };
@@ -16,6 +17,7 @@ struct ModelReference
 struct PointLight
 {
   glm::vec3 position = glm::vec3(0, 0, 0);
+  float pad;
   glm::vec3 color = glm::vec3(1, 1, 1);
   float attenuation = 1;
 };
@@ -23,7 +25,9 @@ struct PointLight
 struct DirectionalLight
 {
   glm::vec3 direction = glm::vec3(0, 1, 0);
+  float pad;
   glm::vec3 color = glm::vec3(1, 1, 1);
+  float pad1;
 };
 
 class Scene
@@ -37,9 +41,11 @@ public:
 
   void traceRay(Ray ray, Payload& payload, const float tmin, const float tmax) const noexcept;
 
-private:
+protected:
+  std::vector<ModelReference> refs;
   std::vector<glm::vec3> positionPool;
   std::vector<glm::vec2> texCoordsPool;
+  std::vector<glm::vec3> normalsPool;
   std::vector<glm::uvec3> indicesPool;
   std::vector<glm::vec3> edgesPool;
   std::vector<glm::vec3> faceNormalsPool;
@@ -60,9 +66,12 @@ private:
   PNode hierarchy;
   std::vector<PModel> models;
 
+  virtual void createRayTracingHierarchy();
+
   // tests if a ray intersects any geometry, no hit information, for shadow rays
   bool testIntersection(const PNode& currentNode, const Ray ray, const float tmin, const float tmax) const noexcept;
   IntersectionInfo generateIntersections(const PNode& currentNode, const Ray ray, const float tmin, const float tmax) const noexcept;
   bool testModel(const ModelReference& reference, const Ray ray, const float tmin, const float tmax) const noexcept;
   IntersectionInfo intersectModel(const ModelReference& reference, const Ray ray, const float tmin, const float tmax) const noexcept;
+  friend class GPURenderer;
 };

src/util/Camera.h

@@ -6,7 +6,7 @@ struct Camera
     glm::vec3 position;
     glm::vec3 target;
     glm::vec2 sensorSize = glm::vec2(0.036, 0.024);
-    float S_O = 6.9;
+    float S_O = 20;
     float f = 0.7;
     float A = 0.35;
 };

src/util/Model.cpp

@@ -7,6 +7,11 @@ void Model::transform(glm::mat4 matrix)
     pos = glm::vec3(matrix * glm::vec4(pos, 1));
   }
 
+  for (auto& nor : normals)
+  {
+    nor = glm::mat3(matrix) * nor;
+  }
+
   boundingBox.transform(matrix);
 
   for (int i = 0; i < indices.size(); i++)

src/util/Model.h

@@ -16,6 +16,7 @@ public:
   AABB boundingBox;
   std::vector<glm::vec3> positions;
   std::vector<glm::vec2> texCoords;
+  std::vector<glm::vec3> normals;
   std::vector<glm::uvec3> indices;
   std::vector<glm::vec3> edges;
   std::vector<glm::vec3> faceNormals;

src/util/ModelLoader.cpp

@@ -8,7 +8,7 @@
 std::vector<PModel> ModelLoader::loadModel(std::string_view filename)
 {
   Assimp::Importer importer;
-  const aiScene* scene = importer.ReadFile(std::string(filename), aiProcess_Triangulate);
+  const aiScene* scene = importer.ReadFile(std::string(filename), aiProcess_Triangulate | aiProcess_GenNormals);
   std::cout << importer.GetErrorString() << std::endl;
   std::vector<PModel> result;
   for (int m = 0; m < scene->mNumMeshes; ++m)
@@ -20,7 +20,15 @@ std::vector<PModel> ModelLoader::loadModel(std::string_view filename)
     {
       auto aiVert = mesh->mVertices[v];
       model->positions.push_back(glm::vec3(aiVert.x, aiVert.y, aiVert.z));
-      model->texCoords.push_back(glm::vec2(mesh->mTextureCoords[0][v].x, mesh->mTextureCoords[0][v].y));
+      if (mesh->HasTextureCoords(0))
+      {
+        model->texCoords.push_back(glm::vec2(mesh->mTextureCoords[0][v].x, mesh->mTextureCoords[0][v].y));
+      }
+      else
+      {
+        model->texCoords.push_back(glm::vec2(0, 0));
+      }
+      model->normals.push_back(glm::vec3(mesh->mNormals[v].x, mesh->mNormals[v].y, mesh->mNormals[v].z));
       aabb.adjust(model->positions.back());
     }
     for (int i = 0; i < mesh->mNumFaces; ++i)

src/util/Ray.h

@@ -3,6 +3,7 @@
 
 struct Payload
 {
+  glm::vec3 rnd01;
   glm::vec3 accumulatedRadiance = glm::vec3(0);
   glm::vec3 accumulatedMaterial = glm::vec3(1);
   uint32_t depth = 0;

vcpkg.json

@@ -5,6 +5,7 @@
       "features": [ "glfw-binding", "opengl3-binding" ]
     },
     "vulkan",
+    "vulkan-memory-allocator",
     "assimp",
     "ktx",
     "glfw3",