diff --git a/projects/zenvdb/VolumeBox.cpp b/projects/zenvdb/VolumeBox.cpp
index 904a3cb7d6..34e1f401de 100644
--- a/projects/zenvdb/VolumeBox.cpp
+++ b/projects/zenvdb/VolumeBox.cpp
@@ -1,3 +1,4 @@
+#include <string>
 #include <zeno/zeno.h>
 #include <zeno/VDBGrid.h>
 #include <zeno/types/PrimitiveObject.h>
@@ -22,6 +23,8 @@ struct CreateVolumeBox : zeno::INode {
         auto scale = get_input2<zeno::vec3f>("scale");
         auto rotate = get_input2<zeno::vec3f>("rotate");
 
+        auto bounds = get_input2<std::string>("Bounds:");
+
         auto order = get_input2<std::string>("EulerRotationOrder:");
         auto orderTyped = magic_enum::enum_cast<EulerAngle::RotationOrder>(order).value_or(EulerAngle::RotationOrder::YXZ);
 
@@ -93,6 +96,7 @@ struct CreateVolumeBox : zeno::INode {
         prim->quads->push_back(zeno::vec4i(3, 2, 6, 7));
 
         primWireframe(prim.get(), true);
+        prim->userData().set2("bounds", bounds);
     
         auto transform_ptr = glm::value_ptr(transform);
             
@@ -122,7 +126,8 @@ ZENDEFNODE(CreateVolumeBox, {
     {"prim"},
     {
         {"enum " + EulerAngle::RotationOrderListString(), "EulerRotationOrder", "XYZ"},
-        {"enum " + EulerAngle::MeasureListString(), "EulerAngleMeasure", "Degree"}
+        {"enum " + EulerAngle::MeasureListString(), "EulerAngleMeasure", "Degree"},
+        {"enum Box Sphere HemiSphere", "Bounds", "Box"}
     },
     {"create"}
 });
diff --git a/zenovis/src/optx/RenderEngineOptx.cpp b/zenovis/src/optx/RenderEngineOptx.cpp
index 3a52826c72..48848fbb72 100644
--- a/zenovis/src/optx/RenderEngineOptx.cpp
+++ b/zenovis/src/optx/RenderEngineOptx.cpp
@@ -473,7 +473,18 @@ struct GraphicsManager {
                     memcpy(transform_ptr+8, row2.data(), sizeof(float)*4);  
                     memcpy(transform_ptr+12, row3.data(), sizeof(float)*4);
 
-                    OptixUtil::preloadVolumeBox(key, mtlid, vbox_transform);
+                    auto bounds = ud.get2<std::string>("bounds");
+                    
+                    uint8_t boundsID = [&]() {
+                        if ("Box" == bounds)
+                            return 0;
+                        if ("Sphere" == bounds)
+                            return 1;
+                        if ("HemiSphere" == bounds)
+                            return 2;
+                    } ();
+
+                    OptixUtil::preloadVolumeBox(key, mtlid, boundsID, vbox_transform);
                     return;
                 }
 
diff --git a/zenovis/xinxinoptix/OptiXStuff.h b/zenovis/xinxinoptix/OptiXStuff.h
index 90b8ab5ae3..1187c5d099 100644
--- a/zenovis/xinxinoptix/OptiXStuff.h
+++ b/zenovis/xinxinoptix/OptiXStuff.h
@@ -642,20 +642,21 @@ inline std::map<std::string, std::pair<uint, uint>> g_vdb_indice_visible;
 
 inline std::map<uint, std::vector<std::string>> g_vdb_list_for_each_shader;
 
-inline std::vector<std::tuple<std::string, glm::mat4>> volumeTrans;
+inline std::vector<std::tuple<std::string, uint8_t, glm::mat4>> volumeTrans;
 inline std::vector<std::tuple<std::string, std::shared_ptr<VolumeWrapper>>> volumeBoxs;
 
-inline bool preloadVolumeBox(std::string& key, std::string& matid, glm::mat4& transform) {
+inline bool preloadVolumeBox(std::string& key, std::string& matid, uint8_t bounds, glm::mat4& transform) {
 
-    volumeTrans.push_back( {matid, transform} );
+    volumeTrans.push_back( {matid, bounds, transform} );
     return true;
 }
 
 inline bool processVolumeBox() {
 
     volumeBoxs.clear();
-    for (auto& [key, val] : volumeTrans) {
+    for (auto& [key, bounds, val] : volumeTrans) {
         auto volume_ptr = std::make_shared<VolumeWrapper>();
+        volume_ptr->bounds = bounds;
         volume_ptr->transform = val;
         buildVolumeAccel(volume_ptr->accel, *volume_ptr, context);
         volumeBoxs.emplace_back( std::tuple{ key, volume_ptr } );
diff --git a/zenovis/xinxinoptix/optixPathTracer.cpp b/zenovis/xinxinoptix/optixPathTracer.cpp
index 7e114534fc..621568cf9d 100644
--- a/zenovis/xinxinoptix/optixPathTracer.cpp
+++ b/zenovis/xinxinoptix/optixPathTracer.cpp
@@ -168,6 +168,7 @@ using Vertex = float3;
 struct PathTracerState
 {
     raii<CUdeviceptr> auxHairBuffer;
+    raii<CUdeviceptr> volumeBoundsBuffer;
 
     OptixTraversableHandle         rootHandleIAS;
     raii<CUdeviceptr>              rootBufferIAS;
@@ -1022,6 +1023,8 @@ void updateRootIAS()
         optix_instances.push_back( opinstance );
     }
 
+    auto firstVolumeInstance = optix_instance_idx+1;
+    std::vector<uint8_t> volumeBounds(OptixUtil::volumeBoxs.size());
     // process volume
 	for (uint i=0; i<OptixUtil::volumeBoxs.size(); ++i) {
         
@@ -1040,7 +1043,9 @@ void updateRootIAS()
 		optix_instance.sbtOffset = sbt_offset;
 		optix_instance.visibilityMask = VolumeMatMask; //VOLUME_OBJECT;
 		optix_instance.traversableHandle = val->accel.handle;
-        
+
+        volumeBounds[i] = val->bounds;
+
 		getOptixTransform( *(val), optix_instance.transform );
         
         if ( OptixUtil::g_vdb_list_for_each_shader.count(shader_index) > 0 ) {
@@ -1071,10 +1076,7 @@ void updateRootIAS()
 
                 auto dummy = glm::transpose(dirtyMatrix);
                 auto dummy_ptr = glm::value_ptr( dummy );
-                for (size_t i=0; i<12; ++i) {   
-                    //optix_instance.transform[i] = mat3r4c[i];
-                    optix_instance.transform[i] = dummy_ptr[i];
-                }
+                memcpy(optix_instance.transform, dummy_ptr, sizeof(float)*12);
             }
         } // count  
 
@@ -1085,6 +1087,17 @@ void updateRootIAS()
 		optix_instances.push_back( optix_instance );
 	}
 
+    state.volumeBoundsBuffer.reset();
+
+    if (!volumeBounds.empty()) {
+        size_t byte_size = sizeof(volumeBounds[0]) * volumeBounds.size();
+        state.volumeBoundsBuffer.resize(byte_size);
+        cudaMemcpy((void*)state.volumeBoundsBuffer.handle, volumeBounds.data(), byte_size, cudaMemcpyHostToDevice);
+
+        state.params.volumeBounds = (void*)state.volumeBoundsBuffer.handle;
+        state.params.firstVolumeOffset = firstVolumeInstance;
+    }
+
     auto op_index = optix_instances.size();
 
     uint32_t MAX_INSTANCE_ID;
@@ -2439,20 +2452,6 @@ struct DrawDat {
 };
 static std::map<std::string, DrawDat> drawdats;
 
-struct DrawDat2 {
-  std::vector<std::string>  mtlidList;
-  std::string mtlid;
-  std::string instID;
-  std::vector<zeno::vec3f> pos;
-  std::vector<zeno::vec3f> nrm;
-  std::vector<zeno::vec3f> tang;
-  std::vector<zeno::vec3f> clr;
-  std::vector<zeno::vec3f> uv;
-  std::vector<int> triMats;
-  std::vector<zeno::vec3i> triIdx;
-};
-static std::map<std::string, DrawDat2> drawdats2;
-
 std::set<std::string> uniqueMatsForMesh() {
 
     std::set<std::string> result;
diff --git a/zenovis/xinxinoptix/optixPathTracer.h b/zenovis/xinxinoptix/optixPathTracer.h
index 5e18214d6e..051db86886 100644
--- a/zenovis/xinxinoptix/optixPathTracer.h
+++ b/zenovis/xinxinoptix/optixPathTracer.h
@@ -188,6 +188,9 @@ struct Params
     unsigned long long triangleLightCoordsBuffer;
     unsigned long long triangleLightNormalBuffer;
 
+    uint32_t firstVolumeOffset;
+    void* volumeBounds;
+
     uint32_t firstSoloSphereOffset;
     void* sphereInstAuxLutBuffer;
 
diff --git a/zenovis/xinxinoptix/volume.cu b/zenovis/xinxinoptix/volume.cu
index 2131272c85..7b2b3dcaec 100644
--- a/zenovis/xinxinoptix/volume.cu
+++ b/zenovis/xinxinoptix/volume.cu
@@ -15,12 +15,36 @@
 using DataTypeNVDB0 = nanovdb::Fp32;
 using GridTypeNVDB0 = nanovdb::NanoGrid<DataTypeNVDB0>;
 
-__inline__ __device__ bool rayBox(const float3& ray_ori, const float3& ray_dir, const nanovdb::BBox<nanovdb::Vec3f>& box, 
-                        float& t0, float& t1) {
+__inline__ __device__ bool rayHit(const float3& ray_ori, const float3& ray_dir, const nanovdb::BBox<nanovdb::Vec3f>& box, 
+                                uint8_t bounds, float& t0, float& t1) {
 
     auto iray = nanovdb::Ray<float>( reinterpret_cast<const nanovdb::Vec3f&>( ray_ori ),
                                      reinterpret_cast<const nanovdb::Vec3f&>( ray_dir ), t0, t1 );
-    return iray.intersects( box, t0, t1 );
+    if (bounds == 0)
+        return iray.intersects( box, t0, t1 );
+    if (bounds == 1)
+        return iray.intersects({0,0,0}, box.max()[1], t0, t1);
+    if (bounds == 2) {
+
+        iray.intersects( box, t0, t1 );
+
+        auto b_t0 = t0;
+        auto b_t1 = t1;
+
+        iray.setEye( iray.eye() * nanovdb::Vec3f{2,1,2} );
+        iray.setDir( iray.dir() * nanovdb::Vec3f{2,1,2} );
+        bool hit = iray.intersects({0,box.min()[1],0}, box.max()[1]*2, t0, t1);
+
+        auto v0 = iray.dir() * t0 * nanovdb::Vec3f{0.5,1,0.5};
+        t0 = v0.length();
+        auto v1 = iray.dir() * t1 * nanovdb::Vec3f{0.5,1,0.5};
+        t1 = v1.length();
+
+        t0 = max(b_t0, t0);
+        t1 = min(b_t1, t1);
+
+        return hit;
+    }
 }
 
 extern "C" __global__ void __intersection__volume()
@@ -52,8 +76,11 @@ extern "C" __global__ void __intersection__volume()
         t1 = t1 * dirlen;
     } 
 
+    auto offsetId = optixGetInstanceId()-params.firstVolumeOffset;
+    auto bounds = reinterpret_cast<uint8_t*>(params.volumeBounds)[offsetId];
+
     //bool inside = box.isInside(reinterpret_cast<const nanovdb::Vec3f&>(ray_ori));
-    auto hitted = rayBox( ray_ori, ray_dir, box, t0, t1 );
+    auto hitted = rayHit( ray_ori, ray_dir, box, bounds, t0, t1 );
     if (!hitted) { return; }
 
     RadiancePRD *prd = getPRD<RadiancePRD>();
@@ -135,8 +162,8 @@ extern "C" __global__ void __closesthit__radiance_volume()
     testPRD.done = false;
     testPRD.seed = prd->seed;
     testPRD.depth == 0;
-    testPRD._tmin_ = t0;
-    testPRD.maxDistance = t1;
+    testPRD._tmin_ = 0;
+    testPRD.maxDistance = NextFloatUp(t1);
     testPRD.test_distance = true;
     
     uint8_t _mask_ = EverythingMask ^ VolumeMatMask;
@@ -146,7 +173,7 @@ extern "C" __global__ void __closesthit__radiance_volume()
     } while(testPRD.test_distance && !testPRD.done);
 
     bool surface_inside = false;
-    if(testPRD.maxDistance < t1)
+    if(testPRD.maxDistance <= t1 && testPRD.maxDistance >= t0)
     {
         t1 = testPRD.maxDistance;
         surface_inside = true;
@@ -162,7 +189,7 @@ extern "C" __global__ void __closesthit__radiance_volume()
         if (surface_inside) { // Hit other material
 
             prd->_mask_ = _mask_;
-            prd->_tmin_ = NextFloatDown(t1);
+            //prd->_tmin_ = NextFloatDown(t1);
 
         } else { // Volume edge
 
diff --git a/zenovis/xinxinoptix/volume/optixVolume.h b/zenovis/xinxinoptix/volume/optixVolume.h
index 4a224bf9a2..ae9a5ac6ca 100644
--- a/zenovis/xinxinoptix/volume/optixVolume.h
+++ b/zenovis/xinxinoptix/volume/optixVolume.h
@@ -142,6 +142,7 @@ struct VolumeAccel
 struct VolumeWrapper
 {
 	//openvdb::math::Transform::Ptr transform; // openvdb::math::Mat4f::identity();
+	uint8_t bounds;
 	glm::mat4 transform; 
 	std::vector<std::string> selected;