Skip to content

[Warp] Part 2/8: Write the PhysX tendon mask setters through the masks - #8824

Open
hujc7 wants to merge 7 commits into
isaac-sim:mainfrom
hujc7:jichuanh/physx-tendon-mask-setters
Open

hujc7 wants to merge 7 commits into
isaac-sim:mainfrom
hujc7:jichuanh/physx-tendon-mask-setters

Conversation

@hujc7

@hujc7 hujc7 commented Oct 26, 2026 •

Copy link
Copy Markdown
Collaborator

Review Map

PR Status Depends on Exact changes
#8821 Part 1/8: shared id and mask helpers, mask writers without host sync — —
📌 #8824 Part 2/8: PhysX tendon mask setters through the masks #8821 link
#6646 Part 3/8: mask-correct partial resets #8821 link
#8783 Part 4/8: shared Warp RNG state #6646 link
#6607 Part 5/8: per-term capture, stable stepping parity #8783 link
#8228 Part 6/8: Warp commands, Rough tasks, live term changes #6607 link
#6582 Part 7/8: Shadow Warp twin, Allegro twin fixes #6607 link
#7012 Part 8/8: --video on the Warp frontend — —
---
config:
  look: neo
  flowchart:
    curve: basis
---
flowchart TD
  P1("<b>#8821</b> · Part 1<br/>id/mask helpers")
  P2("📌 <b>#8824</b> · Part 2<br/>PhysX tendon masks")
  P3("<b>#6646</b> · Part 3<br/>masked resets")
  P4("<b>#8783</b> · Part 4<br/>WarpRng")
  P5("<b>#6607</b> · Part 5<br/>per-term capture")
  P6("<b>#8228</b> · Part 6<br/>commands, Rough")
  P7("<b>#6582</b> · Part 7<br/>Shadow twin")
  P8("<b>#7012</b> · Part 8<br/>--video, independent")
  P1 --> P2
  P1 --> P3 --> P4 --> P5
  P5 --> P6
  P5 --> P7
  style P2 stroke-width:3px
Loading

Summary

The 11 PhysX tendon set_*_mask setters no longer turn their masks into ids with a GPU→host wait: they fill the
tendon buffers through the masks, 83-91 → 18.5 µs per call (5x) with half the environments selected.

Stacked on #8821 (shared id and mask helpers). Until it merges this PR also shows #8821's commits; the change
here is the last commit, Write the PhysX tendon mask setters through the masks.

Description

The fixed and spatial tendon mask setters of the PhysX Articulation only fill IsaacLab's tendon buffers (stiffness,
damping, limit stiffness, position limit, rest length, offset, position target); the push to PhysX happens later in
write_*_tendon_properties_to_sim_*, which stays index based because the PhysX tensor API takes indices. The setters
still turned both masks into ids first (torch.nonzero, which waits for the GPU) and then called the index setters.

Each setter now resolves its masks, checks the shape and launches a masked kernel, the same steps as the Newton tendon
mask setters.

Performance

Desk RTX PRO 6000, PhysX contract test articulation (4096 envs), host µs per call including the GPU wait, median of 7
batches, 2 alternated runs per side.

Setter Selection #8821 (µs) this PR (µs) Speed-up
stiffness / damping / limit stiffness / rest length / offset (fixed and spatial) half envs, 2 tendons 83-91 18.5-19 5x
position limit half envs, 2 tendons 111-112 22-27 4.5x
position target half envs, 2 tendons 174-176 34-42 4.5x
any of the above all 19 17-19 same

No task in the repository calls these mask setters per step, so task FPS does not change.

Type of change

  • New feature (non-breaking change which adds functionality)

Checklist

Docker and GPU tests run on demand. Push the commits you want tested, then
comment run-ci on the pull request.

  • I have read and understood the contribution guidelines
  • I have run the pre-commit checks with uv run isaaclab --format
  • I have made corresponding changes to the documentation
  • My changes generate no new warnings
  • I have validated the changed behavior with focused coverage and applied the test-audit criteria to any test changes
  • I have added a changelog fragment under source/<pkg>/changelog.d/ for every touched package (do not edit CHANGELOG.rst or bump package versions — CI handles that)
  • I have added my name to the CONTRIBUTORS.md or my name already exists there

hujc7 added 7 commits October 26, 2026 02:48
…pers

resolve_ids resolves ids or a boolean mask to ids on a device, as resolve_1d_mask does to a mask. resolve_1d_mask takes an optional scratch mask, copies a Warp mask on another device, and accepts a boolean torch mask on its device inside a graph capture. The experimental package now re-exports it instead of keeping a copy.
…red helpers

SensorBase, Camera and WrenchComposer resolve env_ids, env_mask and body_ids through resolve_ids and resolve_1d_mask instead of their own copies. The test doubles that build these objects without their initializer get the env masks a real sensor creates.
The nodal mask writers of BaseDeformableObject, which the OVPhysX deformable object inherits, and of the PhysX deformable object called wp.nonzero on a Warp array, which raised a RuntimeError. They now resolve the mask to ids with resolve_ids and call the index writers.
…sk writers through masks

