Rework compute kernel
This commit is contained in:
Dynamitos
Vendored
+3
@@ -0,0 +1,3 @@
|
||||
{
|
||||
"cmake.generator": "Ninja"
|
||||
}
|
||||
+24
-13
@@ -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
|
||||
$<TARGET_FILE_DIR:RayTracer>
|
||||
)
|
||||
|
||||
endif()
|
||||
add_subdirectory(src/)
|
||||
add_subdirectory(src/)
|
||||
|
||||
@@ -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);
|
||||
}
|
||||
}
|
||||
+168
-64
@@ -1,81 +1,185 @@
|
||||
import Common;
|
||||
|
||||
struct HitInfo
|
||||
{
|
||||
float3 position;
|
||||
float3 normal;
|
||||
float3 barycentricCoords;
|
||||
uint instanceIndex;
|
||||
uint primitiveIndex;
|
||||
};
|
||||
|
||||
HitInfo get_hit_info(RayQuery<RAY_FLAG_NONE> 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<RAY_FLAG_NONE> 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<RAY_FLAG_NONE> 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<RAY_FLAG_NONE> 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);
|
||||
}
|
||||
|
||||
@@ -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;
|
||||
}
|
||||
@@ -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);
|
||||
}
|
||||
+3
-3
@@ -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,
|
||||
|
||||
@@ -4,4 +4,5 @@ target_sources(RayTracer
|
||||
MetalRenderer.mm
|
||||
MetalScene.h
|
||||
MetalScene.mm
|
||||
Compute.metal)
|
||||
ComputeKernel.metal
|
||||
)
|
||||
Binary file not shown.
@@ -1,323 +0,0 @@
|
||||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
#include <metal_numeric>
|
||||
|
||||
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<float>::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<typename T, typename IndexType>
|
||||
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<float, access::read_write> accumulator [[texture(0)]],
|
||||
texture2d<float, access::read_write> 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<triangle_data, instancing> i;
|
||||
i.assume_geometry_type(geometry_type::triangle);
|
||||
i.force_opacity(forced_opacity::opaque);
|
||||
|
||||
typename intersector<triangle_data, instancing>::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<float> 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);
|
||||
}
|
||||
|
||||
Binary file not shown.
@@ -0,0 +1,357 @@
|
||||
#include <metal_stdlib>
|
||||
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<float, access::write> accumulator [[ texture(0) ]],
|
||||
texture2d<float, access::write> 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<float, access::sample> tex [[texture(0)]]) {
|
||||
sampler s(mag_filter::linear, min_filter::linear);
|
||||
return tex.sample(s, in.uv);
|
||||
}
|
||||
@@ -28,6 +28,7 @@ struct SampleParams
|
||||
uint samplesPerPixel;
|
||||
uint numDirectionalLights;
|
||||
uint numPointLights;
|
||||
uint numModels;
|
||||
};
|
||||
|
||||
class MetalRenderer : public Renderer
|
||||
|
||||
+105
-63
@@ -2,14 +2,23 @@
|
||||
#include "metal/MetalScene.h"
|
||||
#include "scene/Renderer.h"
|
||||
#include "util/Camera.h"
|
||||
#include <Foundation/Foundation.h>
|
||||
#include <GLFW/glfw3.h>
|
||||
#include <filesystem>
|
||||
#include <fstream>
|
||||
#include <imgui.h>
|
||||
#include <imgui_impl_glfw.h>
|
||||
#include <imgui_impl_metal.h>
|
||||
#include <iostream>
|
||||
#include <mutex>
|
||||
#include <sstream>
|
||||
#include <string>
|
||||
#include <vector>
|
||||
|
||||
NSWindow* window;
|
||||
NSWindow* window;
|
||||
CAMetalLayer* metalLayer;
|
||||
id<CAMetalDrawable> drawable;
|
||||
id<MTLTexture> texToDraw;
|
||||
|
||||
id<MTLDevice> device;
|
||||
id<MTLLibrary> library;
|
||||
@@ -18,38 +27,49 @@ id<MTLFunction> function;
|
||||
id<MTLComputePipelineState> computePipeline;
|
||||
id<MTLTexture> accumulator = nullptr;
|
||||
id<MTLTexture> resultTexture = nullptr;
|
||||
|
||||
|
||||
static std::mutex g_metal_mtx;
|
||||
|
||||
MTLRenderPassDescriptor* renderPass;
|
||||
id<MTLRenderCommandEncoder> renderEncoder;
|
||||
id<MTLCommandBuffer> renderCmd;
|
||||
id<MTLRenderPipelineState> 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<MTLTexture> tex;
|
||||
{
|
||||
std::lock_guard<std::mutex> 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<MTLCommandBuffer> cmdBuffer = [queue commandBuffer];
|
||||
id<MTLComputeCommandEncoder> 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<MTLCommandBuffer> _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];
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@@ -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;
|
||||
|
||||
@@ -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<ModelReference> refs;
|
||||
|
||||
Reference in New Issue
Block a user