Skip to content

Commit

Permalink
Refactoring
Browse files Browse the repository at this point in the history
  • Loading branch information
@szellmann
szellmann committed Jun 25, 2024
1 parent 1f4fd05 commit bdb68d0
Show file tree
Hide file tree
Showing 9 changed files with 194 additions and 143 deletions.
212 changes: 140 additions & 72 deletions DeviceCopyableObjects.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -1052,7 +1052,7 @@ VSNRAY_FUNC inline void split_primitive(

// BLS primitives //

struct BLS
struct BLSBase
{
enum Type {
Triangle,
Expand All @@ -1070,6 +1070,10 @@ struct BLS
};
Type type{Unknown};
unsigned blsID{UINT_MAX};
};

struct BLS : BLSBase
{
#ifdef WITH_CUDA
union {
cuda_index_bvh<basic_triangle<3,float>>::bvh_ref asTriangle;
Expand Down Expand Up @@ -1107,32 +1111,50 @@ struct BLS
};

// only world BLS's have instances
struct WorldBLS : BLS
struct WorldBLS : BLSBase
{
// asInstance:
// TODO: it would be even better if the bvh_ref
// was bumped into the base classes union; this
// could potentially be achieved via CRTP?
int instID;
#ifdef WITH_CUDA
cuda_index_bvh<BLS>::bvh_ref theBVH;
union {
cuda_index_bvh<basic_triangle<3,float>>::bvh_ref asTriangle;
cuda_index_bvh<basic_triangle<3,float>>::bvh_ref asQuad;
cuda_index_bvh<basic_sphere<float>>::bvh_ref asSphere;
cuda_index_bvh<dco::Cone>::bvh_ref asCone;
cuda_index_bvh<basic_cylinder<float>>::bvh_ref asCylinder;
cuda_index_bvh<dco::BezierCurve>::bvh_ref asBezierCurve;
cuda_index_bvh<dco::ISOSurface>::bvh_ref asISOSurface;
cuda_index_bvh<dco::Volume>::bvh_ref asVolume;
cuda_index_bvh<BLS>::bvh_ref asInstance;
};
#elif defined(WITH_HIP)
hip_index_bvh<BLS>::bvh_ref theBVH;
union {
hip_index_bvh<basic_triangle<3,float>>::bvh_ref asTriangle;
hip_index_bvh<basic_triangle<3,float>>::bvh_ref asQuad;
hip_index_bvh<basic_sphere<float>>::bvh_ref asSphere;
hip_index_bvh<dco::Cone>::bvh_ref asCone;
hip_index_bvh<basic_cylinder<float>>::bvh_ref asCylinder;
hip_index_bvh<dco::BezierCurve>::bvh_ref asBezierCurve;
hip_index_bvh<dco::ISOSurface>::bvh_ref asISOSurface;
hip_index_bvh<dco::Volume>::bvh_ref asVolume;
hip_index_bvh<BLS>::bvh_ref asInstance;
};
#else
index_bvh<BLS>::bvh_ref theBVH;
#endif
union {
struct {
mat3 affineInv;
vec3 transInv;
} asTransform;
struct {
mat3 *affineInv;
vec3 *transInv;
unsigned len;
box1 time;
} asMotionTransform;
index_bvh<basic_triangle<3,float>>::bvh_ref asTriangle;
index_bvh<basic_triangle<3,float>>::bvh_ref asQuad;
index_bvh<basic_sphere<float>>::bvh_ref asSphere;
index_bvh<dco::Cone>::bvh_ref asCone;
index_bvh<basic_cylinder<float>>::bvh_ref asCylinder;
index_bvh<dco::BezierCurve>::bvh_ref asBezierCurve;
index_bvh<dco::ISOSurface>::bvh_ref asISOSurface;
index_bvh<dco::Volume>::bvh_ref asVolume;
index_bvh<BLS>::bvh_ref asInstance;
};
#endif
int instID;
mat3 *affineInv;
vec3 *transInv;
unsigned len;
box1 time;
};

VSNRAY_FUNC
Expand Down Expand Up @@ -1193,11 +1215,56 @@ inline aabb get_bounds(const BLS &bls)
VSNRAY_FUNC
inline aabb get_bounds(const WorldBLS &bls)
{
if (bls.type == BLS::Transform && bls.theBVH.num_nodes()) {
#ifdef WITH_HIP
// with HIP we currenlty assume that TLSs are built on the host:
bvh_node hip_root;
if (bls.type == BLS::Triangle && bls.asTriangle.num_nodes())
HIP_SAFE_CALL(hipMemcpy(
&hip_root, bls.asTriangle.nodes(), sizeof(hip_root), hipMemcpyDefault));
if (bls.type == BLS::Quad && bls.asQuad.num_nodes())
HIP_SAFE_CALL(hipMemcpy(
&hip_root, bls.asQuad.nodes(), sizeof(hip_root), hipMemcpyDefault));
else if (bls.type == BLS::Sphere && bls.asSphere.num_nodes())
HIP_SAFE_CALL(hipMemcpy(
&hip_root, bls.asSphere.nodes(), sizeof(hip_root), hipMemcpyDefault));
else if (bls.type == BLS::Cone && bls.asCone.num_nodes())
HIP_SAFE_CALL(hipMemcpy(
&hip_root, bls.asCone.nodes(), sizeof(hip_root), hipMemcpyDefault));
else if (bls.type == BLS::Cylinder && bls.asCylinder.num_nodes())
HIP_SAFE_CALL(hipMemcpy(
&hip_root, bls.asCylinder.nodes(), sizeof(hip_root), hipMemcpyDefault));
else if (bls.type == BLS::BezierCurve && bls.asBezierCurve.num_nodes())
HIP_SAFE_CALL(hipMemcpy(
&hip_root, bls.asBezierCurve.nodes(), sizeof(hip_root), hipMemcpyDefault));
else if (bls.type == BLS::ISOSurface && bls.asISOSurface.num_nodes())
HIP_SAFE_CALL(hipMemcpy(
&hip_root, bls.asISOSurface.nodes(), sizeof(hip_root), hipMemcpyDefault));
else if (bls.type == BLS::Volume && bls.asVolume.num_nodes())
HIP_SAFE_CALL(hipMemcpy(
&hip_root, bls.asVolume.nodes(), sizeof(hip_root), hipMemcpyDefault));
#else
if (bls.type == BLS::Triangle && bls.asTriangle.num_nodes())
return bls.asTriangle.node(0).get_bounds();
if (bls.type == BLS::Quad && bls.asQuad.num_nodes())
return bls.asQuad.node(0).get_bounds();
else if (bls.type == BLS::Sphere && bls.asSphere.num_nodes())
return bls.asSphere.node(0).get_bounds();
else if (bls.type == BLS::Cone && bls.asCone.num_nodes())
return bls.asCone.node(0).get_bounds();
else if (bls.type == BLS::Cylinder && bls.asCylinder.num_nodes())
return bls.asCylinder.node(0).get_bounds();
else if (bls.type == BLS::BezierCurve && bls.asBezierCurve.num_nodes())
return bls.asBezierCurve.node(0).get_bounds();
else if (bls.type == BLS::ISOSurface && bls.asISOSurface.num_nodes())
return bls.asISOSurface.node(0).get_bounds();
else if (bls.type == BLS::Volume && bls.asVolume.num_nodes())
return bls.asVolume.node(0).get_bounds();
#endif
else if (bls.type == BLS::Transform && bls.asInstance.num_nodes()) {

aabb bound = bls.theBVH.node(0).get_bounds();
mat3f rot = inverse(bls.asTransform.affineInv);
vec3f trans = -bls.asTransform.transInv;
aabb bound = bls.asInstance.node(0).get_bounds();
mat3f rot = inverse(bls.affineInv[0]);
vec3f trans = -bls.transInv[0];
auto verts = compute_vertices(bound);
aabb result;
result.invalidate();
Expand All @@ -1206,23 +1273,27 @@ inline aabb get_bounds(const WorldBLS &bls)
result.insert(v);
}
return result;
} else if (bls.type == BLS::MotionTransform && bls.asMotionTransform.len) {
} else if (bls.type == BLS::MotionTransform && bls.len) {
aabb result;
result.invalidate();
for (unsigned i = 0; i < bls.asMotionTransform.len; ++i) {
aabb bound = bls.theBVH.node(0).get_bounds();
mat3f rot = inverse(bls.asMotionTransform.affineInv[i]);
vec3f trans = -bls.asMotionTransform.transInv[i];
for (unsigned i = 0; i < bls.len; ++i) {
aabb bound = bls.asInstance.node(0).get_bounds();
mat3f rot = inverse(bls.affineInv[i]);
vec3f trans = -bls.transInv[i];
auto verts = compute_vertices(bound);
for (vec3 v : verts) {
v = rot * v + trans;
result.insert(v);
}
}
return result;
} else {
return get_bounds((const BLS &)bls);
}

#ifdef WITH_HIP
return hip_root.get_bounds();
#else
return {};
#endif
}

VSNRAY_FUNC
Expand Down Expand Up @@ -1252,45 +1323,54 @@ VSNRAY_FUNC
inline hit_record<Ray, primitive<unsigned>> intersect(
const Ray &ray, const WorldBLS &bls)
{
const bool isInstance = bls.type == BLS::Transform ||
bls.type == BLS::MotionTransform;

if (!isInstance)
return intersect(ray, (const BLS &)bls);

mat3 affineInv;
vec3 transInv;

if (bls.type == BLS::Transform) {
affineInv = bls.asTransform.affineInv;
transInv = bls.asTransform.transInv;
if (bls.type == BLS::Triangle && (ray.intersectionMask & Ray::Triangle))
return intersect(ray,bls.asTriangle);
else if (bls.type == BLS::Quad && (ray.intersectionMask & Ray::Quad))
return intersect(ray,bls.asQuad);
else if (bls.type == BLS::Sphere && (ray.intersectionMask & Ray::Sphere))
return intersect(ray,bls.asSphere);
else if (bls.type == BLS::Cone && (ray.intersectionMask & Ray::Cone))
return intersect(ray,bls.asCone);
else if (bls.type == BLS::Cylinder && (ray.intersectionMask & Ray::Cylinder))
return intersect(ray,bls.asCylinder);
else if (bls.type == BLS::BezierCurve && (ray.intersectionMask & Ray::BezierCurve))
return intersect(ray,bls.asBezierCurve);
else if (bls.type == BLS::ISOSurface && (ray.intersectionMask & Ray::ISOSurface))
return intersect(ray,bls.asISOSurface);
else if (bls.type == BLS::Volume && (ray.intersectionMask & Ray::Volume))
return intersect(ray,bls.asVolume);

// Instance types:
else if (bls.type == BLS::Transform) {
affineInv = bls.affineInv[0];
transInv = bls.transInv[0];
} else if (bls.type == BLS::MotionTransform) {
float rayTime = clamp(ray.time,
bls.asMotionTransform.time.min,
bls.asMotionTransform.time.max);
float rayTime = clamp(ray.time, bls.time.min, bls.time.max);

float time01 = rayTime - bls.asMotionTransform.time.min
/ (bls.asMotionTransform.time.max - bls.asMotionTransform.time.min);
float time01 = rayTime - bls.time.min / (bls.time.max - bls.time.min);

unsigned ID1 = unsigned(float(bls.asMotionTransform.len-1) * time01);
unsigned ID2 = min(bls.asMotionTransform.len-1, ID1+1);
unsigned ID1 = unsigned(float(bls.len-1) * time01);
unsigned ID2 = min(bls.len-1, ID1+1);

float frac = time01 * (bls.asMotionTransform.len-1) - ID1;
float frac = time01 * (bls.len-1) - ID1;

affineInv = lerp(bls.asMotionTransform.affineInv[ID1],
bls.asMotionTransform.affineInv[ID2],
affineInv = lerp(bls.affineInv[ID1],
bls.affineInv[ID2],
frac);

transInv = lerp(bls.asMotionTransform.transInv[ID1],
bls.asMotionTransform.transInv[ID2],
transInv = lerp(bls.transInv[ID1],
bls.transInv[ID2],
frac);
}

Ray xfmRay(ray);
xfmRay.ori = affineInv * (xfmRay.ori + transInv);
xfmRay.dir = affineInv * xfmRay.dir;

auto hr = intersect(xfmRay,bls.theBVH);
auto hr = intersect(xfmRay,bls.asInstance);
hr.inst_id = hr.hit ? bls.instID : ~0u;
return hr;
}
Expand Down Expand Up @@ -1353,24 +1433,12 @@ struct Instance
#else
index_bvh<BLS>::bvh_ref theBVH;
#endif
union {
struct {
mat4 xfm;
mat3 normalXfm;
mat3 affineInv;
vec3 transInv;
} asTransform;
struct {
// TODO: use arrays, but that needs to be trivially
// constructible for that!
mat4 *xfms;
mat3 *normalXfms;
mat3 *affineInv;
vec3 *transInv;
size_t len;
box1 time;
} asMotionTransform;
};
mat4 *xfms;
mat3 *normalXfms;
mat3 *affineInv;
vec3 *transInv;
size_t len;
box1 time;
};

// Surface //
Expand Down
19 changes: 8 additions & 11 deletions renderer/common.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -607,23 +607,20 @@ VSNRAY_FUNC
inline mat3 getNormalTransform(const dco::Instance &inst, const Ray &ray)
{
if (inst.type == dco::Instance::Transform) {
return inst.asTransform.normalXfm;
return inst.normalXfms[0];
} else if (inst.type == dco::Instance::MotionTransform) {

float rayTime = clamp(ray.time,
inst.asMotionTransform.time.min,
inst.asMotionTransform.time.max);
float rayTime = clamp(ray.time, inst.time.min, inst.time.max);

float time01 = rayTime - inst.asMotionTransform.time.min
/ (inst.asMotionTransform.time.max - inst.asMotionTransform.time.min);
float time01 = rayTime - inst.time.min / (inst.time.max - inst.time.min);

unsigned ID1 = unsigned(float(inst.asMotionTransform.len-1) * time01);
unsigned ID2 = min((unsigned)inst.asMotionTransform.len-1, ID1+1);
unsigned ID1 = unsigned(float(inst.len-1) * time01);
unsigned ID2 = min((unsigned)inst.len-1, ID1+1);

float frac = time01 * (inst.asMotionTransform.len-1) - ID1;
float frac = time01 * (inst.len-1) - ID1;

return lerp(inst.asMotionTransform.normalXfms[ID1],
inst.asMotionTransform.normalXfms[ID2],
return lerp(inst.normalXfms[ID1],
inst.normalXfms[ID2],
frac);
}

Expand Down
40 changes: 17 additions & 23 deletions scene/Instance.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -40,12 +40,22 @@ Instance *Instance::createInstance(
void Instance::commit()
{
m_id = getParam<uint32_t>("id", ~0u);
m_xfm = getParam<mat4>("transform", mat4::identity());
m_xfmInvRot = inverse(top_left(m_xfm));
m_group = getParamObject<Group>("group");
mat4 xfm = getParam<mat4>("transform", mat4::identity());

if (!m_group)
reportMessage(ANARI_SEVERITY_WARNING, "missing 'group' on ANARIInstance");

m_xfms.resize(1);
m_normalXfms.resize(1);
m_affineInv.resize(1);
m_transInv.resize(1);

m_xfms[0] = xfm;
m_affineInv[0] = inverse(top_left(m_xfms[0]));
m_transInv[0] = -m_xfms[0](3).xyz();
m_normalXfms[0] = inverse(transpose(m_affineInv[0]));

dispatch();
}

Expand All @@ -54,21 +64,6 @@ uint32_t Instance::id() const
return m_id;
}

const mat4 &Instance::xfm() const
{
return m_xfm;
}

const mat3 &Instance::xfmInvRot() const
{
return m_xfmInvRot;
}

bool Instance::xfmIsIdentity() const
{
return xfm() == mat4::identity();
}

const Group *Instance::group() const
{
return m_group.ptr;
Expand All @@ -89,13 +84,12 @@ void Instance::visionarayGeometryUpdate()
m_instance[0].userID = m_id;
m_instance[0].groupID = group()->visionarayScene()->m_groupID;

// set xfm
m_instance[0].theBVH = group()->visionarayScene()->refBVH();
m_instance[0].asTransform.xfm = m_xfm;
m_instance[0].asTransform.affineInv = inverse(top_left(m_xfm));
m_instance[0].asTransform.transInv = -m_xfm(3).xyz();
m_instance[0].asTransform.normalXfm
= inverse(transpose(m_instance[0].asTransform.affineInv ));
m_instance[0].xfms = m_xfms.devicePtr();
m_instance[0].normalXfms = m_normalXfms.devicePtr();
m_instance[0].affineInv = m_affineInv.devicePtr();
m_instance[0].transInv = m_transInv.devicePtr();
m_instance[0].len = m_xfms.size();

vgeom.primitives.data = m_instance.devicePtr();
vgeom.primitives.len = m_instance.size();
Expand Down
Loading

0 comments on commit bdb68d0

Please sign in to comment.