|
| 1 | +# SPDX-FileCopyrightText: Copyright (c) 2026 The Newton Developers |
| 2 | +# SPDX-License-Identifier: Apache-2.0 |
| 3 | + |
| 4 | +"""Shared Warp kernels for :class:`~newton.controllers.ControllerJointImpedance`.""" |
| 5 | + |
| 6 | +import numpy as np |
| 7 | +import warp as wp |
| 8 | + |
| 9 | + |
| 10 | +def _idx_max(idx: wp.array[wp.uint32]) -> int: |
| 11 | + """Return the minimum flat-array size needed to hold all indices.""" |
| 12 | + return int(np.max(idx.numpy())) + 1 |
| 13 | + |
| 14 | + |
| 15 | +@wp.kernel |
| 16 | +def _pd_term_kernel( |
| 17 | + joint_q: wp.array2d[wp.float32], # (robot_count, max_dofs) |
| 18 | + joint_qd: wp.array2d[wp.float32], # (robot_count, max_dofs) |
| 19 | + joint_q_des: wp.array2d[wp.float32], # (robot_count, max_dofs) |
| 20 | + joint_qd_des: wp.array2d[wp.float32], # (robot_count, max_dofs) |
| 21 | + stiffness: wp.array2d[wp.float32], # (robot_count, max_dofs) |
| 22 | + damping: wp.array2d[wp.float32], # (robot_count, max_dofs) |
| 23 | + dofs_per_robot: wp.array[wp.int32], # (robot_count,) |
| 24 | + out: wp.array2d[wp.float32], # (robot_count, max_dofs) |
| 25 | +): |
| 26 | + robot, dof = wp.tid() |
| 27 | + if dof >= dofs_per_robot[robot]: |
| 28 | + return |
| 29 | + out[robot, dof] = stiffness[robot, dof] * (joint_q_des[robot, dof] - joint_q[robot, dof]) + damping[robot, dof] * ( |
| 30 | + joint_qd_des[robot, dof] - joint_qd[robot, dof] |
| 31 | + ) |
| 32 | + |
| 33 | + |
| 34 | +@wp.kernel |
| 35 | +def _add_term_kernel( |
| 36 | + term: wp.array2d[wp.float32], # (robot_count, max_dofs) |
| 37 | + dofs_per_robot: wp.array[wp.int32], # (robot_count,) |
| 38 | + tau: wp.array2d[wp.float32], # (robot_count, max_dofs) |
| 39 | +): |
| 40 | + robot, dof = wp.tid() |
| 41 | + if dof >= dofs_per_robot[robot]: |
| 42 | + return |
| 43 | + tau[robot, dof] = tau[robot, dof] + term[robot, dof] |
| 44 | + |
| 45 | + |
| 46 | +@wp.kernel |
| 47 | +def _mass_matrix_multiply_kernel( |
| 48 | + M: wp.array3d[wp.float32], # (robot_count, max_dofs, max_dofs) |
| 49 | + vec: wp.array2d[wp.float32], # (robot_count, max_dofs) |
| 50 | + dofs_per_robot: wp.array[wp.int32], # (robot_count,) |
| 51 | + out: wp.array2d[wp.float32], # (robot_count, max_dofs) |
| 52 | +): |
| 53 | + robot, dof = wp.tid() |
| 54 | + if dof >= dofs_per_robot[robot]: |
| 55 | + return |
| 56 | + acc = float(0.0) |
| 57 | + for col in range(dofs_per_robot[robot]): |
| 58 | + acc = acc + M[robot, dof, col] * vec[robot, col] |
| 59 | + out[robot, dof] = acc |
| 60 | + |
| 61 | + |
| 62 | +@wp.kernel |
| 63 | +def _gather_dof_flat_kernel( |
| 64 | + src: wp.array[wp.float32], # flat sim array |
| 65 | + indices: wp.array[wp.uint32], # (total_dofs,) — concatenated per-robot, no padding |
| 66 | + dst: wp.array[wp.float32], # flat output (total_dofs,) |
| 67 | +): |
| 68 | + flat = wp.tid() |
| 69 | + dst[flat] = src[indices[flat]] |
| 70 | + |
| 71 | + |
| 72 | +@wp.kernel |
| 73 | +def _gather_dof_kernel( |
| 74 | + src: wp.array[wp.float32], # flat sim array |
| 75 | + dof_indices: wp.array[wp.uint32], # (total_dofs,) — concatenated per-robot indices |
| 76 | + dof_offsets: wp.array[wp.int32], # (robot_count,) — start of each robot in dof_indices |
| 77 | + dofs_per_robot: wp.array[wp.int32], # (robot_count,) |
| 78 | + dst: wp.array2d[wp.float32], # (robot_count, max_dofs) |
| 79 | +): |
| 80 | + robot, dof = wp.tid() |
| 81 | + if dof >= dofs_per_robot[robot]: |
| 82 | + return |
| 83 | + dst[robot, dof] = src[dof_indices[dof_offsets[robot] + dof]] |
| 84 | + |
| 85 | + |
| 86 | +@wp.kernel |
| 87 | +def _scatter_dof_kernel( |
| 88 | + src: wp.array2d[wp.float32], # (robot_count, max_dofs) |
| 89 | + dof_indices: wp.array[wp.uint32], # (total_dofs,) — concatenated per-robot indices |
| 90 | + dof_offsets: wp.array[wp.int32], # (robot_count,) — start of each robot in dof_indices |
| 91 | + dofs_per_robot: wp.array[wp.int32], # (robot_count,) |
| 92 | + dst: wp.array[wp.float32], # flat sim output |
| 93 | +): |
| 94 | + robot, dof = wp.tid() |
| 95 | + if dof >= dofs_per_robot[robot]: |
| 96 | + return |
| 97 | + dst[dof_indices[dof_offsets[robot] + dof]] = src[robot, dof] |
0 commit comments