diff --git a/zeno/src/nodes/mtl/ShaderMath.cpp b/zeno/src/nodes/mtl/ShaderMath.cpp
index 30dc603d44..417d9154bf 100644
--- a/zeno/src/nodes/mtl/ShaderMath.cpp
+++ b/zeno/src/nodes/mtl/ShaderMath.cpp
@@ -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() {
@@ -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 + ")");
         }
     }
diff --git a/zenovis/xinxinoptix/CMakeLists.txt b/zenovis/xinxinoptix/CMakeLists.txt
index a39951e1bc..85cf69dfeb 100644
--- a/zenovis/xinxinoptix/CMakeLists.txt
+++ b/zenovis/xinxinoptix/CMakeLists.txt
@@ -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
diff --git a/zenovis/xinxinoptix/CallableVolume.cu b/zenovis/xinxinoptix/CallableVolume.cu
index 75e169c781..102ceb250a 100644
--- a/zenovis/xinxinoptix/CallableVolume.cu
+++ b/zenovis/xinxinoptix/CallableVolume.cu
@@ -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);
diff --git a/zenovis/xinxinoptix/DefaultFallback.cu b/zenovis/xinxinoptix/DefaultFallback.cu
deleted file mode 100644
index 7070aaaeb5..0000000000
--- a/zenovis/xinxinoptix/DefaultFallback.cu
+++ /dev/null
@@ -1,24 +0,0 @@
-#include <optix.h>
-#include "TraceStuff.h"
-
-extern "C" __global__ void __anyhit__shadow_cutout()
-{
-    RadiancePRD* prd = getPRD();
-    prd->shadowAttanuation = {};
-    prd->attenuation = {};
-    optixTerminateRay();
-    return;
-}
-
-extern "C" __global__ void __closesthit__radiance()
-{
-    RadiancePRD* prd = getPRD();
-    prd->done = true;
-    prd->depth += 1;
-    return;
-}
-
-extern "C" __global__ void __closesthit__occlusion()
-{
-    setPayloadOcclusion( true );
-}
diff --git a/zenovis/xinxinoptix/DeflMatShader.cu b/zenovis/xinxinoptix/DeflMatShader.cu
index 86213d57c2..51ec757c5e 100644
--- a/zenovis/xinxinoptix/DeflMatShader.cu
+++ b/zenovis/xinxinoptix/DeflMatShader.cu
@@ -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_)
@@ -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;
             }
@@ -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
@@ -256,7 +255,7 @@ extern "C" __global__ void __anyhit__shadow_cutout()
             }
         }
 
-        prd->shadowAttanuation = vec3(0);
+        prd->attanuation = vec3(0);
         optixTerminateRay();
         return;
     }
@@ -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();
@@ -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 = {};
@@ -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);
diff --git a/zenovis/xinxinoptix/Light.cu b/zenovis/xinxinoptix/Light.cu
index b4e84ea052..2f90a3237b 100644
--- a/zenovis/xinxinoptix/Light.cu
+++ b/zenovis/xinxinoptix/Light.cu
@@ -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;
     }
     
@@ -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();
@@ -272,9 +272,7 @@ extern "C" __global__ void __anyhit__shadow_cutout()
     }
 
     if (visible) {
-        prd->shadowAttanuation = {};
-        prd->attenuation2 = {};
-        prd->attenuation = {};
+        prd->attanuation = {};
         optixTerminateRay();
     }
 
diff --git a/zenovis/xinxinoptix/Light.h b/zenovis/xinxinoptix/Light.h
index 3cfa0f5d07..a283d6aca9 100644
--- a/zenovis/xinxinoptix/Light.h
+++ b/zenovis/xinxinoptix/Light.h
@@ -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;
@@ -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); 
@@ -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;
 
@@ -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;
 
diff --git a/zenovis/xinxinoptix/TraceStuff.h b/zenovis/xinxinoptix/TraceStuff.h
index ffafb52f72..30971c3b09 100644
--- a/zenovis/xinxinoptix/TraceStuff.h
+++ b/zenovis/xinxinoptix/TraceStuff.h
@@ -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?
@@ -65,7 +86,7 @@ struct RadiancePRD
     unsigned int flags;
     int          countEmitted;
     int          done;
