Move all SolverMuJoCo Warp kernels to kernels.py file #1027#1040
Conversation
…1027 Signed-off-by: Viktor Reutskyy <vreutskyy@nvidia.com>
📝 WalkthroughWalkthroughA new Changes
Sequence Diagram(s)sequenceDiagram
autonumber
participant Solver as SolverMuJoCo
participant Kernels as kernels.py
participant Warp as Warp Arrays
Note over Solver,Kernels: Kernel implementations moved from solver to kernels.py
Solver->>Kernels: convert_body_xforms_to_warp_kernel(inputs)
Kernels-->>Warp: write body transforms, quaternions, DOF indices
Solver->>Kernels: eval_articulation_fk(articulation_state)
Kernels-->>Warp: propagate joint transforms & velocities
Solver->>Kernels: convert_newton_contacts_to_mjwarp_kernel(contacts)
Kernels-->>Warp: write contact entries (distance, normals, frames, params)
Note over Warp: Warp arrays consumed by subsequent physics kernels
Estimated code review effort🎯 3 (Moderate) | ⏱️ ~25 minutes
Possibly related PRs
Suggested reviewers
Pre-merge checks and finishing touches❌ Failed checks (1 warning)
✅ Passed checks (4 passed)
✨ Finishing touches
🧪 Generate unit tests (beta)
Thanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out. Comment |
Codecov Report✅ All modified and coverable lines are covered by tests. 📢 Thoughts on this report? Let us know! |
Signed-off-by: Viktor Reutskyy <vreutskyy@nvidia.com>
Signed-off-by: Viktor Reutskyy <vreutskyy@nvidia.com>
There was a problem hiding this comment.
Actionable comments posted: 0
🧹 Nitpick comments (1)
newton/_src/solvers/mujoco/kernels.py (1)
196-586: Consider removing unused parameters in a future cleanup.Several kernel functions have unused parameters that were carried over from the original implementation:
bodies_per_world(line 196) inconvert_newton_contacts_to_mjwarp_kernelnworld_in(line 212) inconvert_newton_contacts_to_mjwarp_kernelup_axisin multiple kernels (lines 317, 393, 524)body_q(line 525) inapply_mjc_body_f_kernelbody_com(line 546) andjoint_q_start(line 548) inapply_mjc_qfrc_kernelThese parameters can be removed to simplify the function signatures, though this would require updating all call sites.
📜 Review details
Configuration used: Path: .coderabbit.yml
Review profile: CHILL
Plan: Pro
📒 Files selected for processing (2)
newton/_src/solvers/mujoco/kernels.py(1 hunks)newton/_src/solvers/mujoco/solver_mujoco.py(1 hunks)
🧰 Additional context used
🧠 Learnings (12)
📓 Common learnings
Learnt from: shi-eric
Repo: newton-physics/newton PR: 521
File: newton/examples/example_cloth_hanging.py:36-36
Timestamp: 2025-08-12T05:17:34.423Z
Learning: The Newton migration guide (docs/migration.rst) is specifically for documenting how to migrate existing warp.sim functionality to Newton equivalents. New Newton-only features that didn't exist in warp.sim do not need migration documentation.
Learnt from: preist-nvidia
Repo: newton-physics/newton PR: 579
File: newton/examples/example_mujoco.py:350-354
Timestamp: 2025-08-19T12:27:30.629Z
Learning: In Newton examples, Example.use_mujoco is a high-level attribute that controls whether to use MuJoCo solver vs other solvers (like XPBD), while SolverMuJoCo.use_mujoco_cpu controls the backend within MuJoCo (CPU vs Warp). These are separate concepts serving different purposes - the PR rename only applies to the solver parameter, not the Example class attributes.
📚 Learning: 2025-08-12T05:17:34.423Z
Learnt from: shi-eric
Repo: newton-physics/newton PR: 521
File: newton/examples/example_cloth_hanging.py:36-36
Timestamp: 2025-08-12T05:17:34.423Z
Learning: The Newton migration guide (docs/migration.rst) is specifically for documenting how to migrate existing warp.sim functionality to Newton equivalents. New Newton-only features that didn't exist in warp.sim do not need migration documentation.
Applied to files:
newton/_src/solvers/mujoco/kernels.pynewton/_src/solvers/mujoco/solver_mujoco.py
📚 Learning: 2025-09-22T21:08:31.901Z
Learnt from: dylanturpin
Repo: newton-physics/newton PR: 806
File: newton/examples/ik/example_ik_franka.py:121-123
Timestamp: 2025-09-22T21:08:31.901Z
Learning: In the newton physics framework, when creating warp arrays for IK solver joint variables using wp.array(self.model.joint_q, shape=(1, coord_count)), the resulting array acts as a reference/pointer to the original model's joint coordinates, so updates from the IK solver automatically reflect in the model's joint_q buffer used for rendering, despite the general warp documentation suggesting copies are made by default.
Applied to files:
newton/_src/solvers/mujoco/kernels.pynewton/_src/solvers/mujoco/solver_mujoco.py
📚 Learning: 2025-08-12T18:04:06.577Z
Learnt from: nvlukasz
Repo: newton-physics/newton PR: 519
File: newton/_src/solvers/featherstone/kernels.py:75-75
Timestamp: 2025-08-12T18:04:06.577Z
Learning: The Newton physics framework requires nightly Warp builds, which means compatibility concerns with older stable Warp versions (like missing functions such as wp.spatial_adjoint) are not relevant for this project.
Applied to files:
newton/_src/solvers/mujoco/kernels.py
📚 Learning: 2025-08-19T12:27:30.629Z
Learnt from: preist-nvidia
Repo: newton-physics/newton PR: 579
File: newton/examples/example_mujoco.py:350-354
Timestamp: 2025-08-19T12:27:30.629Z
Learning: In Newton examples, Example.use_mujoco is a high-level attribute that controls whether to use MuJoCo solver vs other solvers (like XPBD), while SolverMuJoCo.use_mujoco_cpu controls the backend within MuJoCo (CPU vs Warp). These are separate concepts serving different purposes - the PR rename only applies to the solver parameter, not the Example class attributes.
Applied to files:
newton/_src/solvers/mujoco/kernels.pynewton/_src/solvers/mujoco/solver_mujoco.py
📚 Learning: 2025-08-12T17:51:37.474Z
Learnt from: nvlukasz
Repo: newton-physics/newton PR: 519
File: newton/geometry.py:16-22
Timestamp: 2025-08-12T17:51:37.474Z
Learning: In the Newton public API refactor, common geometry symbols like ParticleFlags and ShapeFlags are now exposed at the top level in newton/__init__.py rather than through newton.geometry. The newton/geometry.py module is intentionally focused only on broad-phase collision detection classes (BroadPhaseAllPairs, BroadPhaseExplicit, BroadPhaseSAP).
Applied to files:
newton/_src/solvers/mujoco/kernels.py
📚 Learning: 2025-09-22T21:08:31.901Z
Learnt from: dylanturpin
Repo: newton-physics/newton PR: 806
File: newton/examples/ik/example_ik_franka.py:121-123
Timestamp: 2025-09-22T21:08:31.901Z
Learning: In the newton physics framework, when creating warp arrays for IK solver joint variables using wp.array(self.model.joint_q, shape=(1, coord_count)), the resulting array acts as a reference/pointer to the original model's joint coordinates, so updates from the IK solver automatically reflect in the model's joint_q buffer used for rendering.
Applied to files:
newton/_src/solvers/mujoco/kernels.pynewton/_src/solvers/mujoco/solver_mujoco.py
📚 Learning: 2025-08-18T15:56:26.587Z
Learnt from: adenzler-nvidia
Repo: newton-physics/newton PR: 552
File: newton/_src/solvers/mujoco/solver_mujoco.py:0-0
Timestamp: 2025-08-18T15:56:26.587Z
Learning: In Newton's MuJoCo solver, when transforming joint axes from Newton's internal frame to MuJoCo's expected frame, use wp.quat_rotate(joint_rot, axis) not wp.quat_rotate_inv(joint_rot, axis). The joint_rot represents rotation from joint-local to body frame, so forward rotation is correct.
Applied to files:
newton/_src/solvers/mujoco/kernels.pynewton/_src/solvers/mujoco/solver_mujoco.py
📚 Learning: 2025-09-25T16:14:22.033Z
Learnt from: dylanturpin
Repo: newton-physics/newton PR: 806
File: newton/_src/sim/ik/ik_objectives.py:0-0
Timestamp: 2025-09-25T16:14:22.033Z
Learning: In NVIDIA Warp's Newton physics library, wp.transform supports direct numerical indexing (e.g., body_tf[0], body_tf[1], body_tf[2] for position and body_tf[3], body_tf[4], body_tf[5], body_tf[6] for quaternion components) to access its elements. This is the standard API used throughout the codebase.
Applied to files:
newton/_src/solvers/mujoco/kernels.py
📚 Learning: 2025-08-19T12:27:30.629Z
Learnt from: preist-nvidia
Repo: newton-physics/newton PR: 579
File: newton/examples/example_mujoco.py:350-354
Timestamp: 2025-08-19T12:27:30.629Z
Learning: In Newton examples, there's a distinction between solver parameters and Example class attributes. The Example class can have its own use_mujoco attribute for controlling example-level behavior (like CUDA graphs, rendering logic), while the solver uses use_mujoco_cpu for backend selection. These serve different purposes and should not be conflated during API renames.
Applied to files:
newton/_src/solvers/mujoco/solver_mujoco.py
📚 Learning: 2025-08-25T11:07:47.818Z
Learnt from: gdaviet
Repo: newton-physics/newton PR: 631
File: newton/examples/mpm/example_anymal_c_walk_on_sand.py:164-169
Timestamp: 2025-08-25T11:07:47.818Z
Learning: In Newton's MuJoCo solver implementation, setting TARGET_POSITION mode on floating base DOFs (first 6 DOFs) does not overconstrain the base - the solver handles this appropriately and maintains expected behavior for floating robots.
Applied to files:
newton/_src/solvers/mujoco/solver_mujoco.py
📚 Learning: 2025-07-14T03:57:29.670Z
Learnt from: Kenny-Vilella
Repo: newton-physics/newton PR: 398
File: newton/examples/example_mujoco.py:352-352
Timestamp: 2025-07-14T03:57:29.670Z
Learning: The use_mujoco option in newton/examples/example_mujoco.py is currently unsupported and causes crashes. The code automatically disables this option with a warning message when users attempt to enable it. This is intentionally kept as a placeholder for future implementation.
Applied to files:
newton/_src/solvers/mujoco/solver_mujoco.py
🧬 Code graph analysis (2)
newton/_src/solvers/mujoco/kernels.py (2)
newton/_src/sim/joints.py (2)
JointMode(81-95)JointType(20-44)newton/_src/sim/collide_unified.py (1)
write_contact(53-149)
newton/_src/solvers/mujoco/solver_mujoco.py (1)
newton/_src/solvers/mujoco/kernels.py (18)
apply_mjc_body_f_kernel(523-538)apply_mjc_control_kernel(505-519)apply_mjc_qfrc_kernel(542-585)convert_body_xforms_to_warp_kernel(762-780)convert_mj_coords_to_warp_kernel(313-385)convert_mjw_contact_to_warp_kernel(460-501)convert_newton_contacts_to_mjwarp_kernel(175-309)convert_warp_coords_to_mj_kernel(389-456)eval_articulation_fk(717-758)repeat_array_kernel(862-869)update_axis_properties_kernel(873-916)update_body_inertia_kernel(812-858)update_body_mass_ipos_kernel(784-808)update_dof_properties_kernel(920-937)update_geom_properties_kernel(1060-1133)update_joint_transforms_kernel(941-1008)update_model_properties_kernel(1049-1056)update_shape_mappings_kernel(1012-1045)
🪛 Ruff (0.14.3)
newton/_src/solvers/mujoco/kernels.py
196-196: Unused function argument: bodies_per_world
(ARG001)
212-212: Unused function argument: nworld_in
(ARG001)
317-317: Unused function argument: up_axis
(ARG001)
393-393: Unused function argument: up_axis
(ARG001)
524-524: Unused function argument: up_axis
(ARG001)
525-525: Unused function argument: body_q
(ARG001)
546-546: Unused function argument: body_com
(ARG001)
548-548: Unused function argument: joint_q_start
(ARG001)
828-828: Ambiguous variable name: I
(E741)
⏰ Context from checks skipped due to timeout of 900000ms. You can increase the timeout in your CodeRabbit configuration to a maximum of 15 minutes (900000ms). (2)
- GitHub Check: Run GPU Benchmarks (Pull Request)
- GitHub Check: Run GPU Unit Tests on AWS EC2 (Pull Request)
🔇 Additional comments (2)
newton/_src/solvers/mujoco/kernels.py (1)
1-38: LGTM!The module header, imports, and type definitions are well-structured. The custom vector types (vec5, vec10f) and constant (MJ_MINVAL) are appropriate for MuJoCo integration.
newton/_src/solvers/mujoco/solver_mujoco.py (1)
46-65: LGTM! Refactoring successfully completed.The kernel imports from
.kernelsare correct and complete. All 18 kernels that were moved tokernels.pyare now properly imported and used throughout the solver implementation. The refactoring achieves its goal of improving code organization while preserving all functionality.
adenzler-nvidia
left a comment
There was a problem hiding this comment.
Thank you, looks straightforward!
Signed-off-by: Viktor Reutskyy <vreutskyy@nvidia.com>
adenzler-nvidia
left a comment
There was a problem hiding this comment.
approving again, thanks!
There was a problem hiding this comment.
Actionable comments posted: 1
🧹 Nitpick comments (3)
newton/_src/solvers/mujoco/kernels.py (3)
770-775: Remove commented debug code.Lines 772-774 contain commented-out quaternion conversion attempts that should be removed to keep the codebase clean.
Apply this diff:
# convert from wxyz to xyzw quat = wp.quat(quat[1], quat[2], quat[3], quat[0]) - # quat = wp.quat(quat[3], quat[0], quat[1], quat[2]) - # quat = wp.quat_identity() - # quat = wp.quat_inverse(quat) body_q[wbi] = wp.transform(pos, quat)
823-823: Consider using a more descriptive variable name.While
Iis standard physics notation for inertia tensor, Python style guides discourage single-letter uppercase variable names to avoid confusion with built-ins. Consider usinginertia_tensororinertiafor better clarity.Apply this diff:
- I = body_inertia[tid] + inertia_tensor = body_inertia[tid] # Calculate eigenvalues and eigenvectors - eigenvectors, eigenvalues = wp.eig3(I) + eigenvectors, eigenvalues = wp.eig3(inertia_tensor)
312-312: Several kernels contain unused parameters.Multiple kernels have unused parameters flagged by static analysis:
up_axisinconvert_mj_coords_to_warp_kernel(line 312)up_axisinconvert_warp_coords_to_mj_kernel(line 388)up_axisandbody_qinapply_mjc_body_f_kernel(lines 519-520)body_comandjoint_q_startinapply_mjc_qfrc_kernel(lines 541, 543)These parameters should either be used in the implementation, removed, or prefixed with underscore if kept for API consistency with other kernels.
Also applies to: 388-388, 519-520, 541-543
📜 Review details
Configuration used: Path: .coderabbit.yml
Review profile: CHILL
Plan: Pro
📒 Files selected for processing (1)
newton/_src/solvers/mujoco/kernels.py(1 hunks)
🧰 Additional context used
🧠 Learnings (7)
📓 Common learnings
Learnt from: shi-eric
Repo: newton-physics/newton PR: 521
File: newton/examples/example_cloth_hanging.py:36-36
Timestamp: 2025-08-12T05:17:34.423Z
Learning: The Newton migration guide (docs/migration.rst) is specifically for documenting how to migrate existing warp.sim functionality to Newton equivalents. New Newton-only features that didn't exist in warp.sim do not need migration documentation.
Learnt from: preist-nvidia
Repo: newton-physics/newton PR: 579
File: newton/examples/example_mujoco.py:350-354
Timestamp: 2025-08-19T12:27:30.629Z
Learning: In Newton examples, Example.use_mujoco is a high-level attribute that controls whether to use MuJoCo solver vs other solvers (like XPBD), while SolverMuJoCo.use_mujoco_cpu controls the backend within MuJoCo (CPU vs Warp). These are separate concepts serving different purposes - the PR rename only applies to the solver parameter, not the Example class attributes.
📚 Learning: 2025-08-12T05:17:34.423Z
Learnt from: shi-eric
Repo: newton-physics/newton PR: 521
File: newton/examples/example_cloth_hanging.py:36-36
Timestamp: 2025-08-12T05:17:34.423Z
Learning: The Newton migration guide (docs/migration.rst) is specifically for documenting how to migrate existing warp.sim functionality to Newton equivalents. New Newton-only features that didn't exist in warp.sim do not need migration documentation.
Applied to files:
newton/_src/solvers/mujoco/kernels.py
📚 Learning: 2025-09-22T21:08:31.901Z
Learnt from: dylanturpin
Repo: newton-physics/newton PR: 806
File: newton/examples/ik/example_ik_franka.py:121-123
Timestamp: 2025-09-22T21:08:31.901Z
Learning: In the newton physics framework, when creating warp arrays for IK solver joint variables using wp.array(self.model.joint_q, shape=(1, coord_count)), the resulting array acts as a reference/pointer to the original model's joint coordinates, so updates from the IK solver automatically reflect in the model's joint_q buffer used for rendering, despite the general warp documentation suggesting copies are made by default.
Applied to files:
newton/_src/solvers/mujoco/kernels.py
📚 Learning: 2025-08-12T18:04:06.577Z
Learnt from: nvlukasz
Repo: newton-physics/newton PR: 519
File: newton/_src/solvers/featherstone/kernels.py:75-75
Timestamp: 2025-08-12T18:04:06.577Z
Learning: The Newton physics framework requires nightly Warp builds, which means compatibility concerns with older stable Warp versions (like missing functions such as wp.spatial_adjoint) are not relevant for this project.
Applied to files:
newton/_src/solvers/mujoco/kernels.py
📚 Learning: 2025-08-19T12:27:30.629Z
Learnt from: preist-nvidia
Repo: newton-physics/newton PR: 579
File: newton/examples/example_mujoco.py:350-354
Timestamp: 2025-08-19T12:27:30.629Z
Learning: In Newton examples, Example.use_mujoco is a high-level attribute that controls whether to use MuJoCo solver vs other solvers (like XPBD), while SolverMuJoCo.use_mujoco_cpu controls the backend within MuJoCo (CPU vs Warp). These are separate concepts serving different purposes - the PR rename only applies to the solver parameter, not the Example class attributes.
Applied to files:
newton/_src/solvers/mujoco/kernels.py
📚 Learning: 2025-09-22T21:08:31.901Z
Learnt from: dylanturpin
Repo: newton-physics/newton PR: 806
File: newton/examples/ik/example_ik_franka.py:121-123
Timestamp: 2025-09-22T21:08:31.901Z
Learning: In the newton physics framework, when creating warp arrays for IK solver joint variables using wp.array(self.model.joint_q, shape=(1, coord_count)), the resulting array acts as a reference/pointer to the original model's joint coordinates, so updates from the IK solver automatically reflect in the model's joint_q buffer used for rendering.
Applied to files:
newton/_src/solvers/mujoco/kernels.py
📚 Learning: 2025-08-18T15:56:26.587Z
Learnt from: adenzler-nvidia
Repo: newton-physics/newton PR: 552
File: newton/_src/solvers/mujoco/solver_mujoco.py:0-0
Timestamp: 2025-08-18T15:56:26.587Z
Learning: In Newton's MuJoCo solver, when transforming joint axes from Newton's internal frame to MuJoCo's expected frame, use wp.quat_rotate(joint_rot, axis) not wp.quat_rotate_inv(joint_rot, axis). The joint_rot represents rotation from joint-local to body frame, so forward rotation is correct.
Applied to files:
newton/_src/solvers/mujoco/kernels.py
🧬 Code graph analysis (1)
newton/_src/solvers/mujoco/kernels.py (2)
newton/_src/sim/joints.py (2)
JointMode(81-95)JointType(20-44)newton/_src/sim/collide_unified.py (1)
write_contact(53-149)
🪛 Ruff (0.14.3)
newton/_src/solvers/mujoco/kernels.py
191-191: Unused function argument: bodies_per_world
(ARG001)
207-207: Unused function argument: nworld_in
(ARG001)
312-312: Unused function argument: up_axis
(ARG001)
388-388: Unused function argument: up_axis
(ARG001)
519-519: Unused function argument: up_axis
(ARG001)
520-520: Unused function argument: body_q
(ARG001)
541-541: Unused function argument: body_com
(ARG001)
543-543: Unused function argument: joint_q_start
(ARG001)
823-823: Ambiguous variable name: I
(E741)
⏰ Context from checks skipped due to timeout of 900000ms. You can increase the timeout in your CodeRabbit configuration to a maximum of 15 minutes (900000ms). (2)
- GitHub Check: Run GPU Benchmarks (Pull Request)
- GitHub Check: Run GPU Unit Tests on AWS EC2 (Pull Request)
🔇 Additional comments (7)
newton/_src/solvers/mujoco/kernels.py (7)
27-28: Custom vector type definitions look good.Both
vec5andvec10use consistent definition patterns withwp.float32. While a past review mentioned alignment, the current definitions are already consistent in structure and dtype.
335-344: Quaternion conversions are correctly implemented.The coordinate conversion kernels properly handle quaternion format differences between MuJoCo (wxyz) and Warp (xyzw). The conversions are consistent across FREE and BALL joint types in both directions.
Also applies to: 359-369, 410-420, 436-440
353-357: Angular velocity transformations are correct.For FREE joints, the angular velocities are properly transformed between coordinate frames using
wp.quat_rotate(MuJoCo→Warp) andwp.quat_rotate_inv(Warp→MuJoCo), maintaining consistency with the quaternion rotations.Also applies to: 429-433
583-708: Forward kinematics implementation is comprehensive and correct.The
eval_single_articulation_fkfunction properly handles all joint types (PRISMATIC, REVOLUTE, BALL, FREE, DISTANCE, D6) with correct transform composition and velocity propagation through the kinematic tree. The D6 joint implementation correctly separates linear and angular DOF processing.
169-304: Contact conversion kernel correctly transforms Newton contacts to MuJoCo format.The kernel properly handles coordinate transformations, contact distance computation with thickness, contact frame construction, and parameter derivation. The use of the midpoint for contact position (line 249) is consistent with XPBD solver conventions.
831-842: Eigenvalue sorting logic is correct.The bubble sort implementation correctly sorts 3 eigenvalues and their corresponding eigenvectors in descending order. For a fixed size of 3 elements, this approach is simple and efficient.
1-24: File structure and organization are well-designed.The file has a clean layout with proper license header, logical grouping of custom types, constants, utility functions, and kernels. The separation of
@wp.funchelpers from@wp.kernelimplementations is clear and maintainable.
…1027 (newton-physics#1040) Signed-off-by: Viktor Reutskyy <vreutskyy@nvidia.com>
…1027 (newton-physics#1040) Signed-off-by: Viktor Reutskyy <vreutskyy@nvidia.com>
Description
This PR refactors the SolverMuJoCo implementation by extracting all Warp kernels and functions into a separate
kernels.pyfile. This improves code organization and maintainability by separating the GPU kernel implementations from the main solver logic.Changes:
newton/_src/solvers/mujoco/kernels.pycontaining:@wp.kernel)@wp.func)vec5,vec10f)MJ_MINVAL)solver_mujoco.pyto import kernels from the new modulesolver_mujoco.pysolver_mujoco.pyThis is a pure refactoring change with no functional modifications. All existing functionality remains intact.
Closes #1027
Newton Migration Guide
Please ensure the migration guide for warp.sim users is up-to-date with the changes made in this PR.
docs/migration.rstis up-to dateBefore your PR is "Ready for review"
newton/tests/test_examples.py)pre-commit run -aSummary by CodeRabbit