""" triton_svd3.py — Fused Triton SVD kernel for batched M×3 matrices. Target use case: CIFAR-sized images (32×32=1024 pixels, 3 channels) where cuSOLVER overhead dominates because the "thin" dimension is only 3. Architecture: 1. Each program handles one (M×3) matrix from the batch 2. Compute 3×3 Gram matrix G = A^T A (6 unique values, in registers) 3. Diagonalize G via Jacobi rotations (3×3 converges in ≤6 sweeps) 4. S = sqrt(eigenvalues), V = eigenvectors 5. U = A @ V @ diag(1/S) (tiled reduction over M) The entire 3×3 eigensolver lives in scalar registers — zero shared memory, zero global memory round-trips. The only bandwidth cost is loading A and writing back U, S, Vh. Author: AbstractPhil / Claude """ import triton import triton.language as tl import torch import math # ============================================================================ # Core kernel: batched SVD for (B, M, 3) tensors # ============================================================================ @triton.jit def _svd3_kernel( # Pointers A_ptr, # (B, M, 3) input U_ptr, # (B, M, 3) left singular vectors (thin) S_ptr, # (B, 3) singular values Vh_ptr, # (B, 3, 3) right singular vectors transposed # Dimensions M: tl.constexpr, # spatial dim (1024 for CIFAR) BLOCK_M: tl.constexpr, # tile size for M-dimension loads JACOBI_ITERS: tl.constexpr, # number of full Jacobi sweeps EPS: tl.constexpr, # numerical floor ): """ One program instance = one (M, 3) matrix in the batch. """ bid = tl.program_id(0) # batch index # --------------------------------------------------------------- # Stage 1: Compute 3×3 Gram matrix G = A^T @ A # # G is symmetric, so we only need 6 accumulators: # g00 g01 g02 # g11 g12 # g22 # --------------------------------------------------------------- g00 = tl.zeros([], dtype=tl.float32) g01 = tl.zeros([], dtype=tl.float32) g02 = tl.zeros([], dtype=tl.float32) g11 = tl.zeros([], dtype=tl.float32) g12 = tl.zeros([], dtype=tl.float32) g22 = tl.zeros([], dtype=tl.float32) base = bid * M * 3 # Tiled accumulation over the spatial dimension for block_start in range(0, M, BLOCK_M): offs = tl.arange(0, BLOCK_M) # [0, 1, ..., BLOCK_M-1] row_idx = block_start + offs # actual row indices mask = row_idx < M # Load 3 columns for this tile # A[bid, row, c] = A_ptr[bid*M*3 + row*3 + c] ptr0 = base + row_idx * 3 + 0 ptr1 = base + row_idx * 3 + 1 ptr2 = base + row_idx * 3 + 2 a0 = tl.load(A_ptr + ptr0, mask=mask, other=0.0).to(tl.float32) a1 = tl.load(A_ptr + ptr1, mask=mask, other=0.0).to(tl.float32) a2 = tl.load(A_ptr + ptr2, mask=mask, other=0.0).to(tl.float32) # Accumulate outer products g00 += tl.sum(a0 * a0) g01 += tl.sum(a0 * a1) g02 += tl.sum(a0 * a2) g11 += tl.sum(a1 * a1) g12 += tl.sum(a1 * a2) g22 += tl.sum(a2 * a2) # --------------------------------------------------------------- # Stage 2: 3×3 Jacobi eigendecomposition (all in scalars) # # Cyclic Jacobi: rotate pairs (0,1), (0,2), (1,2) each sweep. # For 3×3 symmetric PSD, 4-6 sweeps is overkill-level convergence. # # We maintain: # - g_ij: the evolving matrix entries (symmetric, 6 values) # - v_ij: the eigenvector matrix (9 values, starts as I) # --------------------------------------------------------------- # Eigenvector accumulator V (row-major: v_rc = V[r,c]) v00 = 1.0; v01 = 0.0; v02 = 0.0 v10 = 0.0; v11 = 1.0; v12 = 0.0 v20 = 0.0; v21 = 0.0; v22 = 1.0 for _sweep in range(JACOBI_ITERS): # --- Rotation (p=0, q=1): zero out g01 --- # # Jacobi rotation angle: # if |g_pq| < eps: skip # tau = (g_qq - g_pp) / (2 * g_pq) # t = sign(tau) / (|tau| + sqrt(1 + tau^2)) # c = 1/sqrt(1+t^2), s = t*c # -- pair (0, 1) -- off_diag = g01 diag_diff = g11 - g00 abs_off = tl.abs(off_diag) # Compute rotation tau_01 = tl.where(abs_off > EPS, diag_diff / (2.0 * off_diag), 0.0) abs_tau = tl.abs(tau_01) t_01 = tl.where( abs_off > EPS, tl.where(tau_01 >= 0, 1.0, -1.0) / (abs_tau + tl.sqrt(1.0 + tau_01 * tau_01)), 0.0, ) c01 = 1.0 / tl.sqrt(1.0 + t_01 * t_01) s01 = t_01 * c01 # Apply Givens to G (symmetric update for p=0, q=1) new_g00 = c01*c01*g00 - 2.0*s01*c01*g01 + s01*s01*g11 new_g11 = s01*s01*g00 + 2.0*s01*c01*g01 + c01*c01*g11 new_g01 = 0.0 # This is the point new_g02 = c01*g02 - s01*g12 new_g12 = s01*g02 + c01*g12 # g22 unchanged g00 = new_g00; g11 = new_g11; g01 = new_g01 g02 = new_g02; g12 = new_g12 # Apply to V columns: V[:, 0], V[:, 1] nv00 = c01*v00 - s01*v01; nv01 = s01*v00 + c01*v01 nv10 = c01*v10 - s01*v11; nv11 = s01*v10 + c01*v11 nv20 = c01*v20 - s01*v21; nv21 = s01*v20 + c01*v21 v00 = nv00; v01 = nv01; v10 = nv10; v11 = nv11; v20 = nv20; v21 = nv21 # -- pair (0, 2) -- off_diag = g02 diag_diff = g22 - g00 abs_off = tl.abs(off_diag) tau_02 = tl.where(abs_off > EPS, diag_diff / (2.0 * off_diag), 0.0) abs_tau = tl.abs(tau_02) t_02 = tl.where( abs_off > EPS, tl.where(tau_02 >= 0, 1.0, -1.0) / (abs_tau + tl.sqrt(1.0 + tau_02 * tau_02)), 0.0, ) c02 = 1.0 / tl.sqrt(1.0 + t_02 * t_02) s02 = t_02 * c02 new_g00 = c02*c02*g00 - 2.0*s02*c02*g02 + s02*s02*g22 new_g22 = s02*s02*g00 + 2.0*s02*c02*g02 + c02*c02*g22 new_g02 = 0.0 new_g01 = c02*g01 - s02*g12 new_g12_b = s02*g01 + c02*g12 g00 = new_g00; g22 = new_g22; g02 = new_g02 g01 = new_g01; g12 = new_g12_b nv00 = c02*v00 - s02*v02; nv02 = s02*v00 + c02*v02 nv10 = c02*v10 - s02*v12; nv12 = s02*v10 + c02*v12 nv20 = c02*v20 - s02*v22; nv22 = s02*v20 + c02*v22 v00 = nv00; v02 = nv02; v10 = nv10; v12 = nv12; v20 = nv20; v22 = nv22 # -- pair (1, 2) -- off_diag = g12 diag_diff = g22 - g11 abs_off = tl.abs(off_diag) tau_12 = tl.where(abs_off > EPS, diag_diff / (2.0 * off_diag), 0.0) abs_tau = tl.abs(tau_12) t_12 = tl.where( abs_off > EPS, tl.where(tau_12 >= 0, 1.0, -1.0) / (abs_tau + tl.sqrt(1.0 + tau_12 * tau_12)), 0.0, ) c12 = 1.0 / tl.sqrt(1.0 + t_12 * t_12) s12 = t_12 * c12 new_g11 = c12*c12*g11 - 2.0*s12*c12*g12 + s12*s12*g22 new_g22 = s12*s12*g11 + 2.0*s12*c12*g12 + c12*c12*g22 new_g12 = 0.0 new_g01 = c12*g01 - s12*g02 new_g02_b = s12*g01 + c12*g02 g11 = new_g11; g22 = new_g22; g12 = new_g12 g01 = new_g01; g02 = new_g02_b nv01 = c12*v01 - s12*v02; nv02 = s12*v01 + c12*v02 nv11 = c12*v11 - s12*v12; nv12 = s12*v11 + c12*v12 nv21 = c12*v21 - s12*v22; nv22 = s12*v21 + c12*v22 v01 = nv01; v02 = nv02; v11 = nv11; v12 = nv12; v21 = nv21; v22 = nv22 # --------------------------------------------------------------- # Stage 2b: Extract eigenvalues, sort descending # Diagonal of G now holds eigenvalues of A^T A # --------------------------------------------------------------- eig0 = tl.maximum(g00, EPS) eig1 = tl.maximum(g11, EPS) eig2 = tl.maximum(g22, EPS) s0 = tl.sqrt(eig0) s1 = tl.sqrt(eig1) s2 = tl.sqrt(eig2) # Sorting network for 3 elements (descending) # We need to sort S and permute V columns accordingly. # Approach: compare-and-swap on the scalar eigenvalues + V columns # # 3-element sorting network: (0,1), (0,2), (1,2) # swap(0, 1) if s0 < s1 do_swap = s0 < s1 s0, s1 = tl.where(do_swap, s1, s0), tl.where(do_swap, s0, s1) # Swap V columns 0 and 1 tv00 = v00; tv10 = v10; tv20 = v20 v00 = tl.where(do_swap, v01, v00); v01 = tl.where(do_swap, tv00, v01) v10 = tl.where(do_swap, v11, v10); v11 = tl.where(do_swap, tv10, v11) v20 = tl.where(do_swap, v21, v20); v21 = tl.where(do_swap, tv20, v21) # swap(0, 2) if s0 < s2 do_swap = s0 < s2 s0, s2 = tl.where(do_swap, s2, s0), tl.where(do_swap, s0, s2) tv00 = v00; tv10 = v10; tv20 = v20 v00 = tl.where(do_swap, v02, v00); v02 = tl.where(do_swap, tv00, v02) v10 = tl.where(do_swap, v12, v10); v12 = tl.where(do_swap, tv10, v12) v20 = tl.where(do_swap, v22, v20); v22 = tl.where(do_swap, tv20, v22) # swap(1, 2) if s1 < s2 do_swap = s1 < s2 s1, s2 = tl.where(do_swap, s2, s1), tl.where(do_swap, s1, s2) tv01 = v01; tv11 = v11; tv21 = v21 v01 = tl.where(do_swap, v02, v01); v02 = tl.where(do_swap, tv01, v02) v11 = tl.where(do_swap, v12, v11); v12 = tl.where(do_swap, tv11, v12) v21 = tl.where(do_swap, v22, v21); v22 = tl.where(do_swap, tv21, v22) # --------------------------------------------------------------- # Stage 2c: Write S and Vh # --------------------------------------------------------------- s_base = bid * 3 tl.store(S_ptr + s_base + 0, s0) tl.store(S_ptr + s_base + 1, s1) tl.store(S_ptr + s_base + 2, s2) # Vh = V^T (Vh[i,j] = V[j,i]) vh_base = bid * 9 tl.store(Vh_ptr + vh_base + 0, v00) # Vh[0,0] = V[0,0] tl.store(Vh_ptr + vh_base + 1, v10) # Vh[0,1] = V[1,0] tl.store(Vh_ptr + vh_base + 2, v20) # Vh[0,2] = V[2,0] tl.store(Vh_ptr + vh_base + 3, v01) # Vh[1,0] = V[0,1] tl.store(Vh_ptr + vh_base + 4, v11) # Vh[1,1] = V[1,1] tl.store(Vh_ptr + vh_base + 5, v21) # Vh[1,2] = V[2,1] tl.store(Vh_ptr + vh_base + 6, v02) # Vh[2,0] = V[0,2] tl.store(Vh_ptr + vh_base + 7, v12) # Vh[2,1] = V[1,2] tl.store(Vh_ptr + vh_base + 8, v22) # Vh[2,2] = V[2,2] # --------------------------------------------------------------- # Stage 3: Recover U = A @ V @ diag(1/S) # # U[:, c] = (1/S[c]) * A @ V[:, c] # Tiled over M to keep memory pressure bounded. # --------------------------------------------------------------- inv_s0 = 1.0 / (s0 + EPS) inv_s1 = 1.0 / (s1 + EPS) inv_s2 = 1.0 / (s2 + EPS) for block_start in range(0, M, BLOCK_M): offs = tl.arange(0, BLOCK_M) row_idx = block_start + offs mask = row_idx < M ptr0 = base + row_idx * 3 + 0 ptr1 = base + row_idx * 3 + 1 ptr2 = base + row_idx * 3 + 2 a0 = tl.load(A_ptr + ptr0, mask=mask, other=0.0).to(tl.float32) a1 = tl.load(A_ptr + ptr1, mask=mask, other=0.0).to(tl.float32) a2 = tl.load(A_ptr + ptr2, mask=mask, other=0.0).to(tl.float32) # U[:, 0] = (A @ V[:, 0]) / s0 u0 = (a0 * v00 + a1 * v10 + a2 * v20) * inv_s0 # U[:, 1] = (A @ V[:, 1]) / s1 u1 = (a0 * v01 + a1 * v11 + a2 * v21) * inv_s1 # U[:, 2] = (A @ V[:, 2]) / s2 u2 = (a0 * v02 + a1 * v12 + a2 * v22) * inv_s2 u_base = bid * M * 3 tl.store(U_ptr + u_base + row_idx * 3 + 0, u0, mask=mask) tl.store(U_ptr + u_base + row_idx * 3 + 1, u1, mask=mask) tl.store(U_ptr + u_base + row_idx * 3 + 2, u2, mask=mask) # ============================================================================ # Python wrapper # ============================================================================ def batched_svd3( A: torch.Tensor, block_m: int = 128, jacobi_iters: int = 6, ) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor]: """ Batched thin SVD for (B, M, 3) float32 tensors. Args: A: Input tensor of shape (B, M, 3). M can be anything (1024 for CIFAR). block_m: Tile size for the spatial dimension. 128 is good for M=1024. jacobi_iters: Number of cyclic Jacobi sweeps. 6 is overkill for 3×3. Returns: U: (B, M, 3) — thin left singular vectors S: (B, 3) — singular values, descending Vh: (B, 3, 3) — right singular vectors transposed """ assert A.ndim == 3 and A.shape[2] == 3, f"Expected (B, M, 3), got {A.shape}" assert A.is_cuda, "Input must be on CUDA" B, M, _ = A.shape A_f32 = A.contiguous().float() U = torch.empty((B, M, 3), dtype=torch.float32, device=A.device) S = torch.empty((B, 3), dtype=torch.float32, device=A.device) Vh = torch.empty((B, 3, 3), dtype=torch.float32, device=A.device) _svd3_kernel[(B,)]( A_f32, U, S, Vh, M=M, BLOCK_M=block_m, JACOBI_ITERS=jacobi_iters, EPS=1e-12, ) return U, S, Vh # ============================================================================ # Benchmark & validation harness # ============================================================================ def _test_correctness(B=256, M=1024): """Validate against torch.linalg.svd.""" A = torch.randn(B, M, 3, device="cuda", dtype=torch.float32) U, S, Vh = batched_svd3(A) # Reference: torch.linalg.svd on 3D is batched U_ref, S_ref, Vh_ref = torch.linalg.svd(A, full_matrices=False) # Singular values should match (tight — these are just 3 values) s_err = (S - S_ref).abs().max().item() print(f"[correctness] S max error: {s_err:.2e}") # Reconstruction: A ≈ U @ diag(S) @ Vh # Compare against torch's own recon error as the floor — # f32 accumulation over M=1024 rows means ~1e-3 is physical. recon = torch.bmm(U * S.unsqueeze(1), Vh) recon_ref = torch.bmm(U_ref * S_ref.unsqueeze(1), Vh_ref) r_err = (A - recon).abs().max().item() r_ref_err = (A - recon_ref).abs().max().item() print(f"[correctness] Recon error: {r_err:.2e} (ref: {r_ref_err:.2e})") # Orthogonality: U^T U ≈ I UtU = torch.bmm(U.transpose(1, 2), U) eye = torch.eye(3, device="cuda").expand(B, -1, -1) orth_err = (UtU - eye).abs().max().item() print(f"[correctness] U orthog error: {orth_err:.2e}") # S should be tight; recon only needs to match torch's own floor assert s_err < 1e-3, f"Singular value error {s_err} > 1e-3" assert r_err < r_ref_err * 2.0 + 1e-5, ( f"Recon error {r_err:.2e} > 2x torch ref {r_ref_err:.2e}" ) assert orth_err < 1e-3, f"Orthogonality error {orth_err} > 1e-3" print("[correctness] PASSED\n") def _cuda_timer(fn, warmup=25, iters=100): """ CUDA-event-timed benchmark. Returns (mean_ms, std_ms) over `iters` runs. Uses per-iteration events for proper distribution stats. """ # Warmup — let triton autotuner / cublas handle settle for _ in range(warmup): fn() torch.cuda.synchronize() starts = [torch.cuda.Event(enable_timing=True) for _ in range(iters)] ends = [torch.cuda.Event(enable_timing=True) for _ in range(iters)] for i in range(iters): starts[i].record() fn() ends[i].record() torch.cuda.synchronize() times = [starts[i].elapsed_time(ends[i]) for i in range(iters)] t = torch.tensor(times) return t.mean().item(), t.std().item(), t.median().item(), t.min().item(), t.max().item() def _profile_sweep(): """ Full profiling sweep: batch sizes × spatial dims × block sizes. Compares Triton SVD3 vs torch.linalg.svd with CUDA event timing. """ import json device_name = torch.cuda.get_device_name(0) print(f"{'='*72}") print(f" Triton SVD3 Profiling — {device_name}") print(f"{'='*72}\n") # ------------------------------------------------------------------ # Sweep 1: Batch scaling at fixed M=1024 (CIFAR spatial) # ------------------------------------------------------------------ print(f"{'─'*72}") print(f" SWEEP 1: Batch scaling (M=1024, BLOCK_M=128)") print(f"{'─'*72}") print(f" {'B':>6} {'Triton ms':>12} {'±std':>8} {'Torch ms':>12} {'±std':>8} {'Speedup':>8}") print(f" {'─'*62}") batch_sizes = [32, 64, 128, 256, 512, 1024, 2048, 4096, 8192, 16384] M = 1024 batch_results = [] for B in batch_sizes: A = torch.randn(B, M, 3, device="cuda", dtype=torch.float32) tri_mean, tri_std, tri_med, tri_min, tri_max = _cuda_timer( lambda: batched_svd3(A, block_m=128) ) tch_mean, tch_std, tch_med, tch_min, tch_max = _cuda_timer( lambda: torch.linalg.svd(A, full_matrices=False) ) speedup = tch_mean / (tri_mean + 1e-9) print(f" {B:>6} {tri_mean:>10.3f}ms {tri_std:>6.3f} {tch_mean:>10.3f}ms {tch_std:>6.3f} {speedup:>7.2f}x") batch_results.append({ "B": B, "M": M, "triton_mean_ms": round(tri_mean, 4), "triton_std_ms": round(tri_std, 4), "triton_median_ms": round(tri_med, 4), "triton_min_ms": round(tri_min, 4), "torch_mean_ms": round(tch_mean, 4), "torch_std_ms": round(tch_std, 4), "torch_median_ms": round(tch_med, 4), "torch_min_ms": round(tch_min, 4), "speedup": round(speedup, 3), }) del A torch.cuda.empty_cache() # ------------------------------------------------------------------ # Sweep 2: Spatial dim scaling at fixed B=1024 # ------------------------------------------------------------------ print(f"\n{'─'*72}") print(f" SWEEP 2: Spatial scaling (B=1024, BLOCK_M=128)") print(f"{'─'*72}") print(f" {'M':>6} {'Triton ms':>12} {'±std':>8} {'Torch ms':>12} {'±std':>8} {'Speedup':>8}") print(f" {'─'*62}") spatial_dims = [64, 256, 512, 1024, 2048, 4096] # 8×8 to 64×64 B = 1024 spatial_results = [] for M in spatial_dims: A = torch.randn(B, M, 3, device="cuda", dtype=torch.float32) tri_mean, tri_std, tri_med, tri_min, tri_max = _cuda_timer( lambda: batched_svd3(A, block_m=128) ) tch_mean, tch_std, tch_med, tch_min, tch_max = _cuda_timer( lambda: torch.linalg.svd(A, full_matrices=False) ) speedup = tch_mean / (tri_mean + 1e-9) equiv_hw = int(M**0.5) tag = f"~{equiv_hw}x{equiv_hw}" if equiv_hw * equiv_hw == M else f" {M}" print(f" {tag:>6} {tri_mean:>10.3f}ms {tri_std:>6.3f} {tch_mean:>10.3f}ms {tch_std:>6.3f} {speedup:>7.2f}x") spatial_results.append({ "B": B, "M": M, "triton_mean_ms": round(tri_mean, 4), "torch_mean_ms": round(tch_mean, 4), "speedup": round(speedup, 3), }) del A torch.cuda.empty_cache() # ------------------------------------------------------------------ # Sweep 3: BLOCK_M tuning at fixed B=4096, M=1024 # ------------------------------------------------------------------ print(f"\n{'─'*72}") print(f" SWEEP 3: BLOCK_M tuning (B=4096, M=1024)") print(f"{'─'*72}") print(f" {'BLOCK_M':>8} {'Triton ms':>12} {'±std':>8} {'tiles/img':>10}") print(f" {'─'*44}") block_sizes = [32, 64, 128, 256, 512, 1024] B, M = 4096, 1024 A = torch.randn(B, M, 3, device="cuda", dtype=torch.float32) block_results = [] for bm in block_sizes: tri_mean, tri_std, tri_med, tri_min, tri_max = _cuda_timer( lambda: batched_svd3(A, block_m=bm) ) n_tiles = (M + bm - 1) // bm print(f" {bm:>8} {tri_mean:>10.3f}ms {tri_std:>6.3f} {n_tiles:>10}") block_results.append({ "block_m": bm, "triton_mean_ms": round(tri_mean, 4), "triton_std_ms": round(tri_std, 4), "n_tiles": n_tiles, }) del A torch.cuda.empty_cache() # ------------------------------------------------------------------ # Sweep 4: Throughput — images/sec at peak batch # ------------------------------------------------------------------ print(f"\n{'─'*72}") print(f" SWEEP 4: Throughput (images/sec)") print(f"{'─'*72}") for B in [4096, 16384]: A = torch.randn(B, 1024, 3, device="cuda", dtype=torch.float32) tri_mean, *_ = _cuda_timer(lambda: batched_svd3(A, block_m=128)) tch_mean, *_ = _cuda_timer(lambda: torch.linalg.svd(A, full_matrices=False)) tri_ips = B / (tri_mean / 1000) tch_ips = B / (tch_mean / 1000) print(f" B={B:>5}: Triton {tri_ips:>12,.0f} img/s | Torch {tch_ips:>12,.0f} img/s") del A torch.cuda.empty_cache() # ------------------------------------------------------------------ # Memory: peak allocation comparison # ------------------------------------------------------------------ print(f"\n{'─'*72}") print(f" MEMORY: Peak allocation (B=4096, M=1024)") print(f"{'─'*72}") B, M = 4096, 1024 A = torch.randn(B, M, 3, device="cuda", dtype=torch.float32) torch.cuda.reset_peak_memory_stats() _ = batched_svd3(A) torch.cuda.synchronize() tri_peak = torch.cuda.max_memory_allocated() / 1024**2 torch.cuda.reset_peak_memory_stats() _ = torch.linalg.svd(A, full_matrices=False) torch.cuda.synchronize() tch_peak = torch.cuda.max_memory_allocated() / 1024**2 print(f" Triton: {tri_peak:.1f} MB") print(f" Torch: {tch_peak:.1f} MB") print(f" Ratio: {tch_peak / (tri_peak + 1e-9):.2f}x\n") # ------------------------------------------------------------------ # Dump JSON for further analysis # ------------------------------------------------------------------ report = { "device": device_name, "batch_sweep": batch_results, "spatial_sweep": spatial_results, "block_m_sweep": block_results, } with open("svd3_profile.json", "w") as f: json.dump(report, f, indent=2) print(f" Full results written to svd3_profile.json\n") if __name__ == "__main__": _test_correctness() _profile_sweep()