PyC CUDA Kernels
Twenty CUDA .cu prototype files from the PyC kernel lab. The collection
covers baseline and experimental matmul, Ada FP32/Tensor Core variants, Hopper
WMMA BF16/FP16 variants, asynchronous staging, CTA/warp/K-stage experiments,
and cuBLASLt control lanes.
These are source artifacts for inspection, compilation, and performance engineering—not precompiled model weights and not a claim of universal benchmark superiority. The 19 documented lab lanes are:
| Kernel | Short description | Source |
|---|---|---|
matrix_mult |
Compile-only CUDA matmul baseline. | kernels/prototypes/baseline/matmul/kernel.cu |
tokenizer_kernel |
Experimental tokenizer/matrix kernel. | kernels/prototypes/experimental/tokenizer_matmul/kernel.cu |
ada_gemm |
Ada FP32 shared-memory GEMM. | kernels/prototypes/ada/gemm/kernel.cu |
ada_gemm_k64_warp32_async |
Ada FP32 GEMM with cp.async double buffering. |
kernels/prototypes/ada/gemm_k64_warp32_async/kernel.cu |
ada_tensor_core_fp16 |
Ada WMMA FP16 GEMM. | kernels/prototypes/ada/tensor_core/kernel.cu |
ada_tensor_core_bf16 |
Ada WMMA BF16 GEMM. | kernels/prototypes/ada/tensor_core/kernel.cu |
hopper_tensor_core_fp16 |
Hopper 64×64×16 WMMA FP16 GEMM. | kernels/prototypes/hopper/tensor_core/kernel.cu |
hopper_tensor_core_bf16 |
Hopper 64×64×16 WMMA BF16 GEMM. | kernels/prototypes/hopper/tensor_core/kernel.cu |
hopper_tensor_core_bf16_perf |
Large-shape Hopper BF16 throughput lane. | kernels/prototypes/hopper/tensor_core/kernel.cu |
hopper_tensor_core_bf16_warp2n |
Hopper lane with two N fragments per warp. | kernels/prototypes/hopper/tensor_core/kernel.cu |
hopper_tensor_core_bf16_k32 |
Hopper lane with 32-wide K staging. | kernels/prototypes/hopper/tensor_core/kernel.cu |
hopper_tensor_core_bf16_warp2n_k32 |
Combined warp2n and K32 experiment. | kernels/prototypes/hopper/tensor_core/kernel.cu |
hopper_cublaslt_bf16 |
cuBLASLt BF16 hardware-ceiling control. | kernels/prototypes/hopper/cublaslt_bf16/kernel.cu |
hopper_cublaslt_bf16_check |
cuBLASLt BF16 correctness lane. | kernels/prototypes/hopper/cublaslt_bf16/kernel.cu |
hopper_tensor_core_bf16_async |
Hopper BF16 WMMA with async double buffering. | kernels/prototypes/hopper/tensor_core_async/kernel.cu |
hopper_tensor_core_bf16_async_wide |
Async WMMA with a wider 64×128 CTA. | kernels/prototypes/hopper/tensor_core_async/kernel.cu |
hopper_tensor_core_bf16_async_wide_k64 |
Wide async lane with K64 stages. | kernels/prototypes/hopper/tensor_core_async/kernel.cu |
hopper_tensor_core_bf16_async_square |
Async WMMA with a 128×128 CTA. | kernels/prototypes/hopper/tensor_core_async/kernel.cu |
hopper_tensor_core_bf16_async_square_k64 |
Square async WMMA with 128×128×64 tiles. | kernels/prototypes/hopper/tensor_core_async/kernel.cu |
The source files are organized by prototype family; several documented lanes
share one parameterized .cu implementation. See the companion
performance Space
for the optimization story and selected measurements.
Individual Kernel Repositories
Each documented lane is also available as its own Hugging Face Kernel repository:
- matrix_mult
- tokenizer_kernel
- ada_gemm
- ada_gemm_k64_warp32_async
- ada_tensor_core_fp16
- ada_tensor_core_bf16
- hopper_tensor_core_fp16
- hopper_tensor_core_bf16
- hopper_tensor_core_bf16_perf
- hopper_tensor_core_bf16_warp2n
- hopper_tensor_core_bf16_k32
- hopper_tensor_core_bf16_warp2n_k32
- hopper_cublaslt_bf16
- hopper_cublaslt_bf16_check
- hopper_tensor_core_bf16_async
- hopper_tensor_core_bf16_async_wide
- hopper_tensor_core_bf16_async_wide_k64
- hopper_tensor_core_bf16_async_square
- hopper_tensor_core_bf16_async_square_k64
- Downloads last month
- -
cuda
gpu-optimization
kernels
gemm
hpc
apache-2.0