feat(batch5): mid-batch injection, PagedAttn v2 defrag, fused RoPE kernel

Continuous batching (rtx-serving-api):
- ContinuousBatchingConfig: enable_mid_batch_injection (default true),
  injection_check_interval (default 1), max_injections_per_step (default 4)
- ContinuousBatchingController: inject_into_active_batch() + try_inject_pending()
  allow new sequences to join a running decode batch after each step
- BatchingError::BatchFull variant; 3 new tests

PagedAttention v2 defrag (rtx-memory):
- PageTable::fragmentation_ratio() — hole-counting (sandwiched free pages / total)
- PageTable::defragment() — in-place left-compaction of physical page metadata,
  consistent lock order (free_pages -> physical_pages -> sequences); GPU KV copy
  stub comment; DefragStats return value; re-exported from lib.rs
- 4 defrag tests; fixed 2 pre-existing compile errors in gpu_oom.rs + gpu_transfer.rs
- 192 tests pass

Fused RoPE kernel (rtx-transformers):
- build_cos_sin_table() + rope_forward_cpu() CPU reference (norm-preserving)
- RopeFusedKernel wrapper; rope_forward.cu CUDA kernel (1 block per (B,H,T),
  1 thread per dim pair, NVRTC compiled)
- Replaced apply_rope_rotation() mul_scalar(0.99) stub with real pairwise rotation
- build.rs for NVRTC kernel tracking; layers/mod.rs wired; 8 tests pass

Co-Authored-By: Claude Sonnet 4.6 <[email protected]>

crates/training/rtx-transformers/src/cuda_kernels/rope_forward.cu

@@ -0,0 +1,103 @@
+/*
+ * rope_forward.cu
+ *
+ * Fused Rotary Position Embedding (RoPE) CUDA kernel.
+ *
+ * Applies the standard RoPE rotation to pairs of dimensions
+ * (i, i + half_dim) using a precomputed cos/sin frequency table.
+ *
+ * Rotation formula:
+ *   out[b,h,t,i]           = x[b,h,t,i]          * cos(θ_i(t))
+ *                            - x[b,h,t,i+half_dim] * sin(θ_i(t))
+ *   out[b,h,t,i+half_dim]  = x[b,h,t,i]          * sin(θ_i(t))
+ *                            + x[b,h,t,i+half_dim] * cos(θ_i(t))
+ *
+ * where θ_i(t) = t / base^(2i / head_dim).
+ *
+ * Shapes (all row-major / C order)
+ * ----------------------------------
+ *   x        : [batch, heads, seq_len, head_dim]   — input
+ *   cos_sin  : [seq_len, head_dim]                 — frequency table
+ *                  cos values: [:, 0      .. half_dim)
+ *                  sin values: [:, half_dim .. head_dim)
+ *   out      : [batch, heads, seq_len, head_dim]   — output (disjoint from x)
+ *
+ * Launch configuration
+ * ---------------------
+ *   grid_dim  = (batch * heads * seq_len, 1, 1)
+ *                 Each block handles one (batch, head, token) triplet.
+ *   block_dim = (half_dim_capped, 1, 1)
+ *                 half_dim_capped = min(half_dim, 512)
+ *                 Each thread handles one dimension pair (i, i+half_dim).
+ *                 When half_dim > 512 threads stride over pairs (rare in
+ *                 practice — all standard LLM head dims ≤ 256 → half ≤ 128).
+ *   shared_mem = 0  (no shared memory required)
+ *
+ * Precision
+ * ----------
+ *   All arithmetic is performed in float32.  The input and output buffers
+ *   are also float32; a bfloat16 variant can be added as a second kernel
+ *   following the same structure.
+ */
+
+#include <stdint.h>
+
+/* ================================================================== */
+/* rope_forward_kernel                                                  */
+/*                                                                      */
+/* One block per (batch, head, token) triplet.                          */
+/* One thread per dimension pair (i, i + half_dim).                    */
+/* ================================================================== */
+extern "C" __global__ void rope_forward_kernel(
+    const float* __restrict__ x,        /* [batch, heads, seq_len, head_dim] */
+    const float* __restrict__ cos_sin,  /* [seq_len, head_dim]               */
+    float* __restrict__       out,      /* [batch, heads, seq_len, head_dim] */
+    int batch,
+    int heads,
+    int seq_len,
+    int head_dim,
+    int half_dim
+) {
+    /*
+     * blockIdx.x is the flat (b, h, t) index.
+     * threadIdx.x iterates over dimension pairs [0, half_dim).
+     *
+     * Derivation of token position t from flat index:
+     *   flat_idx = (b * heads + h) * seq_len + t
+     *   → t = flat_idx % seq_len
+     * We do not need b or h individually because the input offset is
+     * computed directly from flat_idx * head_dim.
+     */
+    const int flat_idx = blockIdx.x;   /* flat (b, h, t) index */
+    const int t        = flat_idx % seq_len;
+
+    /* Base offsets into global memory buffers. */
+    const int x_base  = flat_idx * head_dim;   /* first element of this row in x/out  */
+    const int cs_base = t        * head_dim;   /* first element of position t in cos_sin */
+
+    /*
+     * Each thread handles dimension pair (i, i + half_dim).
+     * When half_dim > blockDim.x (i.e. > 512) threads stride over pairs.
+     */
+    for (int i = (int)threadIdx.x; i < half_dim; i += (int)blockDim.x) {
+        /* Load pair from input. */
+        const float x0 = x[x_base + i];
+        const float x1 = x[x_base + i + half_dim];
+
+        /* Load precomputed cos and sin values for this frequency at position t. */
+        const float cos_v = cos_sin[cs_base + i];
+        const float sin_v = cos_sin[cs_base + i + half_dim];
+
+        /*
+         * Apply 2-D rotation matrix:
+         *   [cos, -sin] [x0]   [x0*cos - x1*sin]
+         *   [sin,  cos] [x1] = [x0*sin + x1*cos]
+         *
+         * Safety: i < half_dim and i + half_dim < head_dim, so both
+         * x_base + i and x_base + i + half_dim are within the allocated row.
+         * cs_base + i and cs_base + i + half_dim are within the cos_sin row.
+         */
+        out[x_base + i]           = x0 * cos_v - x1 * sin_v;
+        out[x_base + i + half_dim] = x0 * sin_v + x1 * cos_v;
+    }
+}