# GPU GEMM cp.async 64×64 CTA WMMA contract (SM 8.0+)
# POSITIVE RESULT: +120% over baseline via async DMA
#
# Experiment result (2026-04-05): 8-byte cp.async per thread per K-tile.
# True async: WMMA runs while cp.async transfers complete in background.
# NO register pressure from loaded values (unlike ld.global→st.shared).
#
# Grounding:
#   [25] Kerr et al., "CUTLASS," 2023 — async pipeline
#   [48] NVIDIA, "CUTLASS 3.0," 2024 — cp.async + smem double-buffer
#   [66] NVIDIA PTX ISA 8.5, 2024 — cp.async specification
kind: PerformanceContract
name: cgp-gpu-gemm-cpasync
version: "1.0.0"
kernel: gemm_cta64_cpasync_fp16

description: |
  64×64 CTA WMMA with cp.async async global→shared DMA copy.
  Each thread does ONE 8-byte cp.async per K-tile (4 FP16 elements).
  Double-buffered (8KB smem = 2× 4KB). WMMA compute runs concurrently
  with cp.async transfers on separate hardware pipelines.

hardware:
  gpu: "NVIDIA GeForce RTX 4090"
  compute_capability: "8.9"
  minimum_sm: "8.0"  # cp.async requires SM 8.0+ (Ampere)
  smem_per_sm_kb: 100
  tensor_core: "FP16 → FP32 WMMA 16×16×16"

requirements:
  - "K must be multiple of 4 (8-byte alignment for cp.async)"
  - "N must be multiple of 4 (same reason)"
  - "PtxModule target must be sm_80+"

bounds:
  - size: [128, 128, 128]
    current_us: 5.1
    current_tflops: 0.8
    speedup_vs_cta64: 1.39
  - size: [256, 256, 256]
    current_us: 8.4
    current_tflops: 4.0
    speedup_vs_cta64: 1.47
  - size: [512, 512, 512]
    current_us: 16.0
    current_tflops: 16.8
    speedup_vs_cta64: 1.43
  - size: [1024, 1024, 1024]
    current_us: 53.0
    current_tflops: 40.5
    target_tflops: 55.0
    target_vs_cublas: 0.50
    speedup_vs_cta64: 1.38
    note: "+120% from baseline (18.4 → 40.5). Target 0.5x cuBLAS."

design:
  approach: |
    1. 16 warps × 4 elements/thread cooperative load
    2. ONE 8-byte cp.async per thread (4 consecutive FP16 values)
    3. Threads 0-255 load A, threads 256-511 load B
    4. cp.async.commit_group after all threads issue their copies
    5. WMMA compute runs on compute buffer while cp.async completes
    6. cp.async.wait_group 0 + bar_sync before next iteration
  overlap: |
    - Global DMA and tensor core execute on independent pipelines
    - No register pressure from loaded values (DMA direct to smem)
    - Buffer toggle avoids compute/store conflicts

falsification:
  - name: FALSIFY-CPASYNC-001
    description: "cp.async kernel must achieve >= 35 TFLOP/s at 1024"
    check: "tflops >= 35.0"
  - name: FALSIFY-CPASYNC-002
    description: "cp.async must be >= 1.3x faster than CTA64 single-buffer"
    check: "cta64_time / cpasync_time >= 1.3"
  - name: FALSIFY-CPASYNC-003
    description: "PTX must contain cp.async instructions"
    check: "ptx_contains_cp_async == true"
