Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
3 changes: 2 additions & 1 deletion RayTracing/src/BVH.h
Original file line number Diff line number Diff line change
Expand Up @@ -2,6 +2,7 @@

#include "Scene.h"
#include "Ray.h"
#include "Constants.h"

#include <glm/glm.hpp>

Expand Down Expand Up @@ -104,7 +105,7 @@ class BVH
int resultSphere = -1;

// Stack-based traversal (max 64 levels for a ~2^64 sphere scene)
int stack[64];
int stack[kBVHMaxStackDepth];
int stackPtr = 0;
stack[stackPtr++] = 0; // root node

Expand Down
2 changes: 1 addition & 1 deletion RayTracing/src/CPUBackend.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -294,7 +294,7 @@ glm::vec4 CPUBackend::PerPixel(uint32_t x, uint32_t y) const
if (i > 2)
{
// Use luminance (BT.709) for survival probability: fairer than max(channel)
const float p = 0.2126f * contribution.r + 0.7152f * contribution.g + 0.0722f * contribution.b;
const float p = kLuminanceR * contribution.r + kLuminanceG * contribution.g + kLuminanceB * contribution.b;
if (p < kRussianRouletteThreshold || (p < 1.0f && PathTracerCore::RandomFloat(seed) > p))
break;
contribution /= p;
Expand Down
4 changes: 2 additions & 2 deletions RayTracing/src/CUDARenderer.cu
Original file line number Diff line number Diff line change
Expand Up @@ -159,7 +159,7 @@ int CUDARenderer_Init(CUDARenderState* state)
}
std::printf("[CUDA] Using device: %s (Compute %d.%d, %zu MB)\n",
prop.name, prop.major, prop.minor,
prop.totalGlobalMem / (1024 * 1024));
prop.totalGlobalMem / (static_cast<size_t>(kBytesPerMB)));

state->initialized = true;
err = cudaStreamCreate(&state->uploadStream);
Expand Down Expand Up @@ -226,7 +226,7 @@ void CUDARenderer_OnResize(CUDARenderState* state, uint32_t width, uint32_t heig

std::printf("[CUDA] Resized to %ux%u (%.2f MB device memory)\n",
width, height,
static_cast<float>(state->pixelCount * (2 * sizeof(float4) + sizeof(uint32_t) + sizeof(float3))) / (1024.0f * 1024.0f));
static_cast<float>(state->pixelCount * (2 * sizeof(float4) + sizeof(uint32_t) + sizeof(float3))) / static_cast<float>(kBytesPerMB));
}

