File size: 21,827 Bytes
fe51a63
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
306
307
308
309
310
311
312
313
314
315
316
317
318
319
320
321
322
323
324
325
326
327
328
329
330
331
332
333
334
335
336
337
338
339
340
341
342
343
344
345
346
347
348
349
350
351
352
353
354
355
356
357
358
359
360
361
362
363
364
365
366
367
368
# SPDX-License-Identifier: Apache-2.0
"""Host pre-processing of ``autoware_lidar_frnet`` (research/frnet/SPEC.md section 3; PLAN.md 2.8), numpy only.

Two halves, both deterministic:

1. :func:`autoware_preprocess` -- the node's pre-processing, ``PKG/lib/preprocess/preprocess_kernel.cu`` +
   ``PKG/lib/lidar_frnet.cpp:196-262`` @ universe ``9ceaccf``, producing the four network inputs exactly as
   the deployed ONNX takes them (unpadded: the graph sizes its voxel tensors as ``max(inverse_map) + 1``,
   S:frnet:255-262, so the ONNX Runtime / torch goldens must never see padding):

   - **ego crop** (``projectPoints_kernel`` L211-228): with the cached static ``sensor -> base_link`` transform as 12
     float32 (row-major R, then t), a point is dropped iff it lies *strictly* inside the box on all three axes; the
     reference-frame coordinates are evaluated in float32 as the kernel writes them
     (``((r0 * x + r1 * y) + r2 * z) + t``; FMA contraction by nvcc could move a point within one ulp of a face);
   - **projection** (``project2d`` L151-163) in float32: ``depth = sqrtf(x^2+y^2+z^2)``, ``yaw = -atan2f(y, x)``,
     ``pitch = asinf(z / depth)``, ``u = 0.5 (yaw / pi + 1) W``, ``v = (1 - (pitch + |down|) / total) H``, NaN -> 0
     (``fmaxf``), clamp to ``[0, W - 1]`` / ``[0, H - 1]``, truncate. float64 math moves 0.33 % of the points to
     another cell (S:frnet:190-193); numpy's float32 ``arctan2`` / ``arcsin`` can still differ from CUDA's
     ``atan2f`` / ``asinf`` by an ulp (unavoidable, documented);
   - **interpolation grid** (L232-258): one point per pixel stored as 4 x fp16 (x, y, z, intensity); the
     **farthest** point wins, by ``depth^2`` computed in float32 from the fp16-rounded x, y, z (exact: products of
     fp16 values are exact in float32, so nvcc's FMA contraction cannot change it). The CUDA race between the
     first plain store and the CAS loop, and the order among equal depths, are fixed here: the farthest point
     wins, ties go to the later input point (the research reference's policy);
   - **interpolation** (``interpolatePoints_kernel`` L332-379): an empty pixel with ``1 <= x <= W - 2`` whose left
     and right neighbours are filled gets ``float(left + right) * 0.5`` (fp16 sum, no azimuth wrap-around), appended
     after the raw points in row-major pixel order (CUDA appends with ``atomicAdd``: non-deterministic there);
   - **frustum cells and unique** (``generateUniqueCoors`` L403-468): ``coors = (0, row, col)``, key
     ``row * W + col``; ``voxel_coors`` = the occupied cells in ascending key order, ``inverse_map`` = the rank of
     each point's cell -- here a counting sort over the ``H * W`` cells (S:frnet:460-466);
   - **profile checks** (``lidar_frnet.cpp:138-149, 227-254``): the input count, ``N`` after interpolation and ``M``
     must lie in the ``num_points`` / ``num_unique_coors`` profiles. The CUDA appends are not bounds-checked (an
     over-profile frame overflows the device buffers before the ``N`` check, S:frnet:201-205); here every count is
     checked after it is known and an out-of-profile frame is refused (:func:`profile_errors`).

2. :func:`build_plan` -- the tables of the TT device graph (PLAN.md 2.8; PORT_LOG.md "trace layout"): the points
   **sorted by cell** (stable, so a cell's points keep their input order) and padded to the capacity ``N_cap`` =
   ``num_points.max``, the per-cell segments (CSR offsets, last row) of the sorted points, the fp32 8-d point
   features with ``pre_norm`` folded in, and the gather indices that replace every scatter of the graph:
   ``point2pix[s]`` (pixel of each sorted point at stride s: pixel2point) and ``pix2vox[s]`` (cell -> voxel row
   or the zero sentinel row ``M_cap``: frustum2pixel; at s > 1 several cells fall on one pixel and the **last
   voxel in voxel order** wins, which is what ORT's sequential ScatterND does, S:frnet:285-287). The network is
   permutation-equivariant in the points, so the host un-permutes the device output with ``plan.order``.

Nothing here imports torch, onnx or ttnn.
"""
from __future__ import annotations

