L-BFGS kernel: 'too many resources requested for launch' on Blackwell (sm_120) — register pressure, not shared memory

[InfraestruturaRobotec]

Environment

GPU	NVIDIA GeForce RTX 5090 (Blackwell, sm_120, compute capability 12.0)
Driver	580.126.09
torch	2.7.0+cu128 (arch_list includes sm_120, compute_120)
cuRobo	v0.7.6
CUDA toolkit	12.8

Summary

On sm_120, cuRobo's L-BFGS kernel fails to launch with

CUDA error: too many resources requested for launch

from curobo/opt/newton/lbfgs.py:167 → curobo/curobolib/opt.py:58.

This is register pressure, not shared memory, and -maxrregcount works
around it — but the real fix is a __launch_bounds__ on the kernel (or a block
size that adapts to the device's register budget), because the workaround costs
occupancy and spill traffic on every architecture it is applied to.

Reproduce

Build cuRobo for Blackwell and run any MotionGen.plan_single:

TORCH_CUDA_ARCH_LIST="12.0+PTX" pip install --no-build-isolation -e .

(Compiling for sm_89 only produces a different, earlier error —
no kernel image is available for execution on the device — which is the
expected arch mismatch and not what this issue is about.)

Evidence that it is registers, not shared memory

lbfgs_step_kernel.cu requests dynamic shared memory and also has a
use_shared_buffers fallback path that uses a much smaller smemsize:

const int smemsize               = history_m * v_dim * sizeof(float);
const int shared_buffer_smemsize = (((3 * v_dim) + 1) * history_m + 32) * sizeof(float);

Forcing LBFGSOptConfig.use_shared_buffers_kernel = False routes to the small
path — and it fails identically. Since "too many resources" covers shared
memory and registers, and the shared-memory request has been eliminated as a
variable, registers are the remaining candidate.

The launch is:

int threadsPerBlock = v_dim;      // lbfgs_step_kernel.cu:815
int blocksPerGrid   = batch_size;

so regs_per_thread * v_dim must stay under the 64K registers-per-block limit.
sm_120 codegen allocates more registers per thread than sm_89 did for the
same kernel, which pushes it over for the v_dim values used by a dual-arm
humanoid config.

Workaround

Capping registers at compile time makes every kernel launch successfully:

# setup.py, extra_cuda_args["nvcc"]
"-maxrregcount=64",

Verified afterwards with cuobjdump --list-elf that all five extensions carry
sm_120, and MotionGen.plan_single returns success=True with a valid
trajectory. A full Isaac Sim data-collection pipeline then runs with zero CUDA
errors.

Suggested fix

__launch_bounds__(maxThreadsPerBlock, minBlocksPerMultiprocessor) on
lbfgs_update_buffer_and_step*, so nvcc bounds registers for that kernel
specifically rather than requiring a translation-unit-wide -maxrregcount that
penalises the other kernels. Alternatively, query
cudaDeviceGetAttribute(cudaDevAttrMaxRegistersPerBlock) and clamp
threadsPerBlock accordingly.

Happy to test a patch on this hardware.