[OVPHYSX] Articulation rewrite (data class + asset class + kernels)

[AntoineRichard]

Description

Drastic rewrite of OVPhysX Articulation and ArticulationData so they follow the same shape as the post-refactor OVPhysX RigidObject from #5426, with the API surface mirroring Newton Articulation and behavior parity with PhysX Articulation. Single-PR atomic rewrite, clean break (no deprecation aliases for OVPhysX-introduced renames; framework-inherited deprecated shims kept).

The OVPhysX articulation diverged significantly from the rest of the framework conventions. This PR brings it back in line: same docstring template, same section ordering, same naming, same internal patterns, same lifecycle.

Important

Stacked on #5426 ([OVPHYSX] RigidObject + RigidObjectData asset). Review only the 16 articulation-specific commits at the tip of this branch — every commit before that lands via #5426. Once #5426 merges to develop, this PR will rebase cleanly onto develop.

Fixes # (none — internal refactor; no associated issue)

Architectural changes

OVPhysX RigidObject is the design template. It has navigated the hybrid OVPhysX surface — Newton-style mask+index dual API + PhysX-style CPU-only bindings via pinned-host staging à la #5329 + pull-to-refresh binding.read(target):

• Eager TimestampedBufferWarp allocation in _create_buffers (single source of truth — no _invalidate_caches / _ensure_*_buffers machinery).
• Pinned-host CPU staging buffers for every CPU-only binding (mass, COM, inertia, all DOF properties).
• _binding_read / _binding_write / _stage_to_pinned_cpu helpers route CPU-only types through pinned-host memory.
• Every public property returns a ProxyArray (warp + torch dual view); raw wp.array for one-shot config buffers.
• Counts and names (num_instances, num_bodies, num_joints, body_names, joint_names, ...) demoted from @property to plain instance attributes.
• Dual mask+index API on every writer/setter (*_index accepts partial data; *_mask accepts full data with a wp.bool mask).
• All write_* / set_* parameters are kwarg-only after *,. No positional. No full_data flag anywhere.
• The deprecated _write_body_state plumbing layer is removed; deprecated state-writer shims (write_root_state_to_sim, etc.) call the public write_*_to_sim_index methods directly, mirroring RigidObject.

Articulation-specific surface mirrors Newton 1-to-1: joint-state writers, joint-property writers (CPU-only), body-property setters (multi-body shape), joint-command target setters, external-wrench setters via WrenchComposer, fixed/spatial tendon setters, deprecated state-writer shims, full actuator pipeline (compute, _apply_actuator_model, _process_actuators_cfg).

Files changed

• source/isaaclab_ovphysx/isaaclab_ovphysx/assets/articulation/articulation.py — full rewrite (~3863 lines, matches Newton).
• source/isaaclab_ovphysx/isaaclab_ovphysx/assets/articulation/articulation_data.py — full rewrite (~2504 lines).
• source/isaaclab_ovphysx/isaaclab_ovphysx/assets/kernels.py — gained 6 articulation kernels migrated from the stop-gap kernels_old.py (now deleted): _compose_root_com_pose, _compute_heading, _copy_first_body, _projected_gravity, _world_vel_to_body_ang, _world_vel_to_body_lin. Plus 2 new joint-property kernels (write_joint_position_limit_to_buffer_index/mask for trailing-dim-2 limits, write_joint_friction_to_buffer_index/mask for the broadcast-coefficient pattern).
• source/isaaclab_ovphysx/isaaclab_ovphysx/assets/kernels_old.py — deleted.
• source/isaaclab_ovphysx/test/assets/test_articulation.py — verbatim PhysX test mirror (~210 parametrizations) with PhysX-internal root_view.X assertions adapted to the OVPhysX bindings dict and omni.physx.scripts-dependent tests xfailed; mirrors the precedent from the RigidObject test mirror in [OVPHYSX] RigidObject + RigidObjectData asset #5426.

Type of change

• Breaking change (existing functionality will not work without user modification — OVPhysX is at 0.2.x, clean break is acceptable per semver-on-0.x; no deprecation aliases for OVPhysX-introduced renames).
• Code modernization / refactor.

Validation

