Skip to content

Commit e5c3019

Browse files
committed
Added collision filtering feature for VBD
Signed-off-by: Anka Chen <[email protected]>
1 parent c4baa06 commit e5c3019

File tree

4 files changed

+680
-85
lines changed

4 files changed

+680
-85
lines changed

newton/_src/geometry/kernels.py

Lines changed: 113 additions & 25 deletions
Original file line numberDiff line numberDiff line change
@@ -2343,26 +2343,49 @@ def vertex_triangle_collision_detection_kernel(
23432343
vertex_colliding_triangles_count[v_index] = vertex_num_collisions
23442344
vertex_colliding_triangles_min_dist[v_index] = min_dis_to_tris
23452345

2346+
@wp.func
2347+
def _binary_search_contains_int32(sorted_arr: wp.array(dtype=wp.int32),
2348+
start: int, end: int, key: int) -> bool:
2349+
# Search [start, end) — end is exclusive
2350+
lo = wp.int32(start)
2351+
hi = wp.int32(end)
2352+
while lo < hi:
2353+
mid = (lo + hi) // wp.int32(2)
2354+
val = sorted_arr[mid]
2355+
if val < key:
2356+
lo = mid + 1
2357+
else:
2358+
hi = mid
2359+
return (lo < wp.int32(end)) and (sorted_arr[lo] == key)
2360+
23462361

23472362
@wp.kernel
2348-
def vertex_triangle_collision_detection_no_triangle_buffers_kernel(
2349-
query_radius: float,
2363+
def vertex_triangle_collision_detection_kernel(
2364+
max_query_radius: float,
2365+
min_query_radius: float,
23502366
bvh_id: wp.uint64,
23512367
pos: wp.array(dtype=wp.vec3),
23522368
tri_indices: wp.array(dtype=wp.int32, ndim=2),
23532369
vertex_colliding_triangles_offsets: wp.array(dtype=wp.int32),
23542370
vertex_colliding_triangles_buffer_sizes: wp.array(dtype=wp.int32),
2371+
triangle_colliding_vertices_offsets: wp.array(dtype=wp.int32),
2372+
triangle_colliding_vertices_buffer_sizes: wp.array(dtype=wp.int32),
2373+
vertex_triangle_filtering_list: wp.array(dtype=wp.int32),
2374+
vertex_triangle_filtering_list_offsets: wp.array(dtype=wp.int32),
2375+
min_distance_filtering_ref_pos: wp.array(dtype=wp.vec3),
23552376
# outputs
23562377
vertex_colliding_triangles: wp.array(dtype=wp.int32),
23572378
vertex_colliding_triangles_count: wp.array(dtype=wp.int32),
23582379
vertex_colliding_triangles_min_dist: wp.array(dtype=float),
2380+
triangle_colliding_vertices: wp.array(dtype=wp.int32),
2381+
triangle_colliding_vertices_count: wp.array(dtype=wp.int32),
23592382
triangle_colliding_vertices_min_dist: wp.array(dtype=float),
23602383
resize_flags: wp.array(dtype=wp.int32),
23612384
):
23622385
"""
23632386
This function applies discrete collision detection between vertices and triangles. It uses pre-allocated spaces to
2364-
record the collision data. Unlike `vertex_triangle_collision_detection_kernel`, this collision detection kernel
2365-
works only in one way, i.e., it only records vertices' colliding triangles to `vertex_colliding_triangles`.
2387+
record the collision data. This collision detector works both ways, i.e., it records vertices' colliding triangles to
2388+
`vertex_colliding_triangles`, and records each triangles colliding vertices to `triangle_colliding_vertices`.
23662389
23672390
This function assumes that all the vertices are on triangles, and can be indexed from the pos argument.
23682391
@@ -2375,14 +2398,20 @@ def vertex_triangle_collision_detection_no_triangle_buffers_kernel(
23752398
23762399
Attributes:
23772400
bvh_id (int): the bvh id you want to collide with
2378-
query_radius (float): the contact radius. vertex-triangle pairs whose distance are less than this will get detected
2401+
max_query_radius (float): the upper bound of collision distance.
2402+
min_query_radius (float): the lower bound of collision distance. This distance is evaluated based on min_distance_filtering_ref_pos
23792403
pos (array): positions of all the vertices that make up triangles
2380-
vertex_colliding_triangles (array): flattened buffer of vertices' collision triangles, every two elements records
2381-
the vertex index and a triangle index it collides to
2382-
vertex_colliding_triangles_count (array): number of triangles each vertex collides
23832404
vertex_colliding_triangles_offsets (array): where each vertex' collision buffer starts
23842405
vertex_colliding_triangles_buffer_sizes (array): size of each vertex' collision buffer, will be modified if resizing is needed
23852406
vertex_colliding_triangles_min_dist (array): each vertex' min distance to all (non-neighbor) triangles
2407+
triangle_colliding_vertices_offsets (array): where each triangle's collision buffer starts
2408+
triangle_colliding_vertices_buffer_sizes (array): size of each triangle's collision buffer, will be modified if resizing is needed
2409+
min_distance_filtering_ref_pos (array): the position that minimal collision distance evaluation uses.
2410+
vertex_colliding_triangles (array): flattened buffer of vertices' collision triangles
2411+
vertex_colliding_triangles_count (array): number of triangles each vertex collides
2412+
triangle_colliding_vertices (array): positions of all the triangles' collision vertices, every two elements
2413+
records the vertex index and a triangle index it collides to
2414+
triangle_colliding_vertices_count (array): number of triangles each vertex collides
23862415
triangle_colliding_vertices_min_dist (array): each triangle's min distance to all (non-self) vertices
23872416
resized_flag (array): size == 3, (vertex_buffer_resize_required, triangle_buffer_resize_required, edge_buffer_resize_required)
23882417
"""
@@ -2392,21 +2421,34 @@ def vertex_triangle_collision_detection_no_triangle_buffers_kernel(
23922421
vertex_buffer_offset = vertex_colliding_triangles_offsets[v_index]
23932422
vertex_buffer_size = vertex_colliding_triangles_offsets[v_index + 1] - vertex_buffer_offset
23942423

2395-
lower = wp.vec3(v[0] - query_radius, v[1] - query_radius, v[2] - query_radius)
2396-
upper = wp.vec3(v[0] + query_radius, v[1] + query_radius, v[2] + query_radius)
2424+
lower = wp.vec3(v[0] - max_query_radius, v[1] - max_query_radius, v[2] - max_query_radius)
2425+
upper = wp.vec3(v[0] + max_query_radius, v[1] + max_query_radius, v[2] + max_query_radius)
23972426

23982427
query = wp.bvh_query_aabb(bvh_id, lower, upper)
23992428

24002429
tri_index = wp.int32(0)
24012430
vertex_num_collisions = wp.int32(0)
2402-
min_dis_to_tris = query_radius
2431+
min_dis_to_tris = max_query_radius
24032432
while wp.bvh_query_next(query, tri_index):
24042433
t1 = tri_indices[tri_index, 0]
24052434
t2 = tri_indices[tri_index, 1]
24062435
t3 = tri_indices[tri_index, 2]
2436+
24072437
if vertex_adjacent_to_triangle(v_index, t1, t2, t3):
24082438
continue
24092439

2440+
if vertex_triangle_filtering_list:
2441+
fl_start = vertex_triangle_filtering_list_offsets[v_index]
2442+
fl_end = vertex_triangle_filtering_list_offsets[v_index + 1] # start of next vertex slice (end exclusive)
2443+
2444+
if fl_end > fl_start:
2445+
# Optional fast-fail using first/last elements (remember end is exclusive)
2446+
first_val = vertex_triangle_filtering_list[fl_start]
2447+
last_val = vertex_triangle_filtering_list[fl_end - 1]
2448+
if (tri_index >= first_val) and (tri_index <= last_val):
2449+
if _binary_search_contains_int32(vertex_triangle_filtering_list, fl_start, fl_end, tri_index):
2450+
continue
2451+
24102452
u1 = pos[t1]
24112453
u2 = pos[t2]
24122454
u3 = pos[t3]
@@ -2415,7 +2457,17 @@ def vertex_triangle_collision_detection_no_triangle_buffers_kernel(
24152457

24162458
dist = wp.length(closest_p - v)
24172459

2418-
if dist < query_radius:
2460+
if min_distance_filtering_ref_pos and min_query_radius > 0.:
2461+
closest_p_ref, _, __ = triangle_closest_point(min_distance_filtering_ref_pos[t1],
2462+
min_distance_filtering_ref_pos[t2],
2463+
min_distance_filtering_ref_pos[t3],
2464+
min_distance_filtering_ref_pos[v_index])
2465+
dist_ref = wp.length(closest_p_ref - min_distance_filtering_ref_pos[v_index])
2466+
2467+
if dist_ref < min_query_radius:
2468+
continue
2469+
2470+
if dist < max_query_radius:
24192471
# record v-f collision to vertex
24202472
min_dis_to_tris = wp.min(min_dis_to_tris, dist)
24212473
if vertex_num_collisions < vertex_buffer_size:
@@ -2426,21 +2478,35 @@ def vertex_triangle_collision_detection_no_triangle_buffers_kernel(
24262478

24272479
vertex_num_collisions = vertex_num_collisions + 1
24282480

2429-
wp.atomic_min(triangle_colliding_vertices_min_dist, tri_index, dist)
2481+
if triangle_colliding_vertices:
2482+
wp.atomic_min(triangle_colliding_vertices_min_dist, tri_index, dist)
2483+
tri_buffer_size = triangle_colliding_vertices_buffer_sizes[tri_index]
2484+
tri_num_collisions = wp.atomic_add(triangle_colliding_vertices_count, tri_index, 1)
2485+
2486+
if tri_num_collisions < tri_buffer_size:
2487+
tri_buffer_offset = triangle_colliding_vertices_offsets[tri_index]
2488+
# record v-f collision to triangle
2489+
triangle_colliding_vertices[tri_buffer_offset + tri_num_collisions] = v_index
2490+
else:
2491+
resize_flags[TRI_COLLISION_BUFFER_OVERFLOW_INDEX] = 1
24302492

24312493
vertex_colliding_triangles_count[v_index] = vertex_num_collisions
24322494
vertex_colliding_triangles_min_dist[v_index] = min_dis_to_tris
24332495

24342496

24352497
@wp.kernel
24362498
def edge_colliding_edges_detection_kernel(
2437-
query_radius: float,
2499+
max_query_radius: float,
2500+
min_query_radius: float,
24382501
bvh_id: wp.uint64,
24392502
pos: wp.array(dtype=wp.vec3),
24402503
edge_indices: wp.array(dtype=wp.int32, ndim=2),
24412504
edge_colliding_edges_offsets: wp.array(dtype=wp.int32),
24422505
edge_colliding_edges_buffer_sizes: wp.array(dtype=wp.int32),
24432506
edge_edge_parallel_epsilon: float,
2507+
edge_filtering_list: wp.array(dtype=wp.int32),
2508+
edge_filtering_list_offsets: wp.array(dtype=wp.int32),
2509+
min_distance_filtering_ref_pos: wp.array(dtype=wp.vec3),
24442510
# outputs
24452511
edge_colliding_edges: wp.array(dtype=wp.int32),
24462512
edge_colliding_edges_count: wp.array(dtype=wp.int32),
@@ -2449,7 +2515,8 @@ def edge_colliding_edges_detection_kernel(
24492515
):
24502516
"""
24512517
bvh_id (int): the bvh id you want to do collision detection on
2452-
query_radius (float):
2518+
max_query_radius (float): the upper bound of collision distance.
2519+
min_query_radius (float): the lower bound of collision distance. This distance is evaluated based on min_distance_filtering_ref_pos
24532520
pos (array): positions of all the vertices that make up edges
24542521
edge_colliding_triangles (array): flattened buffer of edges' collision edges
24552522
edge_colliding_edges_count (array): number of edges each edge collides
@@ -2469,32 +2536,52 @@ def edge_colliding_edges_detection_kernel(
24692536
lower = wp.min(e0_v0_pos, e0_v1_pos)
24702537
upper = wp.max(e0_v0_pos, e0_v1_pos)
24712538

2472-
lower = wp.vec3(lower[0] - query_radius, lower[1] - query_radius, lower[2] - query_radius)
2473-
upper = wp.vec3(upper[0] + query_radius, upper[1] + query_radius, upper[2] + query_radius)
2539+
lower = wp.vec3(lower[0] - max_query_radius, lower[1] - max_query_radius, lower[2] - max_query_radius)
2540+
upper = wp.vec3(upper[0] + max_query_radius, upper[1] + max_query_radius, upper[2] + max_query_radius)
24742541

24752542
query = wp.bvh_query_aabb(bvh_id, lower, upper)
24762543

24772544
colliding_edge_index = wp.int32(0)
24782545
edge_num_collisions = wp.int32(0)
2479-
min_dis_to_edges = query_radius
2546+
min_dis_to_edges = max_query_radius
24802547
while wp.bvh_query_next(query, colliding_edge_index):
24812548
e1_v0 = edge_indices[colliding_edge_index, 2]
24822549
e1_v1 = edge_indices[colliding_edge_index, 3]
24832550

24842551
if e0_v0 == e1_v0 or e0_v0 == e1_v1 or e0_v1 == e1_v0 or e0_v1 == e1_v1:
24852552
continue
24862553

2554+
if edge_filtering_list:
2555+
fl_start = edge_filtering_list_offsets[e_index]
2556+
fl_end = edge_filtering_list_offsets[e_index + 1] # start of next vertex slice (end exclusive)
2557+
2558+
if fl_end > fl_start:
2559+
# Optional fast-fail using first/last elements (remember end is exclusive)
2560+
first_val = edge_filtering_list[fl_start]
2561+
last_val = edge_filtering_list[fl_end - 1]
2562+
if (colliding_edge_index >= first_val) and (colliding_edge_index <= last_val):
2563+
if _binary_search_contains_int32(edge_filtering_list, fl_start, fl_end, colliding_edge_index):
2564+
continue
2565+
# else: key is out of range, cannot be present -> skip_this remains False
2566+
24872567
e1_v0_pos = pos[e1_v0]
24882568
e1_v1_pos = pos[e1_v1]
24892569

2490-
st = wp.closest_point_edge_edge(e0_v0_pos, e0_v1_pos, e1_v0_pos, e1_v1_pos, edge_edge_parallel_epsilon)
2491-
s = st[0]
2492-
t = st[1]
2493-
c1 = e0_v0_pos + (e0_v1_pos - e0_v0_pos) * s
2494-
c2 = e1_v0_pos + (e1_v1_pos - e1_v0_pos) * t
2570+
std = wp.closest_point_edge_edge(e0_v0_pos, e0_v1_pos, e1_v0_pos, e1_v1_pos, edge_edge_parallel_epsilon)
2571+
dist = std[2]
24952572

2496-
dist = wp.length(c1 - c2)
2497-
if dist < query_radius:
2573+
if min_distance_filtering_ref_pos and min_query_radius > 0.:
2574+
e0_v0_pos_ref, e0_v1_pos_ref, e1_v0_pos_ref, e1_v1_pos_ref = (min_distance_filtering_ref_pos[e0_v0],
2575+
min_distance_filtering_ref_pos[e0_v1],
2576+
min_distance_filtering_ref_pos[e1_v0],
2577+
min_distance_filtering_ref_pos[e1_v1])
2578+
std_ref = wp.closest_point_edge_edge(e0_v0_pos_ref, e0_v1_pos_ref, e1_v0_pos_ref, e1_v1_pos_ref, edge_edge_parallel_epsilon)
2579+
2580+
dist_ref = std_ref[2]
2581+
if dist_ref < min_query_radius:
2582+
continue
2583+
2584+
if dist < max_query_radius:
24982585
edge_buffer_offset = edge_colliding_edges_offsets[e_index]
24992586
edge_buffer_size = edge_colliding_edges_offsets[e_index + 1] - edge_buffer_offset
25002587

@@ -2512,6 +2599,7 @@ def edge_colliding_edges_detection_kernel(
25122599
edge_colliding_edges_min_dist[e_index] = min_dis_to_edges
25132600

25142601

2602+
25152603
@wp.kernel
25162604
def triangle_triangle_collision_detection_kernel(
25172605
bvh_id: wp.uint64,

0 commit comments

Comments
 (0)