Skip to content

refactor(cake_dense_projection_gemm): deterministic names for the dense projection GEMM package - #8046

Merged
yyihuang merged 5 commits into
flashinfer-ai:mainfrom
yyihuang:dense-projection-det-gemm-names
Oct 31, 2026
Merged

refactor(cake_dense_projection_gemm): deterministic names for the dense projection GEMM package#8046
yyihuang merged 5 commits into
flashinfer-ai:mainfrom
yyihuang:dense-projection-det-gemm-names

Conversation

@yyihuang

@yyihuang yyihuang commented Oct 30, 2026 •
edited by coderabbitai Bot
Loading

Copy link
Copy Markdown
Collaborator

The Cake dense_projection_gemm package in flashinfer.experimental has 127 programs built from 124 templates. Most are BF16-operand dense GEMMs that write BF16 or FP32 output, and there is also an FP32 router GEMM that runs through BF16x3 emulation. This PR renames the package to flashinfer.experimental.dense_projection_det_gemm, where "det" means repeated runs with the same inputs, hardware and plan give bitwise-identical results. Kernel machine code and plans do not change. The baseline is the previous Cake version on the same GPU, all on CUDA 13.4. Across 294 shapes the geomean speedup is 1.000 on NVIDIA GB200 (sm_100a), NVIDIA GB300 (sm_103a) and Rubin (sm_107a). Per-row ratios are 0.929–1.024, 0.974–1.019 and 0.990–1.021 respectively, all within run-to-run variation.

The table below gives per-shape times in µs for the previous version and this PR on each GPU (GB200 and GB300 with 152 SMs, Rubin with 216 SMs), with speedup = previous / this PR.

