From fb6b84a7b7b8ea2b765d8f228053ff407a6f388a Mon Sep 17 00:00:00 2001 From: Dynamitos Date: Sun, 2 Aug 2026 07:56:02 +0200 Subject: [PATCH] Rework compute kernel --- .vscode/settings.json | 3 + CMakeLists.txt | 37 ++-- res/shaders/ClosestHit.slang | 125 ----------- res/shaders/ComputeKernel.slang | 232 +++++++++++++++------ res/shaders/Miss.slang | 8 - res/shaders/RayGen.slang | 57 ----- src/main.cpp | 6 +- src/metal/CMakeLists.txt | 3 +- src/metal/Compute.air | Bin 14128 -> 0 bytes src/metal/Compute.metal | 323 ----------------------------- src/metal/Compute.metallib | Bin 19045 -> 0 bytes src/metal/ComputeKernel.metal | 357 ++++++++++++++++++++++++++++++++ src/metal/MetalRenderer.h | 1 + src/metal/MetalRenderer.mm | 168 +++++++++------ src/metal/MetalScene.mm | 2 +- src/scene/Scene.h | 1 + 16 files changed, 665 insertions(+), 658 deletions(-) create mode 100644 .vscode/settings.json delete mode 100644 res/shaders/ClosestHit.slang delete mode 100644 res/shaders/Miss.slang delete mode 100644 res/shaders/RayGen.slang delete mode 100644 src/metal/Compute.air delete mode 100644 src/metal/Compute.metal delete mode 100644 src/metal/Compute.metallib create mode 100644 src/metal/ComputeKernel.metal diff --git a/.vscode/settings.json b/.vscode/settings.json new file mode 100644 index 0000000..bcc8af4 --- /dev/null +++ b/.vscode/settings.json @@ -0,0 +1,3 @@ +{ + "cmake.generator": "Ninja" +} \ No newline at end of file diff --git a/CMakeLists.txt b/CMakeLists.txt index 8c0071e..68eb682 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -17,6 +17,7 @@ find_package(glfw3 CONFIG REQUIRED) find_package(glm CONFIG REQUIRED) find_package(Ktx CONFIG REQUIRED) find_package(imgui CONFIG REQUIRED) +find_package(slang CONFIG REQUIRED) add_executable(RayTracer "") target_include_directories(RayTracer PUBLIC src/) @@ -26,19 +27,29 @@ 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) +target_link_libraries(RayTracer PUBLIC slang::slang) + +if(APPLE) + target_include_directories(RayTracer PUBLIC ${VCPKG_INSTALLED_DIR}/arm64-osx/include) + set(CMAKE_OSX_DEPLOYMENT_TARGET 26.0) + target_link_libraries(RayTracer PUBLIC + "-framework Metal" + "-framework MetalKit" + "-framework AppKit" + "-framework Foundation" + "-framework QuartzCore" + ) +endif() + 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) -elseif(APPLE) -target_include_directories(RayTracer PUBLIC ${VCPKG_INSTALLED_DIR}/arm64-osx/include) -SET(CMAKE_OSX_DEPLOYMENT_TARGET 26.0) -target_link_libraries(RayTracer PUBLIC - "-framework Metal" - "-framework MetalKit" - "-framework AppKit" - "-framework Foundation" - "-framework QuartzCore" + target_include_directories(RayTracer PUBLIC ${VCPKG_INSTALLED_DIR}/x64-windows/include) + target_link_libraries(RayTracer PUBLIC ${VCPKG_INSTALLED_DIR}/x64-windows/lib/slang.lib) +endif() + +add_custom_command(TARGET RayTracer POST_BUILD + COMMAND ${CMAKE_COMMAND} -E copy_directory + ${CMAKE_CURRENT_SOURCE_DIR}/res + $ ) -endif() -add_subdirectory(src/) \ No newline at end of file +add_subdirectory(src/) diff --git a/res/shaders/ClosestHit.slang b/res/shaders/ClosestHit.slang deleted file mode 100644 index 3ad9fe2..0000000 --- a/res/shaders/ClosestHit.slang +++ /dev/null @@ -1,125 +0,0 @@ -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 = pParams.materialData[m.materialIndex]; - float3 emissive = mat.emissive_type.xyz; - - 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 + emissive; - // 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); - } -} \ No newline at end of file diff --git a/res/shaders/ComputeKernel.slang b/res/shaders/ComputeKernel.slang index ed9b834..a159f28 100644 --- a/res/shaders/ComputeKernel.slang +++ b/res/shaders/ComputeKernel.slang @@ -1,81 +1,185 @@ import Common; +struct HitInfo +{ + float3 position; + float3 normal; + float3 barycentricCoords; + uint instanceIndex; + uint primitiveIndex; +}; + +HitInfo get_hit_info(RayQuery q) +{ + HitInfo info; + // In Slang for Metal/Vulkan, these are the standard names for ray query results + info.instanceIndex = q.CommittedInstanceID(); + info.primitiveIndex = q.CommittedPrimitiveIndex(); + float2 baryCenter = q.CommittedRayBarycentrics(); + info.barycentricCoords = float3(1.0f - baryCenter.x - baryCenter.y, baryCenter.x, baryCenter.y); + return info; +} + +Vertex interpolate_vertex(uint vertexIdx0, uint vertexIdx1, uint vertexIdx2, float3 bary) +{ + Vertex v0 = loadVertex(vertexIdx0); + Vertex v1 = loadVertex(vertexIdx1); + Vertex v2 = loadVertex(vertexIdx2); + + Vertex vert; + vert.position = v0.position * bary.x + v1.position * bary.y + v2.position * bary.z; + vert.texCoords = v0.texCoords * bary.x + v1.texCoords * bary.y + v2.texCoords * bary.z; + vert.normal = v0.normal * bary.x + v1.normal * bary.y + v2.normal * bary.z; + return vert; +} + [shader("compute")] [numthreads(8, 8, 1)] -void computeKernel(uint2 threadId [[thread_position_in_grid]]) +void computeKernel(uint2 threadId: SV_DispatchThreadID) { - if (threadId.x >= pParams.cam.width || threadId.y >= pParams.cam.height) - return; + if (threadId.x >= pParams.cam.width || threadId.y >= pParams.cam.height) + return; - uint pass = pSamps.pass; - uint samplesPerPixel = pSamps.samplesPerPixel; - if (pass == samplesPerPixel) return; + uint pass = pSamps.pass; + uint samplesPerPixel = pSamps.samplesPerPixel; + if (pass == samplesPerPixel) + return; - uint2 pix = threadId; - uint imgWidth = pParams.cam.width; - uint imgHeight = pParams.cam.height; + uint2 pix = threadId; + uint imgWidth = pParams.cam.width; + uint imgHeight = pParams.cam.height; - //-- define cam - float3 camPos = pParams.cam.cameraPosition; - float3 camForward = pParams.cam.cameraForward; - float f = pParams.cam.f; - float S_O = pParams.cam.S_O; - float3 fogEmm = pParams.cam.fogEmm; - float ks = pParams.cam.ks; - float A = pParams.cam.A; - float ka = pParams.cam.ka; - float2 sensorSize = pParams.cam.sensorSize; + // -- Camera setup -- + float3 camPos = pParams.cam.cameraPosition; + float3 camForward = pParams.cam.cameraForward; + float f = pParams.cam.f; + float S_O = pParams.cam.S_O; + float3 fogEmm = pParams.cam.fogEmm; + float ks = pParams.cam.ks; + float A = pParams.cam.A; + float ka = pParams.cam.ka; + float2 sensorSize = pParams.cam.sensorSize; - float3 cx = -normalize(cross(camForward, abs(camForward.y) < 0.9 ? float3(0, 1, 0) : float3(0, 0, 1))); - float3 cy = cross(camForward, cx); - const float2 sdim = sensorSize; + float3 cx = -normalize(cross(camForward, abs(camForward.y) < 0.9 ? float3(0, 1, 0) : float3(0, 0, 1))); + float3 cy = cross(camForward, cx); + const float2 sdim = sensorSize; - float S_I = (S_O * f) / (S_O - f); + float S_I = (S_O * f) / (S_O - f); - //-- sample sensor - float3 rnd = rand01(uint3(pix, pass)); - float2 rnd2 = 2.0f * float2(rnd.xy); // tent filter - 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 = ((float2(pix) + 0.5f * (0.5f + float2((pass / 2) % 2, pass % 2) + tent)) / float2(imgWidth, imgHeight) - 0.5f) * sdim; - - float3 lc = camPos + camForward * 0.035f; // sample on 3d sensor plane - float3 spos = camPos + cx * s.x + cy * s.y; - float3 rayDir = normalize(lc - spos); + // -- Sample sensor -- + float3 rnd = rand01(uint3(pix, pass)); + float2 rnd2 = 2.0f * float2(rnd.xy); // tent filter + 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 = ((float2(pix) + 0.5f * (0.5f + float2((pass / 2) % 2, pass % 2) + tent)) / float2(imgWidth, imgHeight) - 0.5f) * sdim; - //-- setup lens (simplified) - float3 lensSample = lc; // for now, just use camera position slightly offset if needed? - // Actually let's do it properly based on A parameter - float3 lensN = -camForward; - float3 lensX = cross(lensN, float3(0, 1, 0)); - float3 lensY = cross(lensN, lensX); - float2 rnd01 = rand01(uint3(pix, pass)).xy; - lensSample = lc + rnd01.x * A * lensX + rnd01.y * A * lensY; + float3 lc = camPos + camForward * 0.035f; // sample on 3d sensor plane + float3 spos = camPos + cx * s.x + cy * s.y; + float3 rayDir = normalize(lc - spos); - float focalPoint = camPos + (S_O + S_I) * camForward; - float t_focus = dot(focalPoint - lensSample, lensN) / dot(rayDir, lensN); - float3 focus = lensSample + t_focus * rayDir; - - float3 rayOrg = lensSample; - float3 rayDirFinal = normalize(focus - lensSample); + // -- Lens (Aperture) -- + float3 lensN = -camForward; + float3 lensX = cross(lensN, float3(0, 1, 0)); + float3 lensY = cross(lensN, lensX); + float2 rndL = rand01(uint3(pix, pass + 100)).xy; + float3 lensSample = lc + (rndL.x - 0.5) * A * lensX + (rndL.y - 0.5) * A * lensY; - // Ray Tracing Loop - RayPayload payload; - payload.light = float3(0); - payload.emissive = 1.0f; - payload.depth = 1; - payload.hit = false; - payload.anyHit = false; + float3 focalPoint = camPos + (S_O + S_I) * camForward; - // Note: We are using the compute-based intersection loop because it's easier to implement in a single kernel - // and we have access to common helper functions. In a full RT pipeline we would use dedicated shaders. - - // Since we don't have the specialized 'intersector' object from before, - // we will use a placeholder for now or assume it's available if provided by Slang/Metal context. - // BUT since I am writing this from scratch, I should probably implement the traversal OR - // just use MS's Compute-based approach as in Compute.metal which worked. - - // Wait! To keep it simple and "lazy", I will just copy the logic from Compute.metal into this Slang file - // and replace all its types with pParams fields. + // Simple ray construction + float3 rayOrg = lensSample; + float3 rayDirFinal = normalize(focalPoint - lensSample); + // -- Path Tracing Loop -- + float3 accumulatedRadiance = float3(0.0); + float3 throughput = float3(1.0); + + for (int bounce = 0; bounce < 4; ++bounce) + { + RayQuery q; + RayDesc rayDesc; + rayDesc.Origin = rayOrg; + rayDesc.Direction = rayDirFinal; + rayDesc.TMin = 0.001; + rayDesc.TMax = 1e20; + q.TraceRayInline(pParams.scene, RAY_FLAG_NONE, 0xff, rayDesc); + + if (q.Proceed()) + { + HitInfo hit = get_hit_info(q); + + ModelReference m = pParams.modelData[hit.instanceIndex]; + uint indexOffset = m.indicesOffset; + uint vertexOffset = m.positionOffset; + + uint v0 = vertexOffset + pParams.indexBuffer[indexOffset + 3 * hit.primitiveIndex + 0]; + uint v1 = vertexOffset + pParams.indexBuffer[indexOffset + 3 * hit.primitiveIndex + 1]; + uint v2 = vertexOffset + pParams.indexBuffer[indexOffset + 3 * hit.primitiveIndex + 2]; + + Vertex vert = interpolate_vertex(v0, v1, v2, hit.barycentricCoords); + MaterialParameter mat = pParams.materialData[m.materialIndex]; + + accumulatedRadiance += throughput * mat.emissive_type.xyz; + + // --- Direct Lighting (NEE) --- + float3 directLight = float3(0); + for (uint i = 0; i < pSamps.numDirectionalLights; ++i) + { + float3 lDir = -pParams.directionalLights[i].direction.xyz; + RayQuery sq; + RayDesc rayDesc; + rayDesc.Origin = vert.position + vert.normal * 0.001; + rayDesc.Direction = lDir; + rayDesc.TMin = 0.001; + rayDesc.TMax = 1e20; + sq.TraceRayInline(pParams.scene, RAY_FLAG_NONE, 0xff, rayDesc); + if (!sq.Proceed()) + { + directLight += mat.shade(vert.normal, -rayDirFinal, lDir, pParams.directionalLights[i].color); + } + } + for (uint i = 0; i < pSamps.numPointLights; ++i) + { + float3 lVec = pParams.pointLights[i].position - vert.position; + float3 lDir = normalize(lVec); + RayQuery sq; + RayDesc rayDesc; + rayDesc.Origin = vert.position + vert.normal * 0.001; + rayDesc.Direction = lDir; + rayDesc.TMin = 0.001; + rayDesc.TMax = 1e20; + sq.TraceRayInline(pParams.scene, RAY_FLAG_NONE, 0xff, rayDesc); + if (sq.Proceed() == false || sq.CommittedRayT() > length(lVec)) + { + directLight += mat.shade(vert.normal, -rayDirFinal, lDir, pParams.pointLights[i].color); + } + } + accumulatedRadiance += throughput * directLight; + + // --- Indirect Lighting (Cosine-weighted sampling) --- + float3 rnd = rand01(uint3(pix, pass + bounce + 200)); + float r1 = 2.0 * PI * rnd.x; + float r2 = rnd.y; + float r2s = sqrt(r2); + + float3 w = vert.normal; + float3 u = normalize(cross(abs(w.x) > 0.1 ? float3(0, 1, 0) : float3(1, 0, 0), w)); + float3 v = cross(w, u); + float3 nextDir = normalize(u * cos(r1) * r2s + v * sin(r1) * r2s + w * sqrt(1.0 - r2)); + + throughput *= mat.albedo_alpha.xyz; + + rayOrg = vert.position + vert.normal * 0.001; + rayDirFinal = nextDir; + + if (length(throughput) < 0.01) + break; + } + else + { + accumulatedRadiance += throughput * float3(0.05, 0.05, 0.1); + break; + } + } + + pParams.image[threadId] = float4(accumulatedRadiance, 1.0); } diff --git a/res/shaders/Miss.slang b/res/shaders/Miss.slang deleted file mode 100644 index 39d092a..0000000 --- a/res/shaders/Miss.slang +++ /dev/null @@ -1,8 +0,0 @@ -import Common; - -[shader("miss")] -void miss(inout RayPayload p) -{ - p.light = float3(0.05, 0.05, 0.1); // Dark blueish background instead of black - p.hit = false; -} \ No newline at end of file diff --git a/res/shaders/RayGen.slang b/res/shaders/RayGen.slang deleted file mode 100644 index 0a77fa5..0000000 --- a/res/shaders/RayGen.slang +++ /dev/null @@ -1,57 +0,0 @@ -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); -} diff --git a/src/main.cpp b/src/main.cpp index 656895c..f2fe199 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -12,12 +12,12 @@ int main() .color = glm::vec3(1, 1, 1), }); renderer->addPointLight(PointLight{}); - renderer->addModels(ModelLoader::loadModel("../res/models/cube.fbx"), + renderer->addModels(ModelLoader::loadModel("../res/models/stanford-bunny.obj"), 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))); - renderer->generate(); + renderer->generate(); Camera camera = Camera{ - .position = glm::vec3(5, 1, 2), + .position = glm::vec3(2, 1, 2), .target = glm::vec3(0, 0, 0), .f = 0, .A = 0, diff --git a/src/metal/CMakeLists.txt b/src/metal/CMakeLists.txt index e89c2a0..465b0fb 100644 --- a/src/metal/CMakeLists.txt +++ b/src/metal/CMakeLists.txt @@ -4,4 +4,5 @@ target_sources(RayTracer MetalRenderer.mm MetalScene.h MetalScene.mm - Compute.metal) \ No newline at end of file + ComputeKernel.metal +) \ No newline at end of file diff --git a/src/metal/Compute.air b/src/metal/Compute.air deleted file mode 100644 index 1fc6a942120e69d16a30a2be329b74166551eda3..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 14128 zcmb7r30%`xw*M~+$qxw0k3|U*;1}Y8dlC>qwuB&}qK(UR=rScS1fgXKNkH7QAc=|+ zTIG@U5;Us@X1y>j2qBF~!TD>0q zK_ZOk;?wxK1QCOgv)6^r$M4fj(1oGpS)v$%wlwD9pY*lU5j8|K)N*4S=y?6sDdK67lRIc7}8d)7R&Kp4@j z7WA28`s4t{UK)zIrk*(}kNr|D7;S|nf{m^iXxr(E>24L|wF+*i1tY?kL38YNdF)1c zHc8m!by7kcLdi7Ql>53Vm)s*G?ECDdz^c~^UyonY+-@KLjfAcuhB)*SKC*ekuLJni z(3{=3NnV$Pb4F_O5{|CDar*nZ+_^{AtXwE!#$KM4OPUs*ryH1#P?hB6sNuUo?e=0O zW-2Sj1l+nr@G0_*s2BbadAox%)p>Dz{C;GU`<55$p!2p zY64?s+hxr9F3F7BIt)|V3!6+}b!v#XAK9CdRvr<5{{d6{adysRQZB}cZ?UK>v*$N@ zhnuZ$ugQ$Xnad6hDKlQcnHh&VH_u_Xdk9RXG?KCaViRV?O=3HNTXDPCuEX8H6nt+3 zAN(6K(O%Iwnin*2fRks+!e*Q3KbUZKt%<$8mc6TlVQ6P>LV#7jI`*nlFxVP$RVWyx#b7%ewjBv< zlg_rSgs}~0Z)<*{JlBR|uByOW1lL*xy{#c*AkR?DjnDTv&)RbUFKBd`-4uBiK9;k{Ts&C3|%VSz2C%mHs~7cx7~lvBq@cUsfW zGcwPbQ!iW78WhpZjP$eev^qvc3q!`cqvU;#5jRNQkca!ZhvoHhZw_+Dl-xU59V4xS zAv>#BaZVB4p-4TZNa;{S4=7R_7#ZiS>1SOTCm5+7#tN*8WY=Pxuf5!Il6%*|y(QxT zhi@?gAr9`>D(<%))(wnT?%@u}c%zudntoZ4iq)l^S45vvq?~i7gJ4cY^jSp;;xstr+i$U(^GH!Vb?{kv(jhI)i;(9%-uL$m#iHl|QDl(2RWar#zXK9d6 zaWc|RSksO@&Pu`?Be1eIZa$yUft|oM%Vo?` zEM&-mIvQ5;Zm77ROZgyDaLCEMNr0`&M!EC?Y$NuLP;d!GUncm{CFlnSYYjnOM+ydn zF`epIQ%KB4ELavKTUMu7(W{8ID>9lvR2U00GA4}I2}MdjqoZ{#M@7q{wb$$NkW^Le zYyASi3RA|2+qPrL#y;92Y(Dl0!M*vYHPFBj+*=OrkeKJJuaKJp$bgnvox(?5 z)QM&j1l3DI!42AgM-qsXRVq?WAbC4H@vo|J`(o zFBUoJ4!ZLHGFmJP8Lj>QFxvODpU_mmn(2dVOn~jh^00t0S}4~``NKFkr;_)znENEx zYPP1=Dbm0hobEJmEe~kK0LwVJ*hMr=%AMTLbzB(4aux4$F?Y<1hWI-NcLe#llO{&G zgDDg%Iv5$yyx-@m@)hT-sr^rKFlx#Sa))60{M#whuUK)+ns$Vdde)o?Lx28%nLaR; zhA^HN=Ae%I*zD^{_b5^t+$+w~(tF)Vuc=A>(7-`QqcBfvG2$-BL!g=R1QT_rvyLO{ z#+j0_D;Ns$SYW#;lkB9$qToi0gGl=&T1lT2Ck+mpt^`H4+FFQ{YzV^JG4QAX8hI5W zc5BE$tKfoAFwh!ut5wiVt-n_XA)ce4ZAx)4(lHjAaWDb%*flLE(r)*~-aPg;u?;-` zUjn9Y3AiQ{z;f{KmH_MaJexsh)04Jj51KXq%W?oql~6E9i-k4i!sDw+GkXVat17T1 zQ!Cw$1Y1>$ZQCQ^0y<*vPGDm^?0785(+QFPNmfecU&0p2{^)wm7vzT@7!snuCEHWRAdjXg zj0fxl&wwyzV%NeJ0WH0@s#@E2Cp)!+vD^fI){y_)fEI%tDz>qOT_aRVwxlf?WeeG1^}kNH&CgJrp11aO2V18;{oB0}cbd-2i7+WtxY*Ep}`kLwjxPjPfczEc`^7@8|C|Xi>^RX8U-;aN?LX$#TdilRe3Xb&qZT25};{T2j)bi zqT`9L2Mn{Dgs-j>8-X>jN0d-OQ!F^;xCLz<>CqVaDUCqqG!sQ7V{+bJBRb6>sqTOh zDnR&8bFJ*h1e?{E$1_n1A&n`)R>nDEFd7Mk2ZH60X&7Aw5K^5IY{sFyT^Q2@1rokK z!4h8^e_XKb7FG+!G=T^tqC)b&qmwW;Ve~;%LJ(5(laG||V2_c#@zJw194R0}L>Mxx zvF$H(0tICoVyQxV^QMdX?z z*3)J@ZJzP0YlccDWJP4liR1X17NW0pSCF!W3@T!_kc?nn#Mv4NXlKb6SXmqI8Kg`>;dUmd87< z{5wpigO>^hSmmCXGdRc7!@@R%Y-IMRaI-2QN1gNYEd+?kIx~CSuz>sbg!NnaD_(%Aj#k$JU3Zqu9F4Zq7SJxI6X=_R~+QKc>#l=AHM>Uld2BW&d z_&9DX(QDO3TlJOIx__zC)*7qz+G!SuuyC8c)Tk{qK2~K=m+Q)=dh>BC)KwZvjUa7d zX+_~yeQ6OAM6KUiXxy&T7FK}V)gZOhk0QX=a5Y8glOEOja+JELSXQY~Pq8i1)*=Ln zMX@O=0$~k>2Ax`?MS#JmuhtlS9hO!UR|5W$wy3bUtWs@E^ovK8ii<1t+tm7^O@&Vt zS8iQdUjC9n{*s!ik{YxXhD!Y=ur?BSTWOK8q(ob~wZusE=~oDi#v_2btaNJyCG{%h z&lI4YUag^ib!vmbXD)+Mt5=rRYRf9B%X3QgS`E@mU6zlS8vs+HarKopQTA1;VZ5gr zPbq_BN-V7aQ)E-)?Nr-5<(P+WBskWOapMqHdrqI@c&#KE-tXQ`>U;bnoH03nf9IZj4FVz_<^`+D#1d}3P z@fQjo?x4`;wa~+p(e$ZSr4?HCo+$fZ(oTWWqI7DY7LgCC5U1*L7%C)U zxq2(*%6_+?sFM({8m%6f6joLnb=AhL+6rV7ZBc@>x+bSaT39^4y5=FO(xPfZNu}Oc zV0^bDs7ma*y zCg4;LmF2q9GVPKgEpl{2VNt084eG;H#cG4G5WJ$KP+RdYbtt$_MW14=QS&g3uOkS; zMzDZf2eFT&e1QZLp%6_bO2yK|B`JwwGJ&cn)-IlgmPqHXLyoibX9l2NTC^R4U#YRu zuymt(`(_9eTK!UkUbB>Hw^RX2R2#KRs8U&3O*ursO5Kw3YJ*XyuiT0vUS)Yjsjd=c zfRRKi!=_Rci%G@8h2-XJxg7RLxueux4-;EVM_f34@rh2HzJ(>id!rgivINi*fF4J{ zEC9;H7ANaNoPG%ado%!*K;Z<6Fg5@J01NCjedmE+HTe zfV}{ipxdw@01_I=ewW0eWcT8h&<2tih{#^w)y+|4Fz<=oIvzc5M|1V z^18RM{ub3Q`-y=_TWFbxR6ff~72C~uRLfcYvQONMyNa{EmTv?^jVYoA-JPs;(f!z| zuJE7S8F`yB9Ok?uw9G{np7Smak(qum3X2ov+cymaMU!Kc8dCW0Bd9L~7p~d!-pS#0{s9u%ja@J3X+lai4 zome~R;Z1h1cX@co7;%oAtOB|~(XWf#|A5)+-|`Sk z)m>}eghCr60>-+xfI2UILncMNY>N8ef~4Py758}YX5Ng2&D5}*?L<=5nHRQD#q*YY zRHQBH7yqs+{I*ZGYB_5@GD%mqDA*vj#)q9DKJ3cK3tpV`N3m5J{w@KQzt(99Ubm#b z#Z$L9=Q|zmKRX7q+kz9*i_|kqnr6OuYG&EcOr0>&sE*ubj@&MB5Armvk%tbXge-h)aO-ULSE4-13oMl6iEfR%~3>q4wNwI?9B{J`b6*BPtZ^`HsiCY|f#vIP(n)oed z@ZSHQKUz<#Ir`i~f3h92H66704%f_Lb7V0sevfO(JLa%G!pJ@9dGFA8_f)(Sa=1eY zDscu0Vw&l{;eM&&8D(riPb+N|yTw^0i0RE!S)^MuXI?Ut?iBr2v(O;X1gHHfdd^3I zZwv(fdNlCO@xTuJ;O5RtBhO#y8@bjukQ+D_|7u54YUH1#sqab8Z=ief<}O-y;lunT zw@&`9|4_w|!6bdq+E)xwJ72P{)hzsRhqZ0(PQ$OFwX#E(m4}WXMmIXI4i5Bj@8P;AjcQ3U~$uXQG)&pQ*TCDB-BZgG*xg{X)qd>oA$w+Y_jh zV66>KG3XprpgZPHIZK0E=Zqt6`b%)A6kLNxPo!W(jebX|CrBQ666HI*s9iK3y6#o- z(1Q{9T!gy-7NqQc>X`@}2;p&$H60Fz4mXPSY7)a{LmJLhEJD`?aGiqHwbRx#EF*C# zLwtq6odt@tj~DS+U!|W(=b!zy|Bgy#JH;ZH@(H?}=vqo;FzgJo({4UhM;46RhJ(s{(mC#gUx8seX&Iyv;l>hQ0wzm|0IYn17p^}?(d3cq=4=(tV`Z|>ZKu-araQB@~}Minmjfa>z;*3c-ZJNx_Fj&rusta zKlj`?cINA?{>-Wm;ZpUs7{e}s?c!JUIyPkI{|U0#oBMw8OYJj9wa*;Rk;5)$F`Plv zhQ$U=rA})|(D*hIYqk0cZQ1#3B=5jgtULbvF9bvDu$V@J#8Mg+~g{*Mq77o-#og{|alAq?>35|RA)RyMP@pFe} zFm`vbD^Bvi6lLl)AFASRw?;3l|Iyy~j{?cys$Jgh1H!D~w&hQAeh7_w|J0U&#_@4z zvbT#}dy@D;l>9V)KQua;{{29Uw`b4jsy%}}y9ZY7_Vny35ZgIR&9~)gC8rXDaK-^%rpY5qv z=j3!{B+nie#-|A5zY)fJh4G)CQp-X^UHe05d&`s~nyfKVmiM#FTWy)6n#^R5C@lWV zQ^F`-s4FmZJ1^AC3!UE>qBv;Bca`B^Ypn5c6Q|;SXo?$G$K933eS1nc&=_)Hm)Evu zguUmQZTDsN?*5+o)|ENTdCANlPRW~i_Xip$oQ*$rH2%=i`27Ge*uz?xIZk!@wL9}m zYvv7&LKPPs@%>=qPacZ0o%cg?V~Bh=$Y#yDr^vd)$O1|vn_xr&3uW|;(D*M-sRx^g z)W#6T0ekPxGW^Cvf>p)c+YmQK!!oib%;{s6bR|X_I7%5WBio*xVPeE6F}{JLP_Z+z zkI*w4tWbt}g$c?notfv{nN;~1573{qWOXUxeDq$ZKbM{7aYGr=6G0+YPnbnD6jzId z%f(w%i-h>h`)YPlXuqUEt_rhqXi+9QYu|XIklFsAIwh2|idiWe?=f~oIpl+J8r-!& z6>6X4cMb+=aQng*!k!f-AOBc#UM>#T;P^VK(xM4(XAUTg(Qu<}_WcCr zP*#tO`DgJK9o4y+{bDFn#OkJ`su`T=5-8%BDoHOfZeBvM$tJmc<#1NLxPifXWBfCs zZc@OQ&nz)_&AtzPKdb84pT$LeN6^y94XokD{JH)yA6tfgsY z1lg@~6AsrR>tKMDBa!od@XTZ6JG!4!n$W<=!p<<=vVh|YvL&3&v+FuI2}hXtF)S>l zS8VFaI!1T{NTEA~U1E1Fa?7UG4cd{hAqi|AEb8X}`}Q0Pd5qOl02DQS~-^WT>*WJe0vs4NV}hM(lV|&6+GOS zjomK*_nG#bMPUz=!$iRTDQ3L1moNnsg6Dh?FqlGEs)qzUojHI{>q|NPw$D2YX3{D)O{> z19_isqQQ7|bXE+)Oc1zVU}i$s-hc{KxIJqYZCteCEDV^$lq;TV<;Y9MSoLxej4?I) z$a!*p17ARm?x5=!+_|<8j!A*FlvxT~j&;`#M#0pP3&525Fy2#y=3g^i{#{Va_-b*F zXBO=k==J;7*a9Y+BHY^-=P~S_k6r_>bce{BP4YK6BK^E&*|aU3<;)^FtAzCz;r+bW zLhzYIj2X6=gn1`IHVcmGv#-iNlqdHEmNI7Gi{}HloLQ0$lc0ZQ1E0vpEwqGVeRkTF zI+EYq`8Gdb=P~yFu~kI$@k?{d;V*lhaX9Y&#kpQuTw5Sbe2q9C2pg}mb>i2>g0{dM z$?HmdK5_gd7C&&jj^1_7Wa8i#7cvere+P1n^|gqs2SMJyNWFQt<0Hu2|KhATJsucf zuX53%h~tWSabU$hy0cl+-bns@MG+YuzbXN1NUawRG;yOYj0}VZVVc8ge8&hMmy$U) zPFa{eIKU+7!QyFuGWL4sGGCjDopWQPy&tLPMqu3nH_`IZFQSTMc5c+yvr8f=4l6C& zZaU;o*_XFs;v4*Wa{Y>;(N{$4eSI~*BB@L1_@#uI+7TW>_pXuzcJ4hqluO?oRfOy6 z)_d`;y$%|0SJ!%PFy!Ps=g4|bP*Ftkl{>@$HsU7dUl)9-^P><`*s%rz>v0gxFiuDVvZG&Nv2rVio7`9Hy%? z%R5+ZQFS>BvWZ0#Ztjh_%^Kp>mNHh9WyAL~I6=4LBq>J;X7Hqw@m}B+bxM5mjLxhw zZ1N@^PfTdewZ?4BKah$5@wnv?KBs-qA7c{fv;MJTg>z8F2pylJvaW>zc7E>xqPc38U89hOCGl@-P*R~ zru20JgV}J^{zkVAB)uNKZN&7#@7%ltQ^-LZUwXUk8krhq=)_i-A&Vz zY5(J2v|8Dx!@K?}Q2dG6p4dw?AM90aSTfvbChwQL8a%EcO_Ph{%am1kgdo3tAD0PJ zFO3G1>0$(DvtSSLkKE+xh4%wZqy{e$r}?984VmfE<>;1L_U5nZTKgXJzNDTQ=;vog z4{#u2H44wmFF*)OwL({UXd!J`YW{|F!AfkQXYyxHEw(Gk3N#tu}EwWR8VD z2$`lV2$>%-zXTOpIG_3!+m2-4VtX!x)$Chrt$$_SNvzlP^H|5mMe(q}W}^jtrD}Qe z5%YN4?fqFnX!RxJ*B`7_f&04!)(gy^fbp&Ve+))Id{lgVj)sRwnVD_Nr*-YfZii`_ z>_WKu@5|9iK}KFco8GH{B+Pk2a$F1gK0OR;Cogs|%mh;nYggMc`u6!$oj zUBSSu{FN0D5ci&uXN%Xiglr}!=6#Z#tgLd>M+83;R5yBDzRf}$A5@Av_Iu+bWHZ5J zD92kv87BkfV)4E>Dcfw!wQYNGUS4#}g4Qtj+ajEKI_AIU3ooa>aj_2LSK#Y$aQTi! zewTMzlNKBZWoQEB9bR5qL0ikg`r8|d>~yD7NbY_94RU&>jX4iluDqtf-!L_>qUaZI zzcJmYy{+i5*=R|eRfMmZw@YmA`rsh`-P^T!hZ3+OsjSO^1MLUx27Gyt|d|5%ewjP$(SXKbv;xBf0IAU)1$cVk1BKCx&9p+#n znuBF6!-TltoHFD(POp!qhtsUju>KjuTEN|XD!(_XJ!O|mD%{#KjPd?jcg!5b6f6qY zvlisiCKlZ`6h|Fh%pXlDuuGC=KKnzFy>;Vzq98&u6-=z3zBQPe+vtxi)|+o(Hq(SH zI=}m{Q+2fUGfODL0yb9$6wHz+j}c7k)22C*K{fWh-jYMP;dIr}mgx3chB^E%<%e5$ zhZTSh*H-QG1orNA2u*KJf*mL3)q8N~-Y#-`Ho_8Zl1zoEuu$LM#XBzeDvZht=dR?M0<=aD-lN1cB+q4zzC;rrZyK`(D4 z>)nX1p}5|?W%CcavmBEH8xw4M$;nf1m2f9z_cd#2G55_AJx4_mmQPbA7~eN8?nz*Z z^yBP1anBvzi06$^90b25j`UJYQdL566U9e=IolM!>HxXmv=(Cr&U;r-5qffprcYGT z+FniJm+B9bPn|3msYcZp61${dl(Mxu73Y zikYoVi&yRuBw9SZM?13)N6AZy)`rISG>}QB-y)kdxx_4B*eem?hsx#`j5LSV0mG$w zs`-Jq)EAkjnilVC7A&(2^d3E)e|VO>tZ1!y?y5s%^69qlL^rz@oa`)M5_~o-mjM?U=)&c8Bb|D^ z!wYYE;CtI>W3O9v0Y9$|C?E&mza=p2PsH(H)^lE~$%4U>&M9^X#e6{+%oWr8Z zIGQdTAW7ELJ(92)-D9DTn#8a;Nd0o4_snwO5(!-(!MzZ3M+4^RVC7o2m<#;TVf zLRjqaHrOhX3X|BtaY@U8wu;2C07#`=BEF9P4zB?F$;5W@?c`x}1BYQ+(C}?Y!x%?W z5FG=)H=zpsQ%OF25mfxMxIgvdsemT{?oa)=8t}P*`%^z|13Ut7f9l771$Y$T{?w13 z1$-9Z{?w0O033~rKlS4PP5TS7E5ebYD39d#rxXrl^!Wu0*da*4SGO`DEqpd#BCV=c z7r{#>M)VRuVIkZC7C!Rev+&tuc(7h1mgLkV7QP78OW-XN^pJQ-@sgUv;za3^YDh~H zskb~HRQW0Ze&kB>vF6gL=2D7Iw;hS0t%a-1?YUDAhfxFS`K{hig}U^0Lm|cLErv&c zhbf9GjXvq785F~7GmnsKDj#cJTJcD|u5#Ownxsd1L@cHuc*I>;QC(IBÆSx{I! z^=bh;1fO~gJzed87XY5>FM*d6)D>IHw1q`#qZ-MGxKYpQr;LvCgWIAJ;1> diff --git a/src/metal/Compute.metal b/src/metal/Compute.metal deleted file mode 100644 index d81e8a5..0000000 --- a/src/metal/Compute.metal +++ /dev/null @@ -1,323 +0,0 @@ -#include -#include -#include - -using namespace metal; - -using namespace raytracing; - - -enum class MaterialType -{ - BlinnPhong -}; - -struct GPUCamera -{ - packed_float3 position; - float f; - packed_float3 forward; - float S_O; - packed_float3 fogEmm; - float ks; - float A; - float ka; - float2 sensorSize; - uint width; - uint height; -}; - -struct SampleParams -{ - uint pass; - uint samplesPerPixel; - uint numDirectionalLights; - uint numPointLights; -}; - -struct Payload -{ - float3 rnd01; - float3 accumulatedRadiance = float3(0); - float3 accumulatedMaterial = float3(1); - uint depth = 0; - float emissive = 1; -}; - -struct ModelReference -{ - uint positionOffset = 0; - uint numPositions = 0; - uint indicesOffset = 0; - uint numIndices = 0; -}; - -struct HitInfo -{ - float t = numeric_limits::max(); - float3 position; - float3 normal; - // not entirely sure what that does - // its the normal being flipped based on some dot product - float3 normalLight; - float2 texCoords; -}; - -struct BRDF -{ - float3 albedo = float3(1, 1, 1); - float alpha = 1; - float3 specularColor = float3(1, 1, 1); - float shininess = 0.004; - float3 emissive = float3(0, 0, 0); - MaterialType materialType; - float3 evaluate(HitInfo hit, float3 viewDir, float3 lightDir, float3 lightColor) - { - float3 normal = hit.normal; - float diffuse = max(dot(normal, lightDir), 0.0f); - float3 h = normalize(lightDir + viewDir); - float specular = pow(clamp(dot(normal, h), 0.0f, 1.0f), shininess); - - return (albedo * diffuse * lightColor) + float3(0.03, 0.03, 0.03); - } -}; - - -struct PointLight -{ - float3 position = float3(0, 0, 0); - packed_float3 color = float3(1, 1, 1); - float attenuation = 1; -}; - -struct DirectionalLight -{ - float3 direction = float3(0, 1, 0); - float3 color = float3(1, 1, 1); - }; - -float3 rand01(uint3 x) -{ // pseudo-random number generator - for (int i = 3; i-- > 0;) - x = ((x >> 8U) ^ uint3(x.y, x.z, x.x)) * 1103515245U; - return float3(x) * (1.0f / float(0xffffffffU)); -} - -template -inline T interpolateVertexAttribute(constant T *attributes, - uint offset, - IndexType i0, - IndexType i1, - IndexType i2, - float2 uv) { - // Look up value for each vertex. - const T T0 = attributes[offset + i0]; - const T T1 = attributes[offset + i1]; - const T T2 = attributes[offset + i2]; - - // Compute the sum of the vertex attributes weighted by the barycentric coordinates. - // The barycentric coordinates sum to one. - return (1.0f - uv.x - uv.y) * T2 + uv.x * T0 + uv.y * T1; -} - -kernel void computeKernel( - uint2 threadId [[thread_position_in_grid]], - constant GPUCamera& camera, - constant SampleParams& sample, - constant packed_uint3* indexBuffer [[buffer(0)]], - constant packed_float3* positions [[buffer(1)]], - constant packed_float2* texCoords [[buffer(2)]], - constant packed_float3* normals [[buffer(3)]], - constant ModelReference* modelRefs [[buffer(4)]], - constant BRDF* materials [[buffer(5)]], - constant DirectionalLight* directionalLights [[buffer(6)]], - constant PointLight* pointLights [[buffer(7)]], - constant MTLAccelerationStructureInstanceDescriptor* instances [[buffer(8)]], - instance_acceleration_structure accelerationStructure [[buffer(9)]], - texture2d accumulator [[texture(0)]], - texture2d image [[texture(1)]] -) -{ - Payload payload; - ray cam; - cam.origin = camera.position; - cam.direction = normalize(camera.forward); - cam.max_distance = INFINITY; - float3 cx = - normalize(cross(cam.direction, abs(cam.direction.y) < 0.9 ? float3(0, 1, 0) : float3(0, 0, 1))), - cy = cross(cx, cam.direction); - const float2 sdim = camera.sensorSize; // sensor size (36 x 24 mm) - - float S_I = (camera.S_O * camera.f) / (camera.S_O - camera.f); - - //-- sample sensor - uint2 pix = threadId; - if(pix.x >= camera.width || pix.y >= camera.height) - return; - - payload.rnd01 = rand01(uint3(pix, sample.pass)); - float2 rnd2 = 2.0f * float2(payload.rnd01.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 = - ((float2(pix) + 0.5f * (0.5f + float2((sample.pass / 2) % 2, sample.pass % 2) + tent)) / float2(camera.width, camera.height) - - 0.5f) * - sdim; - float3 spos = cam.origin + cx * s.x + cy * s.y, lc = cam.origin + cam.direction * 0.035f; // sample on 3d sensor plane - cam.origin = lc; - cam.direction = normalize(lc - spos); // construct ray - - //-- setup lens - float3 lensP = lc; - float3 lensN = -cam.direction; - float3 lensX = cross(lensN, float3(0, 1, 0)); // the exact vector doesnt matter - float3 lensY = cross(lensN, lensX); - - float3 lensSample = lensP + payload.rnd01.x * camera.A * lensX + payload.rnd01.y * camera.A * lensY; - - float3 focalPoint = cam.origin + (camera.S_O + S_I) * cam.direction; - float t = dot(focalPoint - cam.origin, lensN) / dot(cam.direction, lensN); - float3 focus = cam.origin + t * cam.direction; - cam.origin = lensSample; - cam.direction = normalize(focus - lensSample); // TODO: Fix lens - - intersector i; - i.assume_geometry_type(geometry_type::triangle); - i.force_opacity(forced_opacity::opaque); - - typename intersector::result_type intersection; - while(payload.depth < 12) { - i.accept_any_intersection(false); - - intersection = i.intersect(cam, accelerationStructure, 0xff); - - if(intersection.type == intersection_type::none) - break; - - uint instanceId = intersection.instance_id; - constant ModelReference& ref = modelRefs[instanceId]; - - HitInfo info; - info.t = intersection.distance; - const auto indices = indexBuffer[ref.indicesOffset + intersection.primitive_id]; - info.position = interpolateVertexAttribute(positions, ref.positionOffset, indices.x, indices.y, indices.z, intersection.triangle_barycentric_coord); - info.texCoords = interpolateVertexAttribute(texCoords, ref.positionOffset, indices.x, indices.y, indices.z, intersection.triangle_barycentric_coord); - info.normal = normalize(interpolateVertexAttribute(normals, ref.positionOffset, indices.x, indices.y, indices.z, intersection.triangle_barycentric_coord)); - info.normalLight = dot(info.normal, cam.direction) < 0 ? info.normal : -info.normal; - - BRDF brdf; - brdf.albedo = float3(0, 1, 0); - - float p = max(max(brdf.albedo.x, brdf.albedo.y), brdf.albedo.z); - if (payload.depth > 5) - { - if (payload.rnd01.z >= p) - break; - else - payload.accumulatedMaterial /= p; - } - // emissive - payload.accumulatedRadiance += payload.accumulatedMaterial * brdf.emissive * payload.emissive; - payload.accumulatedMaterial *= brdf.albedo; - - // direct lighting - for (uint l = 0; l < sample.numDirectionalLights; ++l) - { - // if there is an intersection, the light is occluded so no lighting - ray shadowRay; - shadowRay.origin = info.position + info.normal * 1e-3f; - shadowRay.direction = -directionalLights[l].direction; - shadowRay.max_distance = INFINITY; - i.accept_any_intersection(true); - intersection = i.intersect(shadowRay, accelerationStructure, 0xff); - if(intersection.type == intersection_type::none) - { - payload.accumulatedRadiance += brdf.evaluate(info, -cam.direction, normalize(shadowRay.direction), directionalLights[l].color); - } - } - for (uint l = 0; l < sample.numPointLights; ++l) - { - float3 lightDir = pointLights[l].position - info.position; - ray shadowRay; - shadowRay.origin = info.position + info.normal * 1e-3f; - shadowRay.direction = lightDir; - shadowRay.max_distance = 1; - i.accept_any_intersection(true); - intersection = i.intersect(shadowRay, accelerationStructure, 0xff); - if (intersection.type == intersection_type::none) - { - float d = length(lightDir); - float illuminance = max(1 - d / pointLights[l].attenuation, 0.0f); - - payload.accumulatedRadiance += illuminance * brdf.evaluate(info, -cam.direction, normalize(lightDir), pointLights[l].color); - } - } - - // TODO: Next Event Estimation for mesh lights - - // indirect lighting - float r1 = 2 * M_PI_F * payload.rnd01.x; - float r2 = payload.rnd01.y; - float r2s = sqrt(r2); - float3 w = info.normalLight; - float3 u = normalize(cross(abs(w.x) > 0.1 ? float3(0, 1, 0) : float3(1, 0, 0), w)); - float3 v = cross(w, u); - cam.origin = info.position; - cam.direction = normalize(u * cos(r1) * r2s + v * sin(r1) * r2s + w * sqrt(1 - r2)); - payload.emissive = 0; - payload.depth++; - } - float resolver = float(sample.samplesPerPixel) / float(sample.pass+1); - float4 previous = float4(0); - if(sample.pass != 0) - { - previous = accumulator.read(threadId); - } - float4 result = previous + float4(payload.accumulatedRadiance / float(sample.samplesPerPixel), 0); - accumulator.write(result, threadId); - image.write(pow(max(result * resolver, 0), float4(0.45f)), threadId); -} - - -// Screen filling quad in normalized device coordinates. -constant float2 quadVertices[] = { - float2(-1, -1), - float2(-1, 1), - float2( 1, 1), - float2(-1, -1), - float2( 1, 1), - float2( 1, -1) -}; - -struct CopyVertexOut { - float4 position [[position]]; - float2 uv; -}; - -// Simple vertex shader that passes through NDC quad positions. -vertex CopyVertexOut copyVertex(unsigned short vid [[vertex_id]]) { - float2 position = quadVertices[vid]; - - CopyVertexOut out; - - out.position = float4(position, 0, 1); - out.uv = position * 0.5f + 0.5f; - - return out; -} - -// Simple fragment shader that copies a texture and applies a simple tonemapping function. -fragment float4 copyFragment(CopyVertexOut in [[stage_in]], - texture2d tex) -{ - constexpr sampler sam(min_filter::nearest, mag_filter::nearest, mip_filter::none); - - float3 color = tex.sample(sam, in.uv).xyz; - - // Apply a simple tonemapping function to reduce the dynamic range of the - // input image into a range which the screen can display. - color = color / (1.0f + color); - - return float4(color, 1.0f); -} - diff --git a/src/metal/Compute.metallib b/src/metal/Compute.metallib deleted file mode 100644 index 06e6ad876a9972747592c7393ad3146e8f74530a..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 19045 zcmeHv3s_TEw(!Y=h)HY)94wIu<9PK%(2!7zn^95Kv%@?^m5(-D)|8`$` zJl&U}%=U#02rTd}moHny!wM_QwbcgAQjM-cQ-O52~^elZfi~Zqy=XbWf^ugqBRl9v05IkR9J`MyUHTy$Jnh)@E4v-Z5 z()iDY7n~EvA$no(#)IVVz>}eHUYG*_)UgE|NxV@I(2erNAekjW7&eYB2FYXC1bEq| zxb5<8<`R}bmb!yh$C(s>hetDiTVNC|qG#+Fm(eefg{9!ar9m>CEnC|qj>`G3Th@g(>O zpTencv#WGPYl--EX285Q2@bwwRc7Gx{MTl&C7H~&dI>-*E{6c; z#0u%lPlv7O?B9=lHzPAB6N{S^E5MnoV1b-Xx8To7Y$QHSBFrqQQ`f0m;>c#2`L{`n znk7H}<=iu=8#%E}^8R2OCpMe@1Vv>-b|xNrG0C|y_Ly{1-*sW8gW1?N^_9Ah9*%1J z-9ud|uO-dLEpp~kCA~AJR>HqY@V}SvZhBeYD0w#pylW2jrkYFTdP@U6jz3z<$L zW12i_W!*&SB9P#8Kx>(I@25My-lngV}1b*&9k2 zH3?Qd&aNq8Je6Rra#%Omt;vV127AKw~*vG>2|s@fEL_FOR7_( zbt@umiUlo})Mm^41EzUL6p?j|d2Nam52K@XIX%Ma)M>Hk3az(C%2kqS&6a?d*;Pw7yk)w$*!1||eh0{jKhMeWDNkwN5~Ul7`QArhM77lf8@5={p_8Hv4yWs@;w#J^y9 z$tgV)&^Q%)78|A!w6Gp_`ImDr^|D8-oTc+=ei0oU&Z|8p??=reHa=K?X^m z=HMq2XAq1J6hfe^?M>`WZe2C@IjjF42u`zOsCdvVlwb(hMrJNlEqMneW3WKTdd`>C z!}4X32dOM$G8XUP!!*X0z>5Z)*g6MW56b{V zAbBr^EBEGJxi|L~P#8n%7J+IXC3WlGvbRXYz$kCG?jvQVND&C93VwT7KwJmRkNGWE z@vH2JTb&XiqhCNH*wxUMF~5LT4W%6HyT24QEyZ3H#@3ODEK0(BqpAREz;=X~a~9`) zrwo;ctWeBQOmqB*QHm2%MldWwmHl~10md{!oPZa@dN9SCiS}*`H-d^uQXmJdA$Ox> zJ57MJ4f3ySYZX^>Sp=WFolT~rX`nzJy9CQ{w50x+Cv7tZVdW^Gw3aEki;Qi1ho>~ zDZ3m2h?RXH)mI@*#Eqh^K2dQPZGlaZQs-{WFtRt+vh^)hoGTdr8&F}4yhKh$bf+uo zWUHuH5jI)vf>?XOK6j zc9CMqmU%}k^LiAi2N?5@Smwiq;Ta68QL&#YVb3jrNzI;Xv})~EEo|5c z?3$qfAiKe6&4-yTwCY-{8(T~>U~l-=%lqERyCz{2bfb;Rw_So`pQ5RfM&(c6G5P8H=gKAc8y(Putp4XvR00J5AoZqjQXM?1> z@n{gV)Vcxq0Zk%-t!pu@y#(8#lXqFk`%V()f_=KbZ(?s>)CFK5JA7V;Wx)YUWSu*u zQIYC#&p+f!?RG;O=X}Vpo9x^Q_^2ImWxD;(z014T4T8;lqA*OUbJc@GgN` zD&c?UWnK01`jxyPN!-P2Y7y+Ly?qm}iKDKFqq@XVpa%?yV~UC2!%n~85CeAoR8QU! z=zVDHCg3yMx1__Hp_aX2s9q1sTssJg zjmG9i0*qOh9lGdjLSf^C2kK+|Nn7+OIeSvJsKb_-Ih^hB z0bw&wWJGnnepD`F6)0!&-VMW6WiBPa*3_3)R%|FNQS0*crJFRxI(4~brna)Q!l2Qq zOLeo#)wTIWnwru=O@2XjaWUZgp@o$ddV{*ca6jBoqSL60*6Av%wLi%*329Q&5})B zhL91L$e=GR)66Q;R8;E9_4!4mdUZjWW^7imT5rfNR~t(5H5EV0L$dHSI=@(BC>)F9 zYh=E*QeSES0p^!hXPsCt90y3pXOu(YDM65#7KMft^Lm1;wxpFJv7 zTwJN!pw<Jz{u6X_29%L{qx1#6Z>QZ;+q8 z(H5%9O4n6T<7Tn)5e2kPr!J&^wQ9ZI*MoYcMyD*T)s$6KmuHshG=->L>arX}Tn{iM z8fw1ODyqwsY7qNK;gQB5o)SwdKol9&5cD?-31PIcG8jRM%1T`kHM$Us)fLrxO%dvP zb(!8@=HcREy~aRQ?CUx`5QfpMAI(_eOE0S|(v+>#z=+pW6l%&xQhJ|Y)nx^mqDpm{ zwnVMhY6`2%)H+3FS*1>2Qd&`30Ue|%FV*WyYc#+)j4q9?R9%(}V&IHS3KVW|xUO+_^|F767AnyEhV8FcCjJ@hfq2N5q*8+4_$m0A?nm#g*bffJZM zKKW-=S5W?&$Z!9)d7Xlkh`mX@|GH%Iol3P5LWR~)sIU!FB_j_|NvDewmZBV=mUoLPc{Znw1|?I-D&iv- zr)Pok(KMU8-lT{)V2NmCB(X!HJlIx-W{V|dtuUocJ`bV3n9EMg93oeHDGH@p-jfoZ zJt}*Hk)25XNUr|D2kxQ3lP~-EF-80Mx#r`C7USdRri&Ajs{A5_s&!XrsgWuclf7AD zd4V@z=Q#vH14t@cI_EhQX~xckE+nwUL7*F6#e!aSAR)b`4NA=Uid;S9Ytr&I*TQyM z8re>~N{{w9+8AkVjD&W28dZLphq%trT&CJnv|u~y?+(@gw$FSfC~^K(iezNk0YzHL zAG|yuynDKb*vHUp_hF}xV83r>j!P`M+QYl#IGTa_3e>@FO1Pd*DoLAplS-cH~VUcV%6x_(dO zBy!8z&eQ#Cubt?>a{9vgtiYAAZ*(N3PCO`icmI1G``?-UC)we{2j;%Jg1)nR>WpP4 zHP_za&f5CcwZCz<&HCn(@``5lq`X%qtqffHynglc=k?PIBPU6ybLBU{ZkPmX#JC_` z0_fxHO?lcD&;i|``KI=`%~wRnTDfrk2Ic~&ig0KKdhUdzODf)1N>HEq{YZ)Ce+37o z9Yzy-V*;h!*IGd{MY?UCcDFm_01am0f@U{;Js1k2^Ke#1i2Bv&cYr$b>Vcy(pF*m3 z(a<4mpOO#b5Lm)afyIVwEw*QVn`Pc^Q1TfI9G3Z@@Y~&JMN<;pxCR~r3Romn*J)z!9liS0M`x;2>Sj4%F z#Tz9hU;XZ~%klelyV#a}EJA2YWw$OnJ-rd@Q887G7*j5|a||a1%QLc zV2hNlKe_s+$hh8rTT$#%H9F>ewfl#52}80AGZXk`0yhuNrJ`mU4ggkRAOg)c5_so_ zWYZ}laKWPJ12-EcFFWwmy3?6|LXpVzwWYPsy#GSOpa1+6iU1na-1Xr;mVhZArbBc# zibSl>cK^lte$tU=dY7U|DoZb5>xIx(gq@B#dT=KQqb~KUy+8dK?zdpXX@eVMx#%-=MHJ= zec3ZXD@l<43_3>j9ZSrIE&CZdMMUyIaOwl0a~}vzejs$t1EEO|geE=^Dw|qaUQw#8 z1ZBlQAiH8!DbnN##q{Y!ZiZZ*|5)y-)KPU2rKGa4nqOI2O*tsQmD*Y5)p~XDqpVG1c z&56+Z5!y|J_7p-}m?w6y9@{Hi$(VDx;HV>_y>-I$hR?~Zf6SA1u%7HV`dsTA<(`?3 za^5yYRO~CL<({ei?RZ9+gSB?Auz)e=V!=@&vRyoJdc#?A>)||c&#{er1y5d@|0sU2 zDmv19rq|(VDC=UC9cfs9fVF;a!;@wB*9C5J;^BQ0BUxYeY<<61qJ0m=SE0D}Y<+?+jgm+zT7RN$6GzYmnnXL}uO_wz7V;Pyry366-|ab#fY`vct1 z2DW}S!2O_C(zdvA_pt`&oaBTzq>=9gM|NC~c*_;>ht`NlI-+yz-nIx&@0ok6Kp}{X zY(CS%J!{1vln-AE`G+!QUvonm;oI49tmDCERbNb%F+?sYP>IC&gxgYf zLa;~LC?|Lp4lVo|ja5I?BxXA9W={#`EM``cL&ptW;dZ$Ggq^COLfkgh zA+V)~$cH|a9+4073UPcsm1&OSw=*v`DFnCmoGffy_Kb^7*7(~A%)a#FBy+!{Kuc8~ z%q|FK#<7l3Tou-D;(h_)=2G%;{$aSWa=>+1*`hKs-Xu z#ilv|;P$6C*hA)?;wm2UG9Ok72sc+jvi5UK6C1?`>J}K>Jt6aH6T_Itv~GdjHN)Z! zuyCZ;`QP^hGjfdek1GYJW1Rf`OgFj6I6Xtk$+gkNGZ;Y$H2Rcl44#bhlY>8n`Lt_c zx7J-1Q6TwPB5>JERdqz((D8)!2BRz0c85@Sc{5D%`5@Aa!3pkQdAi`X^kb!P6j5}( z!Dvc1-kISt-6req_{)n1vW&68bpV#yr9{1I=(vcH&fFkvaCo}3nW~Vw%^n%C8c$V;bw7U*J8R`ZWbBx693SG54rh8JVI1e3513#YT7$tLq2=dcFa{R{+m90pviGPk zBRHZEKeiEu?H!TgG3HfMxA=Cp;9^q-8mgJb2%F?$`f-B!hO~hU8a&0d)ON9-g$V>B zCjGc_BJCNDj9GIkLKtwdi6%7O5z?o7SkWdJg4Y%~Vxa1L4=dR;v?th@{)%Fr5r%6C z^NzGzP{*oJ37HquM+g3sA$2aBiy8@wMSDLI+Gs(vA3v40WGQ1U99QWO>tE!=ZDAK} zsL%;r?nmriBJhjPW7-Yf@K$g~Hw>VN%ceyHwDyucI!bE98xI8rn55?gjzFQIOA?#d z8Ha|jqGPhkSlknNv8n%hJq#O}@_MgHmvN3fS~zt|8#i};(VpG*WM#1Jc>rn^gOiIS zA3MujHkt{pp@-z{jY8ED+)RsUd~QX@@r`!e#@yEzXULcWe^0$LMH~`qcqis3>5Br) z%{Ber=$hgFg1INqEq##;+Y?xC@0^uZVBD_CO~4z+;oYiC|{eQ`nRYxysPxh+Fyw%&mzcyK7dO8h>83JVDr)+VBtj>$XsqOk7ZBcC{1=&EZ$; z`f|2~x2hypgbj|YZ6?CY2+d~-F z7f`(t+-3I%eB$$N1JPZY0bW!qp6mcu-`Pq z4T)95g}eGrnO8U9r`b8Vi-x7OjKP;DwXmi|->kjicu~_4{|x!sBfwiKc4xJDy(3mT zr=9nv={fd7M@sw3-bJr-{w<`P8_-lh8)^~?vT?I1=GrGV+AbKKgYv)5Fdwe_?S}(P zJEo;vO0IyvRdcKN&;IFR^7k3D_XN@;=LJr#&?8ema(DCDvntHLJn1EAuonF*b3tOa zVEDQBRd(Cv*Ev%^y9w3FM(6CC-*Np*Po*lgIHtol$4$5#ds~QKYEX>q+$UUFBRjX34x3>XGiwQKiAUbz+zUOD1gng%;W*SV=5rRmPFc^-zvC ztq?!{HZ?o3T39_c`ieI6`jC%d`J(fW9^ zF>>%EELDlsciSqii_mjap~Cq(XB17n z_R1R8mw`%~d9lkmxpQ-uk<1Wqh6!sHrz4)}GV*J>MavhHuG8E_v0PHMv>U<)l`cJ=`azI>p4DKGySB<3NG~hg3fCf1%oy#R;3lbP?Pq zKdLuJrM0H-*mJ55Yne2_n$}5rGb>*zEP1(ZD=VD-!&H#ojhG_b`_<+vLi(L0;v_oA zH+z?W$&li&N5<~KcHS!4*7vk96-3;#xK4|A`Byp%r)z57%Ogt@Uq-x6K9j}S0o9Tm zPpS2ELdhiMi<1VE2v#xJV{4xCc5wTX4~BKABfHDXrfzd5+yBzLGQnC-yu0t!65e0P zio&I|sDq~S$9Bbq(Z5V7V;tT(^LPUDgQ{MZE^^JbmH4xlC<>xK2#m)mBB?4d-bB&S zoym|yvo{i#_h~S8Y4Sl)MeyzsoD*@`O^OBf?|WCa9d|`%KUFsMi%W34wBTUaH&q?% z)6r|T{eYjpq~sR?OhRFWH zrKZV?pC)eXdo}O*xP^irdsiO%C^*Ij#V_{T_?b}LpzC7C#;kd6p0Bv|ZRUGb9(G2| znr-?(!=)B}-NJ*Mo3W>k?XpyYe7#$Dh z%nQLa;IL8;VF;r+0gjL;U&;nJR6+>o90{E-<-v;{6o$Zy!yP7iwnL`K7sa>`q8B~> zd=y9c{>Pt>LgYGw#rw)2Ctn0ZA%f!op8|1I)*t%esQ?Rw*B|=fh*#ue;1B(9bncIQ zLj9p1z7b#(;q`}pxE)}@@cKhPycJ-GM}O#tzYQ?Nr$6+=od830{GlIy3Sg7r^@o1= z4**Mm*B|=fF(6|B!2O{gK6?c25B>0TfTMc-p&zaUcqG96p&!0}1nv+0a6rW#hp|x! zA&iZ3J=1@m1Hbz3uX9(GUZbs9rH zl{s=RjOv05Xt)tJ%qpG*Rm`fMRXV03itI?uMU{qGHHl-{(CMpCxiR=-0i!o%{>v-4 zhMEPpY3hn~Wt#jVwLy()M;ujDmzBZQ@7;^F{NjH_NrbF1m*2oya#DVcEM=w|U7XJ>S3^~`vgBEkSrXKN zW1g}8)4kBIsGt8M|LN;6$Ny9Q(_udU>9$e->6QP1|MbeS{?kXhfxj(;Z+^;Rvq$}> z!;t@U*ogo1O5{JSF=H>%gjoV?4tUU-C5tY98eEDgi%LD>&08 z8?76_vDN|ZoosLxHAP{sM1imDIb+m`KIB4j&Ux2Q^;>aNH}ZEiGE|h)Y*o9}paRbk za8>oW(OST(NdM2>Xq{2$is7HS&bC+$DeMh!r%=Lp23$J9H}CHbX!4-_JV(J zpXfd>T-!#6RV%bsmoSjCF|XMBi{2ZzWT8pFbANg2XNgLku^+ENnxc2tNUaA^|k9%y5iGdXX> zUdAv2;%Gy`FdMdoORxVb0K?uPDQ86R&kGylpSN<6l3uUmU0!9R2dp6Cu&yjDp>Y!2 z*IfJ%ZggySI`imok#pMwfWP{p=)$_}IBUb9~Nuzxw9=>YMjJ z**DJvlj4_r^V-#KGRFGllfp9!VN3nh#}2vQ{h#{S`CYbBg-$I>hpN3u0iETa7&A-^(B=$LrrNB zhBYigOCbh6K@L0)Uz%igaN_VFso4=0hd&2_2vRb$BOX#Rvyqv|S>yafVldBj4|4y) zaW6KyIPTw|pLOVGLmm8BZSeC#ql;sbxjVMd!ad}S=y1D>LrYZjcmyelVf?TqpK(}^ z$afh&AaX~p<6cDTuTajgZ-}%?MDer+t z#o-kVDRxV0_9XDCIQn+ZwzoNz%>*C&-|jadg}x7E$|=7IWZ$(FW8I&~6WmE#7h$F< zfshLh1jx~$D?P~Pu)s#n40~vqSBM#Z>ug6$2IV_|Y%a=oAgB&BAq@Kz45t71Y8=&D zd{!{%YEbis;p-I6=0Wx|*x`A#!aFKn1?0fW}^5{V!1^4M6lzoY8bKvb2l%M}tj#O4vtuA^LCN=WX zy!(Odf7$36YwuvZU$%E>VVQU$S?Im*TL1@t6g^7i56F!*S!C#%Ai{$vc^A z;Zw;E7Yu(Uq)BWlW&_bC%Fr{rgKeGJFUCI9GMrCuStovFb3`1*3Kpuj5iN0;Ma}G! zo=AHuW#CwC7@MB!<;HP7m2^rtae_~ABznqTDuFQa)`-1S!cg7)s_Or}y<}sgI{*JL znJi-X$S^n3(TAo2e7a_wGv*UDVJ0?%{G%&@siM8fZM=?97(Fa!1Qw#CJ%joQ*;!}~ zgPqkP;bcnH5*z*3@5la|zaK;P)(+BdZ|wqm>!(%i~s z&R-d}ZIenUy1&|0zuHy*lkF-xjM87Ss}fkdXQIJ8*0x$vBzdlj*p{@Qp(8seIrCTh z>i@>Rf)87YG&(q`*1=h}-G2PN?JivhlK@fW|} z7(;fO?>iiX0biF<|5^mj%pYu$C_#*w*m;`=buOUZUshA?tc8?11|(PDmHR; z03V-#L5N|;R#OWu*m;9s;N>Qo2L*BX!&390Fb@9|0s>Mp^I!s`WM(@v4yQ}p5*Ao* zl4g^NmB?8q0gKKZE+~<>wXA>$k7T#WguvQ11SXzD;Ey6af`mj)p@UT>VeD(6MRts} z`XW1MVV?+wT%5qj?9hNkJF7pjH`8`4i#YJedCm!;Av&_z#Wu;u15x-a;T4IyhqYM} zl!zC2S#w#K&m05?{JUQE;BoJBw7*I_9<{KUf@`wvnyS7cMJ_p|nL zx4o6~?fPxA8s7$|ej`%#|JU1#QP2jzy@<>h<_Y6kE_lbkL?cq@Wc6Me&U{ZV7$;T2 zzcTST$~SeFU~|^c$eO)0FZ#(jOu?y1zvaDuW&rQutKLcdv-*1)lp5C_f)2Z2mnH c7FmA&e@%3EHv7LD9lb34FPnuU9AMbL1I27h +using namespace metal; + +// --- Structs (matching C++ and MetalRenderer.mm) --- + +struct GPUCamera { + float3 cameraPosition; + float f; + float3 cameraForward; + float S_O; + float3 fogEmm; + float ks; + float A; + float ka; + float2 sensorSize; + uint width; + uint height; +}; + +struct MaterialParameter { + float4 albedo_alpha; // xyz: albedo, w: alpha + float4 specularColor_sh; // xyz: specularColor, w: shininess + float4 emissive_type; // xyz: emissive, w: materialType (as float) + + float3 shade(float3 normal, float3 viewDir, float3 lightDir, float3 lightColor) const { + float3 albedo = albedo_alpha.xyz; + float shininess = specularColor_sh.w; + float diffuse = max(dot(normal, lightDir), 0.0); + float3 h = normalize(lightDir + viewDir); + float specular = pow(clamp(dot(normal, h), 0.0, 1.0), shininess); + + return (albedo * diffuse * lightColor) + (specularColor_sh.xyz * specular * lightColor); + } +}; + +struct ModelReference { + uint32_t positionOffset; + uint32_t numPositions; + uint32_t indicesOffset; + uint32_t numIndices; + uint32_t materialIndex; +}; + +struct PointLight { + float3 position; + float pad; + float3 color; + float attenuation; +}; + +struct DirectionalLight { + float3 direction; + float pad; + float3 color; + float pad1; +}; + +struct SampleParams { + uint pass; + uint samplesPerPixel; + uint numDirectionalLights; + uint numPointLights; + uint numModels; +}; + +// --- Ray-triangle intersection (Möller-Trumbore) --- + +// Returns t (distance along ray) if hit, or FLT_MAX if no intersection. +// Writes barycentric coordinates into *bcwOut when it hits. +static float rayTriangle(float3 ro, float3 rd, float3 v0, float3 v1, float3 v2, + thread float* tOut, thread float2* bcwOut) { + const float EPSILON = 1e-8; + float3 edge1 = v1 - v0; + float3 edge2 = v2 - v0; + float3 h = cross(rd, edge2); + float a = dot(edge1, h); + if (a > -EPSILON && a < EPSILON) + return FLT_MAX; // ray parallel to triangle + float f = 1.0 / a; + float3 s = ro - v0; + float u = f * dot(s, h); + if (u < 0.0 || u > 1.0) + return FLT_MAX; + float3 q = cross(s, edge1); + float v = f * dot(rd, q); + if (v < 0.0 || u + v > 1.0) + return FLT_MAX; + // t must be positive and within ray bounds + float tt = f * dot(edge2, q); + if (tt < EPSILON) + return FLT_MAX; + *tOut = tt; + *bcwOut = float2(u, v); + return tt;} + +// Test ray against all triangles of all models. Returns closest hit info or FLT_MAX. +struct HitInfo { + float t; + uint modelIndex; + uint primIndex; + float2 bary; +}; + +static HitInfo intersectAll( + device const uint32_t* indexBuf, + device const float* posBuf, + device const ModelReference* models, + uint numModels, + float3 ro, + float3 rd) +{ + HitInfo hit = { FLT_MAX, 0u, 0u, float2(0.0) }; + + for (uint m = 0; m < numModels; ++m) { + uint idxOff = models[m].indicesOffset; + uint vtxOff = models[m].positionOffset; + uint triCount = models[m].numIndices / 3; + for (uint t = 0u; t < triCount; ++t) { + uint i0 = idxOff + 3 * t + 0; + uint i1 = idxOff + 3 * t + 1; + uint i2 = idxOff + 3 * t + 2; + + float3 v0 = float3(posBuf[vtxOff + 3*i0], posBuf[vtxOff + 3*i0+1], posBuf[vtxOff + 3*i0+2]); + float3 v1 = float3(posBuf[vtxOff + 3*i1], posBuf[vtxOff + 3*i1+1], posBuf[vtxOff + 3*i1+2]); + float3 v2 = float3(posBuf[vtxOff + 3*i2], posBuf[vtxOff + 3*i2+1], posBuf[vtxOff + 3*i2+2]); + + float hitT; + float2 hitBC; + float tt = rayTriangle(ro, rd, v0, v1, v2, &hitT, &hitBC); + if (tt < hit.t) { + hit.t = tt; + hit.modelIndex = m; + hit.primIndex = t; + hit.bary = hitBC; + } + } + } + return hit; +} + +// Shadow variant — returns true if any triangle blocks the ray within distance 'dist'. +static bool isBlocked( + device const uint32_t* indexBuf, + device const float* posBuf, + device const ModelReference* models, + uint numModels, + float3 ro, + float3 rd, + float dist) +{ + for (uint m = 0; m < numModels; ++m) { + uint idxOff = models[m].indicesOffset; + uint vtxOff = models[m].positionOffset; + uint triCount = models[m].numIndices / 3; + for (uint t = 0u; t < triCount; ++t) { + uint i0 = idxOff + 3 * t + 0; + uint i1 = idxOff + 3 * t + 1; + uint i2 = idxOff + 3 * t + 2; + + float3 v0 = float3(posBuf[vtxOff + 3*i0], posBuf[vtxOff + 3*i0+1], posBuf[vtxOff + 3*i0+2]); + float3 v1 = float3(posBuf[vtxOff + 3*i1], posBuf[vtxOff + 3*i1+1], posBuf[vtxOff + 3*i1+2]); + float3 v2 = float3(posBuf[vtxOff + 3*i2], posBuf[vtxOff + 3*i2+1], posBuf[vtxOff + 3*i2+2]); + + float hitT; + float2 hitBC; + float tt = rayTriangle(ro, rd, v0, v1, v2, &hitT, &hitBC); + if (tt < dist) return true; + } + } + return false; +} + +// --- Helpers --- + +float3 rand01(uint3 x) { + for (int i = 3; i > 0; --i) { + x = ((x >> 8u) ^ x.yzx) * 1103515245u; + } + return float3(x) * (1.0 / 4294967295.0); // 1/0xFFFFFFFF +} + +// --- Kernels --- + +kernel void computeKernel( + device uint32_t* indexBuffer [[ buffer(0) ]], + device float* positions [[ buffer(1) ]], + device float2* texCoords [[ buffer(2) ]], + device float3* normals [[ buffer(3) ]], + device ModelReference* modelData [[ buffer(4) ]], + device MaterialParameter* materialData [[ buffer(5) ]], + device DirectionalLight* directionalLights [[ buffer(6) ]], + device PointLight* pointLights [[ buffer(7) ]], + device float* instanceBuffer_dummy [[ buffer(8) ]], + texture2d accumulator [[ texture(0) ]], + texture2d image [[ texture(1) ]], + constant GPUCamera& gpuCam [[ buffer(10) ]], + constant SampleParams& pSamps [[ buffer(11) ]], + uint2 threadId [[thread_position_in_grid]] +) { + if (threadId.x >= gpuCam.width || threadId.y >= gpuCam.height) + return; + + uint pass = pSamps.pass; + uint spp = pSamps.samplesPerPixel; + if (pass >= spp) + return; + + // -- Camera setup -- + float3 camPos = gpuCam.cameraPosition; + float3 camFwd = gpuCam.cameraForward; + float S_O = gpuCam.S_O; + float2 sdim = gpuCam.sensorSize; + + float3 cx = -normalize(cross(camFwd, abs(camFwd.y) < 0.9 ? float3(0,1,0) : float3(0,0,1))); + float3 cy = cross(camFwd, cx); + + float S_I = (S_O * gpuCam.f) / (S_O - gpuCam.f); + + // -- Sample sensor with tent filter + dither -- + float3 rnd = rand01(uint3(threadId.x, threadId.y, pass)); + float2 rnd2 = 2.0 * float2(rnd.xy); // tent filter + 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 = ((float2(threadId) + 0.5 * (0.5 + float2((pass / 2) % 2, pass % 2) + tent)) + / float2(gpuCam.width, gpuCam.height) - 0.5) * sdim; + + float3 lc = camPos + camFwd * 0.035; // sample on 3d sensor plane + [[maybe_unused]] + float3 spos = camPos + cx * s.x + cy * s.y; + + // -- Lens (Aperture) / DOF -- + float3 lensN = -camFwd; + float3 lensX = cross(lensN, float3(0,1,0)); + float3 lensY = cross(lensN, lensX); + float2 rndL = rand01(uint3(threadId.x, threadId.y, pass + 100)).xy; + float3 lensSample = lc + (rndL.x - 0.5) * gpuCam.A * lensX + (rndL.y - 0.5) * gpuCam.A * lensY; + + float3 focalPoint = camPos + (S_O + S_I) * camFwd; + + float3 rayOrg = lensSample; + float3 rayDirF = normalize(focalPoint - lensSample); + + // -- Path Tracing Loop -- + float3 accumulatedRadiance = float3(0.0); + float3 throughput = float3(1.0); + + for (int bounce = 0; bounce < 4; ++bounce) { + // Manual ray-triangle intersection against all scene triangles + HitInfo hit = intersectAll(indexBuffer, positions, modelData, pSamps.numModels, rayOrg, rayDirF); + + if (hit.t < FLT_MAX) { + uint instanceIndex = hit.modelIndex; + uint primitiveIndex = hit.primIndex; + float2 bcw = hit.bary; + float3 bary = float3(1.0 - bcw.x - bcw.y, bcw.x, bcw.y); + + ModelReference m = modelData[instanceIndex]; + uint idxOff = m.indicesOffset; + uint vtxOff = m.positionOffset; + + uint v0_idx = vtxOff + indexBuffer[idxOff + 3*primitiveIndex + 0]; + uint v1_idx = vtxOff + indexBuffer[idxOff + 3*primitiveIndex + 1]; + uint v2_idx = vtxOff + indexBuffer[idxOff + 3*primitiveIndex + 2]; + + // Interpolate position, normal, texCoords (barycentric) + float3 p0 = float3(positions[3*v0_idx], positions[3*v0_idx+1], positions[3*v0_idx+2]); + [[maybe_unused]] float2 t0 = texCoords[v0_idx]; + float3 n0 = normals[v0_idx]; + + float3 p1 = float3(positions[3*v1_idx], positions[3*v1_idx+1], positions[3*v1_idx+2]); + [[maybe_unused]] float2 t1 = texCoords[v1_idx]; + float3 n1 = normals[v1_idx]; + + float3 p2 = float3(positions[3*v2_idx], positions[3*v2_idx+1], positions[3*v2_idx+2]); + [[maybe_unused]] float2 t2 = texCoords[v2_idx]; + float3 n2 = normals[v2_idx]; + + float3 pos = p0 * bary.x + p1 * bary.y + p2 * bary.z; + // tex = t0*bary.x + t1*bary.y + t2*bary.z; // for texture lookups + float3 norm = normalize(n0 * bary.x + n1 * bary.y + n2 * bary.z); + + MaterialParameter mat = materialData[m.materialIndex]; + + // Emissive contribution + accumulatedRadiance += throughput * mat.emissive_type.xyz; + + // --- Direct Lighting (NEE) --- + float3 directLight = float3(0); + + for (uint j = 0; j < pSamps.numDirectionalLights; ++j) { + float3 lDir = -directionalLights[j].direction; + if (!isBlocked(indexBuffer, positions, modelData, pSamps.numModels, + pos + norm * 0.001, lDir, FLT_MAX)) { + directLight += mat.shade(norm, -rayDirF, lDir, directionalLights[j].color); + } + } + + for (uint j = 0; j < pSamps.numPointLights; ++j) { + float3 lVec = pointLights[j].position - pos; + float3 lDir = normalize(lVec); + float dist = length(lVec); + if (!isBlocked(indexBuffer, positions, modelData, pSamps.numModels, + pos + norm * 0.001, lDir, dist)) { + directLight += mat.shade(norm, -rayDirF, lDir, pointLights[j].color) + * (1.0 / (dist * dist + 1.0)); + } + } + accumulatedRadiance += throughput * directLight; + + // --- Indirect Lighting (cosine-weighted hemisphere sampling) --- + float3 rnd_ind = rand01(uint3(threadId.x, threadId.y, pass + bounce + 200)); + float r1 = 2.0 * M_PI_F * rnd_ind.x; + float r2 = rnd_ind.y; + float r2s = sqrt(r2); + + float3 w = norm; + float3 u = normalize(cross(abs(w.x) > 0.1 ? float3(0,1,0) : float3(1,0,0), w)); + float3 v = cross(w, u); + float3 nextDir = normalize(u * cos(r1) * r2s + v * sin(r1) * r2s + w * sqrt(1.0 - r2)); + + throughput *= mat.albedo_alpha.xyz; + + rayOrg = pos + norm * 0.001; + rayDirF = nextDir; + + if (length(throughput) < 0.01) break; + } else { + // No hit — sky color + accumulatedRadiance += throughput * float3(0.05, 0.05, 0.1); + break; + } + } + + image.write(float4(accumulatedRadiance, 1.0), threadId); +} + +// --- Quad rendering shaders (pass-through) --- + +struct VertexOut { + float4 position [[position]]; + float2 uv; +}; + +vertex VertexOut copyVertex(uint vid [[vertex_id]]) { + VertexOut out; + float2 pos = float2((float)((vid << 1) & 2), (float)(vid & 2)); + out.position = float4(pos * 2.0 - 1.0, 0.0, 1.0); + out.uv = pos; + return out; +} + +fragment float4 copyFragment(VertexOut in [[stage_in]], + texture2d tex [[texture(0)]]) { + sampler s(mag_filter::linear, min_filter::linear); + return tex.sample(s, in.uv); +} diff --git a/src/metal/MetalRenderer.h b/src/metal/MetalRenderer.h index 2d1746e..03cecb7 100644 --- a/src/metal/MetalRenderer.h +++ b/src/metal/MetalRenderer.h @@ -28,6 +28,7 @@ struct SampleParams uint samplesPerPixel; uint numDirectionalLights; uint numPointLights; + uint numModels; }; class MetalRenderer : public Renderer diff --git a/src/metal/MetalRenderer.mm b/src/metal/MetalRenderer.mm index 96058e5..bee58e5 100644 --- a/src/metal/MetalRenderer.mm +++ b/src/metal/MetalRenderer.mm @@ -2,14 +2,23 @@ #include "metal/MetalScene.h" #include "scene/Renderer.h" #include "util/Camera.h" +#include #include +#include +#include #include #include #include +#include +#include +#include +#include +#include - NSWindow* window; +NSWindow* window; CAMetalLayer* metalLayer; id drawable; +id texToDraw; id device; id library; @@ -18,38 +27,49 @@ id function; id computePipeline; id accumulator = nullptr; id resultTexture = nullptr; - + +static std::mutex g_metal_mtx; + MTLRenderPassDescriptor* renderPass; id renderEncoder; id renderCmd; id pipelineState; -static void glfw_error_callback(int error, const char* description) -{ - fprintf(stderr, "Glfw Error %d: %s\n", error, description); -} +static void glfw_error_callback(int error, const char* description) { fprintf(stderr, "Glfw Error %d: %s\n", error, description); } MetalRenderer::MetalRenderer() { width = 1920; height = 1080; + texToDraw = nullptr; device = MTLCreateSystemDefaultDevice(); - - library = [device newDefaultLibrary]; - if (!library) { - NSError* error = nil; - NSURL* url = [NSURL fileURLWithPath:@"../src/metal/Compute.metallib"]; - library = [device newLibraryWithURL:url error:&error]; - if (!library) { - fprintf(stderr, "Failed to load library from %s: %s\n", [url path], [[error localizedDescription] UTF8String]); + + NSError* error = nil; + MTLCompileOptions* compileOptions = [[MTLCompileOptions alloc] init]; + std::string shaderSource; + { + // Note: Adjust path based on where you run the binary from + std::ifstream shaderFile("../src/metal/ComputeKernel.metal"); + if (!shaderFile.is_open()) + { + std::cerr << "Failed to open shader file!" << std::endl; + return; } + std::stringstream buffer; + buffer << shaderFile.rdbuf(); + shaderSource = buffer.str(); } + library = [device newLibraryWithSource:[NSString stringWithUTF8String:shaderSource.c_str()] options:compileOptions error:&error]; + if(error) + { + std::cerr << "Failed to compile shader: " << [[error localizedDescription] UTF8String] << std::endl; + return; + } queue = [device newCommandQueue]; scene = new MetalScene(device, queue); function = [library newFunctionWithName:@"computeKernel"]; - NSError* error; computePipeline = [device newComputePipelineStateWithFunction:function error:&error]; - + IMGUI_CHECKVERSION(); ImGui::CreateContext(); ImGuiIO& io = ImGui::GetIO(); @@ -61,10 +81,10 @@ MetalRenderer::MetalRenderer() glfwGetMonitorContentScale(glfwGetPrimaryMonitor(), &xscale, &yscale); glfwWindowHint(GLFW_CLIENT_API, GLFW_NO_API); handle = glfwCreateWindow(width / xscale, height / yscale, "RayTracer", nullptr, nullptr); - + ImGui_ImplGlfw_InitForOpenGL(handle, true); ImGui_ImplMetal_Init(device); - + NSWindow* cocoaWindow = glfwGetCocoaWindow(handle); metalLayer = [CAMetalLayer layer]; metalLayer.device = device; @@ -72,22 +92,30 @@ MetalRenderer::MetalRenderer() [[cocoaWindow contentView] setLayer:metalLayer]; [[cocoaWindow contentView] setWantsLayer:true]; renderPass = [[MTLRenderPassDescriptor alloc] init]; - - MTLRenderPipelineDescriptor *renderDescriptor = [[MTLRenderPipelineDescriptor alloc] init]; - + + MTLRenderPipelineDescriptor* renderDescriptor = [[MTLRenderPipelineDescriptor alloc] init]; + renderDescriptor.vertexFunction = [library newFunctionWithName:@"copyVertex"]; renderDescriptor.fragmentFunction = [library newFunctionWithName:@"copyFragment"]; - + renderDescriptor.colorAttachments[0].pixelFormat = MTLPixelFormatBGRA8Unorm; - + pipelineState = [device newRenderPipelineStateWithDescriptor:renderDescriptor error:&error]; - + + // Create persistent compute-result textures once. + MTLTextureDescriptor* texDesc = [[MTLTextureDescriptor alloc] init]; + [texDesc setWidth:1920]; + [texDesc setHeight:1080]; + [texDesc setPixelFormat:MTLPixelFormatRGBA32Float]; + [texDesc setUsage:MTLTextureUsageShaderWrite | MTLTextureUsageShaderRead]; + resultTexture = [device newTextureWithDescriptor:texDesc]; + accumulator = [device newTextureWithDescriptor:texDesc]; + [texDesc release]; + [renderDescriptor release]; } -MetalRenderer::~MetalRenderer() { - [renderPass release]; -} +MetalRenderer::~MetalRenderer() { [renderPass release]; } void MetalRenderer::addPointLight(PointLight point) { scene->addPointLight(point); } void MetalRenderer::addDirectionalLight(DirectionalLight dir) { scene->addDirectionalLight(dir); } @@ -97,31 +125,42 @@ void MetalRenderer::generate() { scene->generate(); } void MetalRenderer::beginFrame() { - @autoreleasepool { - - glfwPollEvents(); - int w, h; - glfwGetFramebufferSize(handle, &w, &h); - framebufferWidth = width; - framebufferHeight = height; - metalLayer.drawableSize = CGSizeMake(framebufferWidth, framebufferHeight); - drawable = [metalLayer nextDrawable]; + @autoreleasepool + { + + glfwPollEvents(); + int w, h; + glfwGetFramebufferSize(handle, &w, &h); + framebufferWidth = w; + framebufferHeight = h; + metalLayer.drawableSize = CGSizeMake(framebufferWidth, framebufferHeight); + drawable = [metalLayer nextDrawable]; renderCmd = [queue commandBuffer]; renderPass.colorAttachments[0].clearColor = MTLClearColorMake(0, 0, 0, 0); renderPass.colorAttachments[0].texture = drawable.texture; renderPass.colorAttachments[0].loadAction = MTLLoadActionClear; renderPass.colorAttachments[0].storeAction = MTLStoreActionStore; ImGui_ImplMetal_NewFrame(renderPass); + ; ImGui_ImplGlfw_NewFrame(); ImGui::NewFrame(); renderEncoder = [renderCmd renderCommandEncoderWithDescriptor:renderPass]; [renderEncoder setRenderPipelineState:pipelineState]; - [renderEncoder setFragmentTexture:resultTexture atIndex:0]; + id tex; + { + std::lock_guard lock(g_metal_mtx); + tex = resultTexture; + if (tex) + [tex retain]; + } + texToDraw = tex; + + [renderEncoder setFragmentTexture:texToDraw atIndex:0]; // Draw a quad which fills the screen. [renderEncoder drawPrimitives:MTLPrimitiveTypeTriangle vertexStart:0 vertexCount:6]; - + [renderEncoder retain]; [renderCmd retain]; [drawable retain]; @@ -130,7 +169,8 @@ void MetalRenderer::beginFrame() void MetalRenderer::update() { - @autoreleasepool { + @autoreleasepool + { ImGui::Render(); ImGui_ImplMetal_RenderDrawData(ImGui::GetDrawData(), renderCmd, renderEncoder); [renderEncoder endEncoding]; @@ -139,6 +179,12 @@ void MetalRenderer::update() [renderCmd commit]; [renderCmd release]; [drawable release]; + + if (texToDraw) + { + [texToDraw release]; + texToDraw = nullptr; + } } } @@ -146,43 +192,35 @@ void MetalRenderer::render(Camera camera, RenderParameter parameter) { GPUCamera gpuCam = { .cameraPosition = camera.position, - .A = camera.A, - .cameraForward = camera.target - camera.position, .f = camera.f, + .cameraForward = camera.target - camera.position, .S_O = camera.S_O, + .fogEmm = glm::vec3(0, 0, 0), + .ks = 0, + .A = camera.A, + .ka = 0, .sensorSize = camera.sensorSize, .width = parameter.width, .height = parameter.height, }; - if(accumulator != nullptr) - { - [accumulator release]; - [resultTexture release]; - } - MTLTextureDescriptor* texDescriptor = [[MTLTextureDescriptor alloc] init]; - [texDescriptor setWidth:parameter.width]; - [texDescriptor setHeight:parameter.height]; - [texDescriptor setPixelFormat:MTLPixelFormatRGBA32Float]; - [texDescriptor setUsage:MTLTextureUsageShaderWrite | MTLTextureUsageShaderRead]; - accumulator = [device newTextureWithDescriptor:texDescriptor]; - resultTexture = [device newTextureWithDescriptor:texDescriptor]; - [texDescriptor release]; for (uint i = 0; i < parameter.numSamples; ++i) { - if(!running) + if (!running) return; - @autoreleasepool{ + @autoreleasepool + { id cmdBuffer = [queue commandBuffer]; id encoder = [cmdBuffer computeCommandEncoder]; // cmdBuffer->addCompletedHandler([this](MTL::CommandBuffer* cmdBuffer) // { std::memcpy(image.data(), resultTexture->buffer(), image.size() * sizeof(glm::vec3)); }); - + SampleParams sample = { - .pass = i, - .samplesPerPixel = parameter.numSamples, - .numDirectionalLights = scene->getNumDirLights(), - .numPointLights = scene->getNumPointLights(), + .pass = i, + .samplesPerPixel = parameter.numSamples, + .numDirectionalLights = scene->getNumDirLights(), + .numPointLights = scene->getNumPointLights(), + .numModels = scene->getNumModels(), }; [encoder setComputePipelineState:computePipeline]; [encoder setBuffer:scene->indicesBuffer offset:0 atIndex:0]; @@ -198,6 +236,10 @@ void MetalRenderer::render(Camera camera, RenderParameter parameter) { [encoder setBuffer:scene->directionalLightBuffer offset:0 atIndex:6]; } + if (scene->getNumPointLights() > 0) + { + [encoder setBuffer:scene->pointLightBuffer offset:0 atIndex:7]; + } [encoder setBuffer:scene->instanceBuffer offset:0 atIndex:8]; [encoder setAccelerationStructure:scene->accelerationStructure atBufferIndex:9]; [encoder setTexture:accumulator atIndex:0]; @@ -230,7 +272,7 @@ void MetalRenderer::render(Camera camera, RenderParameter parameter) [encoder endEncoding]; [cmdBuffer addCompletedHandler:^(id _Nonnull cmd) { sampleTimes.push_back((cmd.GPUEndTime - cmd.GPUStartTime) * 1000.f); - if(sampleTimes.size() > 200) + if (sampleTimes.size() > 200) { sampleTimes.erase(sampleTimes.begin()); } @@ -238,4 +280,4 @@ void MetalRenderer::render(Camera camera, RenderParameter parameter) [cmdBuffer commit]; } } -} \ No newline at end of file +} diff --git a/src/metal/MetalScene.mm b/src/metal/MetalScene.mm index 13dc1bc..6b732d3 100644 --- a/src/metal/MetalScene.mm +++ b/src/metal/MetalScene.mm @@ -40,7 +40,7 @@ void MetalScene::createRayTracingHierarchy() for (uint i = 0; i < refs.size(); ++i) { MTLAccelerationStructureTriangleGeometryDescriptor* descriptor = [MTLAccelerationStructureTriangleGeometryDescriptor descriptor]; - descriptor.triangleCount = refs[i].numIndices; + descriptor.triangleCount = refs[i].numIndices / 3; descriptor.indexBuffer = indicesBuffer; descriptor.indexBufferOffset = refs[i].indicesOffset * sizeof(glm::uvec3); descriptor.indexType = MTLIndexTypeUInt32; diff --git a/src/scene/Scene.h b/src/scene/Scene.h index 2d868a2..1999e39 100644 --- a/src/scene/Scene.h +++ b/src/scene/Scene.h @@ -41,6 +41,7 @@ public: constexpr uint32_t getNumDirLights() const { return (uint)directionalLights.size(); } constexpr uint32_t getNumPointLights() const { return (uint)pointLights.size(); } + constexpr uint32_t getNumModels() const { return (uint)models.size(); } protected: std::vector refs;