The Newton assets resolve ids and masks through resolve_ids and resolve_1d_mask. The RigidObjectCollection mask writers and the fixed-tendon mask setters write through the masks instead of converting them to ids, which waited for the GPU on every call; the joint-coordinate writes reuse a scratch mask; one generic masked 2-D write kernel replaces two copies.
The OVPhysX assets resolve env, body, joint and tendon selections through resolve_ids and resolve_1d_mask instead of their own copies.
The PhysX assets resolve env, body, joint and tendon selections through resolve_ids and resolve_1d_mask instead of their own copies. The sensor test doubles get the env masks a real sensor creates.
The eleven fixed and spatial tendon mask setters only fill IsaacLab's tendon buffers, yet turned both masks into ids with a host wait first. Each now resolves its masks, checks the shape and launches a masked kernel, as the Newton tendon mask setters do. The push to PhysX stays index based.
@hujc7

hujc7 commented Oct 26, 2026

Copy link
Copy Markdown
Collaborator Author

run-ci

@greptile-apps

greptile-apps Bot commented Oct 26, 2026 •

Copy link
Copy Markdown
Contributor

RetriggerView in GreptileConfidence Score: 4/5

[Medium impact] The changes reach selection and write paths across several simulation backends.

Restore negative-step environment selections before merging; also add exact mask-length checks to the new tendon writers.

Findings

  1. P1 Negative slices stop working ▶
  2. P2 Short masks escape checks ▶

Summary

The PR shares index and mask helpers across backends and replaces tendon mask-to-index writes with direct masked writes.

  • Asset selectors use shared helpers to handle ids and masks consistently.
  • Tendon mask setters update selected buffers without first building ids.
  • Newton mask writers update selected bodies and environments directly.
  • Deformable mask writes reach the selected environments through the backend writer.

Diagram

%%{init: {'theme': 'neutral'}}%%
flowchart LR
    A[Full-size tendon values and masks] --> B[Resolve masks]
    B --> C[Write selected buffer cells]
    C --> D[Cached tendon properties]
    D --> E[Explicit property write]
    E --> F[PhysX]
    G[Target length] --> H[Rest length minus target]
    H --> C
    C --> I[Mark target dirty]
    I --> J[write_data_to_sim]
    J --> E
Loading

Reviews (1) · Last reviewed commit: "Write the PhysX tendon mask setters thro..." · Reviewed by Greptile


"""Warp utility functions and shared kernels for isaaclab_experimental."""

from isaaclab.utils.warp.utils import resolve_1d_mask

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

P1 Negative slices stop working

Re-exporting the shared resolve_1d_mask breaks calls such as ManagerBasedEnvWarp.resolve_env_mask(env_ids=slice(None, None, -2)). The shared helper applies the slice directly to a PyTorch tensor, which rejects negative steps. The removed helper expanded slices with range, so this selection previously worked. Preserve that conversion in the shared helper before replacing the experimental implementation.