void CUDARenderer_UploadScene(
Expand Down
8 changes: 4 additions & 4 deletions RayTracing/src/CUDARenderer.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -94,7 +94,7 @@ __device__ inline float3 SampleGGX_VNDF(float3 V, float a, float r1, float r2)

float3 up = make_float3(0.0f, 0.0f, 1.0f);
float3 T1;
if (Vh.z < 0.9999f) {
if (Vh.z < kONBThreshold) {
T1 = make_float3(-Vh.y, Vh.x, 0.0f);
float t1Len = sqrtf(T1.x*T1.x + T1.y*T1.y);
T1 = make_float3(T1.x/t1Len, T1.y/t1Len, 0.0f);
Expand Down Expand Up @@ -205,7 +205,7 @@ __device__ inline GPUHitPayload TraceRay(
);

// Stack-based BVH traversal (max 64 levels)
int stack[64];
int stack[kBVHMaxStackDepth];
int stackPtr = 0;
stack[stackPtr++] = 0; // root node

Expand Down Expand Up @@ -280,7 +280,7 @@ __device__ inline GPUHitPayload TraceRay(
int rightChild = node.Count;

// Push far child first, then near child (near will be popped first)
if (stackPtr + 2 <= 64)
if (stackPtr + 2 <= kBVHMaxStackDepth)
{
stack[stackPtr++] = rightChild;
stack[stackPtr++] = leftChild;
Expand Down Expand Up @@ -345,7 +345,7 @@ __device__ inline float3 PerPixel(
if (i > 2)
{
// Use luminance (BT.709) for survival probability: fairer than max(channel)
float p = 0.2126f * contribution.x + 0.7152f * contribution.y + 0.0722f * contribution.z;
float p = kLuminanceR * contribution.x + kLuminanceG * contribution.y + kLuminanceB * contribution.z;
if (p < kRussianRouletteThreshold || (p < 1.0f && RandomFloat(seed) > p))
break;
float invP = 1.0f / p;
Expand Down
12 changes: 12 additions & 0 deletions RayTracing/src/Constants.h
Original file line number Diff line number Diff line change
Expand Up @@ -11,3 +11,15 @@ constexpr float kNdotMin = 0.001f; // minimum dot(N,*) for mic
constexpr float kRussianRouletteThreshold = 0.001f; // path termination probability threshold
constexpr float kSpecDenominatorEps = 0.001f; // specular BRDF denominator guard
constexpr float kInSphereEpsilon = 0.000001f; // rejection threshold for RandomInUnitSphere (1e-6)

// ─── Luminance Constants (BT.709) ───
constexpr float kLuminanceR = 0.2126f;
constexpr float kLuminanceG = 0.7152f;
constexpr float kLuminanceB = 0.0722f;

// ─── Algorithmic Constants ───
constexpr float kONBThreshold = 0.9999f; // BuildONB degeneracy check
constexpr int kBVHMaxStackDepth = 64; // BVH traversal stack (supports ~2^64 spheres)

// ─── Unit Conversion ───
constexpr int kBytesPerMB = 1024 * 1024;
4 changes: 3 additions & 1 deletion RayTracing/src/OptiXDenoiser.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -4,6 +4,8 @@

#include <cuda.h> // Driver API for context bridging

#include "Constants.h"

// Define the global OptiX function table (required by optix_stubs.h)
// Must be defined in exactly one translation unit.
OptixFunctionTable g_optixFunctionTable_118 = {};
Expand Down Expand Up @@ -156,7 +158,7 @@ bool OptiXDenoiser::Initialize(uint32_t width, uint32_t height, cudaStream_t str
width, height,
static_cast<float>(m_sizes.withoutOverlapScratchSizeInBytes
+ m_sizes.stateSizeInBytes
+ sizeof(float)) / (1024.0f * 1024.0f));
+ sizeof(float)) / static_cast<float>(kBytesPerMB));
return true;
}

Expand Down
15 changes: 12 additions & 3 deletions RayTracing/src/PathTracer.ispc
Original file line number Diff line number Diff line change
Expand Up @@ -14,6 +14,15 @@ static const float kRussianRouletteThreshold = 0.001f;
static const float kSpecDenominatorEps = 0.001f;
static const float kInSphereEpsilon = 1e-6f;

// ── Luminance Constants (BT.709) ──
static const float kLuminanceR = 0.2126f;
static const float kLuminanceG = 0.7152f;
static const float kLuminanceB = 0.0722f;

// ── Algorithmic Constants ──
static const float kONBThreshold = 0.9999f;
#define kBVHMaxStackDepth 64

// ── RNG: PCG Hash (matches Renderer.cpp Utils::PCG_Hash exactly) ──
// Use unsigned int32: PCG hash relies on unsigned overflow behavior (well-defined).
// ISPC's unsigned int32→float conversion is efficient on modern targets (AVX2/AVX512).
Expand Down Expand Up @@ -100,7 +109,7 @@ static inline void SampleGGX_VNDF(

// T1 = normalize(cross(up, Vh))
varying float T1x, T1y, T1z;
if (Vhz < 0.9999f) {
if (Vhz < kONBThreshold) {
T1x = -Vhy; T1y = Vhx; T1z = 0.0f;
varying float t1Len = sqrt(T1x*T1x + T1y*T1y);
T1x /= t1Len; T1y /= t1Len;
Expand Down Expand Up @@ -163,7 +172,7 @@ static inline float TraceRayBVH(
varying bool negZ = (invZ < 0.0f);

// Per-lane varying stack (ISPC handles varying arrays efficiently)
varying int stack[64];
varying int stack[kBVHMaxStackDepth];
varying int stackPtr = 0;
stack[stackPtr++] = 0; // root node

Expand Down Expand Up @@ -374,7 +383,7 @@ export void ISPCRenderPixels(
// ── Russian roulette (after 3 guaranteed bounces) ──
// Use luminance (BT.709) for survival probability: fairer than max(channel)
if (bounce > 2) {
varying float p = 0.2126f * contribR + 0.7152f * contribG + 0.0722f * contribB;
varying float p = kLuminanceR * contribR + kLuminanceG * contribG + kLuminanceB * contribB;

if (p < kRussianRouletteThreshold)
break;
Expand Down
2 changes: 1 addition & 1 deletion RayTracing/src/PathTracerCore.h
Original file line number Diff line number Diff line change
Expand Up @@ -99,7 +99,7 @@ namespace PathTracerCore
{
glm::vec3 Vh = glm::normalize(glm::vec3(a * V.x, a * V.y, V.z));
glm::vec3 up(0.0f, 0.0f, 1.0f);
glm::vec3 T1 = (Vh.z < 0.9999f) ? glm::normalize(glm::cross(up, Vh)) : glm::vec3(1.0f, 0.0f, 0.0f);
glm::vec3 T1 = (Vh.z < kONBThreshold) ? glm::normalize(glm::cross(up, Vh)) : glm::vec3(1.0f, 0.0f, 0.0f);
glm::vec3 T2 = glm::cross(Vh, T1);

float r = glm::sqrt(r1);
Expand Down
Loading