shape (K→N) op out T GB200 previous µs GB200 this PR µs GB200 speedup GB300 previous µs GB300 this PR µs GB300 speedup Rubin previous µs Rubin this PR µs Rubin speedup
indexer_q (2048→4096) dgrad bf16 16172 174.4 174.3 1.001 143.1 143.2 0.999 78.9 78.9 1.000
indexer_q (2048→4096) dgrad bf16 16231 174.1 173.8 1.001 143.1 142.8 1.002 78.7 78.7 1.000
indexer_q (2048→4096) fwd bf16 16172 169.2 169.4 0.999 139.9 139.6 1.002 78.0 78.0 1.000
indexer_q (2048→4096) fwd bf16 16231 170.6 170.2 1.002 140.5 140.0 1.004 78.1 78.0 1.001
indexer_q (2048→4096) wgrad bf16 16172 179.3 179.4 0.999 148.8 148.4 1.003 78.4 78.4 1.000
indexer_q (2048→4096) wgrad bf16 16231 179.8 180.0 0.999 149.0 150.3 0.991 78.8 78.8 1.000
indexer_q (2048→4096) wgrad f32 16172 177.4 177.1 1.001 147.7 146.8 1.006 79.8 79.8 1.000
indexer_q (2048→4096) wgrad f32 16231 179.2 178.0 1.007 147.4 147.7 0.998 80.2 80.1 1.001
shared_down (2048→6144) dgrad bf16 16172 259.5 258.7 1.003 217.6 217.1 1.002 115.9 116.0 1.000
shared_down (2048→6144) dgrad bf16 16231 256.1 254.8 1.005 215.4 215.8 0.998 116.0 116.0 1.000
shared_down (2048→6144) dgrad f32 16172 257.8 262.7 0.982 219.1 219.0 1.001 117.5 117.6 1.000
shared_down (2048→6144) dgrad f32 16231 258.2 258.7 0.998 219.7 218.8 1.004 117.8 117.8 1.000
shared_down (2048→6144) fwd bf16 16172 252.6 254.2 0.994 212.0 212.5 0.998 110.5 110.5 1.000
shared_down (2048→6144) fwd bf16 16231 250.3 251.2 0.997 212.1 212.7 0.997 110.5 110.4 1.000
shared_down (2048→6144) wgrad bf16 16172 259.4 261.8 0.991 218.7 219.1 0.998 115.2 115.2 1.000
shared_down (2048→6144) wgrad bf16 16231 259.9 261.0 0.996 219.6 219.6 1.000 115.6 115.6 1.000
shared_down (2048→6144) wgrad f32 16172 261.5 261.8 0.999 220.8 220.4 1.002 116.1 116.2 0.999
shared_down (2048→6144) wgrad f32 16231 260.9 262.1 0.995 221.7 221.6 1.001 116.7 116.6 1.000
q_b (2048→16384) dgrad bf16 16172 645.0 644.1 1.001 583.4 582.8 1.001 295.5 295.3 1.001
q_b (2048→16384) dgrad bf16 16231 643.2 643.0 1.000 583.5 582.3 1.002 296.2 296.4 0.999
q_b (2048→16384) dgrad f32 16172 645.1 645.4 0.999 586.0 585.5 1.001 299.4 299.7 0.999
q_b (2048→16384) dgrad f32 16231 646.8 646.0 1.001 584.0 584.2 1.000 300.4 300.1 1.001
q_b (2048→16384) fwd bf16 16172 614.5 615.4 0.999 556.5 555.2 1.002 287.2 287.2 1.000
q_b (2048→16384) fwd bf16 16231 618.1 615.2 1.005 556.3 553.3 1.005 287.6 287.3 1.001
q_b (2048→16384) wgrad bf16 16172 630.8 631.8 0.998 568.1 568.5 0.999 292.8 292.2 1.002
q_b (2048→16384) wgrad bf16 16231 632.4 628.4 1.006 576.7 575.7 1.002 294.0 293.9 1.000
q_b (2048→16384) wgrad f32 16172 641.7 643.5 0.997 582.3 577.9 1.008 293.5 293.4 1.000
q_b (2048→16384) wgrad f32 16231 649.8 645.7 1.006 584.0 583.2 1.001 294.4 294.7 0.999
4096→32 dgrad bf16 16231 24.0 24.0 1.001 23.2 23.2 1.001 16.6 16.6 0.999
4096→32 dgrad bf16 16384 24.1 24.1 0.999 23.4 23.4 1.001 16.8 16.8 0.997
4096→32 dgrad f32 16231 44.6 44.6 1.000 42.8 42.9 0.996 30.3 30.2 1.003
4096→32 dgrad f32 16384 44.9 45.1 0.997 43.1 43.2 0.999 30.3 30.4 0.995
4096→32 fwd bf16 16231 26.4 26.6 0.995 26.9 27.0 0.999 19.2 19.2 0.999
4096→32 fwd bf16 16384 26.2 26.1 1.003 26.6 26.5 1.003 19.6 19.6 0.998
4096→32 wgrad bf16 16231 28.2 28.4 0.995 28.4 28.4 0.997 20.0 20.0 1.000
4096→32 wgrad bf16 16384 28.1 28.0 1.005 28.4 28.4 0.999 20.2 20.2 1.002
4096→32 wgrad f32 16231 28.3 28.3 0.999 28.4 28.3 1.003 26.1 26.1 0.998
4096→32 wgrad f32 16384 28.2 28.1 1.003 28.5 28.5 1.001 26.4 26.4 1.001
4096→48 dgrad bf16 16231 24.5 24.5 1.003 23.0 22.9 1.002 16.8 16.8 1.000
4096→48 dgrad bf16 16384 24.6 24.6 1.001 23.4 23.3 1.003 17.0 17.0 1.000
4096→48 dgrad f32 16231 45.1 44.9 1.003 43.2 43.2 1.000 30.6 30.5 1.001
4096→48 dgrad f32 16384 45.4 45.4 1.001 43.6 43.8 0.997 30.7 30.5 1.004
4096→48 fwd bf16 16231 26.4 26.4 0.999 26.8 27.0 0.995 19.4 19.4 1.000
4096→48 fwd bf16 16384 26.1 26.1 1.001 26.7 26.6 1.004 19.6 19.6 1.000
4096→48 wgrad bf16 16231 28.3 28.4 0.998 28.5 28.6 0.998 20.3 20.3 1.001
4096→48 wgrad bf16 16384 28.4 28.4 1.001 28.5 28.6 0.997 20.5 20.6 0.997
4096→48 wgrad f32 16231 28.4 28.4 1.000 28.6 28.7 0.996 27.3 27.4 0.997
4096→48 wgrad f32 16384 28.5 28.4 1.004 28.8 28.7 1.002 27.6 27.6 1.001
4096→64 dgrad bf16 16231 24.4 24.4 1.000 23.6 23.5 1.004 16.6 16.6 1.001
4096→64 dgrad bf16 16384 24.6 24.6 1.001 23.6 23.6 1.003 16.7 16.8 0.992
4096→64 dgrad f32 16231 45.2 45.0 1.006 43.4 43.3 1.001 30.5 30.3 1.007
4096→64 dgrad f32 16384 45.4 45.5 0.996 43.7 43.5 1.003 30.5 30.6 0.999
4096→64 fwd bf16 16231 26.7 26.6 1.002 27.3 27.2 1.002 19.6 19.6 1.001
4096→64 fwd bf16 16384 26.6 26.5 1.003 27.0 27.0 0.999 19.7 19.7 0.998
4096→64 wgrad bf16 16231 28.7 28.8 0.998 29.0 28.9 1.001 20.2 20.2 1.002
4096→64 wgrad bf16 16384 28.3 28.3 0.999 28.6 28.6 1.003 20.5 20.6 0.997
4096→64 wgrad f32 16231 28.8 28.7 1.004 28.8 29.0 0.994 29.3 29.3 0.999
4096→64 wgrad f32 16384 28.7 28.5 1.006 28.7 28.7 1.002 30.4 30.5 0.999
4096→128 dgrad bf16 16231 25.8 25.7 1.002 24.5 24.5 1.002 18.0 18.0 0.997
4096→128 dgrad bf16 16384 25.8 25.9 0.999 24.5 24.4 1.003 18.1 18.1 1.003
4096→128 dgrad f32 16231 47.4 47.2 1.004 44.8 44.7 1.001 32.1 32.1 1.000
4096→128 dgrad f32 16384 47.6 47.6 1.000 45.1 45.0 1.003 32.2 32.2 0.998
4096→128 fwd bf16 16231 27.5 27.5 1.000 28.0 28.0 0.999 20.0 20.0 0.998
4096→128 fwd bf16 16384 27.2 27.3 0.998 27.4 27.6 0.991 20.0 20.0 1.000
4096→128 wgrad bf16 16231 29.6 29.6 1.001 29.5 29.5 0.998 21.0 20.9 1.001
4096→128 wgrad bf16 16384 29.3 29.2 1.004 29.3 29.4 0.997 21.0 20.9 1.005
4096→128 wgrad f32 16231 29.7 29.5 1.005 29.4 29.4 1.000 20.9 20.9 0.998
4096→128 wgrad f32 16384 29.4 29.4 1.001 29.4 29.5 0.997 21.1 21.2 0.998
4096→512 dgrad bf16 16231 50.7 50.8 0.999 41.6 41.6 0.999 29.1 29.1 0.999
4096→512 dgrad bf16 16384 50.8 50.7 1.002 41.7 41.6 1.001 29.2 29.1 1.002
4096→512 dgrad f32 16231 72.0 72.1 0.999 60.6 60.7 0.997 38.5 38.5 1.000
4096→512 dgrad f32 16384 72.1 72.3 0.998 60.7 60.6 1.002 38.8 38.9 0.997
4096→512 fwd bf16 16231 57.2 57.0 1.003 48.6 48.8 0.997 33.2 33.2 1.000
4096→512 fwd bf16 16384 56.9 56.9 0.999 48.6 48.7 0.999 33.2 33.2 1.000
4096→512 wgrad bf16 16231 51.6 51.7 1.000 43.8 43.6 1.004 36.2 36.2 1.000
4096→512 wgrad bf16 16384 52.4 52.3 1.003 43.9 43.9 0.999 36.4 36.5 0.998
4096→512 wgrad f32 16231 51.6 51.5 1.002 43.5 43.4 1.004 73.8 73.6 1.002
4096→512 wgrad f32 16384 51.9 51.8 1.002 43.7 43.8 0.999 74.0 74.2 0.998
4096→1024 dgrad bf16 16231 93.8 93.8 1.000 73.1 73.2 0.998 44.1 44.0 1.001
4096→1024 dgrad bf16 16384 93.9 93.8 1.001 72.9 73.3 0.994 44.3 44.2 1.002
4096→1024 fwd bf16 16231 96.5 96.3 1.002 75.6 76.1 0.993 45.5 45.4 1.000
4096→1024 fwd bf16 16384 96.8 96.6 1.003 76.0 76.2 0.999 45.5 45.5 1.001
4096→1024 wgrad bf16 16231 95.7 95.6 1.001 74.6 74.7 0.999 52.6 52.6 1.001
4096→1024 wgrad bf16 16384 96.6 96.8 0.998 75.2 75.5 0.996 53.0 53.0 0.999
4096→1024 wgrad f32 16231 97.7 98.4 0.993 76.5 76.5 1.001 74.9 75.1 0.999
4096→1024 wgrad f32 16384 98.9 98.7 1.003 77.2 77.6 0.994 75.5 75.5 1.000
4096→4096 dgrad bf16 16231 316.9 311.5 1.017 274.0 273.9 1.000 149.9 149.9 1.000
4096→4096 dgrad bf16 16384 312.9 312.4 1.002 273.4 273.4 1.000 149.9 149.9 1.000
4096→4096 fwd bf16 16231 313.2 313.9 0.998 274.4 273.6 1.003 149.3 149.2 1.001
4096→4096 fwd bf16 16384 314.0 313.3 1.002 274.4 273.9 1.002 149.5 149.4 1.001
4096→4096 wgrad bf16 16231 326.8 326.2 1.002 284.7 284.8 1.000 176.3 176.2 1.000
4096→4096 wgrad bf16 16384 328.6 328.5 1.000 287.4 287.3 1.000 177.5 177.5 1.000
4096→4096 wgrad f32 16231 329.9 330.0 1.000 289.3 289.2 1.000 233.7 233.7 1.000
4096→4096 wgrad f32 16384 330.9 331.2 0.999 291.7 291.5 1.001 235.4 235.3 1.000
4096→12288 dgrad bf16 16231 924.1 924.7 0.999 842.6 845.2 0.997 452.0 452.0 1.000
4096→12288 dgrad bf16 16384 926.9 925.5 1.002 841.8 847.3 0.994 453.4 453.1 1.001
4096→12288 fwd bf16 16231 896.9 896.5 1.000 823.3 819.2 1.005 439.9 438.5 1.003
4096→12288 fwd bf16 16384 880.7 893.1 0.986 827.5 821.4 1.007 440.6 440.5 1.000
4096→12288 wgrad bf16 16231 910.9 912.2 0.999 845.6 850.1 0.995 460.1 460.0 1.000
4096→12288 wgrad bf16 16384 903.0 924.1 0.977 853.1 854.2 0.999 463.5 463.7 1.000
4096→12288 wgrad f32 16231 923.9 946.5 0.976 857.0 859.3 0.997 471.1 471.1 1.000
4096→12288 wgrad f32 16384 933.9 936.6 0.997 860.1 864.2 0.995 474.7 474.8 1.000
5120→32 dgrad bf16 16231 28.6 28.7 0.997 28.0 27.5 1.017 19.7 19.7 0.999
5120→32 dgrad bf16 16384 29.0 28.8 1.007 27.7 27.6 1.001 19.8 19.9 0.994
5120→32 dgrad f32 16231 53.6 53.6 1.000 51.5 51.6 0.999 36.3 36.4 0.999
5120→32 dgrad f32 16384 54.1 54.2 0.998 52.0 51.9 1.001 36.7 36.6 1.003
5120→32 fwd bf16 16231 31.3 31.4 0.998 31.6 31.6 1.001 23.1 23.1 0.999
5120→32 fwd bf16 16384 32.2 32.2 1.000 32.7 32.8 0.999 23.2 23.2 0.999
5120→32 wgrad bf16 16231 32.9 32.8 1.003 34.4 34.3 1.002 22.6 22.7 0.998
5120→32 wgrad bf16 16384 33.4 33.2 1.004 34.4 34.6 0.994 22.8 22.8 1.003
5120→32 wgrad f32 16231 32.9 32.8 1.001 34.2 34.2 1.001 28.0 28.1 0.999
5120→32 wgrad f32 16384 33.3 33.3 1.001 34.5 34.6 0.998 28.0 28.1 0.998
5120→48 dgrad bf16 16231 29.2 29.1 1.003 27.9 28.0 0.998 20.4 20.2 1.008
5120→48 dgrad bf16 16384 29.7 29.4 1.011 28.0 28.0 0.999 20.5 20.6 0.991
5120→48 dgrad f32 16231 54.2 54.4 0.997 52.2 52.1 1.002 36.9 36.9 0.999
5120→48 dgrad f32 16384 54.9 54.8 1.000 52.5 52.6 1.000 37.1 37.0 1.001
5120→48 fwd bf16 16231 31.5 31.3 1.005 31.8 31.7 1.001 23.0 23.0 1.000
5120→48 fwd bf16 16384 32.2 32.2 1.002 32.8 32.8 1.000 23.1 23.1 1.000
5120→48 wgrad bf16 16231 33.5 33.5 0.999 34.6 34.4 1.005 23.2 22.9 1.012
5120→48 wgrad bf16 16384 34.0 34.0 0.999 34.8 34.8 0.998 23.2 23.3 0.997
5120→48 wgrad f32 16231 33.4 33.6 0.995 34.7 34.6 1.004 29.0 28.9 1.004
5120→48 wgrad f32 16384 34.0 34.0 0.998 34.8 34.9 0.998 29.2 29.2 0.997
5120→64 dgrad bf16 16231 29.2 29.2 1.001 27.9 28.0 0.999 19.7 19.8 0.998
5120→64 dgrad bf16 16384 29.5 29.4 1.002 28.1 28.2 0.998 20.0 19.8 1.007
5120→64 dgrad f32 16231 54.7 54.8 0.999 52.3 52.3 1.001 36.9 36.8 1.003
5120→64 dgrad f32 16384 54.9 54.8 1.001 52.6 52.6 1.000 36.8 36.9 0.997
5120→64 fwd bf16 16231 32.0 32.0 0.999 32.4 32.3 1.003 23.2 23.2 1.000
5120→64 fwd bf16 16384 32.6 32.5 1.002 33.0 33.0 1.000 23.4 23.3 1.005
5120→64 wgrad bf16 16231 33.4 33.3 1.001 34.6 34.8 0.995 22.8 22.9 0.996
5120→64 wgrad bf16 16384 33.8 33.7 1.002 35.0 34.9 1.002 23.0 22.9 1.004
5120→64 wgrad f32 16231 33.4 33.3 1.002 34.8 34.5 1.007 31.1 31.2 0.997
5120→64 wgrad f32 16384 33.8 33.6 1.004 34.9 34.8 1.004 32.3 32.4 0.998
5120→128 dgrad bf16 16231 30.8 30.8 0.999 29.2 29.2 1.001 21.2 21.3 0.996
5120→128 dgrad bf16 16384 30.9 30.8 1.001 29.4 29.4 1.000 21.2 21.2 1.000
5120→128 dgrad f32 16231 56.7 56.9 0.997 54.1 54.2 0.998 39.1 39.1 1.000
5120→128 dgrad f32 16384 57.2 57.5 0.995 54.3 54.2 1.003 39.3 39.2 1.002
5120→128 fwd bf16 16231 33.2 33.2 0.998 33.5 33.6 1.000 23.6 23.6 0.999
5120→128 fwd bf16 16384 32.9 32.9 1.000 33.5 33.4 1.003 23.7 23.6 1.004
5120→128 wgrad bf16 16231 33.9 34.0 0.998 35.9 35.8 1.004 23.3 23.3 0.999
5120→128 wgrad bf16 16384 34.6 34.4 1.005 36.2 36.3 0.998 23.6 23.5 1.002
5120→128 wgrad f32 16231 34.1 33.9 1.005 35.9 35.8 1.000 23.3 23.4 0.999
5120→128 wgrad f32 16384 34.5 34.4 1.000 36.1 36.2 0.997 23.5 23.6 0.999
5120→512 dgrad bf16 16231 63.8 63.8 1.000 50.5 50.3 1.003 33.9 33.9 1.000
5120→512 dgrad bf16 16384 63.6 63.7 0.999 50.7 50.3 1.007 34.0 34.1 0.995
5120→512 dgrad f32 16231 88.3 88.6 0.997 74.1 74.4 0.996 46.8 46.9 0.998
5120→512 dgrad f32 16384 88.5 88.2 1.004 74.5 74.8 0.996 46.8 46.9 0.997
5120→512 fwd bf16 16231 70.9 70.9 1.000 58.6 58.9 0.996 39.3 39.3 1.000
5120→512 fwd bf16 16384 71.1 71.0 1.001 58.8 58.8 0.999 39.4 39.4 1.000
5120→512 wgrad bf16 16231 69.6 69.6 1.001 54.1 54.1 0.999 38.1 38.1 1.000
5120→512 wgrad bf16 16384 70.1 70.1 1.000 54.4 54.4 1.000 38.4 38.3 1.001
5120→512 wgrad f32 16231 71.1 70.8 1.005 56.0 56.0 0.999 74.1 73.9 1.004
5120→512 wgrad f32 16384 71.2 71.3 1.000 55.6 55.6 0.999 74.4 74.4 1.000
5120→1024 dgrad bf16 16231 113.4 113.1 1.003 90.9 91.0 0.999 52.6 52.6 1.001
5120→1024 dgrad bf16 16384 112.9 112.9 1.000 90.6 90.8 0.998 52.7 52.6 1.001
5120→1024 fwd bf16 16231 117.7 117.5 1.001 93.9 93.3 1.006 54.8 54.8 1.001
5120→1024 fwd bf16 16384 117.5 117.5 1.000 93.8 94.4 0.994 54.8 54.8 1.001
5120→1024 wgrad bf16 16231 121.9 122.2 0.997 98.4 98.6 0.998 63.5 63.4 1.001
5120→1024 wgrad bf16 16384 122.3 122.8 0.996 99.7 99.8 0.999 63.9 63.8 1.001
5120→1024 wgrad f32 16231 122.8 123.3 0.996 99.7 100.3 0.994 77.4 77.6 0.996
5120→1024 wgrad f32 16384 123.6 124.0 0.997 100.2 100.5 0.998 77.7 77.8 0.999
5120→5120 dgrad bf16 16231 481.5 480.4 1.002 428.1 428.8 0.998 223.8 223.8 1.000
5120→5120 dgrad bf16 16384 482.5 483.4 0.998 430.5 430.2 1.001 224.2 224.4 0.999
5120→5120 fwd bf16 16231 481.5 481.4 1.000 429.0 429.2 1.000 223.8 224.1 0.999
5120→5120 fwd bf16 16384 482.6 484.6 0.996 430.0 430.1 1.000 224.1 224.4 0.999
5120→5120 wgrad bf16 16231 489.1 489.8 0.999 436.0 434.7 1.003 234.3 234.4 1.000
5120→5120 wgrad bf16 16384 494.2 494.3 1.000 436.8 439.3 0.994 236.1 236.2 1.000
5120→5120 wgrad f32 16231 492.5 491.4 1.002 437.8 439.0 0.997 240.7 240.6 1.000
5120→5120 wgrad f32 16384 496.1 492.3 1.008 439.4 442.2 0.994 242.6 242.5 1.000
5120→17408 dgrad bf16 16231 1630.1 1597.1 1.021 1480.0 1473.4 1.004 812.0 812.5 0.999
5120→17408 dgrad bf16 16384 1593.7 1618.6 0.985 1494.0 1494.8 0.999 822.6 821.0 1.002
5120→17408 fwd bf16 16231 1569.7 1574.8 0.997 1455.1 1494.6 0.974 792.8 794.9 0.997
5120→17408 fwd bf16 16384 1622.0 1584.5 1.024 1473.9 1489.9 0.989 797.8 802.7 0.994
5120→17408 wgrad bf16 16231 1579.4 1564.1 1.010 1466.7 1462.1 1.003 803.3 801.7 1.002
5120→17408 wgrad bf16 16384 1604.1 1594.3 1.006 1473.0 1510.5 0.975 812.0 813.8 0.998
5120→17408 wgrad f32 16231 1621.1 1585.6 1.022 1480.1 1481.8 0.999 817.0 817.1 1.000
5120→17408 wgrad f32 16384 1593.8 1612.7 0.988 1490.6 1488.5 1.001 822.4 823.6 0.999
indexer_hw (6144→32) dgrad bf16 16172 35.1 34.4 1.021 32.9 33.1 0.995 22.4 22.6 0.992
indexer_hw (6144→32) dgrad bf16 16231 34.3 34.4 0.996 32.6 32.6 0.998 22.4 22.5 0.994
indexer_hw (6144→32) dgrad f32 16172 61.3 61.3 1.001 58.2 58.3 0.999 38.4 38.6 0.995
indexer_hw (6144→32) dgrad f32 16231 61.4 61.5 0.998 58.5 58.6 0.998 38.7 38.7 1.001
indexer_hw (6144→32) fwd bf16 16172 37.0 37.1 0.997 37.6 37.5 1.002 26.5 26.6 0.997
indexer_hw (6144→32) fwd bf16 16231 37.0 37.1 0.999 37.4 37.4 1.000 26.4 26.4 0.998
indexer_hw (6144→32) wgrad bf16 16172 37.8 37.9 0.998 39.5 39.4 1.003 26.6 26.7 0.998
indexer_hw (6144→32) wgrad bf16 16231 38.0 38.0 1.001 39.4 39.3 1.002 26.6 26.6 0.999
indexer_hw (6144→32) wgrad f32 16172 37.9 37.9 0.999 39.3 39.4 0.998 26.6 26.6 1.001
indexer_hw (6144→32) wgrad f32 16231 38.0 38.0 1.000 39.4 39.5 0.998 26.5 26.5 1.000
indexer_k (6144→128) dgrad bf16 16172 36.4 36.3 1.004 34.2 34.3 0.999 24.4 24.3 1.002
indexer_k (6144→128) dgrad bf16 16231 36.4 36.4 0.998 34.2 34.2 1.001 24.5 24.6 0.995
indexer_k (6144→128) dgrad f32 16172 66.6 66.4 1.004 62.2 62.1 1.001 41.2 41.2 1.002
indexer_k (6144→128) dgrad f32 16231 66.8 66.5 1.005 62.0 61.9 1.002 41.4 41.5 0.998
indexer_k (6144→128) fwd bf16 16172 38.6 38.6 1.001 38.9 38.9 1.000 26.9 26.8 1.003
indexer_k (6144→128) fwd bf16 16231 38.6 38.5 1.003 39.0 39.0 0.998 27.1 27.3 0.994
indexer_k (6144→128) wgrad bf16 16172 39.0 39.1 0.999 40.7 40.9 0.995 27.6 27.6 1.001
indexer_k (6144→128) wgrad bf16 16231 39.1 38.9 1.005 41.2 41.2 1.000 27.8 27.9 0.997
indexer_k (6144→128) wgrad f32 16172 38.8 38.8 0.998 40.6 40.6 0.999 27.5 27.7 0.994
indexer_k (6144→128) wgrad f32 16231 39.1 38.9 1.003 41.1 41.1 0.998 27.9 27.9 1.001
kv_a (6144→576) dgrad bf16 16172 84.5 84.0 1.005 64.8 64.8 0.999 41.6 41.6 1.000
kv_a (6144→576) dgrad bf16 16231 84.2 83.7 1.006 65.0 64.8 1.002 41.4 41.6 0.997
kv_a (6144→576) dgrad f32 16172 111.6 112.6 0.991 92.3 92.5 0.998 55.5 55.7 0.997
kv_a (6144→576) dgrad f32 16231 112.3 111.5 1.007 93.2 93.3 0.999 55.8 55.9 0.999
kv_a (6144→576) fwd bf16 16172 95.6 95.6 1.000 73.1 73.5 0.995 41.4 41.3 1.001
kv_a (6144→576) fwd bf16 16231 96.1 96.0 1.002 72.9 73.4 0.994 41.3 41.3 1.000
kv_a (6144→576) wgrad bf16 16172 83.4 84.6 0.986 65.8 65.2 1.009 43.3 43.3 1.001
kv_a (6144→576) wgrad bf16 16231 84.3 84.5 0.998 65.9 65.9 1.000 43.5 43.3 1.004
kv_a (6144→576) wgrad f32 16172 84.9 85.2 0.997 66.3 66.2 1.002 44.4 44.3 1.002
kv_a (6144→576) wgrad f32 16231 85.5 85.6 0.999 66.8 66.7 1.002 44.7 44.9 0.996
q_a (6144→2048) dgrad bf16 16172 256.2 255.3 1.003 212.0 211.9 1.001 110.6 110.6 1.000
q_a (6144→2048) dgrad bf16 16231 250.7 252.5 0.993 212.1 211.9 1.001 110.6 110.6 0.999
q_a (6144→2048) dgrad f32 16172 263.6 264.1 0.998 223.7 223.8 1.000 114.6 114.7 1.000
q_a (6144→2048) dgrad f32 16231 263.6 263.9 0.999 224.4 224.0 1.002 114.6 114.5 1.001
q_a (6144→2048) fwd bf16 16172 249.6 249.5 1.000 211.1 210.8 1.002 115.1 114.9 1.001
q_a (6144→2048) fwd bf16 16231 249.8 250.0 0.999 210.6 211.2 0.997 115.1 115.1 1.000
q_a (6144→2048) wgrad bf16 16172 258.8 259.6 0.997 219.2 218.8 1.002 115.1 115.2 0.999
q_a (6144→2048) wgrad bf16 16231 259.5 260.0 0.998 219.7 219.5 1.001 115.5 115.5 1.000
q_a (6144→2048) wgrad f32 16172 261.0 260.7 1.001 220.4 219.9 1.002 116.3 116.5 0.999
q_a (6144→2048) wgrad f32 16231 261.3 261.3 1.000 220.8 220.8 1.000 116.5 116.6 0.999
shared_gate_up (6144→2048) dgrad bf16 16172 251.7 250.9 1.003 212.0 211.8 1.001 110.6 110.6 1.000
shared_gate_up (6144→2048) dgrad bf16 16231 250.9 250.8 1.001 212.5 212.0 1.003 110.6 110.6 1.000
shared_gate_up (6144→2048) dgrad f32 16172 263.8 264.2 0.999 224.5 224.1 1.002 114.5 114.5 1.000
shared_gate_up (6144→2048) dgrad f32 16231 264.1 264.0 1.000 225.0 223.9 1.005 114.7 114.7 1.000
shared_gate_up (6144→2048) fwd bf16 16172 249.1 249.6 0.998 211.8 210.5 1.006 115.0 114.9 1.001
shared_gate_up (6144→2048) fwd bf16 16231 249.2 250.2 0.996 211.5 211.0 1.002 115.4 115.3 1.001
shared_gate_up (6144→2048) wgrad bf16 16172 259.2 259.5 0.999 219.1 219.4 0.999 115.2 115.2 1.001
shared_gate_up (6144→2048) wgrad bf16 16231 260.0 259.8 1.001 220.0 220.2 0.999 115.5 115.4 1.001
shared_gate_up (6144→2048) wgrad f32 16172 260.7 261.3 0.998 220.7 220.5 1.001 116.4 116.4 1.000
shared_gate_up (6144→2048) wgrad f32 16231 260.9 261.5 0.998 221.2 221.4 0.999 116.8 116.7 1.000
dense_gate_up (6144→12288) wgrad bf16 16172 1355.8 1335.3 1.015 1238.0 1236.5 1.001 669.2 668.3 1.001
dense_gate_up (6144→12288) wgrad bf16 16231 1347.4 1356.1 0.994 1246.4 1243.1 1.003 665.7 666.1 0.999
dense_gate_up (6144→12288) wgrad f32 16172 1354.5 1358.4 0.997 1241.3 1244.3 0.998 673.1 672.2 1.001
dense_gate_up (6144→12288) wgrad f32 16231 1345.9 1328.7 1.013 1249.1 1254.6 0.996 676.9 676.2 1.001
8192→128 dgrad bf16 16231 47.8 51.5 0.929 43.9 44.8 0.980 30.6 30.7 0.999
8192→128 dgrad bf16 16384 48.0 47.9 1.001 44.0 44.1 0.998 30.6 30.7 0.999
8192→128 dgrad f32 16231 90.7 90.8 0.998 85.4 85.6 0.998 58.3 58.2 1.002
8192→128 dgrad f32 16384 90.1 90.1 1.000 84.9 84.9 1.000 58.8 58.6 1.003
8192→128 fwd bf16 16231 49.0 49.0 1.000 48.6 48.6 1.000 35.1 35.1 1.000
8192→128 fwd bf16 16384 49.5 49.5 0.999 49.1 49.3 0.997 35.2 35.2 1.000
8192→128 wgrad bf16 16231 44.9 44.9 0.999 45.7 45.5 1.005 34.4 34.4 1.000
8192→128 wgrad bf16 16384 45.3 45.2 1.002 46.0 46.0 0.999 34.8 34.8 1.000
8192→128 wgrad f32 16231 44.9 45.1 0.996 45.7 45.6 1.001 34.3 34.3 1.001
8192→128 wgrad f32 16384 45.1 45.1 1.000 45.8 46.0 0.995 34.7 34.8 0.997
12288→4096 dgrad bf16 16231 900.5 898.4 1.002 826.2 824.2 1.002 438.1 438.1 1.000
12288→4096 dgrad bf16 16384 904.5 903.2 1.001 830.5 828.5 1.002 440.2 440.4 1.000
12288→4096 fwd bf16 16231 919.8 918.0 1.002 835.9 836.3 1.000 452.0 452.1 1.000
12288→4096 fwd bf16 16384 900.1 899.6 1.001 840.9 838.6 1.003 453.3 453.3 1.000
12288→4096 wgrad bf16 16231 903.0 894.3 1.010 832.8 829.5 1.004 457.7 458.0 1.000
12288→4096 wgrad bf16 16384 891.3 892.4 0.999 836.6 840.2 0.996 461.7 461.7 1.000
12288→4096 wgrad f32 16231 932.7 938.1 0.994 850.4 853.3 0.997 469.5 469.5 1.000
12288→4096 wgrad f32 16384 924.7 916.9 1.008 851.7 852.7 0.999 473.3 473.2 1.000
dense_down (12288→6144) wgrad bf16 16172 1321.2 1334.4 0.990 1236.4 1218.6 1.015 665.8 666.1 1.000
dense_down (12288→6144) wgrad bf16 16231 1348.5 1333.9 1.011 1244.8 1243.9 1.001 669.6 668.4 1.002
dense_down (12288→6144) wgrad f32 16172 1338.7 1338.9 1.000 1248.9 1236.7 1.010 672.2 671.9 1.001
dense_down (12288→6144) wgrad f32 16231 1345.8 1340.7 1.004 1252.4 1235.7 1.013 676.2 676.2 1.000
o_proj (16384→6144) dgrad bf16 16172 1813.8 1792.7 1.012 1675.0 1644.2 1.019 904.2 903.9 1.000
o_proj (16384→6144) dgrad bf16 16231 1789.2 1813.1 0.987 1670.9 1665.2 1.003 947.2 936.7 1.011
o_proj (16384→6144) dgrad f32 16172 1856.1 1836.6 1.011 1699.5 1690.6 1.005 921.8 921.7 1.000
o_proj (16384→6144) dgrad f32 16231 1826.5 1818.3 1.005 1713.0 1702.2 1.006 921.2 922.2 0.999
o_proj (16384→6144) fwd bf16 16172 1787.5 1802.2 0.992 1690.2 1687.7 1.001 893.9 893.8 1.000
o_proj (16384→6144) fwd bf16 16231 1807.5 1798.5 1.005 1702.8 1681.1 1.013 908.4 901.8 1.007
o_proj (16384→6144) wgrad bf16 16172 1784.6 1776.0 1.005 1646.6 1652.5 0.996 906.7 898.8 1.009
o_proj (16384→6144) wgrad bf16 16231 1814.9 1772.3 1.024 1681.6 1662.3 1.012 903.4 884.5 1.021
o_proj (16384→6144) wgrad f32 16172 1801.0 1792.6 1.005 1655.4 1658.0 0.998 925.6 924.8 1.001
o_proj (16384→6144) wgrad f32 16231 1795.4 1834.7 0.979 1680.5 1677.0 1.002 928.3 930.1 0.998
17408→5120 dgrad bf16 16231 1593.3 1567.2 1.017 1462.3 1471.2 0.994 799.6 801.3 0.998
17408→5120 dgrad bf16 16384 1589.6 1580.3 1.006 1477.2 1456.8 1.014 809.2 807.8 1.002
17408→5120 fwd bf16 16231 1598.2 1580.2 1.011 1483.5 1465.4 1.012 819.9 824.4 0.994
17408→5120 fwd bf16 16384 1571.5 1607.2 0.978 1468.9 1482.8 0.991 823.4 827.8 0.995
17408→5120 wgrad bf16 16231 1575.8 1567.0 1.006 1464.6 1446.3 1.013 803.5 811.2 0.990
17408→5120 wgrad bf16 16384 1593.8 1600.6 0.996 1470.7 1474.5 0.997 814.2 812.5 1.002
17408→5120 wgrad f32 16231 1555.5 1579.3 0.985 1466.4 1461.9 1.003 819.2 821.4 0.997
17408→5120 wgrad f32 16384 1602.0 1607.5 0.997 1489.2 1469.7 1.013 825.5 826.8 0.999
qabs (None→None) dgrad bf16 16172 241.5 241.5 1.000 232.3 232.0 1.001 152.5 152.7 0.999
qabs (None→None) dgrad bf16 16231 239.2 238.9 1.001 230.8 230.7 1.001 152.9 153.0 0.999
qabs (None→None) fwd bf16 16172 266.5 266.2 1.001 241.3 241.6 0.999 159.2 159.1 1.001
qabs (None→None) fwd bf16 16231 275.3 275.3 1.000 245.7 245.8 1.000 160.2 160.2 1.000
qabs (None→None) wgrad bf16 16172 216.1 216.0 1.001 218.9 218.6 1.002 142.3 142.4 0.999
qabs (None→None) wgrad bf16 16231 217.0 217.3 0.999 219.3 220.1 0.996 142.7 142.9 0.999
vproj (None→None) dgrad bf16 16172 290.0 289.7 1.001 265.4 265.0 1.002 171.8 171.8 1.000
vproj (None→None) dgrad bf16 16231 290.7 290.9 0.999 264.2 264.2 1.000 172.9 173.0 1.000
vproj (None→None) fwd bf16 16172 263.9 264.0 1.000 252.8 252.9 1.000 166.0 165.9 1.000
vproj (None→None) fwd bf16 16231 266.8 266.6 1.001 257.0 256.9 1.001 165.9 165.9 1.000
vproj (None→None) wgrad bf16 16172 235.4 235.4 1.000 236.8 236.9 1.000 155.4 155.5 0.999
vproj (None→None) wgrad bf16 16231 236.4 236.6 0.999 237.6 237.7 0.999 156.2 156.4 0.999
router 4096→256 dgrad f32 16231 188.8 188.7 1.001 173.4 174.4 0.994 105.4 105.5 0.998
router 4096→256 fwd f32 16231 192.3 192.5 0.999 166.1 166.7 0.997 106.0 105.9 1.001
router 4096→256 wgrad f32 16231 188.0 187.5 1.003 163.4 164.4 0.994 98.1 98.3 0.998
router 5120→256 dgrad f32 16231 238.3 239.1 0.996 225.4 225.9 0.998 134.1 134.3 0.999
router 5120→256 fwd f32 16231 227.4 226.8 1.002 196.4 196.4 1.000 124.1 124.3 0.999
router 5120→256 wgrad f32 16231 239.6 239.8 0.999 209.4 210.3 0.996 136.4 136.6 0.999
router 6144→256 dgrad f32 16172 273.3 273.6 0.999 260.4 260.5 1.000 153.9 154.2 0.998
router 6144→256 dgrad f32 16231 275.7 275.9 0.999 262.7 263.4 0.997 155.3 155.5 0.999
router 6144→256 fwd f32 16172 262.2 261.4 1.003 228.3 228.2 1.000 142.9 142.9 1.000
router 6144→256 fwd f32 16231 259.3 259.9 0.998 227.2 227.3 0.999 141.8 142.4 0.996
router 6144→256 wgrad f32 16172 248.7 248.2 1.002 216.4 216.3 1.001 138.0 138.3 0.998
router 6144→256 wgrad f32 16231 248.7 249.5 0.997 216.1 216.7 0.998 138.1 138.2 0.999

