Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Faster PRD #1685

Merged
merged 5 commits into from
Jan 3, 2024
Merged

Faster PRD #1685

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
5331205
Less PRD
iaomw Jan 2, 2024
95a1cff
Cache HenyeyGreenstein
iaomw Jan 2, 2024
4814d5b
Light texture tint
iaomw Jan 2, 2024
6282a2b
Ternary operator
iaomw Jan 2, 2024
eb5979c
Minor updates
iaomw Jan 2, 2024
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
8 changes: 6 additions & 2 deletions zeno/src/nodes/mtl/ShaderMath.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -12,7 +12,7 @@ static const char /* see https://docs.gl/sl4/trunc */
" radians sinh cosh tanh asinh acosh atanh round roundEven floor"
" ceil trunc sign step length normalize hsvToRgb rgbToHsv luminance",
binops[] = "add sub mul div mod pow atan2 min max dot cross distance safepower",
ternops[] = "mix clamp smoothstep add3";
ternops[] = "mix clamp smoothstep add3 ?";


static auto &toHlsl() {
Expand Down Expand Up @@ -58,7 +58,11 @@ struct ShaderTernaryMath : ShaderNodeClone<ShaderTernaryMath> {

if (op == "add3") {
return em->emitCode(in1 + " + " + in2 + " + " + in3);
} else {
}
else if (op == "?") {
return em->emitCode(in1 + " ? " + in2 + " : " + in3);
}
else {
return em->emitCode(em->funcName(op) + "(" + in1 + ", " + in2 + ", " + in3 + ")");
}
}
Expand Down
1 change: 0 additions & 1 deletion zenovis/xinxinoptix/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -112,7 +112,6 @@ set(FILE_LIST
${CMAKE_CURRENT_SOURCE_DIR}/@PTKernel.cu
${CMAKE_CURRENT_SOURCE_DIR}/@CallableDefault.cu
${CMAKE_CURRENT_SOURCE_DIR}/@CallableVolume.cu
${CMAKE_CURRENT_SOURCE_DIR}/@DefaultFallback.cu
${CMAKE_CURRENT_SOURCE_DIR}/@DeflMatShader.cu
${CMAKE_CURRENT_SOURCE_DIR}/@DisneyBRDF.h
${CMAKE_CURRENT_SOURCE_DIR}/@DisneyBSDF.h
Expand Down
5 changes: 2 additions & 3 deletions zenovis/xinxinoptix/CallableVolume.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -235,12 +235,11 @@ static __inline__ __device__ vec2 samplingVDB(const unsigned long long grid_ptr,
return vec2 { nanoSampling<decltype(_acc), DataTypeNVDB, Order>(_acc, pos_indexed, volin), _grid->tree().root().maximum() };
}

extern "C" __device__ VolumeOut __direct_callable__evalmat(const float4* uniforms, VolumeIn& _attrs) {
extern "C" __device__ VolumeOut __direct_callable__evalmat(const float4* uniforms, VolumeIn2& attrs) {

VolumeIn2 attrs{_attrs.pos, _attrs.sigma_t, _attrs.seed, _attrs.sbt_ptr };
auto& prd = attrs;

auto att_pos = attrs.pos;
vec3& att_pos = reinterpret_cast<vec3&>(attrs.pos);
auto att_clr = vec3(0);
auto att_uv = vec3(0);
auto att_nrm = vec3(0);
Expand Down
24 changes: 0 additions & 24 deletions zenovis/xinxinoptix/DefaultFallback.cu
View file
Open in desktop

This file was deleted.

36 changes: 17 additions & 19 deletions zenovis/xinxinoptix/DeflMatShader.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -49,10 +49,9 @@ extern "C" __global__ void __anyhit__shadow_cutout()

HitGroupData* rt_data = (HitGroupData*)optixGetSbtDataPointer();

RadiancePRD* prd = getPRD();
ShadowPRD* prd = getPRD<ShadowPRD>();
MatInput attrs{};


bool sphere_external_ray = false;

#if (_SPHERE_)
Expand Down Expand Up @@ -211,14 +210,14 @@ extern "C" __global__ void __anyhit__shadow_cutout()
optixIgnoreIntersection();
}else{

if(length(prd->shadowAttanuation) < 0.01f){
prd->shadowAttanuation = vec3(0.0f);
if(length(prd->attanuation) < 0.01f){
prd->attanuation = vec3(0.0f);
optixTerminateRay();
return;
}

if(specTrans==0.0f){
prd->shadowAttanuation = vec3(0.0f);
prd->attanuation = vec3(0.0f);
optixTerminateRay();
return;
}
Expand All @@ -231,23 +230,23 @@ extern "C" __global__ void __anyhit__shadow_cutout()
}
if(rnd(prd->seed)<(1-specTrans)||prd->nonThinTransHit>1)
{
prd->shadowAttanuation = vec3(0,0,0);
prd->attanuation = vec3(0,0,0);
optixTerminateRay();
return;
}

float nDi = fabs(dot(N,ray_dir));
vec3 fakeTrans = vec3(1)-BRDFBasics::fresnelSchlick(vec3(1)-basecolor,nDi);
prd->shadowAttanuation = prd->shadowAttanuation * fakeTrans;
prd->attanuation = prd->attanuation * fakeTrans;

#if (_SPHERE_)
if (sphere_external_ray) {
prd->shadowAttanuation *= vec3(1, 0, 0);
prd->attanuation *= vec3(1, 0, 0);
if (nDi < (1.0f-_FLT_EPL_)) {
prd->shadowAttanuation = {};
prd->attanuation = {};
optixTerminateRay(); return;
} else {
prd->shadowAttanuation *= fakeTrans;
prd->attanuation *= fakeTrans;
}
}
#endif
Expand All @@ -256,7 +255,7 @@ extern "C" __global__ void __anyhit__shadow_cutout()
}
}

prd->shadowAttanuation = vec3(0);
prd->attanuation = vec3(0);
optixTerminateRay();
return;
}
Expand All @@ -280,10 +279,9 @@ extern "C" __global__ void __closesthit__radiance()
RadiancePRD* prd = getPRD();
if(prd->test_distance)
{
prd->vol_t1 = optixGetRayTmax();
prd->maxDistance = optixGetRayTmax();
return;
}
prd->test_distance = false;

const OptixTraversableHandle gas = optixGetGASTraversableHandle();
const uint sbtGASIndex = optixGetSbtGASIndex();
Expand Down Expand Up @@ -859,12 +857,12 @@ extern "C" __global__ void __closesthit__radiance()
prd->radiance_t *= radiance;
};

