add sparse gemm optimization tutorial

Author: guqiqi

optimization/data/sparse-gemm-local-rerun-2026-04-07.csv

@@ -0,0 +1,11 @@
+date,gpu_index,gpu_name,problem_m,problem_n,problem_k,dtype_variant,kernel_label,kernel_kind,command,timing_warmup,timing_repeat,avg_ms,tflops,hw_efficiency_percent,notes
+2026-04-07,0,"NVIDIA H800 PCIe",4096,8192,8192,fp16,"baseline swizzle64_128",".co","./choreo -gs -t cute -arch=sm_90a --use-warpspec benchmark/performance/gemm_sp/gemm_sp_f16_dyn_sm90_warpspec_1p1c_swizzle64_128.co -o /tmp/f16_base64.cute.result && CUDA_VISIBLE_DEVICES=0 CHOREO_TIMING_WARMUP=5 CHOREO_TIMING_REPEAT=50 bash /tmp/f16_base64.cute.result --execute",5,50,,,,"runtime assert on current checkout"
+2026-04-07,0,"NVIDIA H800 PCIe",4096,8192,8192,fp16,"prepack swizzle64_128",".co","./choreo -gs -t cute -arch=sm_90a --use-warpspec --use-prepack benchmark/performance/gemm_sp/gemm_sp_f16_dyn_sm90_warpspec_1p1c_swizzle64_128_prepack.co -o /tmp/f16_prepack64.cute.result && CUDA_VISIBLE_DEVICES=0 CHOREO_TIMING_WARMUP=5 CHOREO_TIMING_REPEAT=50 bash /tmp/f16_prepack64.cute.result --execute",5,50,4.21844,130.322,8.61348,"runnable"
+2026-04-07,0,"NVIDIA H800 PCIe",4096,8192,8192,fp16,"prepack swizzle128_128",".co","./choreo -gs -t cute -arch=sm_90a --use-warpspec --use-prepack benchmark/performance/gemm_sp/gemm_sp_f16_dyn_sm90_warpspec_1p1c_swizzle128_128_prepack.co -o /tmp/f16_prepack128.cute.result && CUDA_VISIBLE_DEVICES=0 CHOREO_TIMING_WARMUP=5 CHOREO_TIMING_REPEAT=50 bash /tmp/f16_prepack128.cute.result --execute",5,50,1.5171,362.372,23.9506,"runnable; note this kernel also changes tile shape and packed metadata format"
+2026-04-07,0,"NVIDIA H800 PCIe",4096,8192,8192,fp16,"1p2c + 3-stage",".co","./choreo -gs -t cute -arch=sm_90a --use-warpspec --use-prepack --stmatrix benchmark/performance/gemm_sp/gemm_sp_f16_aitune_2026-03-25_iter120.co -o /tmp/iter120.cute.result && CUDA_VISIBLE_DEVICES=0 CHOREO_TIMING_WARMUP=5 CHOREO_TIMING_REPEAT=50 bash /tmp/iter120.cute.result --execute",5,50,1.03411,531.622,35.1369,"runnable; current CLI did not accept historical README flags --wgmma-split-batch / --wgmma-wait-depth=2"
+2026-04-07,0,"NVIDIA H800 PCIe",4096,8192,8192,fp16,"wider-K async metadata path",".cu","CUDA_VISIBLE_DEVICES=0 CHOREO_TIMING_WARMUP=5 CHOREO_TIMING_REPEAT=50 bash benchmark/performance/gemm_sp/gemm_sp_f16_aitune_2026-03-25_iter143/run.sh",5,50,0.701177,784.047,51.8207,"runnable"
+2026-04-07,0,"NVIDIA H800 PCIe",4096,8192,8192,e4m3,"baseline swizzle32_64",".co","./choreo -gs -t cute -arch=sm_90a --use-warpspec benchmark/performance/gemm_sp/gemm_sp_e4m3_dyn_sm90_warpspec_1p1c_swizzle32_64.co -o /tmp/e4m3_base32_64.cute.result && CUDA_VISIBLE_DEVICES=0 CHOREO_TIMING_WARMUP=5 CHOREO_TIMING_REPEAT=50 bash /tmp/e4m3_base32_64.cute.result --execute",5,50,0.672667,817.278,27.0085,"runnable"
+2026-04-07,0,"NVIDIA H800 PCIe",4096,8192,8192,e4m3,"prepack swizzle32_64",".co","./choreo -gs -t cute -arch=sm_90a --use-warpspec --use-prepack benchmark/performance/gemm_sp/gemm_sp_e4m3_dyn_sm90_warpspec_1p1c_swizzle32_64_prepack.co -o /tmp/e4m3_prepack32_64.cute.result && CUDA_VISIBLE_DEVICES=0 CHOREO_TIMING_WARMUP=5 CHOREO_TIMING_REPEAT=50 bash /tmp/e4m3_prepack32_64.cute.result --execute",5,50,0.79649,690.224,22.8098,"runnable"
+2026-04-07,0,"NVIDIA H800 PCIe",4096,8192,8192,e4m3,"prepack swizzle128_128",".co","./choreo -gs -t cute -arch=sm_90a --use-warpspec --use-prepack benchmark/performance/gemm_sp/gemm_sp_e4m3_dyn_sm90_warpspec_1p1c_swizzle128_128_prepack.co -o /tmp/e4m3_swizzle128_prepack.cute.result && CUDA_VISIBLE_DEVICES=0 CHOREO_TIMING_WARMUP=5 CHOREO_TIMING_REPEAT=50 bash /tmp/e4m3_swizzle128_prepack.cute.result --execute --skip-verify",5,50,1.01794,540.069,17.8476,"runnable only with --skip-verify on current checkout"
+2026-04-07,0,"NVIDIA H800 PCIe",4096,8192,8192,e4m3,"TMA metadata staging",".co","./choreo -gs -t cute -arch=sm_90a --use-warpspec --use-prepack benchmark/performance/gemm_sp/gemm_sp_e4m3_aitune_2026-03-21_iter001.co -o /tmp/e4m3_iter001.cute.result && CUDA_VISIBLE_DEVICES=0 CHOREO_TIMING_WARMUP=5 CHOREO_TIMING_REPEAT=50 bash /tmp/e4m3_iter001.cute.result --execute",5,50,0.510716,1076.44,35.5731,"runnable"
+2026-04-07,0,"NVIDIA H800 PCIe",4096,8192,8192,e4m3,"early empty arrive",".cu","CUDA_VISIBLE_DEVICES=0 CHOREO_TIMING_WARMUP=5 CHOREO_TIMING_REPEAT=50 bash benchmark/performance/gemm_sp/e4m3_aitune_2026-03-21_iter068/run.sh",5,50,0.392372,1401.11,46.3023,"runnable"