The kernel's other Cake variants cover kk/kn/nk/nn layouts, row-major and transposed outputs, 64/128/256-row tiles, batched head projection, split-tail and stream-K forms, and eight-warp CTAs on sm_100a/sm_103a, across 254 build entries: 81 for sm_100a, 86 for sm_103a and 87 for sm_107a.

Summary by CodeRabbit

  • New Features

    • Added deterministic dense-projection GEMM support, including prepared and convenience execution paths for projection and router operations.
    • Added architecture-specific kernel selection for supported GPU architectures.
    • Preserved compatibility with existing package imports and API names through aliases.

  • Documentation

    • Updated backend documentation and benchmark examples to use the deterministic GEMM APIs and describe their execution paths.

@yyihuang

yyihuang commented Oct 30, 2026

Copy link
Copy Markdown
Collaborator Author

@flashinfer-bot run tests/experimental/test_cake_dense_projection_det_gemm.py tests/experimental/test_cake_dense_projection_det_gemm_names.py

@coderabbitai

coderabbitai Bot commented Oct 30, 2026 •
edited
Loading

Copy link
Copy Markdown
Contributor

Review in Change Stack →

No actionable comments were generated in the recent review.

Recent review info
Run configuration
  • Source mode: Auto
  • Repository: Unavailable
  • Branch: Unavailable
  • File: No configuration file from this repository applied
  • Commit: Unavailable
  • Layers: defaults
  • Review profile: CHILL
  • Plan: Advanced
  • Run ID: 20266bd7-b192-4b4d-a323-39f0a13e549c
