# GPU GEMM 64×64 CTA WMMA contract (+62% over 32×32)
# Contract-first: defines target before implementation.
#
# Experiment result: 64×64 CTA tile with 16 warps (4×4 grid of 16×16 WMMA).
# Each A element reused by 4 column warps (vs 2 in 32×32), each B element
# by 4 row warps. Compute-to-load ratio: 32 FLOP/byte (vs 16 for 32×32).
#
# Grounding:
#   [25] Kerr et al., "CUTLASS," 2023 — tile size vs arithmetic intensity
#   [9]  Volkov, "Better Performance at Lower Occupancy," 2010 — register blocking
kind: PerformanceContract
name: cgp-gpu-gemm-cta64
version: "1.0.0"
kernel: gemm_cta64_wmma_fp16

description: |
  64×64 CTA-tiled WMMA GEMM. 16 warps cooperatively compute 64×64 output
  tiles via 4×4 grid of 16×16 WMMA subtiles. 2× compute-to-load ratio vs
  32×32 kernel enables better data reuse, amortizing shared memory bandwidth.

hardware:
  gpu: "NVIDIA GeForce RTX 4090"
  compute_capability: "8.9"
  smem_per_sm_kb: 100
  tensor_core: "FP16 → FP32 WMMA 16×16×16"

bounds:
  - size: [128, 128, 128]
    current_us: 7.1
    current_tflops: 0.6
    note: "Small-size: 32×32 wins (4× more CTAs → better SM occupancy)"

  - size: [512, 512, 512]
    current_us: 22.9
    current_tflops: 11.7
    note: "Crossover: 32×32 still wins (18.9µs)"

  - size: [1024, 1024, 1024]
    current_us: 72.4
    current_tflops: 29.7
    target_tflops: 40.0
    target_vs_cublas: 0.40
    note: "1.62x over 32×32 (117→72us) — target 0.5x cuBLAS via cp.async"

design:
  approach: |
    1. 16 warps (512 threads) in 4×4 grid
    2. Each warp owns one 16×16 WMMA output subtile
    3. Shared memory: A[64×16] + B[16×64] = 4KB (2× 32×32)
    4. Cooperative load: threads 0-255 load A, 256-511 load B
    5. 4 elements/thread (vs 8 in 32×32) — less register pressure
  benefits: |
    - 2× data reuse (A element used by 4 col warps vs 2)
    - Higher arithmetic intensity → better hides memory latency
    - 16 WMMAs per K-tile amortizes buffer management overhead

falsification:
  - name: FALSIFY-CTA64-001
    description: "CTA64 at 1024 must achieve >= 25 TFLOP/s"
    check: "tflops >= 25.0"
  - name: FALSIFY-CTA64-002
    description: "CTA64 must be >= 1.3x faster than CTA32 at 1024"
    check: "cta32_time / cta64_time >= 1.3"
  - name: FALSIFY-CTA64-003
    description: "Double-buffer variant must not regress (amortization works)"
    check: "dbuf_time <= single_time * 1.05"
