From 9f923d5cf93e4527998255d49b658d0510b6af7e Mon Sep 17 00:00:00 2001 From: iaomw Date: Tue, 20 Jun 2023 20:48:07 +0800 Subject: [PATCH 1/4] anisotropy sss --- zenovis/xinxinoptix/DeflMatShader.cu | 4 +++- zenovis/xinxinoptix/IOMat.h | 1 + zenovis/xinxinoptix/PTKernel.cu | 14 ++++++++++---- zenovis/xinxinoptix/TraceStuff.h | 6 ++++-- 4 files changed, 18 insertions(+), 7 deletions(-) diff --git a/zenovis/xinxinoptix/DeflMatShader.cu b/zenovis/xinxinoptix/DeflMatShader.cu index a15332e39..a8a31e988 100644 --- a/zenovis/xinxinoptix/DeflMatShader.cu +++ b/zenovis/xinxinoptix/DeflMatShader.cu @@ -72,6 +72,7 @@ static __inline__ __device__ MatOutput evalMat(cudaTextureObject_t zenotex[], fl float mat_NoL = 1.0f; float mat_LoV = 1.0f; vec3 mat_reflectance = att_reflectance; + auto vol_sample_anisotropy = 0.0f; //GENERATED_END_MARK /** generated code here end **/ MatOutput mats; @@ -108,6 +109,7 @@ static __inline__ __device__ MatOutput evalMat(cudaTextureObject_t zenotex[], fl mats.sssParam = mat_sssParam; mats.scatterStep = mat_scatterStep; mats.smoothness = mat_smoothness; + mats.vol_anisotropy = clamp(vol_sample_anisotropy, -0.99f, 0.99f); return mats; } } @@ -950,7 +952,7 @@ extern "C" __global__ void __closesthit__radiance() // printf("into sss, sigma_t, alpha: %f, %f, %f\n", prd->sigma_t.x, prd->sigma_t.y, prd->sigma_t.z,prd->ss_alpha.x, prd->ss_alpha.y, prd->ss_alpha.z); // } - prd->pushMat(prd->sigma_t, prd->ss_alpha); + prd->pushMat(prd->sigma_t, prd->ss_alpha, mats.vol_anisotropy); } prd->scatterDistance = scatterDistance; diff --git a/zenovis/xinxinoptix/IOMat.h b/zenovis/xinxinoptix/IOMat.h index 07b1f9106..2c4e8ef3d 100644 --- a/zenovis/xinxinoptix/IOMat.h +++ b/zenovis/xinxinoptix/IOMat.h @@ -33,6 +33,7 @@ struct MatOutput { vec3 nrm; vec3 emission; + float vol_anisotropy; }; struct MatInput { diff --git a/zenovis/xinxinoptix/PTKernel.cu b/zenovis/xinxinoptix/PTKernel.cu index 36ea234ca..cc6b54a85 100644 --- a/zenovis/xinxinoptix/PTKernel.cu +++ b/zenovis/xinxinoptix/PTKernel.cu @@ -8,6 +8,7 @@ #include "DisneyBSDF.h" #include "zxxglslvec.h" +#include "volume.h" #include extern "C" { @@ -337,10 +338,8 @@ extern "C" __global__ void __miss__radiance() } vec3 sigma_t, ss_alpha; - //vec3 sigma_t, ss_alpha; prd->readMat(sigma_t, ss_alpha); - vec3 transmittance; if (ss_alpha.x < 0.0f) { // is inside Glass transmittance = DisneyBSDF::Transmission(sigma_t, optixGetRayTmax()); @@ -351,12 +350,19 @@ extern "C" __global__ void __miss__radiance() prd->attenuation *= transmittance;//DisneyBSDF::Transmission(prd->extinction,optixGetRayTmax()); prd->attenuation2 *= transmittance;//DisneyBSDF::Transmission(prd->extinction,optixGetRayTmax()); prd->origin += prd->direction * optixGetRayTmax(); - prd->direction = DisneyBSDF::SampleScatterDirection(prd->seed); + + auto anisotropy = prd->anisotropy_queue[prd->curMatIdx]; + pbrt::HenyeyGreenstein hg(anisotropy); float3 newdir; + + float2 uu = {rnd(prd->seed), rnd(prd->seed)}; + auto prob = hg.Sample_p(prd->direction, newdir, uu); + + prd->direction = newdir; //DisneyBSDF::SampleScatterDirection(prd->seed); vec3 channelPDF = vec3(1.0f/3.0f); prd->channelPDF = channelPDF; if (ss_alpha.x < 0.0f) { // is inside Glass - prd->maxDistance = DisneyBSDF::SampleDistance2(prd->seed, sigma_t, sigma_t, channelPDF); + prd->maxDistance = 1e16f; //DisneyBSDF::SampleDistance2(prd->seed, sigma_t, sigma_t, channelPDF); } else { prd->maxDistance = diff --git a/zenovis/xinxinoptix/TraceStuff.h b/zenovis/xinxinoptix/TraceStuff.h index 3c84ecaba..9bcb8769d 100644 --- a/zenovis/xinxinoptix/TraceStuff.h +++ b/zenovis/xinxinoptix/TraceStuff.h @@ -1,6 +1,7 @@ #pragma once #include +#include #include "zxxglslvec.h" #include "optixPathTracer.h" @@ -98,6 +99,7 @@ struct RadiancePRD float Lweight; vec3 sigma_t_queue[8]; vec3 ss_alpha_queue[8]; + half anisotropy_queue[8]; int curMatIdx; float samplePdf; bool fromDiff; @@ -152,7 +154,7 @@ struct RadiancePRD attenuation *= multiplier; } - int pushMat(vec3 extinction, vec3 ss_alpha = vec3(-1.0f)) + int pushMat(vec3 extinction, vec3 ss_alpha = vec3(-1.0f), half aniso = 0.0f) { vec3 d = abs(sigma_t_queue[curMatIdx] - extinction); float c = dot(d, vec3(1,1,1)); @@ -161,7 +163,7 @@ struct RadiancePRD curMatIdx++; sigma_t_queue[curMatIdx] = extinction; ss_alpha_queue[curMatIdx] = ss_alpha; - + anisotropy_queue[curMatIdx] = aniso; } return curMatIdx; From 0fde3be09a33135f317fea4f34f82cde749d8458 Mon Sep 17 00:00:00 2001 From: iaomw Date: Wed, 21 Jun 2023 18:13:46 +0800 Subject: [PATCH 2/4] anisotropy glass --- zenovis/xinxinoptix/DeflMatShader.cu | 10 ++++++++-- zenovis/xinxinoptix/PTKernel.cu | 11 ++++++++--- 2 files changed, 16 insertions(+), 5 deletions(-) diff --git a/zenovis/xinxinoptix/DeflMatShader.cu b/zenovis/xinxinoptix/DeflMatShader.cu index a8a31e988..0c51095f8 100644 --- a/zenovis/xinxinoptix/DeflMatShader.cu +++ b/zenovis/xinxinoptix/DeflMatShader.cu @@ -935,7 +935,8 @@ extern "C" __global__ void __closesthit__radiance() if (isTrans) { vec3 channelPDF = vec3(1.0f/3.0f); prd->maxDistance = scatterStep>0.5f? DisneyBSDF::SampleDistance2(prd->seed, prd->sigma_t, prd->sigma_t, channelPDF) : 1e16f; - prd->pushMat(extinction); + prd->pushMat(extinction, vec3(-1.0f), mats.vol_anisotropy); + } else { vec3 channelPDF = vec3(1.0f/3.0f); @@ -990,7 +991,12 @@ extern "C" __global__ void __closesthit__radiance() prd->isSS = true; prd->maxDistance = DisneyBSDF::SampleDistance2(prd->seed, vec3(prd->attenuation) * ss_alpha, sigma_t, prd->channelPDF); } - else + else if (ss_alpha.x < 0.0f) + { + prd->isSS = false; + prd->maxDistance = scatterStep>0.5f? DisneyBSDF::SampleDistance2(prd->seed, sigma_t, sigma_t, prd->channelPDF) : 1e16f; + } + else { prd->isSS = false; prd->maxDistance = 1e16; diff --git a/zenovis/xinxinoptix/PTKernel.cu b/zenovis/xinxinoptix/PTKernel.cu index cc6b54a85..0df46fe7e 100644 --- a/zenovis/xinxinoptix/PTKernel.cu +++ b/zenovis/xinxinoptix/PTKernel.cu @@ -355,15 +355,20 @@ extern "C" __global__ void __miss__radiance() pbrt::HenyeyGreenstein hg(anisotropy); float3 newdir; float2 uu = {rnd(prd->seed), rnd(prd->seed)}; - auto prob = hg.Sample_p(prd->direction, newdir, uu); + auto prob = hg.Sample_p(-prd->direction, newdir, uu); prd->direction = newdir; //DisneyBSDF::SampleScatterDirection(prd->seed); vec3 channelPDF = vec3(1.0f/3.0f); prd->channelPDF = channelPDF; if (ss_alpha.x < 0.0f) { // is inside Glass - prd->maxDistance = 1e16f; //DisneyBSDF::SampleDistance2(prd->seed, sigma_t, sigma_t, channelPDF); - } else + if (optixGetRayTmax() == 1e16f) { + prd->maxDistance = 1e16f; + } else { + prd->maxDistance = DisneyBSDF::SampleDistance2(prd->seed, sigma_t, sigma_t, channelPDF); + } + } + else { prd->maxDistance = DisneyBSDF::SampleDistance2(prd->seed, vec3(prd->attenuation) * ss_alpha, sigma_t, channelPDF); From 2ea688d3ae48b7bf3c816c4e143a0b35c88322c4 Mon Sep 17 00:00:00 2001 From: iaomw Date: Sun, 25 Jun 2023 20:51:58 +0800 Subject: [PATCH 3/4] Fix half type casting --- zenovis/xinxinoptix/DeflMatShader.cu | 17 +---------------- zenovis/xinxinoptix/PTKernel.cu | 16 ++++++++-------- zenovis/xinxinoptix/TraceStuff.h | 12 ++++++------ 3 files changed, 15 insertions(+), 30 deletions(-) diff --git a/zenovis/xinxinoptix/DeflMatShader.cu b/zenovis/xinxinoptix/DeflMatShader.cu index 90dd9f342..87c8d16cf 100644 --- a/zenovis/xinxinoptix/DeflMatShader.cu +++ b/zenovis/xinxinoptix/DeflMatShader.cu @@ -923,7 +923,6 @@ extern "C" __global__ void __closesthit__radiance() if (prd->curMatIdx > 0) { vec3 sigma_t, ss_alpha; - //vec3 sigma_t, ss_alpha; prd->readMat(sigma_t, ss_alpha); if (ss_alpha.x < 0.0f) { // is inside Glass prd->attenuation *= DisneyBSDF::Transmission(sigma_t, optixGetRayTmax()); @@ -948,11 +947,6 @@ extern "C" __global__ void __closesthit__radiance() //printf("maxdist:%f\n",prd->maxDistance); prd->channelPDF = channelPDF; // already calculated in BxDF - - // if (idx.x == w/2 && idx.y == h/2) { - // printf("into sss, sigma_t, alpha: %f, %f, %f\n", prd->sigma_t.x, prd->sigma_t.y, prd->sigma_t.z,prd->ss_alpha.x, prd->ss_alpha.y, prd->ss_alpha.z); - // } - prd->pushMat(prd->sigma_t, prd->ss_alpha, mats.vol_anisotropy); } @@ -999,17 +993,8 @@ extern "C" __global__ void __closesthit__radiance() else { prd->isSS = false; - prd->maxDistance = 1e16; + prd->maxDistance = 1e16f; } - - // if (prd->medium != DisneyBSDF::PhaseFunctions::vacuum) { - - // prd->bad = true; - - // printf("%f %f %f %f %f %f %f %f \n matIdx = %d isotropic = %d \n", prd->sigma_t_queue[0].x, prd->sigma_t_queue[1].x, prd->sigma_t_queue[2].x, prd->sigma_t_queue[3].x, prd->sigma_t_queue[4].x, prd->sigma_t_queue[5].x, prd->sigma_t_queue[6].x, prd->sigma_t_queue[7].x, - // prd->curMatIdx, prd->medium); - // printf("matIdx = %d isotropic = %d \n\n", prd->curMatIdx, prd->medium); - // } } }else{ if(prd->medium == DisneyBSDF::PhaseFunctions::isotropic){ diff --git a/zenovis/xinxinoptix/PTKernel.cu b/zenovis/xinxinoptix/PTKernel.cu index 870df2501..4d304cf1c 100644 --- a/zenovis/xinxinoptix/PTKernel.cu +++ b/zenovis/xinxinoptix/PTKernel.cu @@ -351,8 +351,8 @@ extern "C" __global__ void __miss__radiance() prd->attenuation2 *= transmittance;//DisneyBSDF::Transmission(prd->extinction,optixGetRayTmax()); prd->origin += prd->direction * optixGetRayTmax(); - auto anisotropy = prd->anisotropy_queue[prd->curMatIdx]; - pbrt::HenyeyGreenstein hg(anisotropy); float3 newdir; + auto anisotropy = prd->anisotropy_queue[prd->curMatIdx]; + pbrt::HenyeyGreenstein hg{__half2float(anisotropy)}; float3 newdir; float2 uu = {rnd(prd->seed), rnd(prd->seed)}; auto prob = hg.Sample_p(-prd->direction, newdir, uu); @@ -362,17 +362,17 @@ extern "C" __global__ void __miss__radiance() vec3 channelPDF = vec3(1.0f/3.0f); prd->channelPDF = channelPDF; if (ss_alpha.x < 0.0f) { // is inside Glass - if (optixGetRayTmax() == 1e16f) { - prd->maxDistance = 1e16f; - } else { - prd->maxDistance = DisneyBSDF::SampleDistance2(prd->seed, sigma_t, sigma_t, channelPDF); - } + prd->maxDistance = DisneyBSDF::SampleDistance2(prd->seed, sigma_t, sigma_t, channelPDF); } - else + else if (ss_alpha.x > 0.0f) { prd->maxDistance = DisneyBSDF::SampleDistance2(prd->seed, vec3(prd->attenuation) * ss_alpha, sigma_t, channelPDF); prd->channelPDF = channelPDF; + } + else + { + prd->maxDistance = 1e16f; } prd->depth++; diff --git a/zenovis/xinxinoptix/TraceStuff.h b/zenovis/xinxinoptix/TraceStuff.h index 9bcb8769d..fa77a1d04 100644 --- a/zenovis/xinxinoptix/TraceStuff.h +++ b/zenovis/xinxinoptix/TraceStuff.h @@ -97,9 +97,9 @@ struct RadiancePRD float3 LP; float3 Ldir; float Lweight; - vec3 sigma_t_queue[8]; - vec3 ss_alpha_queue[8]; - half anisotropy_queue[8]; + vec3 sigma_t_queue[8]{}; + vec3 ss_alpha_queue[8]{}; + half anisotropy_queue[8]{}; int curMatIdx; float samplePdf; bool fromDiff; @@ -154,7 +154,7 @@ struct RadiancePRD attenuation *= multiplier; } - int pushMat(vec3 extinction, vec3 ss_alpha = vec3(-1.0f), half aniso = 0.0f) + int pushMat(vec3 extinction, vec3 ss_alpha = vec3(-1.0f), float aniso = 0.0f) { vec3 d = abs(sigma_t_queue[curMatIdx] - extinction); float c = dot(d, vec3(1,1,1)); @@ -162,8 +162,8 @@ struct RadiancePRD { curMatIdx++; sigma_t_queue[curMatIdx] = extinction; - ss_alpha_queue[curMatIdx] = ss_alpha; - anisotropy_queue[curMatIdx] = aniso; + ss_alpha_queue[curMatIdx] = ss_alpha; + anisotropy_queue[curMatIdx] = __float2half(aniso); } return curMatIdx; From 202f6036304493558a7833a6d61177262bb33f90 Mon Sep 17 00:00:00 2001 From: iaomw Date: Tue, 27 Jun 2023 16:36:32 +0800 Subject: [PATCH 4/4] fix garbage extinction value for glass --- zenovis/xinxinoptix/DeflMatShader.cu | 34 ++++++++++++---------------- zenovis/xinxinoptix/PTKernel.cu | 9 ++++---- 2 files changed, 18 insertions(+), 25 deletions(-) diff --git a/zenovis/xinxinoptix/DeflMatShader.cu b/zenovis/xinxinoptix/DeflMatShader.cu index 87c8d16cf..50b4b38df 100644 --- a/zenovis/xinxinoptix/DeflMatShader.cu +++ b/zenovis/xinxinoptix/DeflMatShader.cu @@ -880,8 +880,6 @@ extern "C" __global__ void __closesthit__radiance() prd->diffDepth++; } - - prd->passed = false; bool inToOut = false; bool outToIn = false; @@ -898,13 +896,15 @@ extern "C" __global__ void __closesthit__radiance() if (prd->curMatIdx > 0) { vec3 sigma_t, ss_alpha; prd->readMat(sigma_t, ss_alpha); + + float3 trans; if (ss_alpha.x<0.0f) { // is inside Glass - prd->attenuation *= DisneyBSDF::Transmission(sigma_t, optixGetRayTmax()); - prd->attenuation2 *= DisneyBSDF::Transmission(sigma_t, optixGetRayTmax()); + trans = DisneyBSDF::Transmission(sigma_t, optixGetRayTmax()); } else { - prd->attenuation *= DisneyBSDF::Transmission2(sigma_t * ss_alpha, sigma_t, prd->channelPDF, optixGetRayTmax(), true); - prd->attenuation2 *= DisneyBSDF::Transmission2(sigma_t * ss_alpha, sigma_t, prd->channelPDF, optixGetRayTmax(), true); + trans = DisneyBSDF::Transmission2(sigma_t * ss_alpha, sigma_t, prd->channelPDF, optixGetRayTmax(), true); } + prd->attenuation *= trans; + prd->attenuation2 *= trans; }else { prd->attenuation *= 1; } @@ -914,8 +914,7 @@ extern "C" __global__ void __closesthit__radiance() //if(flag == DisneyBSDF::transmissionEvent || flag == DisneyBSDF::diracEvent) { if(istransmission || flag == DisneyBSDF::diracEvent) { if(prd->next_ray_is_going_inside){ - if(thin < 0.5f && mats.doubleSide < 0.5f ) - { + outToIn = true; inToOut = false; @@ -933,7 +932,8 @@ extern "C" __global__ void __closesthit__radiance() prd->channelPDF = vec3(1.0f/3.0f); if (isTrans) { vec3 channelPDF = vec3(1.0f/3.0f); - prd->maxDistance = scatterStep>0.5f? DisneyBSDF::SampleDistance2(prd->seed, prd->sigma_t, prd->sigma_t, channelPDF) : 1e16f; + // prd->sigma_t is only pre-calculated for SSS branch, so it's garbage value here, use extinction value instead + prd->maxDistance = scatterStep>0.5f? DisneyBSDF::SampleDistance2(prd->seed, extinction, extinction, channelPDF) : 1e16f; prd->pushMat(extinction, vec3(-1.0f), mats.vol_anisotropy); } else { @@ -946,16 +946,15 @@ extern "C" __global__ void __closesthit__radiance() //prd->maxDistance = max(prd->maxDistance, 10/min_sg); //printf("maxdist:%f\n",prd->maxDistance); prd->channelPDF = channelPDF; - // already calculated in BxDF + // prd->sigma_t and prd->ss_alpha should be pre-calculated in BxDF for this SSS branch prd->pushMat(prd->sigma_t, prd->ss_alpha, mats.vol_anisotropy); } prd->scatterDistance = scatterDistance; prd->scatterStep = scatterStep; - } - - } - else{ + } + else { + outToIn = false; inToOut = true; @@ -985,16 +984,11 @@ extern "C" __global__ void __closesthit__radiance() prd->isSS = true; prd->maxDistance = DisneyBSDF::SampleDistance2(prd->seed, vec3(prd->attenuation) * ss_alpha, sigma_t, prd->channelPDF); } - else if (ss_alpha.x < 0.0f) + else { prd->isSS = false; prd->maxDistance = scatterStep>0.5f? DisneyBSDF::SampleDistance2(prd->seed, sigma_t, sigma_t, prd->channelPDF) : 1e16f; } - else - { - prd->isSS = false; - prd->maxDistance = 1e16f; - } } }else{ if(prd->medium == DisneyBSDF::PhaseFunctions::isotropic){ diff --git a/zenovis/xinxinoptix/PTKernel.cu b/zenovis/xinxinoptix/PTKernel.cu index 4d304cf1c..3ec0ec004 100644 --- a/zenovis/xinxinoptix/PTKernel.cu +++ b/zenovis/xinxinoptix/PTKernel.cu @@ -135,6 +135,7 @@ extern "C" __global__ void __raygen__rg() prd.direction = ray_direction; prd.curMatIdx = 0; prd.test_distance = false; + prd.sigma_t_queue[0] = vec3(0.0f); prd.ss_alpha_queue[0] = vec3(-1.0f); prd.minSpecRough = 0.01; prd.samplePdf = 1.0f; @@ -363,17 +364,15 @@ extern "C" __global__ void __miss__radiance() prd->channelPDF = channelPDF; if (ss_alpha.x < 0.0f) { // is inside Glass prd->maxDistance = DisneyBSDF::SampleDistance2(prd->seed, sigma_t, sigma_t, channelPDF); + prd->isSS = false; } - else if (ss_alpha.x > 0.0f) + else { prd->maxDistance = DisneyBSDF::SampleDistance2(prd->seed, vec3(prd->attenuation) * ss_alpha, sigma_t, channelPDF); prd->channelPDF = channelPDF; + prd->isSS = true; } - else - { - prd->maxDistance = 1e16f; - } prd->depth++;