Generate CUDA for a robot#

GRiD’s code generator (grid_codegen/ in this repository) produces a robot-specific grid.cuh header; generating the header does not compile or run CUDA. The grid-generate CLI ships in the grid-rbd wheel, so pip install grid-rbd (installation) is enough; only a source checkout also needs its Git submodules populated.

Generate a header#

From the repository root (the sample URDFs under grid_assets/robots/ are in the checkout, not the wheel; substitute your own URDF path):

grid-generate grid_assets/robots/iiwa14.urdf --algorithm-list dynamics -o /tmp/iiwa_grid.cuh

--algorithm-list restricts emission to named algorithms or profiles (next section). Omit it to use the full profile. Add -f for a floating base; use --no-mujoco-kernels to omit the additional MuJoCo-convention kernels. For collision geometry, see Collision code generation.

Profiles#

--algorithm-list (and algorithm_list= / codegen_profile= in Python) accepts algorithm keys (underscored, e.g. inverse_dynamics) and profile names (hyphenated). A profile is a curated set; whatever you name is expanded with its dependencies before emission. The profiles, from grid_codegen/_algo_profiles.py:

Profile

Selects

all (default)

every default algorithm: dynamics, gradients, second order, kinematics, integrators, regressors, centroidal and energy terms (the opt-in keys frame_jacobian, frame_jacobian_dot, osc_inertia, end_effector_pose_runtime, end_effector_pose_gradient_runtime, idsva_so_world_frame, integrator_hessian must be named explicitly)

dynamics

inverse / forward dynamics, minv, aba, crba, both first-order gradients, the second-order pair, and the integrator family

dynamics-core

inverse_dynamics, minv, forward_dynamics

dynamics-gradients

the core three plus inverse_dynamics_gradient, forward_dynamics_gradient, f_ext_gradient

second-order

the core three, both first-order gradients, idsva_so_body_frame, fdsva_so

integrators

the core three, both first-order gradients, integrator, integrator_gradient, integrator_with_gradient

kinematics

end_effector_pose (alias: kinematics-only)

kinematics-derivatives

end_effector_pose with its gradient and Hessian

frame-jacobian

end_effector_pose, minv, frame_jacobian, frame_jacobian_dot, osc_inertia

regressor

inverse_dynamics, inverse_dynamics_regressor and the kinetic / potential energy regressors

energy-regressor

inverse_dynamics, end_effector_pose and the kinetic / potential energy regressors

fd-param-gradient

the core three, inverse_dynamics_regressor, forward_dynamics_parameter_gradient

f-ext-gradient

inverse_dynamics, minv, f_ext_gradient

The grid_plant block (costs, barriers, the plant step and the cost menu) rides on the request: it is emitted for all, for any request that carries a plant algorithm (the integrators, fdsva_so, the end-effector pose derivatives, com, ccrba), baked cost or contact frames, or a multi-target batch, and whenever --plant (emit_plant=True in Python) is given. kinematics and dynamics-core headers carry no plant block unless asked.

Fewer kernels in the header means a smaller header, a faster nvcc and less RAM; see Compile time and footprint below.

Generate from Python#

from URDFParser import URDFParser
from grid_codegen import GRiDCodeGenerator

robot = URDFParser().parse("grid_assets/robots/iiwa14.urdf")
generator = GRiDCodeGenerator(robot)
generator.gen_all_code(
    algorithm_list=["dynamics"],
    output_path="/tmp/iiwa_grid.cuh",
)

GRiDCodeGenerator is the Python class name, not a separate repository or installable module. GRiD’s code generator vendors GLASS device-side linear algebra into the header by default. Set vendor_glass=False when using an external glass.cuh on your C++ include path.

Use the generated code#

Include it in your kernel#

This is the entry point for CUDA applications: the header is self-contained (GLASS device-side linear algebra is vendored in by default), so a kernel that #includes it needs only nvcc. Every algorithm has a _device routine that runs inside your block on shared-memory operands, which is how you fuse GRiD dynamics into an MPC, RL or controls kernel:

