Add interactive PyC kernel performance lab

Documentation-first Gradio Space covering 19 CUDA kernels and performance-engineering tradeoffs.

README.md

@@ -1,13 +1,21 @@
 ---
-title: Pyc Kernel Lab
-emoji: 📊
-colorFrom: indigo
-colorTo: green
+title: PyC CUDA Kernel Lab
+emoji: ⚡
+colorFrom: gray
+colorTo: indigo
 sdk: gradio
-sdk_version: 6.25.0
-python_version: '3.13'
+sdk_version: 5.49.1
 app_file: app.py
 pinned: false
+license: apache-2.0
 ---
 
-Check out the configuration reference at https://huggingface.co/docs/hub/spaces-config-reference
+# PyC CUDA Kernel Lab
+
+An interactive map of 19 CUDA kernel prototypes and the performance-engineering
+ideas behind them: tiling, Tensor Cores, asynchronous staging, CTA geometry,
+and library ceilings.
+
+The Space is documentation-first. It does not compile or execute CUDA kernels.
+The source release and selected evidence are available at
+[`AethronPhantom/pyc-kernels`](https://huggingface.co/datasets/AethronPhantom/pyc-kernels).

app.py

@@ -0,0 +1,75 @@
+import json
+from pathlib import Path
+
+import gradio as gr
+
+
+ROOT = Path(__file__).parent
+manifest = json.loads((ROOT / "kernels.json").read_text())
+kernels = manifest["kernels"]
+
+EVIDENCE = {
+    "hopper_tensor_core_fp16": "H100, 1024³: 0.105 ms, 20.473 TFLOPS (correctness lane).",
+    "hopper_tensor_core_bf16": "H100, 1024³: 0.102 ms, 21.024 TFLOPS (correctness lane).",
+    "hopper_cublaslt_bf16": "H100, 4096³: 0.162 ms, 846.466 TFLOPS (cuBLASLt control lane).",
+    "hopper_tensor_core_bf16_async_square_k64": "H100, 4096³: 0.9252 ms in the captured async square-K64 profile.",
+}
+
+
+def rows():
+    return [
+        [
+            k["name"],
+            ", ".join(k.get("tags", [])),
+            k["source"],
+            EVIDENCE.get(k["name"], "See the published performance evidence; no summary number is asserted here."),
+        ]
+        for k in kernels
+    ]
+
+
+def explain(name):
+    k = next(item for item in kernels if item["name"] == name)
+    tags = ", ".join(k.get("tags", []))
+    return (
+        f"### {k['name']}\n\n"
+        f"{k.get('description', 'No description supplied.')}\n\n"
+        f"**Source:** `{k['source']}`\n\n"
+        f"**Tags:** {tags}\n\n"
+        f"**Evidence:** {EVIDENCE.get(name, 'No compact summary is asserted for this entry.')}\n\n"
+        f"**Compile command**\n```text\n{k.get('compile_cmd', '')}\n```\n\n"
+        f"**Run command**\n```text\n{k.get('run_cmd', '') or '(compile-only)'}\n```"
+    )
+
+
+with gr.Blocks(title="PyC CUDA Kernel Lab", theme=gr.themes.Soft()) as demo:
+    gr.Markdown(
+        "# PyC CUDA Kernel Lab\n"
+        "### 19 kernels · one optimization story\n\n"
+        "This is a documentation and performance-engineering map, not a live CUDA runner. "
+        "The useful question is not only *which kernel is fastest?* but *which bottleneck "
+        "did each change target?*"
+    )
+    gr.Markdown(
+        "## How to read the progression\n"
+        "- **Tiling and reuse:** move data from global memory into shared memory and reuse it.\n"
+        "- **Tensor Cores:** use WMMA fragments for FP16/BF16 matrix math.\n"
+        "- **Feed the math pipe:** `cp.async`, double buffering, padding, and wider CTAs target exposed load/synchronization cost.\n"
+        "- **Shape is a control knob:** warp assignment, CTA M×N geometry, and K-stage depth change work per synchronization.\n"
+        "- **Ceiling/control:** cuBLASLt separates custom-kernel limitations from the hardware/library ceiling."
+    )
+    gr.Markdown("## Kernel catalog")
+    gr.Dataframe(rows(), headers=["Kernel", "Tags", "Source", "Evidence"], interactive=False, wrap=True)
+    selector = gr.Dropdown([k["name"] for k in kernels], value=kernels[0]["name"], label="Inspect a kernel")
+    detail = gr.Markdown(explain(kernels[0]["name"]))
+    selector.change(explain, selector, detail)
+    gr.Markdown(
+        "## Measurement discipline\n"
+        "Treat each number as a receipt: record GPU model, architecture, CUDA version, matrix shape, "
+        "warmup/repeat policy, correctness checks, and whether the reference path was enabled. "
+        "A kernel that wins one shape is not automatically a general winner."
+    )
+
+
+if __name__ == "__main__":
+    demo.launch()

kernels.json

@@ -0,0 +1,156 @@
+{
+  "kernels": [
+    {
+      "name": "matrix_mult",
+      "source": "kernels/prototypes/baseline/matmul/kernel.cu",
+      "description": "CUDA matrix multiplication kernel (compile-only baseline).",
+      "tags": ["cuda", "matmul"],
+      "compile_cmd": "{nvcc} -O3 -c {source} -o {build_dir}/{name}.o",
+      "run_cmd": ""
+    },
+    {
+      "name": "tokenizer_kernel",
+      "source": "kernels/prototypes/experimental/tokenizer_matmul/kernel.cu",
+      "description": "Experimental tokenizer/matrix CUDA kernel file.",
+      "tags": ["cuda", "tokenizer", "experimental"],
+      "compile_cmd": "{nvcc} -O2 -c {source} -o {build_dir}/{name}.o",
+      "run_cmd": ""
+    },
+    {
+      "name": "ada_gemm",
+      "source": "kernels/prototypes/ada/gemm/kernel.cu",
+      "description": "Ada-focused FP32 shared-memory GEMM prototype with standalone correctness and timing harness.",
+      "tags": ["cuda", "matmul", "ada", "sm89", "prototype"],
+      "compile_cmd": "{nvcc} -O3 -std=c++17 -lineinfo -gencode arch=compute_89,code=sm_89 -gencode arch=compute_89,code=compute_89 {source} -o {build_dir}/{name}",
+      "run_cmd": "{build_dir}/{name} 1024 1024 1024 10 50"
+    },
+    {
+      "name": "ada_gemm_k64_warp32_async",
+      "source": "kernels/prototypes/ada/gemm_k64_warp32_async/kernel.cu",
+      "description": "Ada FP32 GEMM winner with 64x64x64 tiles, 32x8 threads, and cp.async double-buffered shared-memory stages.",
+      "tags": ["cuda", "matmul", "ada", "sm89", "prototype", "fp32", "vectorized", "winner", "k64", "warp32", "async", "cpasync"],
+      "compile_cmd": "{nvcc} -O3 -std=c++17 -lineinfo -gencode arch=compute_89,code=sm_89 -gencode arch=compute_89,code=compute_89 {source} -o {build_dir}/{name}",
+      "run_cmd": "{build_dir}/{name} 1024 1024 1024 10 50"
+    },
+    {
+      "name": "ada_tensor_core_fp16",
+      "source": "kernels/prototypes/ada/tensor_core/kernel.cu",
+      "description": "Ada Tensor Core FP16 GEMM prototype with WMMA-backed 32x64x16 correctness and timing harness.",
+      "tags": ["cuda", "matmul", "ada", "sm89", "prototype", "tensor-core", "fp16"],
+      "compile_cmd": "{nvcc} -O3 -std=c++17 -lineinfo -DPYC_ADA_TENSOR_CORE_USE_BF16=0 -gencode arch=compute_89,code=sm_89 -gencode arch=compute_89,code=compute_89 {source} -o {build_dir}/{name}",
+      "run_cmd": "{build_dir}/{name} 1024 1024 1024 10 50"
+    },
+    {
+      "name": "ada_tensor_core_bf16",
+      "source": "kernels/prototypes/ada/tensor_core/kernel.cu",
+      "description": "Ada Tensor Core BF16 GEMM prototype with WMMA-backed 32x64x16 correctness and timing harness.",
+      "tags": ["cuda", "matmul", "ada", "sm89", "prototype", "tensor-core", "bf16"],
+      "compile_cmd": "{nvcc} -O3 -std=c++17 -lineinfo -DPYC_ADA_TENSOR_CORE_USE_BF16=1 -gencode arch=compute_89,code=sm_89 -gencode arch=compute_89,code=compute_89 {source} -o {build_dir}/{name}",
+      "run_cmd": "{build_dir}/{name} 1024 1024 1024 10 50"
+    },
+    {
+      "name": "hopper_tensor_core_fp16",
+      "source": "kernels/prototypes/hopper/tensor_core/kernel.cu",
+      "description": "Hopper Tensor Core FP16 GEMM prototype with a 64x64x16 WMMA tile and larger per-block warpgroup-style work.",
+      "tags": ["cuda", "matmul", "hopper", "sm90", "prototype", "tensor-core", "fp16"],
+      "compile_cmd": "{nvcc} -O3 -std=c++17 -lineinfo -DPYC_HOPPER_TENSOR_CORE_USE_BF16=0 -gencode arch=compute_90,code=sm_90 -gencode arch=compute_90,code=compute_90 {source} -o {build_dir}/{name}",
+      "run_cmd": "{build_dir}/{name} 1024 1024 1024 5 20"
+    },
+    {
+      "name": "hopper_tensor_core_bf16",
+      "source": "kernels/prototypes/hopper/tensor_core/kernel.cu",
+      "description": "Hopper Tensor Core BF16 GEMM prototype with a 64x64x16 WMMA tile and larger per-block warpgroup-style work.",
+      "tags": ["cuda", "matmul", "hopper", "sm90", "prototype", "tensor-core", "bf16"],
+      "compile_cmd": "{nvcc} -O3 -std=c++17 -lineinfo -DPYC_HOPPER_TENSOR_CORE_USE_BF16=1 -gencode arch=compute_90,code=sm_90 -gencode arch=compute_90,code=compute_90 {source} -o {build_dir}/{name}",
+      "run_cmd": "{build_dir}/{name} 1024 1024 1024 5 20"
+    },
+    {
+      "name": "hopper_tensor_core_bf16_perf",
+      "source": "kernels/prototypes/hopper/tensor_core/kernel.cu",
+      "description": "Hopper BF16 baseline perf lane at 4096^3 with reference disabled to isolate steady-state Tensor Core throughput.",
+      "tags": ["cuda", "matmul", "hopper", "sm90", "prototype", "tensor-core", "bf16", "perf", "next-loop"],
+      "compile_cmd": "{nvcc} -O3 -std=c++17 -lineinfo -DPYC_HOPPER_TENSOR_CORE_USE_BF16=1 -gencode arch=compute_90,code=sm_90 -gencode arch=compute_90,code=compute_90 {source} -o {build_dir}/{name}",
+      "run_cmd": "{build_dir}/{name} 4096 4096 4096 3 30 1"
+    },
+    {
+      "name": "hopper_tensor_core_bf16_warp2n",
+      "source": "kernels/prototypes/hopper/tensor_core/kernel.cu",
+      "description": "Hopper BF16 perf lane with 8 warps per CTA and each warp computing two N fragments to increase work per warp.",
+      "tags": ["cuda", "matmul", "hopper", "sm90", "prototype", "tensor-core", "bf16", "perf", "next-loop", "warp2n"],
+      "compile_cmd": "{nvcc} -O3 -std=c++17 -lineinfo -DPYC_HOPPER_TENSOR_CORE_USE_BF16=1 -DPYC_HOPPER_TC_WARP_COL_TILES=2 -DPYC_HOPPER_TC_WARP_ROW_GROUPS=4 -DPYC_HOPPER_TC_WARP_COL_GROUPS=2 -DPYC_HOPPER_TC_SHARED_PAD_A=8 -DPYC_HOPPER_TC_SHARED_PAD_B=8 -gencode arch=compute_90,code=sm_90 -gencode arch=compute_90,code=compute_90 {source} -o {build_dir}/{name}",
+      "run_cmd": "{build_dir}/{name} 4096 4096 4096 3 30 1"
+    },
+    {
+      "name": "hopper_tensor_core_bf16_k32",
+      "source": "kernels/prototypes/hopper/tensor_core/kernel.cu",
+      "description": "Hopper BF16 perf lane with a 64x64x32 shared-memory stage to cut synchronization frequency across K.",
+      "tags": ["cuda", "matmul", "hopper", "sm90", "prototype", "tensor-core", "bf16", "perf", "next-loop", "k32"],
+      "compile_cmd": "{nvcc} -O3 -std=c++17 -lineinfo -DPYC_HOPPER_TENSOR_CORE_USE_BF16=1 -DPYC_HOPPER_TC_TILE_K=32 -DPYC_HOPPER_TC_SHARED_PAD_A=8 -DPYC_HOPPER_TC_SHARED_PAD_B=8 -gencode arch=compute_90,code=sm_90 -gencode arch=compute_90,code=compute_90 {source} -o {build_dir}/{name}",
+      "run_cmd": "{build_dir}/{name} 4096 4096 4096 3 30 1"
+    },
+    {
+      "name": "hopper_tensor_core_bf16_warp2n_k32",
+      "source": "kernels/prototypes/hopper/tensor_core/kernel.cu",
+      "description": "Hopper BF16 perf lane combining 8-warp CTAs, two N fragments per warp, padded shared-memory strides, and 32-wide K staging.",
+      "tags": ["cuda", "matmul", "hopper", "sm90", "prototype", "tensor-core", "bf16", "perf", "next-loop", "warp2n", "k32"],
+      "compile_cmd": "{nvcc} -O3 -std=c++17 -lineinfo -DPYC_HOPPER_TENSOR_CORE_USE_BF16=1 -DPYC_HOPPER_TC_WARP_COL_TILES=2 -DPYC_HOPPER_TC_WARP_ROW_GROUPS=4 -DPYC_HOPPER_TC_WARP_COL_GROUPS=2 -DPYC_HOPPER_TC_TILE_K=32 -DPYC_HOPPER_TC_SHARED_PAD_A=8 -DPYC_HOPPER_TC_SHARED_PAD_B=8 -gencode arch=compute_90,code=sm_90 -gencode arch=compute_90,code=compute_90 {source} -o {build_dir}/{name}",
+      "run_cmd": "{build_dir}/{name} 4096 4096 4096 3 30 1"
+    },
+    {
+      "name": "hopper_cublaslt_bf16",
+      "source": "kernels/prototypes/hopper/cublaslt_bf16/kernel.cu",
+      "description": "Hopper cuBLASLt BF16 GEMM control lane for measuring the hardware-native library ceiling before a custom TMA/WGMMA implementation lands.",
+      "tags": ["cuda", "matmul", "hopper", "sm90", "prototype", "cublaslt", "bf16", "control", "ceiling"],
+      "compile_cmd": "{nvcc} -O3 -std=c++17 -lineinfo -gencode arch=compute_90,code=sm_90 -gencode arch=compute_90,code=compute_90 {source} -lcublasLt -lcublas -o {build_dir}/{name}",
+      "run_cmd": "{build_dir}/{name} 4096 4096 4096 3 30 1"
+    },
+    {
+      "name": "hopper_cublaslt_bf16_check",
+      "source": "kernels/prototypes/hopper/cublaslt_bf16/kernel.cu",
+      "description": "Smaller Hopper cuBLASLt BF16 correctness lane with reference enabled.",
+      "tags": ["cuda", "matmul", "hopper", "sm90", "prototype", "cublaslt", "bf16", "correctness"],
+      "compile_cmd": "{nvcc} -O3 -std=c++17 -lineinfo -gencode arch=compute_90,code=sm_90 -gencode arch=compute_90,code=compute_90 {source} -lcublasLt -lcublas -o {build_dir}/{name}",
+      "run_cmd": "{build_dir}/{name} 512 512 512 2 10 0"
+    },
+    {
+      "name": "hopper_tensor_core_bf16_async",
+      "source": "kernels/prototypes/hopper/tensor_core_async/kernel.cu",
+      "description": "Hopper BF16 async WMMA lane with cp.async double-buffered staging, 64x64x32 tiles, and four fragments per warp to cut exposed feed bubbles.",
+      "tags": ["cuda", "matmul", "hopper", "sm90", "prototype", "tensor-core", "bf16", "async", "cpasync", "double-buffered"],
+      "compile_cmd": "{nvcc} -O3 -std=c++17 -lineinfo -DPYC_HOPPER_TENSOR_CORE_USE_BF16=1 -gencode arch=compute_90,code=sm_90 -gencode arch=compute_90,code=compute_90 {source} -o {build_dir}/{name}",
+      "run_cmd": "{build_dir}/{name} 4096 4096 4096 3 30 1"
+    },
+    {
+      "name": "hopper_tensor_core_bf16_async_wide",
+      "source": "kernels/prototypes/hopper/tensor_core_async/kernel.cu",
+      "description": "Hopper BF16 async WMMA lane widened to a 64x128x32 CTA so each stage amortizes more scheduling overhead before the WGMMA/TMA path lands.",
+      "tags": ["cuda", "matmul", "hopper", "sm90", "prototype", "tensor-core", "bf16", "async", "cpasync", "double-buffered", "wide"],
+      "compile_cmd": "{nvcc} -O3 -std=c++17 -lineinfo -DPYC_HOPPER_TENSOR_CORE_USE_BF16=1 -DPYC_HOPPER_TC_WARP_COL_GROUPS=4 -gencode arch=compute_90,code=sm_90 -gencode arch=compute_90,code=compute_90 {source} -o {build_dir}/{name}",
+      "run_cmd": "{build_dir}/{name} 4096 4096 4096 3 30 1"
+    },
+    {
+      "name": "hopper_tensor_core_bf16_async_wide_k64",
+      "source": "kernels/prototypes/hopper/tensor_core_async/kernel.cu",
+      "description": "Hopper BF16 async WMMA lane keeping the 64x128 CTA but doubling stage depth to 64 so each stage amortizes more feed and loop overhead.",
+      "tags": ["cuda", "matmul", "hopper", "sm90", "prototype", "tensor-core", "bf16", "async", "cpasync", "double-buffered", "wide", "k64", "gap-close"],
+      "compile_cmd": "{nvcc} -O3 -std=c++17 -lineinfo -DPYC_HOPPER_TENSOR_CORE_USE_BF16=1 -DPYC_HOPPER_TC_WARP_COL_GROUPS=4 -DPYC_HOPPER_TC_TILE_K=64 -gencode arch=compute_90,code=sm_90 -gencode arch=compute_90,code=compute_90 {source} -o {build_dir}/{name}",
+      "run_cmd": "{build_dir}/{name} 4096 4096 4096 3 30 1"
+    },
+    {
+      "name": "hopper_tensor_core_bf16_async_square",
+      "source": "kernels/prototypes/hopper/tensor_core_async/kernel.cu",
+      "description": "Hopper BF16 async WMMA lane widened into a 128x128x32 CTA so the square 4096^3 regime gets more work per CTA before switching stages.",
+      "tags": ["cuda", "matmul", "hopper", "sm90", "prototype", "tensor-core", "bf16", "async", "cpasync", "double-buffered", "square", "gap-close"],
+      "compile_cmd": "{nvcc} -O3 -std=c++17 -lineinfo -DPYC_HOPPER_TENSOR_CORE_USE_BF16=1 -DPYC_HOPPER_TC_WARP_ROW_GROUPS=4 -DPYC_HOPPER_TC_WARP_COL_GROUPS=4 -gencode arch=compute_90,code=sm_90 -gencode arch=compute_90,code=compute_90 {source} -o {build_dir}/{name}",
+      "run_cmd": "{build_dir}/{name} 4096 4096 4096 3 30 1"
+    },
+    {
+      "name": "hopper_tensor_core_bf16_async_square_k64",
+      "source": "kernels/prototypes/hopper/tensor_core_async/kernel.cu",
+      "description": "Hopper BF16 async WMMA square lane with a 128x128x64 CTA to test the last cheap work-per-stage lever before the WGMMA/TMA path takes over.",
+      "tags": ["cuda", "matmul", "hopper", "sm90", "prototype", "tensor-core", "bf16", "async", "cpasync", "double-buffered", "square", "k64", "gap-close"],
+      "compile_cmd": "{nvcc} -O3 -std=c++17 -lineinfo -DPYC_HOPPER_TENSOR_CORE_USE_BF16=1 -DPYC_HOPPER_TC_WARP_ROW_GROUPS=4 -DPYC_HOPPER_TC_WARP_COL_GROUPS=4 -DPYC_HOPPER_TC_TILE_K=64 -gencode arch=compute_90,code=sm_90 -gencode arch=compute_90,code=compute_90 {source} -o {build_dir}/{name}",
+      "run_cmd": "{build_dir}/{name} 4096 4096 4096 3 30 1"
+    }
+  ]
+}

requirements.txt

@@ -0,0 +1 @@
+gradio==5.49.1