Skip to content

Commit

Permalink
Improve volume bounds
Browse files Browse the repository at this point in the history
  • Loading branch information
iaomw committed Nov 7, 2024
1 parent 1018171 commit e61e176
Show file tree
Hide file tree
Showing 7 changed files with 80 additions and 33 deletions.
7 changes: 6 additions & 1 deletion projects/zenvdb/VolumeBox.cpp
Original file line number Diff line number Diff line change
@@ -1,3 +1,4 @@
#include <string>
#include <zeno/zeno.h>
#include <zeno/VDBGrid.h>
#include <zeno/types/PrimitiveObject.h>
Expand All @@ -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);

Expand Down Expand Up @@ -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);

Expand Down Expand Up @@ -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"}
});
Expand Down
13 changes: 12 additions & 1 deletion zenovis/src/optx/RenderEngineOptx.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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;
}

Expand Down
9 changes: 5 additions & 4 deletions zenovis/xinxinoptix/OptiXStuff.h
Original file line number Diff line number Diff line change
Expand Up @@ -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 } );
Expand Down
37 changes: 18 additions & 19 deletions zenovis/xinxinoptix/optixPathTracer.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -168,6 +168,7 @@ using Vertex = float3;
struct PathTracerState
{
raii<CUdeviceptr> auxHairBuffer;
raii<CUdeviceptr> volumeBoundsBuffer;

OptixTraversableHandle rootHandleIAS;
raii<CUdeviceptr> rootBufferIAS;
Expand Down Expand Up @@ -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) {

Expand All @@ -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 ) {
Expand Down Expand Up @@ -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

Expand All @@ -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;
Expand Down Expand Up @@ -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;
Expand Down
3 changes: 3 additions & 0 deletions zenovis/xinxinoptix/optixPathTracer.h
Original file line number Diff line number Diff line change
Expand Up @@ -188,6 +188,9 @@ struct Params
unsigned long long triangleLightCoordsBuffer;
unsigned long long triangleLightNormalBuffer;

uint32_t firstVolumeOffset;
void* volumeBounds;

uint32_t firstSoloSphereOffset;
void* sphereInstAuxLutBuffer;

Expand Down
43 changes: 35 additions & 8 deletions zenovis/xinxinoptix/volume.cu
Original file line number Diff line number Diff line change
Expand Up @@ -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()
Expand Down Expand Up @@ -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>();
Expand Down Expand Up @@ -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;
Expand All @@ -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;
Expand All @@ -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

Expand Down
1 change: 1 addition & 0 deletions zenovis/xinxinoptix/volume/optixVolume.h
Original file line number Diff line number Diff line change
Expand Up @@ -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;

Expand Down

0 comments on commit e61e176

Please sign in to comment.