Skip to content

Commit

Permalink
Merge branch 'ershi/extra-left-score' into 'main'
Browse files Browse the repository at this point in the history 
Fix unused "left_score" warnings, typos

See merge request omniverse/warp!1005
  • Loading branch information
@shi-eric
shi-eric committed Jan 21, 2025
2 parents 4d9effa + 0f314ba commit 7801865
Show file tree
Hide file tree
Showing 4 changed files with 15 additions and 31 deletions.
6 changes: 3 additions & 3 deletions warp/native/bvh.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -123,9 +123,9 @@ int TopDownBVHBuilder::partition_median(const vec3* lowers, const vec3* uppers,
return k;
}

struct PartitionPredictateMidPoint
struct PartitionPredicateMidPoint
{
PartitionPredictateMidPoint(const vec3* lowers, const vec3* uppers, int a, float m) : lowers(lowers), uppers(uppers), axis(a), mid(m) {}
PartitionPredicateMidPoint(const vec3* lowers, const vec3* uppers, int a, float m) : lowers(lowers), uppers(uppers), axis(a), mid(m) {}

bool operator()(int index) const
{
Expand All @@ -152,7 +152,7 @@ int TopDownBVHBuilder::partition_midpoint(const vec3* lowers, const vec3* uppers
int axis = longest_axis(edges);
float mid = center[axis];

int* upper = std::partition(indices+start, indices+end, PartitionPredictateMidPoint(lowers, uppers, axis, mid));
int* upper = std::partition(indices+start, indices+end, PartitionPredicateMidPoint(lowers, uppers, axis, mid));

int k = upper-indices;

Expand Down
12 changes: 6 additions & 6 deletions warp/native/bvh.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -18,7 +18,7 @@
#include <cuda_runtime_api.h>

#define THRUST_IGNORE_CUB_VERSION_CHECK
#define REODER_HOST_TREE
#define REORDER_HOST_TREE

#include <cub/cub.cuh>

Expand All @@ -28,7 +28,7 @@ namespace wp
void bvh_create_host(vec3* lowers, vec3* uppers, int num_items, int constructor_type, BVH& bvh);
void bvh_destroy_host(BVH& bvh);

// for LBVH: this will start with some muted leaf nodes, but that is okay, we can still trace up becuase there parents information is still valid
// for LBVH: this will start with some muted leaf nodes, but that is okay, we can still trace up because there parents information is still valid
// the only thing worth mentioning is that when the parent leaf node is also a leaf node, we need to recompute its bounds, since their child information are lost
// for a compact tree such as those from SAH or Median constructor, there is no muted leaf nodes
__global__ void bvh_refit_kernel(int n, const int* __restrict__ parents, int* __restrict__ child_count, int* __restrict__ primitive_indices, BVHPackedNodeHalf* __restrict__ node_lowers, BVHPackedNodeHalf* __restrict__ node_uppers, const vec3* item_lowers, const vec3* item_uppers)
Expand Down Expand Up @@ -456,7 +456,7 @@ void LinearBVHBuilderGPU::build(BVH& bvh, const vec3* item_lowers, const vec3* i
// allocate temporary memory used during building
indices = (int*)alloc_device(WP_CURRENT_CONTEXT, sizeof(int)*num_items*2); // *2 for radix sort
keys = (int*)alloc_device(WP_CURRENT_CONTEXT, sizeof(int)*num_items*2); // *2 for radix sort
deltas = (int*)alloc_device(WP_CURRENT_CONTEXT, sizeof(int)*num_items); // highest differenting bit between keys for item i and i+1
deltas = (int*)alloc_device(WP_CURRENT_CONTEXT, sizeof(int)*num_items); // highest differentiating bit between keys for item i and i+1
range_lefts = (int*)alloc_device(WP_CURRENT_CONTEXT, sizeof(int)*bvh.max_nodes);
range_rights = (int*)alloc_device(WP_CURRENT_CONTEXT, sizeof(int)*bvh.max_nodes);
num_children = (int*)alloc_device(WP_CURRENT_CONTEXT, sizeof(int)*bvh.max_nodes);
Expand Down Expand Up @@ -533,11 +533,11 @@ T* make_device_buffer_of(void* context, T* host_buffer, size_t buffer_size)

void copy_host_tree_to_device(void* context, BVH& bvh_host, BVH& bvh_device_on_host)
{
#ifdef REODER_HOST_TREE
#ifdef REORDER_HOST_TREE


// reorder bvh_host such that its nodes are in the front
// this is enssential for the device refit
// this is essential for the device refit
BVHPackedNodeHalf* node_lowers_reordered = new BVHPackedNodeHalf[bvh_host.max_nodes];
BVHPackedNodeHalf* node_uppers_reordered = new BVHPackedNodeHalf[bvh_host.max_nodes];

Expand Down Expand Up @@ -617,7 +617,7 @@ void copy_host_tree_to_device(void* context, BVH& bvh_host, BVH& bvh_device_on_h
bvh_host.node_lowers = node_lowers_reordered;
bvh_host.node_uppers = node_uppers_reordered;
bvh_host.node_parents = node_parents_reordered;
#endif // REODER_HOST_TREE
#endif // REORDER_HOST_TREE

bvh_device_on_host.num_nodes = bvh_host.num_nodes;
bvh_device_on_host.num_leaf_nodes = bvh_host.num_leaf_nodes;
Expand Down
3 changes: 1 addition & 2 deletions warp/native/bvh.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -161,7 +161,7 @@ struct BVH
// used for fast refits
int* node_parents;
int* node_counts;
// reordered pritimive indices corresponds to the ordering of leaf nodes
// reordered primitive indices corresponds to the ordering of leaf nodes
int* primitive_indices;

int max_depth;
Expand Down Expand Up @@ -542,4 +542,3 @@ void bvh_refit_device(uint64_t id);
#endif

} // namespace wp

25 changes: 5 additions & 20 deletions warp/native/mesh.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -225,12 +225,9 @@ CUDA_CALLABLE inline bool mesh_query_point(uint64_t id, const vec3& point, float
float left_dist_sq = distance_to_aabb_sq(point, vec3(left_lower.x, left_lower.y, left_lower.z), vec3(left_upper.x, left_upper.y, left_upper.z));
float right_dist_sq = distance_to_aabb_sq(point, vec3(right_lower.x, right_lower.y, right_lower.z), vec3(right_upper.x, right_upper.y, right_upper.z));

float left_score = left_dist_sq;
float right_score = right_dist_sq;

wp::vec2i child_indices;
wp::vec2 child_dist;
if (left_dist_sq < right_score)
if (left_dist_sq < right_dist_sq)
{
child_indices = wp::vec2i(right_index, left_index);
child_dist = wp::vec2(right_dist_sq, left_dist_sq);
Expand Down Expand Up @@ -415,12 +412,9 @@ CUDA_CALLABLE inline bool mesh_query_point_no_sign(uint64_t id, const vec3& poin
float left_dist_sq = distance_to_aabb_sq(point, vec3(left_lower.x, left_lower.y, left_lower.z), vec3(left_upper.x, left_upper.y, left_upper.z));
float right_dist_sq = distance_to_aabb_sq(point, vec3(right_lower.x, right_lower.y, right_lower.z), vec3(right_upper.x, right_upper.y, right_upper.z));

float left_score = left_dist_sq;
float right_score = right_dist_sq;

wp::vec2i child_indices;
wp::vec2 child_dist;
if (left_dist_sq < right_score)
if (left_dist_sq < right_dist_sq)
{
child_indices = wp::vec2i(right_index, left_index);
child_dist = wp::vec2(right_dist_sq, left_dist_sq);
Expand Down Expand Up @@ -605,12 +599,9 @@ CUDA_CALLABLE inline bool mesh_query_furthest_point_no_sign(uint64_t id, const v
float left_dist_sq = furthest_distance_to_aabb_sq(point, vec3(left_lower.x, left_lower.y, left_lower.z), vec3(left_upper.x, left_upper.y, left_upper.z));
float right_dist_sq = furthest_distance_to_aabb_sq(point, vec3(right_lower.x, right_lower.y, right_lower.z), vec3(right_upper.x, right_upper.y, right_upper.z));

float left_score = left_dist_sq;
float right_score = right_dist_sq;

wp::vec2i child_indices;
wp::vec2 child_dist;
if (left_dist_sq > right_score)
if (left_dist_sq > right_dist_sq)
{
child_indices = wp::vec2i(right_index, left_index);
child_dist = wp::vec2(right_dist_sq, left_dist_sq);
Expand Down Expand Up @@ -847,12 +838,9 @@ CUDA_CALLABLE inline bool mesh_query_point_sign_normal(uint64_t id, const vec3&
float left_dist_sq = distance_to_aabb_sq(point, vec3(left_lower.x, left_lower.y, left_lower.z), vec3(left_upper.x, left_upper.y, left_upper.z));
float right_dist_sq = distance_to_aabb_sq(point, vec3(right_lower.x, right_lower.y, right_lower.z), vec3(right_upper.x, right_upper.y, right_upper.z));

float left_score = left_dist_sq;
float right_score = right_dist_sq;

wp::vec2i child_indices;
wp::vec2 child_dist;
if (left_dist_sq < right_score)
if (left_dist_sq < right_dist_sq)
{
child_indices = wp::vec2i(right_index, left_index);
child_dist = wp::vec2(right_dist_sq, left_dist_sq);
Expand Down Expand Up @@ -1130,12 +1118,9 @@ CUDA_CALLABLE inline bool mesh_query_point_sign_winding_number(uint64_t id, cons
float left_dist_sq = distance_to_aabb_sq(point, vec3(left_lower.x, left_lower.y, left_lower.z), vec3(left_upper.x, left_upper.y, left_upper.z));
float right_dist_sq = distance_to_aabb_sq(point, vec3(right_lower.x, right_lower.y, right_lower.z), vec3(right_upper.x, right_upper.y, right_upper.z));

float left_score = left_dist_sq;
float right_score = right_dist_sq;

wp::vec2i child_indices;
wp::vec2 child_dist;
if (left_dist_sq < right_score)
if (left_dist_sq < right_dist_sq)
{
child_indices = wp::vec2i(right_index, left_index);
child_dist = wp::vec2(right_dist_sq, left_dist_sq);
Expand Down

0 comments on commit 7801865

Please sign in to comment.