RadiancePRD shadow_prd {};
shadow_prd.seed = prd->seed;
shadow_prd.shadowAttanuation = make_float3(1.0f, 1.0f, 1.0f);
shadow_prd.nonThinTransHit = (mats.thin == false && mats.specTrans > 0) ? 1 : 0;
ShadowPRD shadowPRD {};
shadowPRD.seed = prd->seed;
shadowPRD.attanuation = make_float3(1.0f, 1.0f, 1.0f);
shadowPRD.nonThinTransHit = (mats.thin == false && mats.specTrans > 0) ? 1 : 0;

shadow_prd.origin = rtgems::offset_ray(P, prd->geometryNormal); // camera space
shadowPRD.origin = rtgems::offset_ray(P, prd->geometryNormal); // camera space
auto shadingP = rtgems::offset_ray(P + params.cam.eye, prd->geometryNormal); // world space

prd->radiance = {};
Expand All @@ -876,7 +874,7 @@ extern "C" __global__ void __closesthit__radiance()
dummy_prt = &radianceNoShadow;
}

DirectLighting<true>(prd, shadow_prd, shadingP, ray_dir, evalBxDF, &taskAux, dummy_prt);
DirectLighting<true>(prd, shadowPRD, shadingP, ray_dir, evalBxDF, &taskAux, dummy_prt);
if(mats.shadowReceiver > 0.5f)
{
auto radiance = length(prd->radiance);
Expand Down
8 changes: 3 additions & 5 deletions zenovis/xinxinoptix/Light.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -31,7 +31,7 @@ extern "C" __global__ void __closesthit__radiance()
RadiancePRD* prd = getPRD();
if(prd->test_distance)
{
prd->vol_t1 = optixGetRayTmax();
prd->maxDistance = optixGetRayTmax();
prd->test_distance = false; return;
}

Expand Down Expand Up @@ -231,7 +231,7 @@ extern "C" __global__ void __anyhit__shadow_cutout()
auto instanceId = optixGetInstanceId();
auto isLightGAS = checkLightGAS(instanceId);

RadiancePRD* prd = getPRD();
ShadowPRD* prd = getPRD<ShadowPRD>();

if (params.num_lights == 0 || !isLightGAS) {
optixIgnoreIntersection();
Expand Down Expand Up @@ -272,9 +272,7 @@ extern "C" __global__ void __anyhit__shadow_cutout()
}

if (visible) {
prd->shadowAttanuation = {};
prd->attenuation2 = {};
prd->attenuation = {};
prd->attanuation = {};
optixTerminateRay();
}

Expand Down
21 changes: 10 additions & 11 deletions zenovis/xinxinoptix/Light.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -75,13 +75,13 @@ static __inline__ __device__ vec3 cihouLightEmission(LightSampleRecord &lsr, Gen
auto intensity = (depth == 0 && light.vIntensity >= 0.0f) ? light.vIntensity : light.intensity;

if (light.tex != 0u) {
auto color = texture2D(light.tex, lsr.uv);
float3 color = (vec3)texture2D(light.tex, lsr.uv);
if (light.texGamma != 1.0f) {
color = pow(color, light.texGamma);
}

color = color * light.color;
color = color * intensity;
return *(vec3*)&color;
return color;
}

return light.color * intensity;
Expand Down Expand Up @@ -204,7 +204,7 @@ namespace detail {

template<bool _MIS_, typename TypeEvalBxDF, typename TypeAux = void>
static __forceinline__ __device__
void DirectLighting(RadiancePRD *prd, RadiancePRD& shadow_prd, const float3& shadingP, const float3& ray_dir,
void DirectLighting(RadiancePRD *prd, ShadowPRD& shadowPRD, const float3& shadingP, const float3& ray_dir,
TypeEvalBxDF& evalBxDF, TypeAux* taskAux=nullptr, float3* RadianceWithoutShadow=nullptr) {

const float3 wo = normalize(-ray_dir);
Expand Down Expand Up @@ -399,9 +399,9 @@ void DirectLighting(RadiancePRD *prd, RadiancePRD& shadow_prd, const float3& sha
}

if (lsr.NoL > _FLT_EPL_ && lsr.PDF > _FLT_EPL_) {
shadow_prd.depth = 0;
traceOcclusion(params.handle, shadow_prd.origin, lsr.dir, 0, lsr.dist, &shadow_prd);
light_attenuation = shadow_prd.shadowAttanuation;

traceOcclusion(params.handle, shadowPRD.origin, lsr.dir, 0, lsr.dist, &shadowPRD);
light_attenuation = shadowPRD.attanuation;

if (nullptr==RadianceWithoutShadow && lengthSquared(light_attenuation) == 0.0f) return;

Expand Down Expand Up @@ -458,13 +458,12 @@ void DirectLighting(RadiancePRD *prd, RadiancePRD& shadow_prd, const float3& sha
if (envpdf < __FLT_DENORM_MIN__) {
return;
}
shadow_prd.depth = 0;
//LP = rtgems::offset_ray(LP, sun_dir);
traceOcclusion(params.handle, shadow_prd.origin, sun_dir,
traceOcclusion(params.handle, shadowPRD.origin, sun_dir,
1e-5f, // tmin
1e16f, // tmax,
&shadow_prd);
light_attenuation = shadow_prd.shadowAttanuation;
&shadowPRD);
light_attenuation = shadowPRD.attanuation;

if (nullptr==RadianceWithoutShadow && lengthSquared(light_attenuation) == 0.0f) return;

Expand Down
40 changes: 29 additions & 11 deletions zenovis/xinxinoptix/TraceStuff.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -43,6 +43,27 @@ enum medium{
isotropicScatter
};

struct VolumePRD {
float vol_t0;
float vol_t1;

bool origin_inside :1;
bool surface_inside :1;
};

struct ShadowPRD {
float3 origin;
uint32_t seed;
float3 attanuation;
uint8_t nonThinTransHit;

VolumePRD vol;

float rndf() {
return rnd(seed);
}
};

struct RadiancePRD
{
// TODO: move some state directly into payload registers?
Expand All @@ -65,7 +86,7 @@ struct RadiancePRD
unsigned int flags;
int countEmitted;
int done;
float3 shadowAttanuation;

int medium;
float scatterDistance;
float scatterPDF;
Expand All @@ -74,7 +95,6 @@ struct RadiancePRD
int diffDepth;
bool isSS;
float scatterStep;
int nonThinTransHit;
float pixel_area;
float Lweight;
vec3 sigma_t_queue[8];
Expand Down Expand Up @@ -111,11 +131,8 @@ struct RadiancePRD
float3 tmp_normal {};

// cihou nanovdb
float vol_t0=0, vol_t1=0;

bool test_distance = false;
bool origin_inside_vdb = false;
bool surface_inside_vdb = false;
VolumePRD vol;
bool test_distance = false ;

float _tmin_ = 0;
float3 geometryNormal;
Expand Down Expand Up @@ -193,7 +210,7 @@ static __forceinline__ __device__ void traceRadiance(
float3 ray_direction,
float tmin,
float tmax,
RadiancePRD *prd,
void *prd,
OptixVisibilityMask mask=255u)
{
unsigned int u0, u1;
Expand All @@ -218,7 +235,7 @@ static __forceinline__ __device__ void traceOcclusion(
float3 ray_direction,
float tmin,
float tmax,
RadiancePRD *prd,
void *prd,
OptixVisibilityMask mask=255u)
{
unsigned int u0, u1;
Expand All @@ -236,11 +253,12 @@ static __forceinline__ __device__ void traceOcclusion(
u0, u1);
}

static __forceinline__ __device__ RadiancePRD* getPRD()
template <typename TypePRD = RadiancePRD>
static __forceinline__ __device__ TypePRD* getPRD()
{
const unsigned int u0 = optixGetPayload_0();
const unsigned int u1 = optixGetPayload_1();
return reinterpret_cast<RadiancePRD*>( unpackPointer( u0, u1 ) );
return reinterpret_cast<TypePRD*>( unpackPointer( u0, u1 ) );
}


Expand Down
Loading
Loading