Comment on lines +2805 to +2812
env_mask = self._resolve_mask(env_mask, self._ALL_ENV_MASK)
fixed_tendon_mask = self._resolve_mask(fixed_tendon_mask, self._ALL_FIXED_TENDON_MASK)
self.assert_shape_and_dtype_mask(stiffness, (env_mask, fixed_tendon_mask), wp.float32, "stiffness")
wp.launch(
articulation_kernels.float_data_to_buffer_with_mask
if isinstance(stiffness, float)
else shared_kernels.write_2d_data_to_buffer_with_mask_kernel(wp.float32),
dim=(self.num_instances, self.num_fixed_tendons),

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

P2 Short masks escape checks

assert_shape_and_dtype_mask checks the data against the supplied mask lengths, not the full asset dimensions. A short environment mask with scalar stiffness=1.0 therefore passes, but this launch still visits every environment and reads beyond the mask. The old path converted the mask into valid indices and visited only those entries.

Add exact mask-length checks to the new tendon writers, including the Newton counterparts, so callers get a clear input error instead of an out-of-bounds GPU read.

@isaaclab-review-bot isaaclab-review-bot Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Isaac Lab Review Bot

Needs attention: Design and architecture (1), Compatibility and deprecation (3), Implementation (2), Style consistency (3), Test quality (3). Other reviewed categories look good.

  • PR description: The stack scope is explained, but classify this as a performance refactor and add the benchmark command, warmup, and software versions; hardware, workload, revisions, and results are supplied.
  • Design and architecture: Equivalent masked 2-D kernels are duplicated in PhysX and Newton; move the backend-independent operation to the core Warp kernel owner.
  • Scope and over-engineering: No over-engineering issues.
  • API: API changes look good.
  • Compatibility and deprecation: Breaking changes: camera tensors no longer cast or flatten, experimental direct imports disappear, and ignored ids now fail during captured mask resolution; preserve each old path with warnings and transition coverage.
  • Implementation: The masked target setter launches full-grid indexing without validating mask lengths, allowing out-of-bounds access; resolve masks first and validate them with assert_shape_and_dtype_mask.
  • Style consistency: Newton tendon tips remain stale, added kernel docs omit required units, and repeated comments merely narrate mask resolution calls; correct or remove them.
  • Test quality: Selection-helper tests encode undocumented integer masks, repeat the production nonzero oracle, and hard-code cuda:0; use boolean literals, independent expectations, and configured CUDA devices.

Needs rework. Posted 12 inline findings.

Automated review; human maintainers own approval decisions.

full_data=True,
if isinstance(target, wp.array):
target = wp.to_torch(target)
self.assert_shape_and_dtype(target, (self.num_instances, self.num_fixed_tendons), wp.float32, "target")

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🔴 Critical · Implementation — Validate target masks before launch

The setter validates only target dimensions, so undersized masks pass and the full-grid kernel indexes beyond them. Resolve masks first and use assert_shape_and_dtype_mask, matching sibling setters.

else:
env_ids = np.asarray(env_ids, dtype=np.int32).reshape(-1)
return wp.array(env_ids, dtype=wp.int32, device=self._device)
return resolve_ids(ids=env_ids, all_ids=self._ALL_INDICES, device=self._device, dtype=wp.int32)

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🟡 Warning · Compatibility — Keep camera selector normalization

Camera selectors previously cast floating tensors and flattened multidimensional tensors; resolve_ids(..., dtype=wp.int32) now rejects both. Preserve the existing cast-and-flatten behavior before delegation.

mask = {
"warp": wp.from_torch(expected, dtype=wp.bool),
"torch_bool": expected,
"torch_int": expected.to(torch.int64),

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🟡 Warning · Test Quality — Remove undocumented integer masks

The integer-mask rows codify behavior outside the documented boolean-mask contract. Remove them or change them to boolean tensors.

}[mask_type]
all_ids = _all_ids(n, device)
ids = resolve_ids(ids=[0, 1], mask=mask, all_ids=all_ids, device=device)
assert ids.dtype == torch.int64 and torch.equal(ids, torch.nonzero(expected).flatten())

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🟡 Warning · Test Quality — Use an independent IDs oracle

The expected value repeats the production torch.nonzero transformation, so the test cannot independently detect that transformation failing. Use a small literal mask and literal expected indices.

resolve_1d_mask(ids=[1], all_mask=all_mask, device=device)


@pytest.mark.skipif(not torch.cuda.is_available(), reason="needs a CUDA device")

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🟡 Warning · Test Quality — Use configured CUDA devices

Both CUDA tests hard-code cuda:0, bypassing runtime device masks and multi-GPU routing. Parametrize them with test_devices(DeviceScope.CUDA) and use the supplied device.

stiffness=self._select_full_data(stiffness, env_ids, fixed_tendon_ids),
fixed_tendon_ids=fixed_tendon_ids,
env_ids=env_ids,
env_mask = self._resolve_mask(env_mask, self._ALL_ENV_MASK)

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🟡 Warning · Style Consistency — Update stale mask advice

The three rewritten Newton tendon mask setters no longer convert masks to indices, but their tips still claim they do and recommend index setters. Update those tips.



@wp.kernel
def set_body_com_pose_to_sim_mask(

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🟡 Warning · Style Consistency — Document masked kernel units

The three added body-mask kernels omit units for poses, offsets, spatial velocities, and accelerations. Add the required physical units to their argument documentation.

self.write_body_link_pose_to_sim_index(
body_poses=body_poses, env_ids=env_ids, body_ids=body_ids, full_data=True, skip_forward=skip_forward
env_mask = resolve_1d_mask(mask=env_mask, all_mask=self._ALL_ENV_MASK, device=self.device)
# a body selection given as ids becomes a body mask on the device, so the same mask kernel applies

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🟡 Warning · Style Consistency — Remove repeated mask narration

The same comment merely restates the following resolve_1d_mask call in four writers. Delete these repeated narration comments.

target = wp.to_torch(target)
self.assert_shape_and_dtype(target, (self.num_instances, self.num_fixed_tendons), wp.float32, "target")
# buffer only; ``write_data_to_sim`` pushes the offset with the other tendon properties
offset = self.data.fixed_tendon_rest_length.torch - target

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🟡 Warning · Implementation — Avoid full offset allocation

Every sparse masked target call allocates and computes a full offset tensor. Subtract rest length from target inside a specialized masked kernel.

return _WRITE_2D_DATA_TO_BUFFER_WITH_INDICES_AND_SIM_IDS_DISPATCHER.select(env_ids, item_ids)


@wp.kernel

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🟡 Warning · Design Architecture — Centralize masked writer kernels

PhysX and Newton now maintain equivalent generic masked 2-D writers and specialization factories. Move this backend-independent operation to the existing core Warp kernel owner.

This branch has not been deployed

No deployments
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

isaac-lab Related to Isaac Lab team

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant