Skip to content

Commit 4fd36b1

Browse files
authored
[MISC] Rigid methods set_qpos,set_dofs_position now clear error code. (#2253)
1 parent 10eb7ee commit 4fd36b1

File tree

4 files changed

+41
-11
lines changed

4 files changed

+41
-11
lines changed

genesis/engine/solvers/rigid/collider_decomp.py

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1505,7 +1505,7 @@ def func_broad_phase(
15051505
continue
15061506

15071507
if n_broad == collider_info.max_collision_pairs_broad[None]:
1508-
errno[None] = errno[None] | 0b00000000000000000000000000000001
1508+
errno[i_b] = errno[i_b] | 0b00000000000000000000000000000001
15091509
break
15101510
collider_state.broad_collision_pairs[n_broad, i_b][0] = i_ga
15111511
collider_state.broad_collision_pairs[n_broad, i_b][1] = i_gb
@@ -2230,7 +2230,7 @@ def func_add_contact(
22302230

22312231
collider_state.n_contacts[i_b] = i_c + 1
22322232
else:
2233-
errno[None] = errno[None] | 0b00000000000000000000000000000010
2233+
errno[i_b] = errno[i_b] | 0b00000000000000000000000000000010
22342234

22352235

22362236
@ti.func

genesis/engine/solvers/rigid/constraint_solver_decomp.py

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1369,7 +1369,7 @@ def func_update_qacc(
13691369
dofs_state.force[i_d, i_b] = dofs_state.qf_smooth[i_d, i_b] + constraint_state.qfrc_constraint[i_d, i_b]
13701370
constraint_state.qacc_ws[i_d, i_b] = constraint_state.qacc[i_d, i_b]
13711371
if ti.math.isnan(constraint_state.qacc[i_d, i_b]):
1372-
errno[None] = errno[None] | 0b00000000000000000000000000000100
1372+
errno[i_b] = errno[i_b] | 0b00000000000000000000000000000100
13731373

13741374
ti.loop_config(serialize=static_rigid_sim_config.para_level < gs.PARA_LEVEL.ALL)
13751375
for i_b in range(_B):

genesis/engine/solvers/rigid/rigid_solver_decomp.py

Lines changed: 37 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -1040,9 +1040,14 @@ def substep(self, f):
10401040
self._static_rigid_sim_config,
10411041
)
10421042

1043+
def get_error_envs_mask(self):
1044+
return ti_to_torch(self._errno) > 0
1045+
10431046
def check_errno(self):
1044-
# Zero-copy cannot be used for 0-dim tensors because of torch limitation, so we are better of calling a kernel
1045-
errno = kernel_read_scalar_int32(self._errno)
1047+
if gs.use_zerocopy:
1048+
errno = np.bitwise_or.reduce(ti_to_numpy(self._errno))
1049+
else:
1050+
errno = kernel_bit_reduction(self._errno)
10461051

10471052
if errno & 0b00000000000000000000000000000001:
10481053
max_collision_pairs_broad = self.collider._collider_info.max_collision_pairs_broad[None]
@@ -1576,6 +1581,13 @@ def get_state(self, f=None):
15761581
def set_state(self, f, state, envs_idx=None):
15771582
if self.is_active:
15781583
envs_idx = self._scene._sanitize_envs_idx(envs_idx)
1584+
1585+
if gs.use_zerocopy:
1586+
errno = ti_to_torch(self._errno, copy=False)
1587+
errno[envs_idx] = 0
1588+
else:
1589+
kernel_set_zero(envs_idx, self._errno)
1590+
15791591
kernel_set_state(
15801592
qpos=state.qpos,
15811593
dofs_vel=state.dofs_vel,
@@ -1609,7 +1621,6 @@ def set_state(self, f, state, envs_idx=None):
16091621
self._is_forward_pos_updated = True
16101622
self._is_forward_vel_updated = True
16111623

1612-
self._errno[None] = 0
16131624
self.collider.clear(envs_idx)
16141625
self.constraint_solver.clear(envs_idx)
16151626

@@ -1907,6 +1918,7 @@ def set_geoms_friction_ratio(self, friction_ratio, geoms_idx=None, envs_idx=None
19071918
def set_qpos(self, qpos, qs_idx=None, envs_idx=None, *, skip_forward=False):
19081919
if gs.use_zerocopy:
19091920
data = ti_to_torch(self._rigid_global_info.qpos, transpose=True, copy=False)
1921+
errno = ti_to_torch(self._errno, copy=False)
19101922
qs_mask = indices_to_mask(qs_idx)
19111923
if (
19121924
(not qs_mask or isinstance(qs_mask[0], slice))
@@ -1919,9 +1931,11 @@ def set_qpos(self, qpos, qs_idx=None, envs_idx=None, *, skip_forward=False):
19191931
else:
19201932
qpos = broadcast_tensor(qpos, gs.tc_float, qs_data.shape)
19211933
torch.where(envs_idx[:, None], qpos, qs_data, out=qs_data)
1934+
errno.masked_fill_(envs_idx, 0.0)
19221935
else:
19231936
mask = (0, *qs_mask) if self.n_envs == 0 else indices_to_mask(envs_idx, *qs_mask)
19241937
assign_indexed_tensor(data, mask, qpos)
1938+
errno[envs_idx] = 0
19251939
if mask and isinstance(mask[0], torch.Tensor):
19261940
envs_idx = mask[0].reshape((-1,))
19271941
else:
@@ -1931,6 +1945,7 @@ def set_qpos(self, qpos, qs_idx=None, envs_idx=None, *, skip_forward=False):
19311945
if self.n_envs == 0:
19321946
qpos = qpos[None]
19331947
kernel_set_qpos(qpos, qs_idx, envs_idx, self._rigid_global_info, self._static_rigid_sim_config)
1948+
kernel_set_zero(envs_idx, self._errno)
19341949

19351950
if self.collider is not None:
19361951
self.collider.reset(envs_idx)
@@ -2204,6 +2219,12 @@ def set_dofs_position(self, position, dofs_idx=None, envs_idx=None):
22042219
self._static_rigid_sim_config,
22052220
)
22062221

2222+
if gs.use_zerocopy:
2223+
errno = ti_to_torch(self._errno, copy=False)
2224+
errno[envs_idx] = 0
2225+
else:
2226+
kernel_set_zero(envs_idx, self._errno)
2227+
22072228
self.collider.reset(envs_idx)
22082229
self.constraint_solver.reset(envs_idx)
22092230

@@ -7097,7 +7118,7 @@ def func_copy_next_to_curr(
70977118
for i_q in range(n_qs):
70987119
rigid_global_info.qpos[i_q, i_b] = rigid_global_info.qpos_next[i_q, i_b]
70997120
else:
7100-
errno[None] = errno[None] | 0b00000000000000000000000000001000
7121+
errno[i_b] = errno[i_b] | 0b00000000000000000000000000001000
71017122

71027123

71037124
@ti.func
@@ -8407,13 +8428,22 @@ def kernel_set_geoms_friction(
84078428
static_rigid_sim_config: ti.template(),
84088429
):
84098430
ti.loop_config(serialize=ti.static(static_rigid_sim_config.para_level < gs.PARA_LEVEL.ALL))
8410-
for i_g_ in ti.ndrange(geoms_idx.shape[0]):
8431+
for i_g_ in range(geoms_idx.shape[0]):
84118432
geoms_info.friction[geoms_idx[i_g_]] = friction[i_g_]
84128433

84138434

84148435
@ti.kernel(fastcache=gs.use_fastcache)
8415-
def kernel_read_scalar_int32(tensor: array_class.V_ANNOTATION) -> ti.i32:
8416-
return tensor[None]
8436+
def kernel_bit_reduction(tensor: array_class.V_ANNOTATION) -> ti.i32:
8437+
flag = ti.i32(0)
8438+
for i in range(tensor.shape[0]):
8439+
flag = ti.atomic_or(flag, tensor[i])
8440+
return flag
8441+
8442+
8443+
@ti.kernel(fastcache=gs.use_fastcache)
8444+
def kernel_set_zero(envs_idx: ti.types.ndarray(), tensor: array_class.V_ANNOTATION):
8445+
for i_b in range(envs_idx.shape[0]):
8446+
tensor[i_b] = 0
84178447

84188448

84198449
@ti.func

genesis/utils/array_class.py

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1892,7 +1892,7 @@ def __init__(self, solver):
18921892
self.geoms_state_adjoint_cache = get_geoms_state(solver)
18931893

18941894
self.rigid_adjoint_cache = get_rigid_adjoint_cache(solver)
1895-
self.errno = V_SCALAR_FROM(dtype=gs.ti_int, value=0)
1895+
self.errno = V(dtype=gs.ti_int, shape=(solver._B,))
18961896

18971897

18981898
DofsState = StructDofsState if gs.use_ndarray else ti.template()

0 commit comments

Comments
 (0)