from dataclasses import dataclass, field
from typing import Any, Dict, List, Optional, Sequence, Tuple

import numpy as np

from ..reference.config import STRIDES, FRNetConfig

__all__ = ["NetworkInputs", "FramePlan", "crop_transform_12", "crop_keep_mask", "project2d", "interpolation_grid",
           "interpolate_points", "unique_coors", "autoware_preprocess", "profile_errors", "decorate_points",
           "build_plan", "segmax_rounds", "segmax_shift_tables"]

PI_F32 = np.float32(np.pi)   # M_PIf


# ------------------------------------------------------------------------------------------------- the node

@dataclass
class NetworkInputs:
    """The four ONNX inputs of one frame, unpadded, plus the bookkeeping of how they were made."""

    points: np.ndarray          # (N, 4) float32: kept raw points (input order), then interpolated points
    coors: np.ndarray           # (N, 3) int64 (0, row, col) frustum cell of every point
    voxel_coors: np.ndarray     # (M, 3) int64 occupied cells in ascending key order
    inverse_map: np.ndarray     # (N,) int64 voxel row of every point
    num_points_raw: int         # rows [0, num_points_raw) of `points` are raw input points
    keep_mask: np.ndarray       # (N0,) bool: input rows that survived the ego crop box
    crop_applied: bool = True

    @property
    def num_points(self) -> int:
        return int(self.points.shape[0])

    @property
    def num_voxels(self) -> int:
        return int(self.voxel_coors.shape[0])

    @property
    def num_interpolated(self) -> int:
        return self.num_points - self.num_points_raw

    def onnx_feeds(self) -> Dict[str, np.ndarray]:
        """``{points, coors, voxel_coors, inverse_map}`` as the deployed ONNX binds them (float32 / int64)."""
        return {"points": np.ascontiguousarray(self.points, np.float32),
                "coors": np.ascontiguousarray(self.coors, np.int64),
                "voxel_coors": np.ascontiguousarray(self.voxel_coors, np.int64),
                "inverse_map": np.ascontiguousarray(self.inverse_map, np.int64)}

    def stats(self) -> Dict[str, Any]:
        counts = np.bincount(self.inverse_map) if self.num_points else np.zeros(1, np.int64)
        return {"num_points_input": int(self.keep_mask.shape[0]), "num_points_raw": self.num_points_raw,
                "num_points_interpolated": self.num_interpolated, "num_points": self.num_points,
                "num_unique_coors": self.num_voxels, "cropped_by_ego_box": int((~self.keep_mask).sum()),
                "crop_applied": self.crop_applied, "max_points_per_cell": int(counts.max()) if counts.size else 0}


def crop_transform_12(T_ref_from_sensor: Any) -> np.ndarray:
    """The node's cached crop transform (``lidar_frnet_node.cpp:160-180``): ``static_cast<float>`` of the 3x3
    rotation (row-major) then the translation of the double ``base_link <- sensor`` transform -> (12,) float32."""
    T = np.asarray(T_ref_from_sensor, np.float64)
    if T.shape not in ((4, 4), (3, 4)):
        raise ValueError(f"crop transform must be 4x4 or 3x4, got {T.shape}")
    return np.concatenate([T[:3, :3].reshape(-1), T[:3, 3]]).astype(np.float32)


