Resource-tier system

The resource-tier system lets an inline-CUDA caller choose the shared-memory and register profile of every generated kernel and device routine. It has shipped in every packaged release (since 0.5.0), including the per-tier spill machinery described below. L2 pinning of the spilled workspace is off by default; the measurement is in the spilled-memory section. The global scratch arena is d_workspace (device memory).

This page is written for inline-CUDA users (#include "grid.cuh" from their own kernel). The Python wrappers (grid_rbd.RobotHandle, grid_rbd.jax.JaxRobotHandle) launch each kernel at its per-algorithm baked tier, grid::launch_cfg<GRID_ALGO_*>::TIER, autotuned into grid_assets/launch_configs/<robot>/<gpu>.json and baked at codegen. On a tuned robot many algorithms run at TIER_LITE/TIER_MINIMAL; an untuned robot (or algorithm) falls back to TIER_SHARED.

Note

This page is the reference (definition, memory model, exposed constants, coverage); the design rationale is collected at the end of the page.

Definition

Every emitted __global__ kernel and every inline-callable _device/_inner function takes a non-type template parameter int RESOURCE_TIER (defaulting to TIER_SHARED). The tier picks a (launch_bounds, smem footprint, register cap) profile so an inline-CUDA caller can fit a GRiD primitive into their outer kernel’s resource budget.

The three tiers:

Table 2 Tier semantics

Tier	launch_bounds	Register cap (sm_120)	Smem behavior
TIER_SHARED (default)	MAX_PERF_LEVEL_THREADS (288-512 per robot)	~128-186 regs/thread	Full inner scratch lives in shared memory; best performance.
TIER_LITE	min(2*SUGGESTED, 768)	~85 regs/thread	Picks the lowest spill rung that fits the ~48 KB LITE smem target (cuda_target_lite_shared_mem_bytes), clamped to be ≥ the SHARED rung. On robots and algorithms with multi-rung ladders this is an intermediate spill level, not an alias of MINIMAL.
TIER_MINIMAL	1024 (hardware cap)	~64 regs/thread	Inner scratch routes entirely to d_workspace. Smallest smem footprint; maximum block-size flexibility for tight outer kernels.

The register cap follows from regs_per_thread * max_threads <= 65536 on sm_120: a tighter launch_bounds lets nvcc allocate more registers per thread, a looser one forces it to budget for more threads and use fewer registers each.

Audience and entry points

GRiD is a code generator for power users: people who call hand-tuned, robot-specialized rigid-body-dynamics kernels directly from their own CUDA code and want every cycle and byte of the GPU. Everything below the convenience layer is built for that person.

Using GRiD does not require being that person. We ship a ladder of entry points, from a one-line Python call that needs no GPU knowledge up to the raw block-parallel device routine whose shared memory the caller manages. Pick the rung that matches how much control you need:

Table 3 Entry points, easiest to most powerful

Goal	Entry point	Caller manages
The answer, from Python	grid_rbd.RobotHandle / grid_rbd.jax.JaxRobotHandle	Nothing. Arrays in, arrays out. The tier comes from the per-algorithm baked launch config (TIER_SHARED fallback when untuned).
The answer, from C++/CUDA host code	grid::<algo>(hd_data, ...) host wrapper	Nothing on-device. The wrapper does H2D/D2H copies, picks launch dims, sets shared-mem attributes, launches the kernel.
A kernel to drop into your own launch	grid::<algo>_kernel<T, TIER> (__global__)	The launch (grid/block dims, dynamic-smem bytes, streams); the per-trajectory batch loop is done inside.
A block-parallel routine to call inside your own kernel	grid::<algo>_inner<T, PLACEMENT> / _device (__device__)	Everything: shared-memory arenas, scratch placement, syncs. This is the engine; the layers above are conveniences.

A new user starts at the top of that table; the RobotHandle tutorial covers that rung and the rest of this page is not needed. A user fitting GRiD inside a fused planning, MPC or learning kernel needs the full machinery, which the rest of this page documents.

Memory hierarchy: shared versus global, spills, and L2 pinning

The other half of the inner’s intelligence is the memory hierarchy. An inner’s working set is a mix of:

• inputs and outputs, supplied by the caller, who decides where they live;

• persistent scratch, needed across the whole routine;

• transient scratch, needed only within a phase and freely reused.

On a GPU these can live in shared memory (fast and scarce: ~48 KB default / ~100 KB opt-in per block on sm_120) or global memory (abundant, slower, and L2-pinnable). The inner decides, per buffer, where each goes, and that decision is exposed as a compile-time placement parameter so the caller can pick a profile that fits their outer kernel’s pressure.

Placement is the inner’s job, not the kernel’s. Each inline-callable inner is keyed on a placement template parameter and chooses s_temp (shared) versus d_workspace (global) for each spillable buffer at the top of the function. The caller (kernel, device wrapper, or your own code) is a thin shim: it sizes both arenas from the exposed constants, hands both pointers in, and passes the placement. A surgical-spill change, moving one more buffer to global or splitting a buffer hot/cold, is therefore local to the inner: repoint a sub-buffer and update its size constant, with no kernel edit. This is what keeps the spill machinery tractable to evolve.

Placement parameters currently emitted (all default to “in shared memory” so existing call sites are unchanged):

Table 4 Inner placement parameters

Inner	Placement param	Buffer it routes
minv_inner	bool F_IN_SMEM	the 6·NV² articulated-body F-region
forward_dynamics_inner	bool MINV_F_IN_SMEM	the internal Minv F-region (FD does not take it as a parameter)
aba_inner	bool TEMP_IN_SMEM	the 140·NJ+ recursion scratch band
end_effector_pose_gradient_inner	bool TEMP_IN_SMEM	the double-buffered kinematic-chain workspace
fdsva_so_inner	bool SCRATCH_IN_SMEM	the 4·NV³ contraction scratch
integrator_inner	bool MINV_F_IN_SMEM	the FD inner’s Minv F-region (value path); see “Integrator surgical spill”
integrator_gradient_kernel	per-tier rung (Dqdd / dAB / inner level)	composes inverse_dynamics_gradient selective/global_temp + spills Dqdd & the dAB output
multi_target_position / _gradient kernels	per-tier output placement	at TIER_SHARED the MT kernels stage outputs in static shared memory and copy out; at spilled tiers they write direct-to-global (&d_out[k*out_size]) with per-timestep inner scratch in the d_workspace SO band. The generated hosts pass hd_data->d_workspace and use workspace-clamped launches.
*_device (inverse_dynamics_gradient / forward_dynamics_gradient / idsva_so / end_effector_pose_hessian)	int RESOURCE_TIER	whole inner s_temp arena (via the tier_workspace_expr helper)

Spilled buffers are recursion-hot, and the original design pinned the workspace in persisting L2 (GRID_CUDA_ENABLE_L2_PERSISTING=1) so a spilled access would cost an L2 hit rather than an HBM round-trip. A controlled A/B on an RTX 5090 (spilling algorithms plus shared-tier controls on iiwa14-fixed / g1-floating / h1_2-floating, four interleaved reps, spreads ≤0.7%) found that the pin never helped and hurt 17 of 51 cells, by up to 23% (integrator family), including shared-tier cells that spill nothing. The persisting window is installed on the stream whenever a workspace exists, and its carve evicts more general L2 traffic than it saves; modern L2s already cache the spilled band well on their own. The generated default is therefore 0; opt back in per build with -DGRID_CUDA_ENABLE_L2_PERSISTING=1 if your GPU or workload measures otherwise.

Each algorithm has a fixed menu of spill levels (level 0 = everything in shared memory; higher levels progressively move buffers to d_workspace). Which level a given RESOURCE_TIER maps to is decided at code-generation time, per robot, based on what fits the smem budget for that robot. Small robots (for example iiwa14, go2) keep every tier at level 0: there is nothing to spill, so the tiers share the same spill placement and emit a single kernel body. This does not make them byte-identical SASS: each tier still carries its own __launch_bounds__(tier_max_threads<TIER>()) (SHARED = MAX_PERF_LEVEL_THREADS, LITE = min(2×, 768), MINIMAL = 1024), so ptxas budgets a different register cap per tier and can produce different register counts and SASS even when the placement is identical. For algorithms with no smem spill at all (for example inverse_dynamics) the launch_bounds is the only thing that differs between tiers, which is why a looser-bounds tier can run slower through register starvation (see the autotune A.7 note). Large robots (for example h1_2) map the lower tiers to deeper spill levels. Because the mapping is per-robot, multiple tiers can share a level, and the generated *_IN_SMEM<TIER>() / *_SCRATCH_IN_SMEM<TIER>() constexprs expose which placement each tier resolves to for the robot you generated.

This reconciles two goals that look opposed: the inner stays self-contained and decidable from a single placement flag (good for inline reuse and for evolving spills), while the choice of flag per tier is a per-robot, fit-driven decision made once at codegen time (good for not paying for a spill you do not need).

This is the required standard for every algorithm, not a per-kernel option. New algorithms and refactors must put the spill decision in the inner (a placement template plus if constexpr selecting s_temp versus d_workspace at the top), and never repoint the inner’s scratch from the kernel or host caller. A caller that aliases or reassigns s_temp from the outside must be migrated. Two consequences follow:

• Composing kernels inherit spills. fdsva_so embeds the idsva_so inner; because that inner owns its placement, fdsva_so spills the dominant idsva_so scratch by passing SCRATCH_IN_SMEM=false, with no pointer surgery in fdsva_so, and any future surgical (cold-only) idsva_so spill propagates to fdsva_so automatically.

• The XImats/XmatsHom load helper is a separate scratch user. It is called from the kernel outside the inner and dereferences s_temp for its 2*num_pos sincos scratch, so a spilled (null) s_temp crashes it. Handle this uniformly: either keep the small sincos scratch in smem always, or repoint s_temp at the workspace before the helper call.

The conformance record for this pattern (which inners own their placement, and the deferred surgical de-alias work) lives in docs/design_history/idsva_so_inner_refactor_notes.md in the repository.

Block-wide parallelism in the inners

The inners are written to saturate the whole thread block, not a fixed lane count. Every parallel region is emitted as a block-stride loop of the form:

for (int i = threadIdx.x + threadIdx.y * blockDim.x;
     i < WORK_ITEMS;
     i += blockDim.x * blockDim.y) { ... }

This has three consequences that power users rely on:

• The routine is correct at any block size. The same generated routine runs correctly whether it is launched with 32 threads or 1024. Work items are distributed across however many threads the block has; there are no hard-coded lane assumptions and no out-of-bounds writes when blockDim does not divide the work evenly. Every parallel write in the emitted code is inside a block-stride loop; there are no bare threadIdx-indexed stores.

• The caller chooses the occupancy/latency trade. Because the routine adapts to the launch, block size can be tuned for a fused kernel’s occupancy without regenerating anything. The tier system’s __launch_bounds__ only bounds the maximum threads (to control the register budget); it does not fix the launch.

• Parallelism is maximal by construction. Each algorithm exposes its natural parallel width (per-(i,j,k) tensor elements, per-DOF columns, per-body 6×6 blocks) directly as the loop bound, so a large block fills with useful work rather than idling. Where the recursion structure forces seriality (the BFS sweeps), the parallel regions sit between syncs and still use the full block.

The inner is the unit of parallelism: the caller brings the threads and the inner uses all of them.

Exposed sizing and placement constants

For every spillable inner, codegen emits a matched trio so the caller can allocate correctly and know what codegen chose. Using fdsva_so as the template:

// bytes to reserve in shared memory for this placement
template <typename T, bool SCRATCH_IN_SMEM = true>
constexpr size_t FDSVA_SO_INNER_SMEM_BYTES();

// bytes to reserve in global memory for this placement
template <typename T, bool SCRATCH_IN_SMEM = true>
constexpr size_t FDSVA_SO_INNER_WORKSPACE_BYTES();

// the placement codegen assigned to each tier, for THIS robot
template <int TIER>
constexpr bool FDSVA_SO_SCRATCH_IN_SMEM();

The same trio is emitted for the other converted algorithms, with the placement bool named for the buffer it controls:

• MINV_INNER_{SMEM,WORKSPACE}_BYTES<T, F_IN_SMEM> + MINV_F_IN_SMEM<TIER>

• FD_INNER_{SMEM,WORKSPACE}_BYTES<T, MINV_F_IN_SMEM> + FD_MINV_F_IN_SMEM<TIER>

• ABA_INNER_{SMEM,WORKSPACE}_BYTES<T, TEMP_IN_SMEM> + ABA_TEMP_IN_SMEM<TIER>

• EE_GRAD_INNER_{SMEM,WORKSPACE}_BYTES<T, TEMP_IN_SMEM> + EE_GRAD_TEMP_IN_SMEM<TIER>

The *_device inline entry points (inverse_dynamics_gradient / forward_dynamics_gradient / idsva_so / end_effector_pose_hessian) expose their sizing as *_DEVICE_INLINE_{SMEM,WORKSPACE}_BYTES<T, TIER> (keyed on tier rather than a placement bool); they decide placement internally via the tier_workspace_expr arena helper, and their kernels inline and spill at the kernel level. Converting those kernels to call placement-deciding inners is a tracked follow-up.

For an inline call:

1. Pick a TIER (or call <ALGO>_..._IN_SMEM<TIER>() to see the placement it resolves to for your robot).

2. Reserve ..._INNER_SMEM_BYTES<T, placement>() in your block’s dynamic shared memory for the primitive’s s_temp.

3. cudaMalloc (once) ..._INNER_WORKSPACE_BYTES<T, placement>() per concurrently-resident block for d_workspace (0 when the placement keeps everything in shared).

4. Call <algo>_inner<T, placement>(..., s_temp, d_workspace, ...).

Inline-CUDA entry points

The tier knob is exposed at the kernel level on every emitted *_kernel<T, RESOURCE_TIER>. At the inline-CUDA _device / _inner level, these functions accept RESOURCE_TIER and a caller-provided T *d_workspace argument:

• fdsva_so_inner<T, RESOURCE_TIER>(s_df2, s_idsva_so, s_Minv, s_df_du, s_XImats, s_temp, d_workspace, gravity): 4*nv³ inner scratch routes between s_temp (SHARED) and d_workspace (LITE/MINIMAL).

• forward_dynamics_gradient_device<T, RESOURCE_TIER>(s_df_du, s_q, s_qd, [s_qdd, s_Minv | s_u], d_robotModel, gravity, d_workspace): whole s_temp arena routes per tier.

• inverse_dynamics_gradient_device<T, RESOURCE_TIER>(s_dc_du, s_q, s_qd, [s_qdd], d_robotModel, gravity, d_workspace): whole s_temp arena routes per tier.

• end_effector_pose_hessian_device<T, RESOURCE_TIER> (s_d2eePos, s_deePos, s_q, d_robotModel, d_workspace): s_d2eeTemp slot (the 2*16*num_ees*n² portion) routes per tier; inner_no_d2 stays in smem at all tiers.

• idsva_so_device<T, RESOURCE_TIER>(s_idsva_so, s_q, s_qd, s_qdd, d_robotModel, gravity, d_workspace): codegen-time frame dispatcher, body_frame_inner for fixed-base and world_frame_inner for floating-base. Inner temp arena routes per tier.

Sizing constants, called from host code to allocate the right buffers:

template <typename T, int TIER = TIER_SHARED>
constexpr size_t FDSVA_SO_INNER_SMEM_BYTES();           // bytes for s_temp at TIER
template <typename T, int TIER = TIER_SHARED>
constexpr size_t FDSVA_SO_INNER_WORKSPACE_BYTES();      // bytes for d_workspace at TIER

// Same pattern: FORWARD_DYNAMICS_GRADIENT_DEVICE_INLINE_*, INVERSE_DYNAMICS_GRADIENT_DEVICE_INLINE_*,
//               END_EFFECTOR_POSE_HESSIAN_DEVICE_INLINE_*, IDSVA_SO_DEVICE_INLINE_*

At TIER_SHARED the SMEM_BYTES value matches current behavior (the temp is in shared); at TIER_LITE/TIER_MINIMAL the SMEM_BYTES value drops (temp moved out) and the WORKSPACE_BYTES value covers the moved temp.

Inline-CUDA usage example:

__global__ void my_outer_kernel(...) {
    extern __shared__ unsigned char s_arena[];
    // ... slice s_arena for your own buffers ...
    T *s_grid_temp = /* slice for GRiD primitive */;

    // At launch site we passed sizeof(s_grid_temp) =
    //   grid::FDSVA_SO_INNER_SMEM_BYTES<T, grid::TIER_MINIMAL>()
    //   == 0 at MINIMAL — no smem reserved for GRiD temp.
    // Workspace was malloc'd to:
    //   grid::FDSVA_SO_INNER_WORKSPACE_BYTES<T, grid::TIER_MINIMAL>()
    //   == 4 * NV^3 * sizeof(T) at MINIMAL.

    grid::fdsva_so_inner<T, grid::TIER_MINIMAL>(
        s_df2, s_idsva_so, s_Minv, s_df_du, s_XImats,
        /* s_temp */ nullptr,       // unused at MINIMAL
        /* d_workspace */ workspace, // global mem
        gravity);
}

References

• Tier baseline matrix: test/diagnostics/tier_baseline.py measures the per-(algo, robot) register + smem + spill data that drives tier decisions (the report it writes is a local artifact, not part of the repository); the shipped picks were made from an RTX 5090 (sm_120) run.

• Smoke test: test/diagnostics/tier_instantiation_smoke.py verifies that all 9 single-overload kernels compile at all 3 tiers and that 17 static_asserts validate the per-tier SMEM/WORKSPACE invariants.

• Reusable arena helper: gen_declare_shared_arena (and its tier_workspace_expr keyword) in grid_codegen/helpers/_code_generation_helpers.py.

• Existing bench harness: test/benchmarks/run_multi_version.py (multi-column driver), test/benchmarks/baselines/{grid,pinocchio,frax,mjx}/ (per- baseline runners), test/benchmarks/generate_report.py (output renderer that understands frax_cpu/frax_gpu columns and renders missing cells as —).

Design rationale

One body per algorithm

The math is identical at every tier; if constexpr branches only differ in pointer routing (s_temp versus d_workspace). One body per algorithm, with up to two pointer-routing branches, means less duplicated code and fewer drift bugs than three independent bodies per tier.

No tier knob on the Python surface

The Python wrapper is a sealed product that never touches nvcc, so the tier is a per-algorithm baked choice rather than a runtime argument: the autotuners write {tier, threads} into grid_assets/launch_configs/, codegen bakes it as grid::launch_cfg<GRID_ALGO_*>::TIER, and every binding launch site (the numpy/pybind host-wrapper calls and the jax/torch direct kernel launches) instantiates that tier. Divergent-tier instantiations get their own dynamic-smem registration in init_grid_kernel_attrs; a distinct __global__ per tier is the reason that registration exists. Untuned robots and algorithms fall back to TIER_SHARED via the primary launch_cfg template.

Suggested thread counts

The tuned launch_cfg entries are per GPU and exist only for the robots the autotuners have visited. For everything else, and for a consumer that fuses several algorithms into one kernel, the header ends with one derived constant per algorithm, grid::<ALGO>_SUGGESTED_THREADS (and grid::TRACKING_COST_SUGGESTED_THREADS when the plant preset is emitted). It is the smallest multiple of 32 that covers the widest block-parallel compute loop reachable from the algorithm’s kernel, read back from the emitted code at generation time (grid_codegen/suggested_threads.py); copy and zero-fill loops do not count, because striding them costs nothing against the serial compute of the loops that do arithmetic. The value is independent of the GPU and needs no tuning; a fused kernel takes the maximum over the algorithms it calls instead of MAX_PERF_LEVEL_THREADS. Values are clamped to [32, 1024]; an algorithm whose loops exceed 1024 strides, so 1024 is the honest answer there. launch_cfg<GRID_ALGO_*>::THREADS remains the tuned override where one exists. On iiwa14 the inverse dynamics suggestion is 96 and the inverse-dynamics gradient 352, which is also its MAX_PERF_LEVEL_THREADS.

A separate LITE smem target

With a single codegen-time spill pick, TIER_LITE collapses onto SHARED or MINIMAL on robots with multi-rung spill ladders. Codegen therefore computes three picks per algorithm (select_shared_tier_3way against the ~48 KB cuda_target_lite_shared_mem_bytes target), emits per-tier if constexpr bodies where the picks diverge, and routes the s_temp slot through gen_declare_shared_arena(tier_workspace_expr=...) so that LITE is an intermediate level.

Relation to block size

Removing cuBLASDx from the generated code (before the first packaged release) made every routine block-size-agnostic below MAX_PERF_LEVEL_THREADS (set_threads_per_block); the tier system extends that above MAX_PERF_LEVEL_THREADS via TIER_MINIMAL’s launch_bounds=1024. The emission layers the tier templates attach to (_inner / _device / _kernel) are described in Codegen architecture.