Commits

Reviewing files that changed from the base of the PR and between 8066248 and 70a9f07.

Files selected for processing: 192
  • .pre-commit-config.yaml
  • benchmarks/bench_cake_dense_projection_det_gemm.py
  • flashinfer/experimental/dense_projection_det_gemm/.gitignore
  • flashinfer/experimental/dense_projection_det_gemm/README.md
  • flashinfer/experimental/dense_projection_det_gemm/__init__.py
  • flashinfer/experimental/dense_projection_det_gemm/build.json
  • flashinfer/experimental/dense_projection_det_gemm/cake_backend.py
  • flashinfer/experimental/dense_projection_det_gemm/cake_build.py
  • flashinfer/experimental/dense_projection_det_gemm/cake_jit.py
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/.clang-format
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_011c85b820ec0cf6863d_tpl_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0a71a49da71c7d34ab37_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0b91f29d88de00658fc5_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0b91f29d88de00658fc5_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0bbd5d9f1d72a90ef3ea_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0bbd5d9f1d72a90ef3ea_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0d487a8f71dac54321ef_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0d487a8f71dac54321ef_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0db65f3ddd780ed42d4d_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0dd5420b570cabc144ab_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0dd5420b570cabc144ab_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_13790f025541d7d162a7_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_13790f025541d7d162a7_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_13d9310f37868a18eace_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_14192244441c1003db74_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_20a2e5c04427e93572da_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_20a2e5c04427e93572da_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_23305bdb787b7c6d7378_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_23305bdb787b7c6d7378_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_295aef2f2e0332de2372_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_29d013dfb72e932eaee5_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_29d013dfb72e932eaee5_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_29de51d5e9cdd6d22491_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_29de51d5e9cdd6d22491_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_3088b8056e0b7f9ca32c_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_31be7faa868eabcedada_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_31be7faa868eabcedada_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_3a87cc189b4a7c1a8d7d_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_3a87cc189b4a7c1a8d7d_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_3d7b66f61c710f893b29_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_3d7b66f61c710f893b29_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_438841f42a0743b63fb5_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_446497e750658dfb057b_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_49acb91191bbcced9051_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_4de304a13db1f2a0e850_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_4de304a13db1f2a0e850_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_4eb96fb2eb1aaa14c7f7_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_4eb96fb2eb1aaa14c7f7_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_505fbe05ce8b62f0a589_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_51c81f11e3b490b5aebb_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_51c81f11e3b490b5aebb_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5226b6c854d7db395d7e_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5445e0f16ad84868fa01_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_552070c21b97faa67a71_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_552070c21b97faa67a71_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5568657a526e2040f0bc_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_56be8af43d5d0822da87_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_59fb66ccddd67565e592_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_59fb66ccddd67565e592_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5a06c5e2f33772a33174_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5b4e87b523f799e09384_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5b4e87b523f799e09384_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5e77ef909e88a99c4357_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_624eb3d7232071ddaf91_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_624eb3d7232071ddaf91_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_6579637169bec3c2f1d5_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_6579637169bec3c2f1d5_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_66df6446dd2cca3106d3_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_736ad30e83d88796883f_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_736ad30e83d88796883f_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_75b1dcdaef42319f5ed4_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_75b1dcdaef42319f5ed4_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_765ece06955b70a95b29_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_7a0ce55fd7221e354322_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_7a0ce55fd7221e354322_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_7d403fa7b167188ccc24_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_7ef28d420de555e7a28d_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_7ef28d420de555e7a28d_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_806927d101fe72fc1317_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_82f698f67a1cca6e4291_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_895f9ba7159bea07aedd_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_8e1f6805342f620cca8c_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_9269f6841d7aeb58471c_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_9269f6841d7aeb58471c_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_9475731f174aaefb2e44_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_9475731f174aaefb2e44_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_9aa9aa60b0c544bfab06_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_9cc6363566760a38b4fe_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_9ffc37c198f2d2f42c1f_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_9ffc37c198f2d2f42c1f_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_a24a3d3fd3ceac3e1cdf_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_a24a3d3fd3ceac3e1cdf_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_aa4f7c2fbd0b9a99e535_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_aa4f7c2fbd0b9a99e535_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_b26ac3f3ed56a78cb052_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_b2a9665ee12fd6afea2b_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_b2a9665ee12fd6afea2b_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_b58b50619ef5ed46e109_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_b773cc0e3894fd0ef2e0_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_b773cc0e3894fd0ef2e0_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_bab00fb2b95c93758095_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_bab00fb2b95c93758095_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_bab9e9dfd7c6ca5e17c0_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_bab9e9dfd7c6ca5e17c0_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_bb0e7457db664d6f0f6a_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_bb0e7457db664d6f0f6a_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_c363601373e2dfb8e9af_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_c363601373e2dfb8e9af_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_c63d1400e64208d7865c_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_c63d1400e64208d7865c_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_ca9a846738b1443f7318_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_d2a5e70bd49083888179_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_d2ea37e4d394bf40922a_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_d2ea37e4d394bf40922a_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_d6ebaab2415ea7b00de9_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_ddff2b911137449c1dfc_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_de7492eeb6e24b31687a_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_device_common.cuh
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_dfd65a75564d04d3826f_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_dfd65a75564d04d3826f_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_e2555d533680e08880ae_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_e2555d533680e08880ae_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_e33ac44bfff66daaee64_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_e33ac44bfff66daaee64_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_e98b3ee7b774c63d7544_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_eb2bb538645fe303bcca_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_eb2bb538645fe303bcca_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_ebbb505a3077e8da0a50_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_ebbb505a3077e8da0a50_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_ec7277a904a7211d67ce_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_ec7277a904a7211d67ce_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_ed1f3ab86ddc55345e69_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_ed1f3ab86ddc55345e69_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_eec6f87f8b7f25eb6aa2_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_eec6f87f8b7f25eb6aa2_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_f2e6d4e79cf21105fa1e_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_f31e1dcb1df00973dc0c_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_f589a6aff5e88318c36c_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_f589a6aff5e88318c36c_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_f6c9b40184336f49390d_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_f6c9b40184336f49390d_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_f993d8de6f29f9e8a1ac_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_f993d8de6f29f9e8a1ac_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_fb4bcdf13915dfb96c59_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_fb766b1be68412e6fda7_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_fb766b1be68412e6fda7_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_fdb106e6e73897521cbb_binding.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_fdb106e6e73897521cbb_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_host_common.cuh
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n128_c1__hen_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n128_pro7__hen_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n256_m256_cg_or_pro7__ht_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n256_pro7__cg_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kn_n128_f32_tma1_c1__pd2_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kn_n256_tma1_bg8_pro7__sol_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kn_n256_tma1_pro7__f32_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_hen_t_skx_pro7__f32_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_pro7__f32_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_skx_pro15__f32_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_t_pro7__f32_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n192_hee_t_tma1_pro7__f32_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_f32_tma1_pro7__ht_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_pro7__ht_kernel.cu
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_device_helpers/cake_cp_async_bulk_gmem2smem_5198c9805c6ecc8d.cuh
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_device_helpers/cake_elect_commit_8131fdc67daf4d23.cuh
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_device_helpers/cake_elect_commit_cg2_multicast_18e7b4ad8e69c467.cuh
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_device_helpers/cake_elect_sync_1ee8cd91025d8932.cuh
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_device_helpers/cake_incr_smem_desc_lo_4429b045244f74bc.cuh
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_device_helpers/cake_make_warp_uniform_15c3664d3e11e8cd.cuh
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_device_helpers/cake_mbarrier_arrive_71ef360c3ac78eb7.cuh
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_device_helpers/cake_mbarrier_arrive_expect_tx_54ce59d489f69526.cuh
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_device_helpers/cake_mbarrier_init_15e581aef85ee586.cuh
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_device_helpers/cake_mbarrier_wait_c93bf7e39b218733.cuh
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_device_helpers/cake_mbarrier_wait_cluster_hint_bdb29a0e36dedbee.cuh
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_device_helpers/cake_tcgen05_mma_f16_03fde8c1cbf058c8.cuh
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_device_helpers/cake_tma_3d_gmem2smem_7df1e07811f83fc6.cuh
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_device_helpers/cake_tma_3d_gmem2smem_cta2_dcd29ed146891990.cuh
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_device_helpers/cake_tma_store_3d_a0d668161ec63906.cuh
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_device_helpers/cake_tmem_ld_x16_a9ed61940fdba296.cuh
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_ffi_checks.h
  • flashinfer/experimental/dense_projection_gemm/__init__.py
  • flashinfer/experimental/dense_projection_gemm/build.json
  • flashinfer/experimental/dense_projection_gemm/cake_jit.py
  • pyproject.toml
  • scripts/check_cake_build_manifests.py
  • tests/experimental/test_cake_dense_projection_det_gemm.py
  • tests/experimental/test_cake_dense_projection_det_gemm_names.py
  • tests/test_helpers/cake_package.py