def crop_keep_mask(xyz: np.ndarray, sensor_to_ref_12: np.ndarray, bounds_f32: np.ndarray) -> np.ndarray:
    """``projectPoints_kernel`` L211-228: True = keep (outside the open box ``min < p < max`` on some axis)."""
    T = np.asarray(sensor_to_ref_12, np.float32).reshape(12)
    b = np.asarray(bounds_f32, np.float32).reshape(6)
    x, y, z = (np.asarray(xyz[:, i], np.float32) for i in range(3))
    px = ((T[0] * x + T[1] * y) + T[2] * z) + T[9]
    py = ((T[3] * x + T[4] * y) + T[5] * z) + T[10]
    pz = ((T[6] * x + T[7] * y) + T[8] * z) + T[11]
    inside = (px > b[0]) & (px < b[3]) & (py > b[1]) & (py < b[4]) & (pz > b[2]) & (pz < b[5])
    return ~inside


def project2d(xyz: np.ndarray, scale_x: int, scale_y: int, fov: Tuple[np.float32, np.float32, np.float32]
              ) -> Tuple[np.ndarray, np.ndarray]:
    """``project2d`` (preprocess_kernel.cu L151-163) in float32 -> (col, row) int64, clamped then truncated."""
    _, down, total = fov
    x = np.asarray(xyz[:, 0], np.float32)
    y = np.asarray(xyz[:, 1], np.float32)
    z = np.asarray(xyz[:, 2], np.float32)
    depth = np.sqrt(x * x + y * y + z * z, dtype=np.float32)
    yaw = -np.arctan2(y, x).astype(np.float32)
    with np.errstate(invalid="ignore", divide="ignore"):
        pitch = np.arcsin(z / depth).astype(np.float32)
    proj_x = np.float32(0.5) * (yaw / PI_F32 + np.float32(1.0))
    proj_y = np.float32(1.0) - (pitch + np.float32(abs(down))) / np.float32(total)
    proj_x = proj_x * np.float32(scale_x)
    proj_y = proj_y * np.float32(scale_y)
    proj_x = np.where(np.isnan(proj_x), np.float32(0.0), proj_x)      # fmaxf(NaN, 0) == 0
    proj_y = np.where(np.isnan(proj_y), np.float32(0.0), proj_y)
    px = np.minimum(np.maximum(proj_x, np.float32(0.0)), np.float32(scale_x) - np.float32(1.0))
    py = np.minimum(np.maximum(proj_y, np.float32(0.0)), np.float32(scale_y) - np.float32(1.0))
    return px.astype(np.int64), py.astype(np.int64)


def interpolation_grid(raw: np.ndarray, config: FRNetConfig) -> Tuple[np.ndarray, np.ndarray]:
    """The interpolation grid of ``projectPoints_kernel`` (L232-258) -> ``(grid (IH * IW, 4) float16, filled
    (IH * IW,) bool)``: per pixel the farthest point by fp16 ``depth^2``, ties to the later input point."""
    iw, ih = config.interp_w, config.interp_h
    grid = np.zeros((iw * ih, 4), np.float16)
    filled = np.zeros(iw * ih, bool)
    if not len(raw):
        return grid, filled
    u, v = project2d(raw[:, :3], iw, ih, config.fov)
    pix = v * iw + u
    h = raw.astype(np.float16)                            # half4(x, y, z, intensity): __float2half (RNE)
    hf = h[:, :3].astype(np.float32)
    depth2 = hf[:, 0] * hf[:, 0] + hf[:, 1] * hf[:, 1] + hf[:, 2] * hf[:, 2]   # exact products (fp16 operands)
    order = np.lexsort((np.arange(len(raw)), depth2, pix))   # by pixel, then depth, then input index
    ps = pix[order]
    last = np.r_[ps[1:] != ps[:-1], True]                 # last of each pixel group = farthest, latest on ties
    grid[ps[last]] = h[order[last]]
    filled[ps[last]] = True
    return grid, filled


def interpolate_points(grid: np.ndarray, filled: np.ndarray, config: FRNetConfig) -> np.ndarray:
    """``interpolatePoints_kernel`` (L332-379) -> (K, 4) float32 points in row-major pixel order."""
    iw, ih = config.interp_w, config.interp_h
    f2 = filled.reshape(ih, iw)
    cand = np.zeros_like(f2)
    cand[:, 1:-1] = ~f2[:, 1:-1] & f2[:, :-2] & f2[:, 2:]
    ys, xs = np.nonzero(cand)
    g2 = grid.reshape(ih, iw, 4)
    with np.errstate(over="ignore"):                                 # an fp16 overflow is inf, as __hadd gives
        s16 = (g2[ys, xs - 1] + g2[ys, xs + 1]).astype(np.float16)  # __half operator+ (RNE)
    return s16.astype(np.float32) * np.float32(0.5)


