in extensions/mvpraymarch/utils.h [729:824]
static __forceinline__ __device__ void ray_subset_fixedbvh(
unsigned warpmask,
int K,
float3 raypos,
float3 raydir,
float2 tminmax,
float2 &rtminmax,
int * sortedobjid,
int2 * nodechildren,
float3 * nodeaabb,
const typename PrimTransfT::Data & primtransf_data,
int *hitboxes,
int & num) {
float3 iraydir = 1.0f/raydir;
int stack[64];
int* stack_ptr = stack;
*stack_ptr++ = -1;
int node = 0;
do {
// check if we're in a leaf
if (node >= (K - 1)) {
{
int k = node - (K - 1);
float3 r0, rd;
PrimTransfT::forward2(primtransf_data, k, raypos, raydir, r0, rd);
float3 ird = 1.0f/rd;
float3 t0 = (-1.f - r0) * ird;
float3 t1 = (1.f - r0) * ird;
float3 tmin = fminf(t0,t1), tmax = fmaxf(t0,t1);
float trmin = max_component(tmin);
float trmax = min_component(tmax);
bool intersection = trmin <= trmax;
if (intersection) {
// hit
rtminmax.x = fminf(rtminmax.x, trmin);
rtminmax.y = fmaxf(rtminmax.y, trmax);
}
if (sync) {
intersection = __any_sync(warpmask, intersection);
}
if (intersection) {
if (sortboxes) {
if (num < maxhitboxes) {
int j = num - 1;
while (j >= 0 && hitboxes[j] > k) {
hitboxes[j + 1] = hitboxes[j];
j = j - 1;
}
hitboxes[j + 1] = k;
num++;
}
} else {
if (num < maxhitboxes) {
hitboxes[num++] = k;
}
}
}
}
node = *--stack_ptr;
} else {
int2 children = make_int2(node * 2 + 1, node * 2 + 2);
// check if we're in each child's bbox
float3 * nodeaabb_ptr = nodeaabb + children.x * 2;
bool traverse_l = ray_aabb_hit_ird(nodeaabb_ptr[0], nodeaabb_ptr[1], raypos, iraydir);
bool traverse_r = ray_aabb_hit_ird(nodeaabb_ptr[2], nodeaabb_ptr[3], raypos, iraydir);
if (sync) {
traverse_l = __any_sync(warpmask, traverse_l);
traverse_r = __any_sync(warpmask, traverse_r);
}
// update stack
if (!traverse_l && !traverse_r) {
node = *--stack_ptr;
} else {
node = traverse_l ? children.x : children.y;
if (traverse_l && traverse_r) {
*stack_ptr++ = children.y;
}
}
if (sync) {
__syncwarp(warpmask);
}
}
} while (node != -1);
}