#include "grid.cuh"   // generated: grid-generate robot.urdf -o grid.cuh

template <typename T>
__global__ void my_kernel(T *d_tau, const T *d_q, const T *d_qd, const T *d_qdd,
                          const grid::robotModel<T> *model, T gravity) {
    __shared__ T s_q[grid::NUM_JOINTS], s_qd[grid::NUM_JOINTS],
                 s_qdd[grid::NUM_JOINTS], s_tau[grid::NUM_JOINTS];
    for (int i = threadIdx.x; i < grid::NUM_JOINTS; i += blockDim.x) {
        s_q[i] = d_q[i]; s_qd[i] = d_qd[i]; s_qdd[i] = d_qdd[i];
    }
    __syncthreads();
    grid::inverse_dynamics_device<T>(s_tau, s_q, s_qd, s_qdd, model,
                                     /*d_f_ext=*/nullptr, gravity);
    __syncthreads();   // fuse your own math on s_tau here
    for (int i = threadIdx.x; i < grid::NUM_JOINTS; i += blockDim.x) d_tau[i] = s_tau[i];
}
// host: see "Initialising the header" below for the two init spellings;
// auto *model = grid::init_robotModel<float>();
// my_kernel<float><<<1, 128, grid::INVERSE_DYNAMICS_DEVICE_DYNAMIC_SHARED_MEM_BYTES<float>()>>>(...);

nvcc -arch=sm_90 -std=c++17 my_kernel.cu -o my_kernel   # pick your GPU's arch

Initialising the header: aggregate or per algorithm#

Before the first launch the generated kernels’ dynamic shared-memory attributes must be registered and the host streams created. The header offers two spellings of that:

// (a) aggregate — simplest. Registers EVERY kernel in the header, so the
//     translation unit also compiles every one of them.
cudaStream_t *streams = grid::init_grid<float>();

// (b) per algorithm — register only what you launch; what the test
//     runners and the grid_rbd bindings use.
cudaStream_t *streams = grid::init_grid_streams<float>();
gpuErrchk(grid::init_grid_kernel_attr_inverse_dynamics_checked<float>());
// one init_grid_kernel_attr_<algo>_checked<T>() per algorithm you launch

init_grid<T>() is the shortest path when the header was generated with a trimmed --algorithm-list, because then “every kernel” is the handful you asked for. With a full-profile header it is what makes a small .cu file take minutes to compile: a translation unit compiles every kernel it takes the address of, so compile time and peak RAM scale with the number of kernels in the header, not with the number you call. The per-algorithm twins init_grid_streams<T>() + init_grid_kernel_attr_<algo>_checked<T>() (one per algorithm key, e.g. ..._attr_forward_dynamics_checked, ..._attr_minv_checked; grep your header for init_grid_kernel_attr_) register only what you name, so the rest of the header stays uncompiled. The _checked form returns cudaError_t and never exits; Library-safe initialization and cleanup has the full contract and the matching teardown.

Compile time and footprint#

Three levers decide how long nvcc takes on a generated header and how much RAM it needs:

  • Emit less: --algorithm-list / algorithm_list= with the profiles above (and --no-mujoco-kernels on a floating base, where the MuJoCo-convention twins roughly double a build). The header only contains what you name plus dependencies.

  • Register less: with a full header, the per-algorithm init twins above compile only the kernels you launch.

  • Fit the device: the resource tiers trade shared memory for global workspace per kernel; the tier you instantiate is the one that gets compiled.

Measured cold-build times, RAM and the warm path for the Python wrapper’s .so (which compiles the same header) are tabulated on Fast Robot Setup.

examples/cuda/ holds the full version of this kernel (the easy _device call and the full-control _inner call with explicit workspace), plus build_and_validate.sh, which generates the header, builds, runs and checks the output against the reference implementation; the notebook examples/notebooks/07_inline_cuda.ipynb is the same walk-through run inline. Each _device routine’s dynamic shared-memory requirement is exposed as grid::<ALGO>_DEVICE_DYNAMIC_SHARED_MEM_BYTES<T>(); the resource tiers page covers fitting a routine into a tight outer kernel.