NVIDIA / warp

A Python framework for high performance GPU simulation and graphics
https://nvidia.github.io/warp/
Other
1.75k stars 148 forks source link

Multiple Hashgrids #214

Open steinraf opened 1 month ago

steinraf commented 1 month ago

I am trying to use different hashgrids for spatially separated points but encountering strange values from the hash_grid_point_id function.

I have created an MRE below that shows the issue.

import warp as wp
import numpy as np

# wp.config.mode = "debug"

wp.init()

n_envs = 2
n_points = 2
offset = np.array([10.0, 0.0, 0.0])

r = 0.1

pos = np.random.rand(n_envs, n_points, 3)

grids = []
for i in range(n_envs):
    grid = wp.HashGrid(20, 20, 20)
    pos[i, ...] += i * offset
    grid.build(wp.array(pos[i, ...], dtype=wp.vec3), radius=r)
    grids.append(grid.id)

p = wp.array2d(pos, dtype=wp.vec3)
g = wp.array(grids, dtype=wp.uint64)

print("grids", g)

@wp.kernel
def hashgrid_tester(
    grids: wp.array(dtype=wp.uint64),
    pos: wp.array2d(dtype=wp.vec3),
):
    for env in range(n_envs):
        for tid in range(n_points):

            grid = grids[env]

            i = wp.hash_grid_point_id(grid, tid)
            if i == -1:
                return

            wp.printf(
                "env %d, tid %d, i %d, grid %lld \n",
                env,
                tid,
                i,
                grid,
            )

wp.launch(hashgrid_tester, dim=[1], inputs=[g, p])
wp.synchronize()

print output:

grids [12918456320 12918521344]
env 0, tid 0, i 117, grid 12918456320 
env 0, tid 1, i 479, grid 12918456320 
env 1, tid 0, i 1, grid 12918521344 
env 1, tid 1, i 0, grid 12918521344 

The indices for the second grid get reordered correctly, but the values for the first grid seem to be random.

When using more than 2 points, a Warp CUDA error 716: misaligned address gets thrown, suggesting that the randomness might be from some sort of out of bounds unaligned read.

For the meantime I am using a single HashGrid for the whole domain, but am wondering if multiple instances of the hashgrid are supported and if this is the correct way.

Thanks