def unique_coors(coors: np.ndarray, frustum_w: int, frustum_h: int) -> Tuple[np.ndarray, np.ndarray]:
    """``generateUniqueCoors`` (L403-468) as a counting sort over the frustum cells -> ``(voxel_coors (M, 3) int64,
    inverse_map (N,) int64)``: the occupied cells in ascending key order and each point's cell rank."""
    keys = coors[:, 1] * frustum_w + coors[:, 2]
    counts = np.bincount(keys, minlength=frustum_w * frustum_h)
    occupied = counts > 0
    cell_rank = np.cumsum(occupied) - 1
    cells = np.flatnonzero(occupied)
    voxel_coors = np.stack([np.zeros_like(cells), cells // frustum_w, cells % frustum_w], axis=1).astype(np.int64)
    return voxel_coors, cell_rank[keys].astype(np.int64)


def autoware_preprocess(points_xyzi: np.ndarray, config: FRNetConfig, *,
                        sensor_to_ref_12: Optional[np.ndarray] = None,
                        crop_bounds: Optional[Sequence[float]] = None) -> NetworkInputs:
    """The node's pre-processing of one sensor-frame cloud ``(N0, 4)`` x, y, z, intensity (raw; float32) ->
    :class:`NetworkInputs`. ``crop_bounds`` (6 numbers) overrides the config's box; the crop runs iff a bound is
    non-zero, and then needs ``sensor_to_ref_12`` (:func:`crop_transform_12`): the node skips the frame when its TF
    lookup fails, so a missing transform raises ``ValueError`` here. Profile checks are separate
    (:func:`profile_errors`)."""
    pts = np.ascontiguousarray(points_xyzi, np.float32)
    if pts.ndim != 2 or pts.shape[1] != 4:
        raise ValueError(f"points must be (N, 4) x, y, z, intensity, got {pts.shape}")
    bounds = config.crop_bounds_f32 if crop_bounds is None else np.asarray(crop_bounds, np.float64).astype(np.float32)
    crop = bool(np.any(bounds != 0.0))
    if crop:
        if sensor_to_ref_12 is None:
            raise ValueError("the ego crop box is enabled but no sensor -> base_link transform was given")
        keep = crop_keep_mask(pts[:, :3], sensor_to_ref_12, bounds)
    else:
        keep = np.ones(len(pts), bool)
    raw = pts[keep]
    grid, filled = interpolation_grid(raw, config)
    interp = interpolate_points(grid, filled, config)
    allpts = np.concatenate([raw, interp], axis=0) if len(interp) else raw.copy()
    fx, fy = project2d(allpts[:, :3], config.frustum_w, config.frustum_h, config.fov)
    coors = np.stack([np.zeros_like(fy), fy, fx], axis=1).astype(np.int64)
    voxel_coors, inverse_map = unique_coors(coors, config.frustum_w, config.frustum_h)
    return NetworkInputs(points=np.ascontiguousarray(allpts, np.float32), coors=coors, voxel_coors=voxel_coors,
                         inverse_map=inverse_map, num_points_raw=int(len(raw)), keep_mask=keep, crop_applied=crop)


def profile_errors(num_input_points: int, inputs: Optional[NetworkInputs], config: FRNetConfig) -> List[str]:
    """The node's three profile checks, in its order (``lidar_frnet.cpp:138-149`` input count before the crop,
    L227-237 ``N`` after interpolation, L243-254 ``M``); empty when the frame is accepted."""
    lo, opt, hi = config.num_points
    if not lo <= num_input_points <= hi:
        return [f"Number of points ({num_input_points}) in input pointcloud does not match the profile "
                f"[min: {lo}, opt: {opt}, max: {hi}]"]
    if inputs is None:
        return []
    if not lo <= inputs.num_points <= hi:
        return [f"Number of points ({inputs.num_points}) after interpolation does not match the profile "
                f"[min: {lo}, opt: {opt}, max: {hi}]"]
    ulo, uopt, uhi = config.num_unique
    if not ulo <= inputs.num_voxels <= uhi:
        return [f"Number of unique coors ({inputs.num_voxels}) does not match the profile "
                f"[min: {ulo}, opt: {uopt}, max: {uhi}]"]
    return []


# ------------------------------------------------------------------------------------------- the TT device plan

def decorate_points(points: np.ndarray, inverse_map: np.ndarray, num_voxels: int, *,
                    pre_norm: Optional[Tuple[np.ndarray, np.ndarray]] = None) -> np.ndarray:
    """The frustum encoder's 8-d point features ``[x, y, z, i, |xyz|, x - mx, y - my, z - mz]`` (``m`` = the
    cell mean of x, y, z; AWML ``frustum_encoder.py:79-100``), computed in float64. With ``pre_norm = (scale,
    shift)`` the ``pre_norm`` BatchNorm1d is applied (``x * scale + shift``, float64) -- the port's host fold
    (PLAN.md 2.8: xyz up to +-200 m minus a cell mean needs fp32 or better, never bf16). One rounding to float32."""
    p = np.asarray(points, np.float64)
    n = len(p)
    out = np.empty((n, 8), np.float64)
    out[:, :4] = p
    out[:, 4] = np.sqrt((p[:, :3] ** 2).sum(axis=1))
    if n:
        counts = np.bincount(inverse_map, minlength=num_voxels).astype(np.float64)
        for c in range(3):
            mean = np.bincount(inverse_map, weights=p[:, c], minlength=num_voxels) / np.maximum(counts, 1.0)
            out[:, 5 + c] = p[:, c] - mean[inverse_map]
    if pre_norm is not None:
        scale, shift = (np.asarray(a, np.float64).reshape(8) for a in pre_norm)
        out = out * scale + shift
    return out.astype(np.float32)


def segmax_rounds(max_segment_length: int) -> int:
    """Rounds of the log-step segmented max (Hillis-Steele) that cover segments up to that length:
    ``ceil(log2(L))`` (0 for L <= 1). PLAN.md 1.3: 5 for OT128 (<= 17 points per cell seen), 7 for QT128 (<= 111)."""
    return int(np.ceil(np.log2(max_segment_length))) if max_segment_length > 1 else 0


@dataclass
class FramePlan:
    """The per-frame host tables of the device graph (module docstring, part 2). Row ``i`` of every per-point table
    is the ``i``-th point in **sorted** order; rows ``N .. N_cap - 1`` are padding (one trailing sentinel segment:
    their voxel id is ``M_cap``, their pixel 0, their features 0, and nothing the output keeps reads them)."""

    n: int                                   # real points
    m: int                                   # real voxels (occupied cells)
    n_raw: int
    n_cap: int
    m_cap: int
    frustum_hw: Tuple[int, int]
    order: np.ndarray                        # (N,) int64: sorted row -> network-input point index
    seg_id: np.ndarray                       # (N_cap,) int64 voxel id of each sorted row (padding: M_cap)
    seg_offsets: np.ndarray                  # (M + 1,) int64 CSR offsets of the voxels' sorted rows
    seg_last: np.ndarray                     # (M_cap,) uint32 last sorted row of each voxel (unused voxels: N_cap - 1)
    point2pix: Dict[int, np.ndarray]         # stride -> (N_cap,) uint32 pixel of each sorted row (padding: 0)
    pix2vox: Dict[int, np.ndarray]           # stride -> ((H/s) * (W/s),) uint32 voxel row or M_cap (zero row)
    features: np.ndarray                     # (N_cap, 8) float32 decorated (+ pre_norm) features
    max_segment_length: int = 0
    strides: Tuple[int, ...] = field(default=STRIDES)

    @property
    def rank(self) -> np.ndarray:
        """(N,) int64: network-input point index -> sorted row (the inverse of ``order``)."""
        r = np.empty(self.n, np.int64)
        r[self.order] = np.arange(self.n)
        return r

    def unsort(self, sorted_rows: np.ndarray) -> np.ndarray:
        """Device rows (sorted, padded) -> the network-input point order (the first ``N`` sorted rows only)."""
        return np.asarray(sorted_rows)[: self.n][self.rank]

    def stats(self) -> Dict[str, Any]:
        return {"n": self.n, "m": self.m, "n_cap": self.n_cap, "m_cap": self.m_cap,
                "max_points_per_cell": self.max_segment_length,
                "segmax_rounds": segmax_rounds(self.max_segment_length)}


def build_plan(inputs: NetworkInputs, config: FRNetConfig, *,
               pre_norm: Optional[Tuple[np.ndarray, np.ndarray]] = None,
               n_cap: Optional[int] = None, m_cap: Optional[int] = None,
               strides: Sequence[int] = STRIDES) -> FramePlan:
    """The device tables of one accepted frame (:class:`FramePlan`). ``n_cap`` / ``m_cap`` default to the
    config's capacities (COMPILE: the profile maxima); a frame above them raises ``ValueError`` (the profile check
    refuses it first)."""
    n_cap = config.n_cap if n_cap is None else int(n_cap)
    m_cap = config.m_cap if m_cap is None else int(m_cap)
    n, m = inputs.num_points, inputs.num_voxels
    if n > n_cap or m > m_cap:
        raise ValueError(f"frame N={n} / M={m} exceeds the capacity N_cap={n_cap} / M_cap={m_cap}")
    H, W = config.frustum_h, config.frustum_w
    inv = inputs.inverse_map
    order = np.argsort(inv, kind="stable")              # voxel ids are key ranks: == the stable sort by cell key
    counts = np.bincount(inv, minlength=m).astype(np.int64)
    offsets = np.zeros(m + 1, np.int64)
    np.cumsum(counts, out=offsets[1:])
    seg_id = np.full(n_cap, m_cap, np.int64)
    seg_id[:n] = inv[order]
    seg_last = np.full(m_cap, n_cap - 1, np.uint32)
    seg_last[:m] = (offsets[1:] - 1).astype(np.uint32)
    coors_sorted = inputs.coors[order]
    vox = inputs.voxel_coors
    point2pix: Dict[int, np.ndarray] = {}
    pix2vox: Dict[int, np.ndarray] = {}
    for s in strides:
        hs, ws = H // s, W // s
        p2p = np.zeros(n_cap, np.uint32)
        p2p[:n] = ((coors_sorted[:, 1] // s) * ws + coors_sorted[:, 2] // s).astype(np.uint32)
        point2pix[s] = p2p
        last = np.full(hs * ws, -1, np.int64)
        np.maximum.at(last, (vox[:, 1] // s) * ws + vox[:, 2] // s, np.arange(m, dtype=np.int64))
        pix2vox[s] = np.where(last < 0, m_cap, last).astype(np.uint32)
    feats = np.zeros((n_cap, 8), np.float32)
    feats[:n] = decorate_points(inputs.points[order], inv[order], m, pre_norm=pre_norm)
    return FramePlan(n=n, m=m, n_raw=inputs.num_points_raw, n_cap=n_cap, m_cap=m_cap, frustum_hw=(H, W),
                     order=order, seg_id=seg_id, seg_offsets=offsets, seg_last=seg_last, point2pix=point2pix,
                     pix2vox=pix2vox, features=feats, max_segment_length=int(counts.max()) if m else 0,
                     strides=tuple(strides))


def segmax_shift_tables(seg_id: np.ndarray, rounds: int) -> List[np.ndarray]:
    """The gather tables of the exact log-step segmented max (PLAN.md 1.3, the K1 stock-op fallback): round
    ``r`` uses ``idx[i] = i - 2^r`` if row ``i - 2^r`` is in the same segment, else ``i`` (max is idempotent);
    after ``rounds`` rounds each segment's last row holds the segment max when its length is <= ``2^rounds``.
    -> ``rounds`` x (N_cap,) uint32."""
    seg = np.asarray(seg_id)
    i = np.arange(len(seg), dtype=np.int64)
    out = []
    for r in range(rounds):
        d = 1 << r
        j = i - d
        same = np.zeros(len(seg), bool)
        same[d:] = seg[d:] == seg[:-d]
        out.append(np.where(same, j, i).astype(np.uint32))
    return out