`primitive_narrowphase` kernel hash collision

[aaydinoglu]

Note: I am using warp through the mjx interface, but still wanted to open the issue here as I believe it is related to warp side.

Me and @bmabsout are encountering a hash collision issue that causes objects to fall through the ground. Unfortunately, I can not share the model that causes the issue as it includes things I am not allowed to share (companies property). I have tried to replicate the issue with different models and failed but still wanted to let you know about it (perhaps it could be easier for you to replicate). I apologize for the inconvenience. Let me explain the issue in detail.

Process: First I set warp data and model via:

model = mjx.put_model(mj_model, impl=mjx.Impl.WARP)
data = mjx.make_data(model, impl=mjx.Impl.WARP, nconmax=..., njmax=..., device=...)

and then simply step the environment.

Experiment steps:

1. Clean warp cache.
2. Start a model with a sphere (and a robot I am not allowed the share) and compile the model. You will see that the primitive narrowphase is compiled:

Module _primitive_narrowphase__locals__primitive_narrowphase_112e6761 fa4a16c load on device 'cuda:0' took 506.52 ms  (compiled)

3. Step the environment, everything works well.
4. DO NOT clean the warp cache.
5. Start a model with a cube (instead of a sphere) and compile the model. You will see that primitive narrowphase is cached, even though we switched to a cube from a sphere. I see the following:

Module _primitive_narrowphase__locals__primitive_narrowphase_112e6761 fa4a16c load on device 'cuda:0' took 0.51 ms  (cached)

6. Step the environment and cube falls through the ground.
7. Clean ONLY the module (_primitive_narrowphase__locals__primitive_narrowphase_) mentioned above from the cache.
8. Start a model with a cube and compile the model. It does compile the kernel and everything works well.

I am quite confident issue is with this line. We have a local fix by switching that line with:

def order_independent_hash(items):
    """Collision-free hash for any iterable of ints, order-independent."""
    hash_val = 0
    for item in sorted(items): # sorting ensures order-independence
        # upper sudzik as a collision free hash
        hash_val = max(hash_val, item)**2 + min(hash_val, item)
    return hash_val

types_hash = order_independent_hash(map(order_independent_hash, primitive_collisions_types))
unique_kernel_name = f"primitive_narrowphase_{types_hash}"
@wp.kernel(module=unique_kernel_name, enable_backward=False)

Relevant versions:

mujoco                             3.5.0
mujoco-mjx                         3.5.0
warp-lang                          1.11.0

[thowell]

@aaydinoglu thank you for reporting this issue!

@shi-eric @adenzler-nvidia should the reported issue be resolved in warp-lang >= 1.12.0 with NVIDIA/warp#1231?

[adenzler-nvidia]

This sounds like issues we've been fighting recently, a few things have been fixed in 1.12. So yes, it would be nice if you could try 1.12, the rc versions are on the nvidia index:

pip install -U --pre warp-lang --extra-index-url=https://pypi.nvidia.com/

[aaydinoglu]

This sounds like issues we've been fighting recently, a few things have been fixed in 1.12. So yes, it would be nice if you could try 1.12, the rc versions are on the nvidia index:

pip install -U --pre warp-lang --extra-index-url=https://pypi.nvidia.com/

I have tested with:

Name: warp-lang
Version: 1.13.0.dev20260227

This actually fixed our issue. Specifically, first attempt with sphere compiled its own as expected:

Module _primitive_narrowphase__locals__primitive_narrowphase_1ec26745 1ec2674 load on device 'cuda:0' took 84.03 ms  (compiled)

Then I switched to cube, and this time it did compile again and cube did not fall through the ground:

Module _primitive_narrowphase__locals__primitive_narrowphase_e85bcc45 e85bcc4 load on device 'cuda:0' took 102.32 ms  (compiled)

Note that I double-checked that in the same setup if I roll back to warp-lang==1.11.0, the issue persists. And in both tests I started with a removed cache. So this fixes this particular issue.