-    float3       shadowAttanuation;
+
     int          medium;
     float        scatterDistance;
     float        scatterPDF;
@@ -74,7 +95,6 @@ struct RadiancePRD
     int          diffDepth;
     bool         isSS;
     float        scatterStep;
-    int          nonThinTransHit;
     float        pixel_area;
     float        Lweight;
     vec3         sigma_t_queue[8];
@@ -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;
@@ -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;
@@ -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;
@@ -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 ) );
 }
 
 
diff --git a/zenovis/xinxinoptix/volume.cu b/zenovis/xinxinoptix/volume.cu
index 2ddfece051..7afcf0a5c4 100644
--- a/zenovis/xinxinoptix/volume.cu
+++ b/zenovis/xinxinoptix/volume.cu
@@ -32,18 +32,31 @@ extern "C" __global__ void __intersection__volume()
                                      reinterpret_cast<const nanovdb::Vec3f&>( ray_dir ), t0, t1 );
     // auto fbox = nanovdb::BBox<nanovdb::Vec3f>(nanovdb::Vec3f(dbox.min()), nanovdb::Vec3f(dbox.max()));
 
+    auto flags = optixGetRayFlags();
+    auto anyhit = flags & OPTIX_RAY_FLAG_ENFORCE_ANYHIT;
+
     if( iray.intersects( dbox, t0, t1 )) // t0 >= 0
     {
         // report the entry-point as hit-point
         //auto kind = optixGetHitKind();
         t0 = fmaxf(t0, optixGetRayTmin());
 
-        RadiancePRD* prd = getPRD();
-        prd->vol_t0 = t0;
-        prd->origin_inside_vdb = (t0 == 0);
+        if (anyhit) {
+            ShadowPRD *prd = getPRD<ShadowPRD>();
+
+            prd->vol.vol_t0 = t0;
+            prd->vol.origin_inside = (t0 == 0);
+            prd->vol.vol_t1 = t1; //min(optixGetRayTmax(), t1);
+            prd->vol.surface_inside = (optixGetRayTmax() < t1);
+        } else {
+
+            RadiancePRD* prd = getPRD();
+            prd->vol.vol_t0 = t0;
+            prd->vol.origin_inside = (t0 == 0);
 
-        prd->vol_t1 = t1; //min(optixGetRayTmax(), t1);
-        prd->surface_inside_vdb = (optixGetRayTmax() < t1); // In case triangles were visited before vdb
+            prd->vol.vol_t1 = t1; //min(optixGetRayTmax(), t1);
+            prd->vol.surface_inside = (optixGetRayTmax() < t1); // In case triangles were visited before volume
+        }
 
         if (optixGetRayTmax() > 0) {
             optixReportIntersection(t0, 0);
@@ -54,7 +67,6 @@ extern "C" __global__ void __intersection__volume()
 extern "C" __global__ void __closesthit__radiance_volume()
 {
     RadiancePRD* prd = getPRD();
-    //if(prd->test_distance) { return; }
     
     prd->countEmitted = false;
     prd->radiance = vec3(0);
@@ -67,22 +79,22 @@ extern "C" __global__ void __closesthit__radiance_volume()
     const float3 ray_orig = optixGetWorldRayOrigin();
     const float3 ray_dir  = optixGetWorldRayDirection();
 
-        float t0 = prd->vol_t0; // world space
-        float t1 = prd->vol_t1; // world space
+        float t0 = prd->vol.vol_t0; // world space
+        float t1 = prd->vol.vol_t1; // world space
 
     RadiancePRD testPRD {};
-    testPRD.vol_t1 = _FLT_MAX_;
-    testPRD.test_distance = true;
     testPRD.isSS = false;
-
+    testPRD.maxDistance = _FLT_MAX_;
+    testPRD.test_distance = true;
+    
     uint16_t _mask_ = EverythingMask ^ VolumeMatMask;
 
     traceRadiance(params.handle, ray_orig,ray_dir, 0, _FLT_MAX_, &testPRD, _mask_);
 
-    if(testPRD.vol_t1 < t1)
+    if(testPRD.maxDistance < t1)
     {
-        t1 = testPRD.vol_t1;
-        prd->surface_inside_vdb = true;
+        t1 = testPRD.maxDistance;
+        prd->vol.surface_inside = true;
     }
 
     const float t_max = fmax(0.f, t1 - t0); // world space
@@ -107,7 +119,7 @@ extern "C" __global__ void __closesthit__radiance_volume()
 
         if (t_ele >= t_max) {
 
-            if (prd->surface_inside_vdb) { // Hit other material
+            if (prd->vol.surface_inside) { // Hit other material
 
                 prd->_mask_ = _mask_;
                 prd->_tmin_ = 0;
@@ -141,7 +153,7 @@ extern "C" __global__ void __closesthit__radiance_volume()
             v_density = 0.0f; continue;
         }
 
-        pbrt::HenyeyGreenstein hg { vol_out.anisotropy };
+        pbrt::HenyeyGreenstein hg (vol_out.anisotropy);
         float2 uu = { prd->rndf(), prd->rndf() };
         auto pdf = hg.sample(-ray_dir, new_dir, uu);              
         //auto relative_prob = prob * (CUDART_PI_F * 4);
@@ -176,19 +188,19 @@ extern "C" __global__ void __closesthit__radiance_volume()
 
     scattering = vol_out.albedo;
 
-    RadiancePRD shadow_prd {};
-    shadow_prd.seed = prd->seed;
-    shadow_prd.origin = new_orig; //camera sapce
-    shadow_prd.shadowAttanuation = vec3(1.0f);
+    ShadowPRD shadowPRD {};
+    shadowPRD.seed = prd->seed;
+    shadowPRD.origin = new_orig; //camera sapce
+    shadowPRD.attanuation = vec3(1.0f);
     
     auto evalBxDF = [&](const float3& _wi_, const float3& _wo_, float& thisPDF) -> float3 {
 
-        pbrt::HenyeyGreenstein hg { vol_out.anisotropy };
+        pbrt::HenyeyGreenstein hg(vol_out.anisotropy);
         thisPDF = hg.p(_wo_, _wi_);
         return scattering * thisPDF;
     };
 
-    DirectLighting<true>(prd, shadow_prd, new_orig+params.cam.eye, ray_dir, evalBxDF);
+    DirectLighting<true>(prd, shadowPRD, new_orig+params.cam.eye, ray_dir, evalBxDF);
     
     prd->depth += 1;
     prd->radiance += prd->emission;
@@ -201,11 +213,11 @@ extern "C" __global__ void __anyhit__occlusion_volume()
     const float3 ray_orig = optixGetWorldRayOrigin();
     const float3 ray_dir  = optixGetWorldRayDirection();
 
-    RadiancePRD* prd = getPRD();
+    ShadowPRD* prd = getPRD<ShadowPRD>();
     const HitGroupData* sbt_data = reinterpret_cast<HitGroupData*>( optixGetSbtDataPointer() );
 
-    const float t0 = prd->vol_t0;
-    const float t1 = prd->vol_t1;
+    const float t0 = prd->vol.vol_t0;
+    const float t1 = prd->vol.vol_t1;
 
     const float t_max = t1 - t0; // world space
           float t_ele = 0;
@@ -213,6 +225,9 @@ extern "C" __global__ void __anyhit__occlusion_volume()
     float3 test_point = ray_orig; 
     float3 transmittance = make_float3(1.0f);
 
+    float hgp = 1.0f;
+    pbrt::HenyeyGreenstein hg(9.0f);
+    
     const float sigma_t = sbt_data->vol_extinction;
 
     auto level = sbt_data->vol_depth;
@@ -235,10 +250,13 @@ extern "C" __global__ void __anyhit__occlusion_volume()
         auto prob_scatter = clamp(v_density, 0.0f, 1.0f);
         auto prob_nulling = 1.0f - prob_scatter;
 
-        pbrt::HenyeyGreenstein hg { vol_out.anisotropy };
-        auto prob_continue = hg.p(-ray_dir, ray_dir) * prob_scatter;
-        prob_continue = clamp(prob_continue, 0.0, 1.0f);
-        //printf("prob_continue %f \n", prob_continue);
+        if (vol_out.anisotropy != hg.g ) {
+            hg = pbrt::HenyeyGreenstein(vol_out.anisotropy);
+            hgp = hg.p(-ray_dir, ray_dir);
+        }
+
+        float prob_continue = hgp * prob_scatter;
+        prob_continue = clamp(prob_continue, 0.0, prob_scatter);
 
         auto tr = transmittance * prob_nulling;
         tr += transmittance * prob_continue * vol_out.albedo;
@@ -257,7 +275,7 @@ extern "C" __global__ void __anyhit__occlusion_volume()
         }
     }
 
-    prd->shadowAttanuation *= transmittance;
+    prd->attanuation *= transmittance;
     optixIgnoreIntersection();
     //prd->origin = ray_orig;
     //prd->direction = ray_dir;
diff --git a/zenovis/xinxinoptix/volume.h b/zenovis/xinxinoptix/volume.h
index 420f705046..97da7163e5 100644
--- a/zenovis/xinxinoptix/volume.h
+++ b/zenovis/xinxinoptix/volume.h
@@ -16,6 +16,9 @@ struct VolumeIn {
     uint32_t* seed;
 
     unsigned long long sbt_ptr;
+
+    float3 _local_pos_ = make_float3(CUDART_NAN_F);
+    float3 _uniform_pos_ = make_float3(CUDART_NAN_F);
 };
 
 struct VolumeOut {
@@ -32,36 +35,29 @@ struct VolumeOut {
 namespace pbrt {
 
 struct HenyeyGreenstein {
-    float g;
-    __device__ HenyeyGreenstein(float g) : g(g) {}
+    float g, gg;
+    __device__ HenyeyGreenstein(float g) : g(g), gg(g*g) {}
     
     float p(const float3 &wo, const float3 &wi) const;
     float sample(const float3 &wo, float3 &wi, const float2 &uu) const;
 };
 
-inline float Schlick(float cosTheta, float k)
-{
-    return (1.0f - k * k) / (4.0f * M_PIf * pow(1.0f - k * cosTheta, 2.0f));
-}
-
 // Media Inline Functions
-inline float PhaseHG(float cosTheta, float g) {
-    float gg = g * g;
+inline float PhaseHG(float cosTheta, float g, float gg) {
+
     float denom = 1 + gg + 2 * g * cosTheta;
 
     if (denom < __FLT_EPSILON__) {
         return 1.0f;
     }
 
-    return (0.25f / M_PIf) * (1 - gg) / (denom * sqrtf(denom));
+    auto P = (0.25f / M_PIf) * (1 - gg) / (denom * sqrtf(denom));
+    return clamp(P, 0.0f, 1.0f);
 }
 
 // HenyeyGreenstein Method Definitions
 inline float HenyeyGreenstein::p(const float3 &wo, const float3 &wi) const {
-    return PhaseHG(dot(wo, wi), g);
-    
-    // float k = 1.55f*g - 0.55f*g*g*g;
-    // return Schlick(dot(wo, wi), k);
+    return PhaseHG(dot(wo, wi), g, gg);
 }
 
 inline float HenyeyGreenstein::sample(const float3 &wo, float3 &wi, const float2 &uu) const {
@@ -76,7 +72,6 @@ inline float HenyeyGreenstein::sample(const float3 &wo, float3 &wi, const float2
     if (fabs(g) < 1e-3f)
         cosTheta = 1 - 2 * uu.x;
     else {
-        float gg = g * g;
         float sqrTerm = (1 - gg) / (1 + g - 2 * g * uu.x);
         cosTheta = -(1 + gg - sqrTerm * sqrTerm) / (2 * g);
     }
@@ -90,7 +85,7 @@ inline float HenyeyGreenstein::sample(const float3 &wo, float3 &wi, const float2
     //CoordinateSystem(wo, &v1, &v2);
     wi = SphericalDirection(sinTheta, cosTheta, phi, v1, v2, wo);
 
-    return PhaseHG(cosTheta, g);
+    return PhaseHG(cosTheta, g, gg);
 }
 
 // Spectrum Function Declarations