Three layers, run on GPU and CPU separately (the wheel's process-global device-mode lock makes a single invocation lock to one device):

1. Real-backend port — test_articulation.py (verbatim PhysX mirror). ./scripts/run_ovphysx.sh -m pytest <path> -k 'cuda:0' and ... -k 'cpu'. Expected end state: each pass shows <X> passed, <Y> xfailed, 0 failed. Every xfail carries a reason pointing at the wheel-gaps spec.
2. Cross-backend interface — source/isaaclab/test/assets/test_articulation_iface.py will gain an ovphysx backend, mirroring the rigid-object iface treatment from [OVPHYSX] RigidObject + RigidObjectData asset #5426.
3. API consistency audit — per-method side-by-side checklist comparing Newton, RigidObject (post-refactor), and the rewritten Articulation; verifies method name, kwarg-only signature, parameter order, return type, docstring template, section-header placement.

Status

Active triage — not yet ready for review.

• ✅ Implementation complete (all writers, setters, properties, lifecycle, actuator pipeline).
• ✅ Initial GPU root-cause bug fixed: _read_transform_binding now routes BODY_COM_POSE through _binding_read so the wheel's CPU-only-binding device check is satisfied on a GPU sim.
• ✅ Verbatim PhysX-internals assertions (root_view.max_dofs == shared_metatype.dof_count, link_paths[0] round-trip) adapted to the OVPhysX bindings dict — they now check binding.shape[1] == num_joints / num_bodies for each per-DOF / per-link binding.
• 🔄 In-flight: tendon-init device-routing bug. _read_initial_properties reads FIXED_TENDON_* / SPATIAL_TENDON_* via numpy assuming CPU residency, but the wheel exposes them as GPU-resident (consistent with PhysX's set_fixed_tendon_properties not cloning to CPU). Plan is to remove tendon types from _CPU_ONLY_TYPES and read them directly into the sim-device buffer.
• ⏳ Pending: cross-backend test_articulation_iface.py extension, API consistency audit, CHANGELOG + version bump (0.2.x → 0.3.0).

Checklist

• I have read and understood the contribution guidelines
• I have run the pre-commit checks with ./isaaclab.sh --format
• I have made corresponding changes to the documentation
• My changes generate no new warnings
• I have added tests that prove my fix is effective or that my feature works (the verbatim PhysX test mirror is the contract; bug-fixing in progress)
• I have updated the changelog and the corresponding version in the extension's config/extension.toml file (deferred to final-pass commit)
• I have added my name to the CONTRIBUTORS.md or my name already exists there

[isaaclab-review-bot]

🤖 Isaac Lab Review Bot

Summary

This PR is a massive (~10K+ lines) atomic rewrite of the OVPhysX Articulation and ArticulationData classes to mirror the post-refactor RigidObject shape and align with Newton/PhysX conventions. The changes include dual mask+index API patterns, pinned-host CPU staging for CPU-only bindings, eager buffer allocation, and comprehensive test coverage mirroring the PhysX test suite. The implementation appears architecturally sound but has several correctness issues in kernel dispatch dimensions, missing buffer invalidations, and test fixture lifecycle problems.

Architecture Impact

High impact. This PR touches the core asset abstraction layer for articulations in OVPhysX:

• Articulation and ArticulationData are foundational classes used by every RL environment running on OVPhysX
• The kernels.py module is shared between RigidObject and Articulation
• Test fixtures install session-scoped monkey patches on OvPhysxManager that affect all OVPhysX tests
• The tensor_types.py changes add new constants that downstream code may depend on

Breaking changes are acceptable per semver-0.x, but the deprecation shims must work correctly.

Implementation Verdict

Significant concerns. Several kernel launch dimension bugs and missing cache invalidations could cause incorrect simulation state.

Test Coverage

The test suite is a verbatim port of PhysX tests (~210 parametrizations), which is excellent for API parity. However:

• The session-scoped fixture approach (_ovphysx_session_patches) is fragile and could mask test isolation bugs
• Missing explicit tests for the new tendon property writers
• Missing regression tests for the deprecated state-writer shims

CI Status

No CI checks available yet — cannot verify if the test suite passes.

Findings

🔴 Critical: articulation.py:1432-1438 — write_joint_velocity_to_sim_index syncs _previous_joint_vel with wrong launch dimensions

wp.launch(
    shared_kernels.write_2d_data_to_buffer_with_indices,
    dim=(env_ids.shape[0], joint_ids.shape[0]),
    inputs=[velocity, env_ids, joint_ids],
    outputs=[self._data._previous_joint_vel],
    device=self._device,
)

When joint_ids is a subset (not all joints), this kernel writes velocity (which has shape (len(env_ids), len(joint_ids))) into _previous_joint_vel at the specified indices. However, if _previous_joint_vel was never initialized from the current simulation state, the unwritten columns will contain stale data from the previous step, causing incorrect FD acceleration on the unwritten joints. The same issue exists in write_joint_velocity_to_sim_mask at line 1478. Fix: Either always sync the full _previous_joint_vel from the current binding before the partial overwrite, or document that partial joint writes produce undefined acceleration for unwritten joints.

🔴 Critical: articulation.py:597-601 — _process_cfg calls update_soft_joint_pos_limits before joint_pos_limits is populated

if D > 0:
    wp.launch(
        update_soft_joint_pos_limits,
        dim=(N, D),
        inputs=[self._data.joint_pos_limits, cfg.soft_joint_pos_limit_factor],

At this point in _process_cfg, self._data.joint_pos_limits refers to the property which reads from _joint_pos_limits TimestampedBuffer. However, _read_initial_properties() is called in ArticulationData._create_buffers() (line 233 in articulation_data.py), which populates _joint_pos_limits. The call order in _initialize_impl is: construct data → _create_buffers() → _process_cfg(). So this should be fine, but the timestamp check in joint_pos_limits property (line ~870) may return stale data if _sim_timestamp is still 0.0. Verify that _read_initial_properties sets the timestamp correctly.

🔴 Critical: articulation_data.py:155-156 — _binding_getter is stored but never validated

self._binding_getter = binding_getter

The binding_getter callable is passed from Articulation._initialize_impl and used in _get_binding. However, if binding_getter is None (which happens during mock testing or if the caller forgets to pass it), _get_binding at line 377 will fail:

if self._binding_getter is not None:
    return self._binding_getter(tensor_type)
return self._bindings.get(tensor_type)

This fallback to self._bindings.get() is correct, but the docstring says "lazily creates bindings on first access" which is misleading when no getter is provided. Minor but could cause confusion during debugging.

🟡 Warning: articulation.py:2285-2295 — set_masses_index doesn't invalidate body_mass timestamp

def set_masses_index(
    self,
    *,
    masses: torch.Tensor | wp.array,
    body_ids: Sequence[int] | torch.Tensor | wp.array | None = None,
    env_ids: Sequence[int] | torch.Tensor | wp.array | None = None,
) -> None:

After writing to the binding, the cached _body_mass buffer is updated via the kernel, but self._data._body_mass.timestamp is never set to self._data._sim_timestamp. This means subsequent reads of body_mass property will re-read from the binding (which is correct, but inefficient) rather than using the just-written cached value. The same pattern appears in set_coms_index and set_inertias_index. Fix: Add self._data._body_mass.timestamp = self._data._sim_timestamp after the kernel write.

🟡 Warning: test_articulation.py:140-150 — _patched_warmup_and_load doesn't handle first-run device mismatch

physx = _PHYSX_BY_DEVICE.get(cache_key)
if physx is None:
    _orig_warmup_and_load(OvPhysxManager)
    _PHYSX_BY_DEVICE[cache_key] = OvPhysxManager._physx
    return

If a test runs on cuda:0 first, then another test somehow requests cpu (despite the _ovphysx_skip_other_device fixture), the code will try to create a second PhysX instance which will crash due to the device-mode lock. The skip fixture mitigates this, but if a test doesn't parametrize on device, it will slip through. The fixture at line 210 handles this:

device = callspec.params.get("device") if callspec is not None else None
if device is None:
    return  # <-- allows through

Tests without device param should still work, but could behave unexpectedly if they assume a specific device.

🟡 Warning: kernels.py:785-795 — write_body_inertia_to_buffer_index assumes contiguous inertia layout

@wp.kernel
def write_body_inertia_to_buffer_index(
    inertia: wp.array3d(dtype=wp.float32),  # (K, L, 9)
    env_ids: wp.array(dtype=wp.int32),
    body_ids: wp.array(dtype=wp.int32),
    out: wp.array3d(dtype=wp.float32),  # (N, B, 9)
):

The kernel accesses inertia[i, j, k] where i is the launch index (not env_ids[i]), which is correct. But if inertia is a non-contiguous view (e.g., from slicing a larger tensor), Warp may produce incorrect results. The set_inertias_index method at line 2382 passes user data directly without ensuring contiguity. Fix: Add inertia = wp.array(inertia, copy=True) if the array might be non-contiguous, or document the requirement.

🔵 Improvement: articulation.py:1173-1175 — Duplicate launch for set_root_link_pose_to_sim_index kernel

wp.launch(
    shared_kernels.set_root_link_pose_to_sim_index,
    dim=env_ids.shape[0],
    inputs=[root_pose, env_ids],
    outputs=[self.data.root_link_pose_w],  # <-- This is a ProxyArray property

Using self.data.root_link_pose_w as a kernel output triggers the property getter, which may read from the binding. The kernel then overwrites the buffer. This is inefficient — should use self.data._root_link_pose_w.data directly (the underlying TimestampedBuffer's data). Same pattern in all write_root_* methods.

🔵 Improvement: articulation_data.py:307-310 — Unused _cpu_joint_friction_* buffers allocated unconditionally

self._cpu_joint_friction_static = wp.zeros((N, D), dtype=wp.float32, device="cpu", pinned=True)
self._cpu_joint_friction_dynamic = wp.zeros((N, D), dtype=wp.float32, device="cpu", pinned=True)
self._cpu_joint_friction_viscous = wp.zeros((N, D), dtype=wp.float32, device="cpu", pinned=True)

These buffers are allocated even on CPU device (if dev != "cpu": check is at line 279, but these are inside the block). However, the friction properties are never written via the public API (no write_joint_friction_* methods exist), so these buffers are allocated but never used. Consider making allocation lazy or removing if friction writes aren't supported yet.

[isaaclab-review-bot]

🤖 Isaac Lab Review Bot

Summary

This follow-up review checks the new commits against previously identified concerns. The PR has grown significantly with the addition of RigidObject and RigidObjectData classes and extensive test coverage.

Previous Findings Status

1. write_joint_velocity_to_sim_index FD acceleration issue — Still present. The kernel writes partial velocity to _previous_joint_vel at lines 1432-1438, which can cause incorrect FD acceleration for unwritten joints when joint_ids is a subset.

2. _process_cfg soft limits timing — Addressed. The call order ensures _read_initial_properties() runs before _process_cfg().

3. _binding_getter validation — Acknowledged as minor, behavior is correct with fallback.

4. set_masses_index timestamp invalidation — Still present at line 2285. After writing masses, self._data._body_mass.timestamp is not updated.

5. Test fixture device mismatch — Mitigated by _ovphysx_skip_other_device fixture.

6. write_body_inertia_to_buffer_index contiguity — Still present, no contiguity enforcement added.

7. ProxyArray property as kernel output — Still present throughout write_root_* methods.

8. Unused friction staging buffers — Still allocated but unused.

New Findings

🔴 Critical: rigid_object.py:695-700 — _get_cpu_env_ids returns wrong pinned buffer on partial writes

def _get_cpu_env_ids(self, env_ids: wp.array) -> wp.array:
    if env_ids.ptr == self._ALL_INDICES.ptr:
        return self._cpu_env_ids_all
    return wp.clone(env_ids, device="cpu")

The wp.clone(..., device="cpu") creates an unpinned CPU array. For CPU-only bindings that require pinned memory for async transfers, this will silently work but defeat the performance optimization. Should use wp.zeros(..., pinned=True) + wp.copy() pattern like _get_cpu_env_mask does.

🟡 Warning: rigid_object_data.py:372-378 — body_com_acc_w FD uses stale _previous_body_com_vel after velocity writes

if self._previous_body_com_vel is None:
    self._previous_body_com_vel = wp.clone(self.body_com_vel_w.warp)

After write_root_com_velocity_to_sim_index, the _previous_body_com_vel is not updated, causing the next body_com_acc_w access to compute acceleration from the old velocity. This mirrors the articulation issue but for rigid objects.

Implementation Verdict

Minor fixes needed. The previous critical FD acceleration issues remain unaddressed. The new RigidObject implementation has the same pattern issue.

[isaaclab-review-bot]

🤖 Isaac Lab Review Bot

Summary

This is a follow-up review checking commits since the previous review. The core Articulation and RigidObject implementations remain largely unchanged.

Previous Findings Status

All previously identified issues remain unaddressed:

1. FD acceleration partial-write bug (write_joint_velocity_to_sim_index lines 1432-1438) — Still present
2. set_masses_index missing timestamp update (line 2285) — Still present
3. _get_cpu_env_ids creates unpinned buffer (rigid_object.py:695-700) — Still present
4. body_com_acc_w stale _previous_body_com_vel after velocity writes — Still present
5. ProxyArray properties used as kernel outputs — Still present throughout

New Findings

No new issues introduced by the recent commits. The changes appear to be minor adjustments that don't introduce additional bugs.

Implementation Verdict

Minor fixes still needed. The FD acceleration issues in both Articulation and RigidObject can cause incorrect acceleration reporting after partial velocity writes. Consider syncing _previous_joint_vel / _previous_body_com_vel in the velocity write paths to avoid spurious acceleration spikes on reset.

[isaaclab-review-bot]

🤖 Isaac Lab Review Bot

Summary

This is a follow-up review after new commits. Checking whether previous findings were addressed.

Previous Findings Status

1. FD acceleration partial-write bug — ✅ ADDRESSED. The write_joint_velocity_to_sim_index and write_joint_velocity_to_sim_mask methods now sync _previous_joint_vel via a second kernel launch (lines 1367-1373 and 1410-1416 in articulation.py), preventing spurious acceleration spikes after velocity writes.

2. set_masses_index missing timestamp update — Still present in articulation.py line 2283 (the _body_mass timestamp is not bumped after a write).

3. _get_cpu_env_ids creates unpinned buffer — Still present in rigid_object.py:695-700.

4. body_com_acc_w stale _previous_body_com_vel after velocity writes — Still present in RigidObject.

5. ProxyArray properties used as kernel outputs — Still present throughout.

New Findings

🔵 Improvement: articulation.py:2283 — set_masses_index should update _body_mass.timestamp

The articulation joint-velocity writers now correctly sync the previous-velocity buffer, but the mass writer path still doesn't bump the timestamp on self._data._body_mass after a write, so subsequent reads may return stale cached values until the next update() call.

Implementation Verdict

Minor fixes still needed — but the critical FD acceleration bug for articulations is now fixed. Ship it for articulation; the remaining issues are edge-case correctness improvements rather than blocking bugs.

[isaaclab-review-bot]

🤖 Isaac Lab Review Bot

Summary

Follow-up review after new commits. Checking previous findings and scanning for new issues.

Previous Findings Status

1. FD acceleration partial-write bug — ✅ Previously addressed.

2. set_masses_index missing timestamp update — ✅ NOW ADDRESSED. Line 2287 in articulation.py now updates self._data._body_mass.timestamp = self._data._sim_timestamp after the write.

3. _get_cpu_env_ids creates unpinned buffer — Still present in rigid_object.py:695-700, but this is a minor performance concern, not a correctness bug.

4. body_com_acc_w stale _previous_body_com_vel after velocity writes — Still present in RigidObject, but the articulation side (which is the focus of this PR) handles this correctly.

5. ProxyArray properties used as kernel outputs — This is a pattern choice, not a bug.

New Issues

No new correctness issues introduced by the latest commits.

Implementation Verdict

Ship it. The critical articulation bugs have been addressed. The remaining RigidObject edge cases (_get_cpu_env_ids pinning, acceleration buffer sync) are pre-existing patterns from #5426 and are minor performance/edge-case concerns rather than blocking issues for this articulation rewrite PR.