From 2c7138b3d7e18d78b1e506aec78844bf00d83440 Mon Sep 17 00:00:00 2001 From: Eric Shi Date: Tue, 21 Jan 2025 09:48:16 -0800 Subject: [PATCH 1/2] Fix left_score was declared but never referenced --- warp/native/mesh.h | 25 +++++-------------------- 1 file changed, 5 insertions(+), 20 deletions(-) diff --git a/warp/native/mesh.h b/warp/native/mesh.h index ba8710b5..2f452e73 100644 --- a/warp/native/mesh.h +++ b/warp/native/mesh.h @@ -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); @@ -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); @@ -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); @@ -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); @@ -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); From 0f314bae923fe328bea923f772e2d0f58f96e96e Mon Sep 17 00:00:00 2001 From: Eric Shi Date: Tue, 21 Jan 2025 09:55:55 -0800 Subject: [PATCH 2/2] Fix some typos --- warp/native/bvh.cpp | 6 +++--- warp/native/bvh.cu | 12 ++++++------ warp/native/bvh.h | 3 +-- 3 files changed, 10 insertions(+), 11 deletions(-) diff --git a/warp/native/bvh.cpp b/warp/native/bvh.cpp index 5258194d..a96f8d3d 100644 --- a/warp/native/bvh.cpp +++ b/warp/native/bvh.cpp @@ -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 { @@ -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; diff --git a/warp/native/bvh.cu b/warp/native/bvh.cu index 8f3efb15..15967d60 100644 --- a/warp/native/bvh.cu +++ b/warp/native/bvh.cu @@ -18,7 +18,7 @@ #include #define THRUST_IGNORE_CUB_VERSION_CHECK -#define REODER_HOST_TREE +#define REORDER_HOST_TREE #include @@ -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) @@ -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); @@ -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]; @@ -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; diff --git a/warp/native/bvh.h b/warp/native/bvh.h index ea177900..1629991f 100644 --- a/warp/native/bvh.h +++ b/warp/native/bvh.h @@ -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; @@ -542,4 +542,3 @@ void bvh_refit_device(uint64_t id); #endif } // namespace wp -