Files with no reviewable changes: 24
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/.clang-format
  • flashinfer/experimental/dense_projection_det_gemm/cake_build.py
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_device_common.cuh
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_device_helpers/cake_elect_commit_8131fdc67daf4d23.cuh
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_device_helpers/cake_mbarrier_wait_c93bf7e39b218733.cuh
  • flashinfer/experimental/dense_projection_det_gemm/.gitignore
  • flashinfer/experimental/dense_projection_gemm/build.json
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_device_helpers/cake_tmem_ld_x16_a9ed61940fdba296.cuh
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_device_helpers/cake_elect_commit_cg2_multicast_18e7b4ad8e69c467.cuh
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_device_helpers/cake_mbarrier_arrive_71ef360c3ac78eb7.cuh
  • flashinfer/experimental/dense_projection_gemm/cake_jit.py
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_device_helpers/cake_cp_async_bulk_gmem2smem_5198c9805c6ecc8d.cuh
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_device_helpers/cake_elect_sync_1ee8cd91025d8932.cuh
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_device_helpers/cake_tma_3d_gmem2smem_7df1e07811f83fc6.cuh
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_device_helpers/cake_tma_3d_gmem2smem_cta2_dcd29ed146891990.cuh
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_device_helpers/cake_tma_store_3d_a0d668161ec63906.cuh
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_host_common.cuh
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_device_helpers/cake_mbarrier_init_15e581aef85ee586.cuh
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_device_helpers/cake_mbarrier_wait_cluster_hint_bdb29a0e36dedbee.cuh
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_device_helpers/cake_make_warp_uniform_15c3664d3e11e8cd.cuh
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_device_helpers/cake_tcgen05_mma_f16_03fde8c1cbf058c8.cuh
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_ffi_checks.h
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_device_helpers/cake_incr_smem_desc_lo_4429b045244f74bc.cuh
  • flashinfer/experimental/dense_projection_det_gemm/csrc/cake_device_helpers/cake_mbarrier_arrive_expect_tx_54ce59d489f69526.cuh
