qbhf2's picture
added NvidiaWarp and GarmentCode repos
66c9c8a
# Copyright (c) 2022 NVIDIA CORPORATION. All rights reserved.
# NVIDIA CORPORATION and its licensors retain all intellectual property
# and proprietary rights in and to this software, related documentation
# and any modifications thereto. Any use, reproduction, disclosure or
# distribution of this software and related documentation without an express
# license agreement from NVIDIA CORPORATION is strictly prohibited.
import unittest
import numpy as np
import warp as wp
from warp.tests.unittest_utils import *
wp.init()
num_points = 4096
dim_x = 128
dim_y = 128
dim_z = 128
scale = 150.0
cell_radius = 8.0
query_radius = 8.0
num_runs = 4
print_enabled = False
@wp.kernel
def count_neighbors(grid: wp.uint64, radius: float, points: wp.array(dtype=wp.vec3), counts: wp.array(dtype=int)):
tid = wp.tid()
# order threads by cell
i = wp.hash_grid_point_id(grid, tid)
# query point
p = points[i]
count = int(0)
# construct query around point p
for index in wp.hash_grid_query(grid, p, radius):
# compute distance to point
d = wp.length(p - points[index])
if d <= radius:
count += 1
counts[i] = count
@wp.kernel
def count_neighbors_reference(
radius: float, points: wp.array(dtype=wp.vec3), counts: wp.array(dtype=int), num_points: int
):
tid = wp.tid()
i = tid % num_points
j = tid // num_points
# query point
p = points[i]
q = points[j]
# compute distance to point
d = wp.length(p - q)
if d <= radius:
wp.atomic_add(counts, i, 1)
def test_hashgrid_query(test, device):
wp.load_module(device=device)
rng = np.random.default_rng(123)
grid = wp.HashGrid(dim_x, dim_y, dim_z, device)
for i in range(num_runs):
if print_enabled:
print(f"Run: {i+1}")
print("---------")
points = rng.random(size=(num_points, 3)) * scale - np.array((scale, scale, scale)) * 0.5
def particle_grid(dim_x, dim_y, dim_z, lower, radius, jitter):
points = np.meshgrid(
np.linspace(0, dim_x, dim_x), np.linspace(0, dim_y, dim_y), np.linspace(0, dim_z, dim_z)
)
points_t = np.array((points[0], points[1], points[2])).T * radius * 2.0 + np.array(lower)
points_t = points_t + rng.random(size=points_t.shape) * radius * jitter
return points_t.reshape((-1, 3))
points = particle_grid(16, 32, 16, (0.0, 0.3, 0.0), cell_radius * 0.25, 0.1)
points_arr = wp.array(points, dtype=wp.vec3, device=device)
counts_arr = wp.zeros(len(points), dtype=int, device=device)
counts_arr_ref = wp.zeros(len(points), dtype=int, device=device)
profiler = {}
with wp.ScopedTimer("grid operations", print=print_enabled, dict=profiler, synchronize=True):
with wp.ScopedTimer("brute", print=print_enabled, dict=profiler, synchronize=True):
wp.launch(
kernel=count_neighbors_reference,
dim=len(points) * len(points),
inputs=[query_radius, points_arr, counts_arr_ref, len(points)],
device=device,
)
wp.synchronize()
with wp.ScopedTimer("grid build", print=print_enabled, dict=profiler, synchronize=True):
grid.build(points_arr, cell_radius)
with wp.ScopedTimer("grid query", print=print_enabled, dict=profiler, synchronize=True):
wp.launch(
kernel=count_neighbors,
dim=len(points),
inputs=[grid.id, query_radius, points_arr, counts_arr],
device=device,
)
counts = counts_arr.numpy()
counts_ref = counts_arr_ref.numpy()
if print_enabled:
print(f"Grid min: {np.min(counts)} max: {np.max(counts)} avg: {np.mean(counts)}")
print(f"Ref min: {np.min(counts_ref)} max: {np.max(counts_ref)} avg: {np.mean(counts_ref)}")
print(f"Passed: {np.array_equal(counts, counts_ref)}")
test.assertTrue(np.array_equal(counts, counts_ref))
devices = get_test_devices()
class TestHashGrid(unittest.TestCase):
pass
add_function_test(TestHashGrid, "test_hashgrid_query", test_hashgrid_query, devices=devices)
if __name__ == "__main__":
wp.build.clear_kernel_cache()
unittest.main(verbosity=2, failfast=False)