perf: transpose-quantize backward operands to route all gemms nt

- quantize gains out_layout (0 row-major / 1 transposed / 2 single-read dual-write); modes 1/2 run a new 32x32 smem-tile transpose kernel
- backward feeds g8/w8T and g8T/x8T to trans_b=True gemms, dropping the NN-swap and TT crosswise kernels from training; fp8 weights keep the swap fallback
- a 64x64 tile variant tied on the real step mix and was reverted; noted in the kernel header

Benchmark: NVIDIA L20, 1.2B model, full train step fwd+bwd+CE
- M=8192: fp8 551.8 -> 532.2 ms, 1.21x -> 1.26x vs bf16; M=2048 0.90x -> 0.95x
- kernel-level grad_x +3.7..12.4%, grad_w +13.8..20.8%; layouts byte-exact, fp8 tests 36/36

csrc/kernels/fp8/quantize.cuh

@@ -169,5 +169,89 @@ void launch_fp8_quantize(const FP8QuantizeParams& p, cudaStream_t stream) {
     fp8_quantize_kernel<Fmt, InT><<<blocks, kThreads, 0, stream>>>(p);
 }
 
+// Tiled transpose quantize (out_layout 1/2): reads the [rows][cols]
+// row-major input once and writes the fp8 bytes transposed ([cols][rows],
+// so the contract dim lands K-contiguous for NT GEMM operands) and, in
+// mode 2, the plain row-major copy too. A 32x32 tile stages through shared
+// memory: input-row-major loads and output writes both stay coalesced, and
+// the byte-wide staging is conflict-free — the +4 pad makes the store
+// stride 9 (words) coprime with the 32 banks and the load is a 32-byte
+// broadcast segment. A 64x64 split-half variant (16 elems/thread, paired
+// 2-byte scatter stores) measured +21% on L2-resident shapes but -3..5%
+// on the DRAM-bound ones that carry the training traffic (occupancy and
+// memory-level parallelism, not instruction count, gate the DRAM regime);
+// weighted by the real step's mix the two tie, so the simpler tile stays.
+template <FP8Format Fmt, typename InT>
+__global__ void fp8_quantize_tiled_kernel(FP8QuantizeParams p) {
+    constexpr int kTile = 32;
+    __shared__ uint8_t tile[kTile][kTile + 4];
+    const float mult = *p.scale;
+    const auto* x = static_cast<const InT*>(p.input_ptr);
+    const int r0 = blockIdx.y * kTile;
+    const int c0 = blockIdx.x * kTile;
+    const int r = r0 + threadIdx.y * 4;
+    const int c = c0 + threadIdx.x;
+
+    uint8_t q[4];
+    float local_amax = 0.0f;
+#pragma unroll
+    for (int j = 0; j < 4; ++j) {
+        q[j] = 0;
+        if (r + j < p.rows && c < p.cols) {
+            const float v =
+                quant_in_traits<InT>::to_float(x[(int64_t)(r + j) * p.cols + c]);
+            local_amax = fmaxf(local_amax, fabsf(v));
+            if constexpr (Fmt == FP8Format::E5M2)
+                q[j] = __nv_fp8_e5m2(v * mult).__x;
+            else
+                q[j] = __nv_fp8_e4m3(v * mult).__x;
+        }
+    }
+    if (p.out_layout == 2) {
+        uint8_t* out = static_cast<uint8_t*>(p.output_ptr);
+#pragma unroll
+        for (int j = 0; j < 4; ++j)
+            if (r + j < p.rows && c < p.cols)
+                out[(int64_t)(r + j) * p.cols + c] = q[j];
+    }
+#pragma unroll
+    for (int j = 0; j < 4; ++j) tile[threadIdx.x][threadIdx.y * 4 + j] = q[j];
+    __syncthreads();
+    // Transposed scatter: output element (c, r) lives at c * rows + r; r
+    // tracks threadIdx.x so each warp writes one contiguous run. The read
+    // swaps the staging indices — tile[col][row] was written, so the value
+    // for input (r0+tx, c0+ty*4+j) sits at tile[ty*4+j][tx].
+    uint8_t* out_t = static_cast<uint8_t*>(p.output_transposed_ptr);
+#pragma unroll
+    for (int j = 0; j < 4; ++j) {
+        const int oc = c0 + threadIdx.y * 4 + j;
+        if (oc < p.cols && r0 + threadIdx.x < p.rows)
+            out_t[(int64_t)oc * p.rows + r0 + threadIdx.x] =
+                tile[threadIdx.y * 4 + j][threadIdx.x];
+    }
+    if (p.amax) {
+        local_amax = warp_reduce_max(local_amax);
+        __shared__ float slots[8];
+        // blockDim.x is 32, so warp id == threadIdx.y; only complete warps
+        // exist (blockDim.y == 8).
+        if (threadIdx.x == 0) slots[threadIdx.y] = local_amax;
+        __syncthreads();
+        if (threadIdx.x == 0 && threadIdx.y == 0) {
+            float v = 0.0f;
+            for (int w = 0; w < (int)blockDim.y; ++w) v = fmaxf(v, slots[w]);
+            atomic_max_float(p.amax, v);
+        }
+    }
+}
+
+template <FP8Format Fmt, typename InT>
+void launch_fp8_quantize_tiled(const FP8QuantizeParams& p,
+                               cudaStream_t stream) {
+    const dim3 grid((p.cols + 31) / 32, (p.rows + 31) / 32);
+    if (grid.x == 0 || grid.y == 0) return;
+    fp8_quantize_tiled_kernel<Fmt, InT>
+        <<<grid, dim3(32, 8), 0, stream>>>(p);
+}
+
 }  // namespace fp8
 }  // namespace astrai