static __forceinline__ __device__ void ray_subset_fixedbvh()

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);
}