Dynamic shared memory size at kernel launch

[tetsuo-cpp]

Summary

Allow shared-memory footprint to be chosen at launch time rather than only via compile-time \! shared name T[N] constants.

Motivation

Static shared sizes force kernels to compile for a max problem size (e.g. attention SCORES f64[1024]). Dynamic shared memory (CUDA extern __shared__ / launch sharedMemBytes) is standard for libraries that adapt to runtime seq_len or tile sizes.

Proposed design (sketch)

Language

\! shared SCORES f64[]       \ dynamic, extent from host
\! shared SCORES f64[MAX]    \ static max + optional dynamic use

Or a header directive:

\! dynamic-shared 8192       \ bytes requested by default; host may override

Lowering

• MLIR GPU workgroup memory with dynamic size where supported
• Host API / runner passes shared_mem_bytes at launch
• Kernel uses base pointers as today; bounds are programmer-managed (debug checks optional)

Acceptance criteria

• Documented way to request dynamic shared bytes from host
• At least one kernel uses launch-time shared size successfully on GPU
• Interaction with static \! shared declarations defined (replace vs additional)
• Metadata schema includes dynamic shared requirement
• LIT/GPU tests

Related

• Dynamic-extent params (separate)
• Runner (Replace warpforth-runner C++ CLI with Python cuda-python + JSON protocol #50) and Python launch API
• Attention demo seq-len limit