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 1fc6a94..0000000 Binary files a/src/metal/Compute.air and /dev/null differ 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 06e6ad8..0000000 Binary files a/src/metal/Compute.metallib and /dev/null differ diff --git a/src/metal/ComputeKernel.metal b/src/metal/ComputeKernel.metal new file mode 100644 index 0000000..9586bbc --- /dev/null +++ b/src/metal/ComputeKernel.metal @@ -0,0 +1,357 @@ +#include +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;