Files skipped from review as they are similar to previous changes: 1
  • flashinfer/experimental/dense_projection_det_gemm/init.py

Included review availability: This review used your included allowance. Your plan provides up to 8 included reviews per hour; 6 remain after this review.

Walkthrough

The pull request renames the experimental dense-projection GEMM package, public APIs, JIT registry, generated CUDA bindings, kernels, and templates to use DET GEMM naming. It preserves aliases for former package and API names and updates package metadata, benchmarks, documentation, and tests.

Changes: Dense Projection DET GEMM

Layer / File(s) Summary
Backend API and compatibility
flashinfer/experimental/dense_projection_det_gemm/cake_backend.py, flashinfer/experimental/dense_projection_det_gemm/README.md, benchmarks/bench_cake_dense_projection_det_gemm.py, flashinfer/experimental/dense_projection_gemm/__init__.py		 Dense-projection and router APIs use DET names. Former API and package names remain aliases. The README and benchmark use the canonical package and API names.
JIT registry and package wiring
flashinfer/experimental/dense_projection_det_gemm/cake_jit.py, flashinfer/experimental/dense_projection_det_gemm/build.json, pyproject.toml, scripts/check_cake_build_manifests.py, .pre-commit-config.yaml, tests/experimental/test_cake_dense_projection_det_gemm_names.py, tests/test_helpers/cake_package.py		 Adds architecture-specific JIT registry and loading helpers plus a build manifest. Package metadata, manifest checks, pre-commit exclusions, and naming tests use the canonical package name.
Generated kernels and validation
flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/*, tests/experimental/test_cake_dense_projection_det_gemm.py		 Generated bindings, kernel symbols, namespaces, and compile-time selectors use DET GEMM names. Two bindings exchange TMA box dimensions. The test suite updates existing checks and adds repeated-launch bitwise equality checks.

Estimated code review effort: 4 (Complex) | ~60 minutes

Change: Refactor

Merge Risk: Minimal · up to 70a9f

This change renames the dense projection GEMM package to the DET naming and keeps aliases for the former package and API names. No concrete merge-blocking defect was identified. The author's reported benchmark parity and the requested test runs are the remaining checks.

Pre-merge checks: Passed 3 · Failed 1 · Inconclusive 1

Failed Failed checks: 1 warning · 1 inconclusive

Check name Status Explanation Resolution
Description check Warning The description gives substantial technical detail, benchmark results, and the intended deterministic naming change. However, it omits the repository template sections, including related issues, pre-c… Update the description to use the repository template. Add the required sections and headings, link the experimental tracking issue, complete the pre-commit and test checklists, state test results and any full-scope CI justification, docume…
Docstring Coverage Inconclusive Docstring coverage is 12.50% which is insufficient. The required threshold is 80.00%. Docstring coverage is scoped to functions touched by this diff with cyclomatic complexity of at least 2. Documenta… Write docstrings for the functions missing them to satisfy the coverage threshold.
Passed Passed checks: 3 passed
Check name Status Explanation
Title check Passed The title clearly and concisely describes the main change: renaming the Cake dense projection GEMM package and its symbols to deterministic names.
Linked Issues check Passed Check skipped because no linked issues were found for this pull request.
Out of Scope Changes check Passed Check skipped because no linked issues were found for this pull request.
Full details: Description check

Explanation

The description gives substantial technical detail, benchmark results, and the intended deterministic naming change. However, it omits the repository template sections, including related issues, pre-commit and test checklists, experimental-track declarations, tracking issue details, and the required experimental-tests block.

Resolution

Update the description to use the repository template. Add the required sections and headings, link the experimental tracking issue, complete the pre-commit and test checklists, state test results and any full-scope CI justification, document the experimental API owner, rationale, and graduation plan, and add the required experimental-tests fenced block with valid targets such as tests/experimental/test_cake_dense_projection_det_gemm.py and tests/experimental/test_cake_dense_projection_det_gemm_names.py.

Full details: Docstring Coverage

Explanation

Docstring coverage is 12.50% which is insufficient. The required threshold is 80.00%. Docstring coverage is scoped to functions touched by this diff with cyclomatic complexity of at least 2. Documentation-only edits and deletions within surviving functions are included; unchanged diff context and fully deleted functions are excluded. Analyzed 48 functions across 46 files. (122 skipped: 4 unsupported, 4 too large, 114 over the file limit.)

Resolution

Write docstrings for the functions missing them to satisfy the coverage threshold.

Finishing Touches: 💡 1
🛠️ Fix failing CI checks: 💡
  • Commit to this branch
  • Create a new PR
🧪 Generate unit tests (beta)
  • Create a new PR
  • Autofix · Keep fixing CodeRabbit findings and required CI, and resolving merge conflicts

Thanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out.

❤️ Share

Comment @coderabbitai help to get the list of available commands.

@yyihuang

yyihuang commented Oct 30, 2026

Copy link
Copy Markdown
Collaborator Author

@flashinfer-bot run tests/experimental/test_cake_dense_projection_det_gemm.py tests/experimental/test_cake_dense_projection_det_gemm_names.py

@yyihuang
yyihuang enabled auto-merge (squash) October 30, 2026 20:55
@yyihuang yyihuang mentioned this pull request Oct 30, 2026
Closed
4 of 10 tasks
@averyhNV
refactor(dense_projection_det_gemm): move the package to flashinfer.e…
bff0a3c 
…xperimental.dense_projection_det_gemm

Move the package (registry, host, loader, manifest, shared headers and
device helpers) to its det name. The former import path keeps an
__init__ that maps its cake_backend, cake_jit and cake_build modules to
the renamed ones, so existing imports name the same objects. The
generated units are retired and rewritten under the new names by the
regenerated package in a following commit.

Co-Authored-By: Claude <noreply@anthropic.com>
@averyhNV
refactor(dense_projection_det_gemm): det entry-point names, with the …
2cd0ac8 
…former names as aliases

Rename the public entry points of cake_backend to their det names
(DetGemmPlan, prepare_dense_projection_det_gemm, projection_det_forward,
router_fp32_det_gemm, ...) and the loader of cake_jit; every former name
stays an alias of the same object. Error messages, the README, the test
suite and the benchmark (both renamed) use the det names. A new test
checks the alias identities and the former import path, and a GPU test
holds prepared launches to bitwise-identical reruns.

Co-Authored-By: Claude <noreply@anthropic.com>
@averyhNV
refactor(dense_projection_det_gemm): det kernel, template and macro n…
27a92d2 
…ames in the registry and host planner

The registry keys, kernel symbols and instance templates take the det
spellings (dense_proj_det_gemm_*, router_fp32_det_gemm_*,
split_fp32_det_bf16x3), and the host planner selects them by those
names.  The programs whose builds the registry used to restate keep
them with the generator instead.

Co-Authored-By: Claude <noreply@anthropic.com>
@averyhNV
refactor(dense_projection_det_gemm): regenerate the package units und…
22ad060 
…er the det names

The device units, bindings and build manifest move to
csrc/cake_dense_projection_det_gemm with the det kernel, namespace and
macro spellings; the units of the former directory are removed.  The
compiled programs are the same: each unit is the former unit with the
names changed.

Co-Authored-By: Claude <noreply@anthropic.com>
@averyhNV
style(dense_projection_det_gemm): ruff format the lines the det names…
70a9f07 
… lengthened

The longer det names push three call sites and one error message past the
line limit; this is the ruff format output for those lines only.

Co-Authored-By: Claude <noreply@anthropic.com>
@yyihuang
yyihuang force-pushed the dense-projection-det-gemm-names branch from 8066248 to 70a9f07 Compare October 30, 2026 23:11
@yyihuang

yyihuang commented Oct 30, 2026

Copy link
Copy Markdown
Collaborator Author

@flashinfer-bot run tests/experimental/test_cake_dense_projection_det_gemm.py tests/experimental/test_cake_dense_projection_det_gemm_names.py

@yyihuang yyihuang mentioned this pull request Oct 30, 2026
Closed
6 tasks
@yyihuang
yyihuang merged commit 6b6d2c9 into flashinfer-ai:main Oct 31, 2026
23 of 24 checks passed
@yyihuang yyihuang mentioned this pull request Oct 31, 2026
Open
1 task done
This was referenced Nov 4, 2026
Merged
Merged
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

@yzh119 yzh119 yzh119 approved these changes
@flashinfer-bot flashinfer-bot Awaiting requested review from flashinfer-bot flashinfer-bot is a code owner
@bkryu bkryu Awaiting requested review from bkryu bkryu is a code owner
@aleozlx aleozlx Awaiting requested review from aleozlx aleozlx is a code owner
@YangXu1990uiuc YangXu1990uiuc Awaiting requested review from YangXu1990uiuc YangXu1990uiuc is a code owner
@Anerudhan Anerudhan Awaiting requested review from Anerudhan Anerudhan is a code owner
@Jingfan-Sun Jingfan-Sun Awaiting requested review from Jingfan-Sun Jingfan-Sun is a code owner

Assignees

No one assigned

Labels

run-ci

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

4 participants

@yyihuang @yzh119 @flashinfer-bot @averyhNV