Open Kernel Challenge: campaign status

Team 3xC1ibR. Generated Tue Oct 27 13:38 CT (2026-10-27T18:38:16+00:00). Phase: kernels.
Kernel track closes
8d 13h
Nov 4 23:59 PT
Serving track opens
9d 1h
Nov 5 12:00 PT
Serving track closes
23d 13h
Nov 19 23:59 PT
Overall kernel rank
526
of 1341
Judge runs recorded
60
57 passed
Judge queue
0 waiting
64 running on 72 GPUs

Where we stand

ProblemOur rankOur best#1#10#50Gap to #50
Dense projections
muse_bf16_proj
599 / 11661.220x1.637x1.570x1.533x+25.7%
Gated attention decode
muse_attention_decode
352 / 8821.748x3.032x2.668x2.486x+42.2%
Routed experts (MoE)
muse_moe_w4a16
650 / 9241.358x2.620x2.473x2.313x+70.2%
Whole decoder layer
muse_decode_layer
506 / 8041.381x2.399x2.240x2.089x+51.3%
Prefill layer
muse_prefill_layer
453 / 5221.034x2.025x1.903x1.707x+65.1%

Overall kernel board: rank 526 of 1341. The top 50 overall when the kernel track closes get first place in the serving queue. Snapshot taken Oct 27 13:38 CT.

Judge runs left today (reset 00:00 UTC = 7 PM CDT, 6 PM CST)

ProblemScored runsCheck runs
muse_bf16_proj7 of 1013 of 20
muse_attention_decode8 of 1018 of 20
muse_moe_w4a166 of 1010 of 20
muse_decode_layer3 of 104 of 20
muse_prefill_layer9 of 1012 of 20

Our best runs

ProblemModeSpeedup95% intervalFileJobAgentWhen
muse_bf16_projsubmit1.219x1.201x to 1.238xprivate/kernels/muse_bf16_proj/pooler_final_cta_projection.py7e1e3ca4eae241d9okc-dense-projectionsOct 27 13:13 CT
muse_bf16_projcheck1.217x1.198x to 1.235xprivate/kernels/muse_bf16_proj/pooler_final_cta_projection.py0dbd345400b74d5aokc-dense-projectionsOct 27 13:11 CT
muse_attention_decodesubmit1.748x1.722x to 1.778xprivate/kernels/muse_attention_decode/v3_block_diag_128b.py27abf122aa954f53okc-attention-decodeOct 27 12:49 CT
muse_attention_decodecheck1.749x1.723x to 1.775xprivate/kernels/muse_attention_decode/v3_block_diag_128b.py6e1d80f8da7447abokc-attention-decodeOct 27 12:41 CT
muse_moe_w4a16submit1.358x1.338x to 1.379xprivate/kernels/muse_moe_w4a16/fused_local_route_preparation.py8393702059ca48c4okc-moe-expertsOct 27 12:02 CT
muse_moe_w4a16check1.371x1.351x to 1.392xprivate/kernels/muse_moe_w4a16/six_stage_expert_pipeline.py53a911aac7a04df8okc-moe-expertsOct 27 13:07 CT
muse_decode_layersubmit1.381x1.360x to 1.402xprivate/kernels/muse_decode_layer/combined_six_stage_experts.pyb5aaec708c9145bbokc-decode-layerOct 27 13:25 CT
muse_decode_layercheck1.390x1.369x to 1.411xprivate/kernels/muse_decode_layer/combined_six_stage_experts.py932fe81888db478eokc-decode-layerOct 27 13:22 CT
muse_prefill_layersubmit1.034x1.019x to 1.050xprivate/kernels/muse_prefill_layer/filtered_expert_alignment.pya31ef34796454e30okc-prefill-layerOct 27 11:31 CT
muse_prefill_layercheck1.033x1.017x to 1.048xprivate/kernels/muse_prefill_layer/filtered_expert_alignment.py296645f9d3cb40c0okc-prefill-layerOct 27 11:27 CT

What worked

WhenScopeChangeResultJobAgent
Oct 27 13:22 CTmuse_decode_layerCompose selective FFN with six-stage non-EP experts, tiny-token geometry and filtered local routes. Thirteen nearby/threshold graphs and all12cold-L2 scored cases pass locally with promising decoder gains.1.341x → 1.390x932fe81888db478eokc-decode-layer
Oct 27 13:13 CTmuse_decode_layerRetain locally faster small-token FFN projections and singleton unpooler; use stockcombined FP32 elsewhere to remove regressions. Ten nearby graphs and four cold-L2 cases pass; quantify incremental whole-decoder gain.1.335x → 1.341xe7970d768e064b51okc-decode-layer
Oct 27 13:13 CTmuse_bf16_projPromote checked singleton GEMV and deterministic final-CTA pooler completion; check0dbd345400b74d5a lower bound1.198316 exceeds scored1.197021, full175 nearby cases pass1.217x → 1.219x7e1e3ca4eae241d9okc-dense-projections
Oct 27 13:11 CTmuse_bf16_projFuse fixed-order split-K completion into the final pooler CTA and retain validated singleton GEMV; full175 nearby cases pass and paired pooler gains1.7-4.2percent1.207x → 1.217x0dbd345400b74d5aokc-dense-projections
Oct 27 13:07 CTmuse_moe_w4a16Compose six-stage Marlin expert pipelines with validated tiny-batch launch geometry1.364x → 1.371x53a911aac7a04df8okc-moe-experts
Oct 27 12:58 CTmuse_decode_layerPromote check60de8e30065544b3 at1.328504x CI[1.308576,1.348432], whose lower bound exceeds scored1.280506; exact-source ten nearby graphs and four cold-L2 cases pass.1.329x → 1.335x68bb4291c12646f3okc-decode-layer
Oct 27 12:54 CTmuse_decode_layerCombine both KV heads in128-byte cache loads and deepen the long-context pipeline, preserving FP32partials and BF16before norm; ten nearby graphs and four cold-L2 cases pass locally.1.281x → 1.329x60de8e30065544b3okc-decode-layer
Oct 27 12:53 CTmuse_bf16_projSelective original-precision GEMV at m1 QKV/unpooler/router; direct30-replay wins19/45/42%, full175-case scored+nearby correctness passed, unchanged O/pooler/other batches.1.197x → 1.207x4347e30b97864255okc-dense-projections
Oct 27 12:41 CTmuse_attention_decodeBlock-diagonal [16, 128] single-load 128-byte coalesced decode kernel + 5-stage software pipelining + BF16 split buffers1.399x → 1.749x6e1d80f8da7447abokc-attention-decode
Oct 27 12:40 CTmuse_moe_w4a16Tune small-batch Marlin launch geometry and instantiate narrow NVFP4 down-projection tile1.358x → 1.364xa5c5014655cd4071okc-moe-experts
Oct 27 12:36 CTmuse_decode_layerScore split-KV decoder after passing check CI lower bound exceeds scored best and exact-file nearby reference validation1.276x → 1.281xd4dfbb2c4aef412eokc-decode-layer
Oct 27 12:30 CTmuse_decode_layerIntegrate split-KV FP8 attention with BF16-rounded combine and fused norm gate; retain exact preprocessing and scored dense/router/expert paths1.239x → 1.276x558c47e25d1d4408okc-decode-layer
Oct 27 12:27 CTmuse_bf16_projCheck f24562501f164303 passed1.17864x CI[1.16096,1.19632], lower bound above scored1.1177; unchanged source passed all175 local scored+nearby cases with no severe cliffs.1.194x → 1.197xef5f707514cf4ef3okc-dense-projections
Oct 27 12:19 CTmuse_attention_decodeCoalesced vectorized FP8 split-KV Flash-Decoding with exact 2048 sliding-window bounds, fused prologue, and single-pass fused combine-norm-gate epilogue– → 1.399x6acd36cb62154b0dokc-attention-decode
Oct 27 12:13 CTmuse_bf16_projParallel tree reduction removes serialized split-K router loads, while cuBLAS dispatch recovers larger qkv/O/unpooler shapes; full local scored suite passed at 1.1954x.1.120x → 1.194x11ef2ffcc2484412okc-dense-projections
Oct 27 12:02 CTmuse_moe_w4a16Score fused small-batch route preparation after confident check and nearby validation1.349x → 1.358x8393702059ca48c4okc-moe-experts
Oct 27 11:55 CTmuse_moe_w4a16Fuse small-batch expert mapping, nonlocal filtering, sorting and Marlin padding in one CTA1.315x → 1.349x3558f029424a409aokc-moe-experts
Oct 27 11:47 CTmuse_bf16_projk-major packed weight tiles + per-shape split-K configs; router hi+lo at all m; local geomean ~1.14x0.910x → 1.120x9d3fc93a472246e9okc-dense-projections
Oct 27 11:38 CTmuse_moe_w4a16Score filtered local-only Marlin for M>=64 after confident check and exact-file nearby validation1.314x → 1.315x463b8db9cef04df3okc-moe-experts
Oct 27 11:31 CTmuse_prefill_layerAdopt filtered local-expert alignment after passing check 296645f9d3cb40c0 with CI lower bound 1.01749 and all five nearby-shape reference checks1.033x → 1.034xa31ef34796454e30okc-prefill-layer
Oct 27 11:28 CTmuse_moe_w4a16Filter nonlocal routes and use local-only Marlin scheduling for M>=64 while preserving small batches1.286x → 1.314x905a7ab1b74e4c84okc-moe-experts
Oct 27 11:27 CTmuse_prefill_layerExclude nonlocal routes from expert alignment after fused prefill boundaries and long full-attention cache gathering1.013x → 1.033x296645f9d3cb40c0okc-prefill-layer
Oct 27 11:28 CTmuse_decode_layerCombine validated M16 expert tiles with concatenated pooler-router BF16 GEMM and fused normalization-routing epilogue after ten nearby FP32-reference graph tests pass1.213x → 1.239xfcf8a6d7049f4f0aokc-decode-layer
Oct 27 11:21 CTmuse_decode_layerIntegrate expert owner M16 tiles for token counts at least 64 to reduce large-verify Marlin latency1.158x → 1.213x8cc823437b384958okc-decode-layer
Oct 27 11:21 CTmuse_moe_w4a16Score Marlin M16 large-batch tiles after confident check and nearby correctness pass1.279x → 1.286x7abd4026af8c4e8cokc-moe-experts
Oct 27 11:17 CTmuse_prefill_layerCombine boundary fusion with contiguous FP8 scratch only for long full-attention blocks1.012x → 1.013x4e832c3cc25a43efokc-prefill-layer
Oct 27 11:18 CTmuse_moe_w4a16Use Marlin M16 tiles for tokens>=64 after local graph sweep showed large batch gains1.165x → 1.279x47a3ba9ac45f4537okc-moe-experts
Oct 27 11:17 CTmuse_decode_layerScore full-precision small-token router plus fused decoder boundaries after six nearby CUDA graph correctness passes1.156x → 1.158x85e649964e524d1bokc-decode-layer
Oct 27 11:12 CTmuse_prefill_layerCombine prefill-sized attention preprocessing and expert route-normalization fusions with a wider hi-lo router projection1.005x → 1.012xd285b1c1b9ff472cokc-prefill-layer
Oct 27 11:11 CTmuse_prefill_layerOne wider GEMM computes the BF16 hi and lo router projections with the same FP32 output sum1.004x → 1.005x7c040b9b86604204okc-prefill-layer
Oct 27 11:09 CTmuse_decode_layerReplace small-token FP32 router GEMM with direct FP32 per-expert reductions while preserving original FP32 inputs and weights1.124x → 1.156x4d1aa5b56f724027okc-decode-layer
Oct 27 11:06 CTmuse_prefill_layerGather complete visible KV spans byte-exactly for contiguous FA3 on full and sliding attention1.002x → 1.004x484f86bc7ead4e54okc-prefill-layer
Oct 27 11:05 CTmuse_decode_layerFuse pre-expert RMS normalization and global-to-local route mapping at the pooler-to-expert boundary1.105x → 1.124x39dcce7bba1b4eabokc-decode-layer
Oct 27 11:02 CTmuse_decode_layerIntegrate validated small-batch expert normalization and fuse global-to-local route mapping with QK preprocessing fusion1.016x → 1.105xff47441cbcbf440bokc-decode-layer
Oct 27 11:03 CTmuse_moe_w4a16Score the passing fused post-expert norm after nearby-shape local correctness passed1.165x → 1.165x58ac3238662442baokc-moe-experts
Oct 27 11:01 CTmuse_decode_layerFuse QK normalization, FP8 query conversion, and FP8 cache writes into one launch with exact BF16 rounding0.998x → 1.016x1b6b13404f534141okc-decode-layer
Oct 27 10:59 CTmuse_prefill_layerEstablish baseline prefill profiles and paired noise before changing expert tiling or normalization– → 1.002xf76eedd14fcc4bd3okc-prefill-layer
Oct 27 10:58 CTmuse_moe_w4a16Fuse the small-batch post-expert norm and weighted sum, preserving Marlin GEMMs1.001x → 1.165x520eed439309445bokc-moe-experts
Oct 27 10:57 CTmuse_decode_layerProfile SGLang decoder components and establish paired timing control– → 0.998xd8bf34d101f44f0aokc-decode-layer
Oct 27 10:58 CTmuse_bf16_projCalibrate: plain PyTorch starter; collect SGLang per-shape times from the profile– → 0.910x76ba44dc39f1455aokc-dense-projections
Oct 27 10:46 CTmuse_moe_w4a16Calibration: SGLang's own kernel against itself should score 1.00x; measures judge noise and turnaround (62 s end to end).– → 1.001x63e23c03bb8941cforchestrator
Oct 27 11:01 CTmuse_moe_w4a16Fuse post-expert normalization for all batch sizes (idea I-20261027-7f0d)Check 520eed439309445b passed 5/5; 1.16497x CI [1.14750, 1.18245]. bs1 76.90 -> 50.23 us; bs8 150.88 -> 117.57; bs16 195.10 -> 158.86. Nearby-shape local correctness pending before scored submit.okc-moe-experts
Oct 27 11:21 CTservingModel speculative serving costs and validate reduced-layer development (idea I-20261027-39a3)Documented serving model and compute plan; native reduced-layer BF16 configuration validation passed in judge image. Synthetic server execution and full-model performance remain unvalidated.okc-serving-strategy
Oct 27 11:05 CTmuse_decode_layerFuse QK normalization with FP8 cache stores and query conversion (idea I-20261027-2544)Judge check 1b6b13404f534141 passed all 13 tests at 1.01597x, CI [1.00073,1.03121]; saves about 4-6 us per layer. Nearby correctness pending before scored submission.okc-decode-layer
Oct 27 11:06 CTmuse_decode_layerIntegrate optimized projections attention and routed experts across decoder boundaries (idea I-20261027-7a5a)Whole-layer check ff47441cbcbf440b passed 13/13 at 1.10488x, CI [1.08831,1.12146]. Six nearby graph cases pass SGLang and FP32 reference; submitting scored run.okc-decode-layer
Oct 27 11:08 CTmuse_decode_layerFuse pre-expert normalization with local route mapping (idea I-20261027-47ee)Check 39dcce7bba1b4eab passed 13/13 at 1.12367x CI [1.10682,1.14053], versus integrated predecessor 1.10488x; about 3-5 us saved across shapes.okc-decode-layer
Oct 27 11:12 CTmuse_decode_layerUse direct FP32 router reductions for small token counts (idea I-20261027-7f41)Direct FP32 router check 4d1aa5b56f724027 passed 13/13 at 1.15604x CI [1.13870,1.17338]; verify-one layer improves to 188.9/202.0 us (1.384x/1.351x). Original FP32 inputs and weights retained.okc-decode-layer
Oct 27 11:31 CTmuse_prefill_layerCombine attention preprocessing and expert routing boundary fusions (idea I-20261027-7097)Boundary fusions retained in filtered alignment check 296645f9d3cb40c0: 1.032986x CI [1.017491,1.048481], all five nearby cases pass. Fusion-only check d285b1c1b9ff472c was 1.012347x.okc-prefill-layer
Oct 27 11:21 CTmuse_moe_w4a16Use Marlin row tile16 for medium and large expert batches (idea I-20261027-e78d)Judge check47a3ba9ac45f4537:1.27941x CI[1.26022,1.29860] vs1.16519 scored best; M128337.89us and M256466.70us. Local M1/4/33/257 and M128skewed pass at relL2 0.00535-0.00556.okc-moe-experts
Oct 27 11:32 CTmuse_decode_layerCombine pooler and split-router projection with fused boundary routing (idea I-20261027-f8dc)Integrated combined pooler/router plus M16 checkfcf8a6d7049f4f0a passes13/13 at1.23921x CI[1.22062,1.25780]. Ten nearby graph/routing/FP32-reference cases pass. Lower bound exceeds recovered M16 scored best1.20033x; submitting.okc-decode-layer
Oct 27 11:24 CTmuse_decode_layerIntegrate validated M16 expert tiles for larger verify batches (idea I-20261027-b9b9)Whole decoder check 8cc823437b384958 passed 13 tests at 1.21297x CI [1.19478,1.23117]. Ten nearby graph/reference cases passed, including verify batches15/17/31/33; submit eligible improvement over scored 1.15829x.okc-decode-layer
Oct 27 11:38 CTmuse_moe_w4a16Filter nonlocal routes before Marlin expert scheduling (idea I-20261027-0f51)Judge check905a7ab1b74e4c84 passes1.31442x CI[1.29471,1.33414] versus1.28648 scored. M128311.43us,M256436.54us. Exact merged file local M1/4/33/257/skewed128 all pass with relL2 0.00535-0.00556; scored submit follows.okc-moe-experts
Oct 27 11:31 CTmuse_prefill_layerExclude nonlocal routes from expert alignment (idea I-20261027-847e)296645f9d3cb40c0 passed: 1.013047x combined candidate -> 1.032986x filtered alignment, 95% CI [1.017491,1.048481]; five nearby-shape checks pass with local speedups 1.022-1.054x.okc-prefill-layer
Oct 27 11:52 CTmuse_bf16_projRepack weights into contiguous k-major tiles in prepare() (idea I-20261027-9b24)job 9d41ecbeddc4471f: scored 1.118x CI[1.101,1.134], 9/9 tests; k-major packed tiles + per-shape split-Kokc-dense-projections
Oct 27 11:44 CTservingIntegrate owner kernels behind a serving flag (idea I-20261027-f01b)Development adapter kept off by default. Actual SGLang expert methods passed 16 shape/rank cases and 96 changing-input CUDA graph replays; max relative L2 vs stock 0.000381. Full-model serving gain unmeasured.okc-serving-strategy
Oct 27 12:55 CTmuse_attention_decodeFuse QK-norm, RoPE, FP8 Q/K/V cast, and KV-cache store into one prologue kernel with wave-aligned FA3 scheduler_metadata (idea I-20261027-b1e4)Adopted in v2 (698a5ce2a9ea4b7f, 1.383x) and v3 (27abf122aa954f53, 1.748x): _fused_prologue_kernel fuses Q/K RMSNorm + RoPE + FP8 quant + K/V page scatter into 1.1-2.2 us.okc-attention-decode
Oct 27 12:55 CTmuse_attention_decodeCoalesced 128-byte dual-KV-head split-KV CUDA kernel with fused split-combine, output RMSNorm, and sigmoid gate (idea I-20261027-d98c)Adopted in v3_block_diag_128b.py (check 6e1d80f8da7447ab 1.749x, submit 27abf122aa954f53 1.748x): block-diagonal [16, 128] 128B coalesced decode + 5-stage pipelining + BF16 split buffers cut decode B=16 full 161.4->108.4us and B=32 full 299.1->207.4us.okc-attention-decode
Oct 27 12:02 CTmuse_moe_w4a16Fuse local expert mapping with small-batch route padding (idea I-20261027-9e60)3558f029424a409a: check 1.34917x, 95% CI [1.32893,1.36941] beats scored 1.31497x; 288 arbitrary route cases, nearby FP32 reference cases, and fresh-input CUDA graph checks passed.okc-moe-experts
Oct 27 12:30 CTmuse_bf16_projHybrid dispatch: cuBLAS fallback for shapes where Triton loses (idea I-20261027-7b9a)Scored ef5f707514cf4ef3 passed1.19702x CI[1.17907,1.21498] vs previous1.1177; local175-case nearby correctness passed and severe partial-M cliffs resolved.okc-dense-projections
Oct 27 12:30 CTmuse_bf16_projReduce split-K partials with parallel vector loads (idea I-20261027-6638)Scored ef5f707514cf4ef3 passed1.19702x CI[1.17907,1.21498] vs previous1.1177; local175-case nearby correctness passed and severe partial-M cliffs resolved.okc-dense-projections
Oct 27 12:36 CTmuse_decode_layerIntegrate coalesced split KV attention (idea I-20261027-240f)Decoder check558c47e25d1d4408 passes13 official tests at1.275725x CI[1.256590,1.294861], lower bound beats scored1.228263. Exact file passes10 nearby officialFP32/routing checks and four cold-L2 local shapes. Explicit BF16 attention boundary preserved; exact source eligible for scoring.okc-decode-layer
Oct 27 12:30 CTmuse_bf16_projDispatch inefficient partial M tiles to cuBLAS (idea I-20261027-9607)Scored ef5f707514cf4ef3 passed1.19702x CI[1.17907,1.21498] vs previous1.1177; local175-case nearby correctness passed and severe partial-M cliffs resolved.okc-dense-projections
Oct 27 13:15 CTmuse_bf16_projCompute singleton BF16 projections with one CTA per output row (idea I-20261027-dc1f)Selective singleton QKV/unpooler/router GEMV is now in scored job7e1e3ca4eae241d9=1.21949868x CI[1.20120620,1.23779116], combined with final-CTA pooler. Prior singleton-only check1.20705953; direct paired QKV/unpooler/router m1 gains1.1920/1.4544/1.4242; full175 nearby cases pass.okc-dense-projections
Oct 27 13:00 CTmuse_decode_layerCombine KV heads in block diagonal attention (idea I-20261027-6136)Check60de8e30065544b3 passed13/13 at1.328504x CI[1.308576,1.348432], lower bound above scored1.280506. Exact-file10nearby graphs and4cold-L2 cases passed; scored68bb4291c12646f3 passed13/13 at1.335054x CI[1.315028,1.355080], a4.26% improvement. Retain FP32 split partials and BF16round before norm;4scores/6checks remain.okc-decode-layer
Oct 27 13:15 CTmuse_bf16_projFuse deterministic split-K completion into the final CTA (idea I-20261027-9e3c)Final-CTA completion restricted to pooler is now in scored job7e1e3ca4eae241d9=1.21949868x CI[1.20120620,1.23779116] versus prior scored1.19702062. Check0dbd345400b74d5a=1.21656443 CIlo1.19831597 qualified;175 nearby+9 preliminary passed;768-call counter/reset stress and PTX publication verified. Broad router/QKV variants rejected for latency; tree variant experimental.okc-dense-projections
Oct 27 13:28 CTmuse_decode_layerCompose selective FFN and six stage experts (idea I-20261027-851d)Composition passed13nearby/threshold graphs and12cold-L2 shapes; check932fe81888db478e passed13/13 at1.390097x CI[1.369245,1.410948], qualifying above1.335054. Scored b5aaec708c9145bb passed13/13 at1.381116x CI[1.360399,1.401833], a3.45% gain.3scores/4checks left;reserve3until22UTC.okc-decode-layer

What did not work

WhenScopeChangeEvidence / errorAgent
Oct 27 11:25 CTinfrastructureReproduce paired cold-cache CUDA graph measurements locallyStrict judge absolute-time reproduction was tested and failed5% on bs1/bs8; published tooling is retained for relative screening. All5 complete correctness/self suites pass. Runtime-child and four-copy padding probes did not resolve the gap; fresh input schedule/driver remain unmatched.okc-infrastructure
Oct 27 11:07 CTmuse_moe_w4a16Use grouped Triton GEMM with in-kernel NVFP4 dequantizationCheck c97f4764e0b94363 passes correctness but only 0.35598x CI [0.35064,0.36132]; BM16 BN64 BK64 dequantizing transposed tiles is 250-1561 us vs Marlin 80-618 us. This initial decomposition is rejected; preserve correctness kernel for tile/scheduling redesign.okc-moe-experts
Oct 27 11:19 CTmuse_prefill_layerTune Marlin token tiles for prefill expert groupsCorrected local paired T16384 sliding: M32 0.8388x, M48 0.9396x against defaultM64; T753 preliminary timings also regressed. Preserve M64.okc-prefill-layer
Oct 27 11:06 CTmuse_prefill_layerFuse pre-expert RMS normalization and reduce intermediate traffic8f59fa181cba4944 passed 5/5 at 1.00185x CI [0.98682,1.01688]; no measurable speedup beyond baseline noise. Retained only as component in existing experiments.okc-prefill-layer
Oct 27 11:11 CTmuse_prefill_layerGather the live sliding attention span into contiguous FP8 scratch484f86bc7ead4e54 passed corrected byte-exact gather; sliding T753 speedup1.0001x, T16384 1.0002x. Gather overhead cancels attention benefit.okc-prefill-layer
Oct 27 11:11 CTmuse_prefill_layerGather full-context FP8 pages for contiguous FA3 attention484f86bc7ead4e54 passed5/5 overall1.00357x CI[0.98852,1.01863]. Long full contexts improve1.0172x/1.0354x but not enough aggregate benefit to clear noise; retain as possible component.okc-prefill-layer
Oct 27 11:21 CTmuse_prefill_layerUse grouped BF16 tensor-core expert tiles with in-kernel NVFP4 dequantizationLocal T16384 sweep BM64/128/256/512 and BN64/128/256: best0.9321x (BM256 BN128); deeper pipeline0.9316x. Same routes and0.0037 relativeL2. Dequantized grouped tiles do not beat Marlin.okc-prefill-layer
Oct 27 11:15 CTmuse_moe_w4a16Use split-K NVFP4 matrix-vector kernels for small batchesCheck b73543230a8840f4 passes but M1 GEMV is234.93us versus52us best; aggregate0.9020x. Compile stats BN16 BK1024 warps8:148 registers,0 spills; BN32 causes14 spills. Direct split-K launches tens of thousands of route/tile CTAs including inactive routes, with excessive decode instruction cost.okc-moe-experts
Oct 27 11:16 CTmuse_prefill_layerPack router high and low BF16 weights into one wider projection7c040b9b86604204 passed5/5 at1.00466x CI[0.98959,1.01973]; wider router alone does not clear noise. Retained as component of combined boundary experiment.okc-prefill-layer
Oct 27 11:18 CTmuse_moe_w4a16Size Marlin expert alignment to the local expert bankLocal cold-L2 graph comparison: count32 alignment adds1.5-3.7us at M1-33 and is neutral at M128-257. Forced M16 separately improves M128529->375us, M256660->473us, M257662->536us. Reject count-only change; isolate row tile tuning.okc-moe-experts
Oct 27 11:20 CTmuse_moe_w4a16Loop over tiny-batch routes within persistent NVFP4 GEMV tilesLocal correctness passes but persistent route-loop GEMV gives M1 167.87us vs baseline80.99us and best fused path about54us; M2 229.57us vs85.15. Reduces split-K overhead but remains3x slower than best; no judge run spent.okc-moe-experts
Oct 27 11:21 CTmuse_prefill_layerRepack dense BF16 weights for contiguous projection output columnsLocal T16384 column-major dense layout1.0022x with noisy pairs; no gain beyond combined fusion. No judge run spent.okc-prefill-layer
Oct 27 11:21 CTmuse_prefill_layerAvoid unnecessary split-KV reduction for sliding prefill attentionLocal T753 split1/2/4 gives1.0234/1.0262/1.0274x vs fused auto1.0268x; no incremental gain. Split1 also changes outputs/route choices materially. Keep automatic FA3 scheduling.okc-prefill-layer
Oct 27 11:25 CTmuse_prefill_layerAlign routed tokens using only the local expert bankLocal T16384 replay: local expert count 1.0047x versus combined candidate 1.0058x, identical routes and relative error below 0.00022; no meaningful gain.okc-prefill-layer
Oct 27 11:29 CTmuse_moe_w4a16Decode NVFP4 tiles with packed BF16 arithmeticLocal packed-BF16 dequant M1281084.80us vs baseline374.82us; M2561747.84us vs643.90us. RelativeL2 about0.0061 vs baseline. Packed arithmetic does not fix the grouped layout/pipeline bottleneck; no judge budget spent.okc-moe-experts
Oct 27 11:33 CTmuse_moe_w4a16Transpose grouped GEMM operands to align NVFP4 decoding with Hopper MMAWeight-left variant passes local tests at relL2 0.00448-0.00464 and generates8wgmma.mma_async,0mma.sync instructions,96registers,0spills. Still M1281553.73us vs374.21us baseline; M2562481.95 vs644.93. Native WGMMA alone does not fix weight decoding/data movement pipeline; no judge run spent.okc-moe-experts
Oct 27 11:42 CTmuse_prefill_layerTune Marlin scheduling after filtering nonlocal routesLocal paired graph replay: T16384 both 1.03277x; T753 scheduler variant 1.04273x versus scored candidate 1.04977x. Routes identical and numerical error small; no gain, no judge run spent.okc-prefill-layer
Oct 27 11:50 CTmuse_decode_layerFilter nonlocal expert routes before medium batch GEMMsDecoder check0d67368d3f5b42e4 passes13 tests at1.234535x CI[1.216017,1.253053], below1.228263 scored-best criterion.13 nearby graph/reference cases pass. M256 Marlin~436us vs prior~408us; component gain did not transfer reliably. No scored run spent.okc-decode-layer
Oct 27 12:03 CTmuse_decode_layerIntegrate packed split-K dense projectionsFull packed QKV/O check335a36cc48644883 fails fresh timing routing; narrowed post-router pooler/unpooler check55f0e68d8fb44abc passes13 tests+10nearby cases at1.237084x CI[1.218528,1.255641]. Lower bound does not beat current scored best1.228263; no scored run spent. Retain validated narrowed source for future integration.okc-decode-layer
Oct 27 12:13 CTmuse_bf16_projInterleave router hi/lo into one packed streamCold-L2 local router comparison: interleaved m128 52.63 us versus parallel dual-stream 10.02 us; m256 97.06 versus 12.65; all six router cases slower. JSON: private/kernels/muse_bf16_proj/local/interleaved_router_projection.json. Keep separate hi/lo streams.okc-dense-projections
Oct 27 12:38 CTmuse_moe_w4a16Extend Marlin row tile16 to medium decode batches40 fresh routes per nearby shape in small_expert_composition_20261027.log: rowtile16 atM16 gives0.97820x vs current, M17 gives0.98266x; previous12-route screening gain did not reproduce. Gate geometry on rowtile16 regresses further. Exclude from candidate; no judge run spent.okc-moe-experts
Oct 27 12:18 CTmuse_decode_layerFuse small batch local expert route preparationDecoder check75afe45528d44114 passes13 judge tests at1.237194x CI[1.218636,1.255752], plus624arbitrary route graph cases and10nearby decoder checks. Lower bound does not beat current scored1.228263. Retain validated helper for future composition; no scored run spent. Original result recovered read-only after supervisor exit.okc-decode-layer
Oct 27 12:34 CTmuse_bf16_projCompute small router batches with one CTA per output rowCombined aligned candidate check8f444f50d9074e88 passed1.19380x CI[1.17590,1.21171], below scored1.19702 and not a significant aggregate gain. Local changed-case wins did not transfer to a better judge score; retain component observations, do not submit.okc-dense-projections
Oct 27 12:34 CTmuse_bf16_projFill H200 SMs with narrow output tiles before split-KCombined aligned candidate check8f444f50d9074e88 passed1.19380x CI[1.17590,1.21171], below scored1.19702 and not a significant aggregate gain. Local changed-case wins did not transfer to a better judge score; retain component observations, do not submit.okc-dense-projections
Oct 27 12:37 CTmuse_bf16_projPipeline dense projection loads through Hopper TMA descriptorsBoth TMA layouts passed correctness but lost: raw0.8572x, packed0.8656x local scored vs best1.19702. Compiler emits WGMMA and TMA, no spills,58/106/202-203 registers and98/147/123KB shared. Packing did not rescue it; do not spend judge runs or repeat broad tuning.okc-dense-projections
Oct 27 12:44 CTmuse_bf16_projPretranspose BF16 weights for cuBLAS NN dispatchPretransposed cuBLAS local suite passed correctness but aggregate1.16516x CI[1.14768,1.18264] trails scored1.19702 and local generalized~1.21. No projection range showed a sufficient consistent gain to justify a judge check.okc-dense-projections
Oct 27 13:17 CTmuse_decode_layerPack combined FFN projections with parallel FP32 reductionBroad combined FP32 packed path regresses atM32/128/256; refined selective candidate passes10nearby graphs,4cold-L2 cases and13judge tests. Checke7970d768e064b51=1.340513x CI[1.320405,1.360621], lower bound below scored1.335054. No score spent; retain passing selective source for composition/pre-reset policy.4scores/5checks remain.okc-decode-layer
Oct 27 13:03 CTmuse_bf16_projCompress BF16 weight planes losslessly with exact exceptionsLocal job127df77e: 20/20 exact BF16 reconstruction checks, 9 preliminary and30 timed correctness cases passed, but aggregate0.205600x CI[0.202516,0.209586]. Pooler m1/128/256=180.45/310.80/572.20us versus baseline22.80/27.60/32.96us. Decode/exception path collectively overwhelms traffic savings; specific compiler cause not isolated. No judge budget spent.okc-dense-projections
Oct 27 12:54 CTmuse_moe_w4a16Tune NVFP4 Marlin asynchronous pipeline depthFixed K64/N128/128-thread geometry and3blocks perSM:48 six-batch variants across20 fresh inputs each passed correctness. Both GEMMs at3stages give0.95185-0.95641x vs stock;2stages give0.67935-0.75664x. Four-stage control tracks stock. Lower pipeline depth consistently slows; no judge run spent.okc-moe-experts
Oct 27 13:21 CTmuse_moe_w4a16Choose separate row tiles for expert projectionsForty changing inputs each at9sizes,45 candidate/control configurations passed with max relativeL2 0.000506. Native control ratios1.0014-1.0057. Mixed8/16 or16/8 schedules regress at every tested size once extra alignment included; ratios M128=.8748/.9282 and M256=.7893/.8790. Same8/8 M29/32 shows only small route-sensitive screening gain. No judge run spent;11checks/6scores remain.okc-moe-experts
Oct 27 13:20 CTmuse_bf16_projRetune pooler occupancy after fusing split-K completionCompleted36 cold-L2 paired configurations x10freshreplays with correctness/counter-reset/resource checks. Changed configurations gained atmost1.49% for m8 and0.48% m1, lost atlarger batches; m128 alternatives0.695-0.872x, m2560.558-0.777x. All0spills. Unchanged configs plus tree reduction gained3.20% m16 and3.47% m128, retained as separate completion refinement, not a tile/occupancy win. No additionaljudge spent.okc-dense-projections
Oct 27 13:36 CTmuse_decode_layerComplete split FFN projections in final CTABF16 completion passed448bit-exact changing graphs and counter resets;FP32 combined passed320fresh graphs,maxrelativeL2=5.37e-8,and counter resets. Each variant passed10nearby decoder cases;affected cold-L2 shapes passed. Paired isolated BF16 cost is unchanged,FP32T12–16 is0–3%slower. Across-run decoder changes are not compelling evidence; no judge run spent. Best1.381116x;3scores/4checks remain.okc-decode-layer
Oct 27 13:28 CTmuse_bf16_projReduce register pressure in exact BF16 weight reconstructionAll18 lower-pressure variants passed correctness and reached0spills (71-163registers), yet best direct speedups versus scored implementation were0.2604 atpooler1,0.1644 at128,0.1009 at256; some256variants0.0348. Removing spills doesnot recover viability; decoder/exception/layout overhead remains dominant collectively. No judge budget spent; close this representation line rather than broad tuning.okc-dense-projections
Oct 27 13:34 CTmuse_bf16_projSelect completion reduction by split count and tile sizeFull9preliminary+175scored/nearby cases passed. Localaggregate1.214604x CI[1.196385,1.232823], nearby1.139731x, worst inheritedpooler255 .857x. Prior scored-source localaggregate1.222330x/nearby1.141313x. Small component tree wins didnot establish aggregate gain; no judge check or scored run spent. Retain code as validated ablation, not promoted.okc-dense-projections
Oct 27 11:53 CTmuse_decode_layerIntegrate validated packed split-K BF16 qkv, O, unpooler and small-token pooler; retain current router, combined pooler-router and M16 expertsdoneokc-decode-layer
Oct 27 11:28 CTmuse_decode_layerRecover completed M16 scored job 2acc91eadaed4546 after the local supervisor exited; identical file reuses its resultduplicateokc-decode-layer
Oct 27 11:03 CTmuse_prefill_layerReplace sliding attention page traversal with a complete live-span FP8 gather and contiguous FA3 calldoneokc-prefill-layer

Ideas queue

StatusScopeIdeaExpectedEffortOwnerId
activemuse_attention_decodeFused sliding-window attention megakernel (QK-norm + RoPE + KV write + 2048-window split-KV + norm-gate)
Sliding-window layers (weight 46 vs 16 for full) have a 0.1-3.7 us roofline floor but take 25.4-37.9 us in baseline due to 9 separate kernel launches; account for 67% of decode-1/verify-1 load time and 21-30% of batched load time.
3-5x speedup on sliding-window shapeshoursokc-attention-decodeI-20261027-90fc
activemuse_bf16_projTriton split-K GEMV/GEMM with vectorized 16B weight loads
All scored shapes are memory-bound skinny GEMMs (floor = weight bytes at 4.8 TB/s); cuBLAS bf16 dispatch leaves headroom vs floor. Split-K gives enough CTAs to saturate HBM at n=256..8192.
1.3-1.7x overalldayokc-dense-projectionsI-20261027-cd4a
activemuse_moe_w4a16Bound Marlin registers to enable four-block residency
CUDA attributes isolate six-stage row16 occupancy limit:157 registers per thread, zero local bytes, active3blocks at57088B shared; row8 uses127registers and fits4. Compile an isolated128-register cap and verify actual attributes/occupancy before timing. Compare spill/local-memory cost against extra resident CTAs in changing-input cold-L2 graphs, preserve precision and schedules, and reject if arithmetic or gains fail.
Unmeasured; register spilling may outweigh residency benefithoursokc-moe-expertsI-20261027-fddd
activemuse_moe_w4a16Compose bounded-register small expert launches
The128-register sweep passed90fresh-input configurations. It made row16 four-block residency possible but slowed every row16 shape; row8 gate4/down3 improved M8/M16 about1.8-2.1% vs validated six-stage baseline. Compose capped launches only for non-EP row8 tokens>4 and retain uncapped tiny/row16/EP paths. Validate exact-file FP32 nearby and40fresh graph inputs before deciding on judge.
Unmeasured combined gain; target small-shape gains without large-shape regressionhoursokc-moe-expertsI-20261027-476b
parkedmuse_moe_w4a16Tune Marlin persistent launch geometry for sparse local routes
Fused-route check3558f029424a409a attributes 86.6% to95.5% of GPU time to Marlin. The local C++ selector ranks two tile configurations by theoretical occupancy and fixes one to four persistent blocks per SM; the Python wrapper hides thread_k/thread_n. Sparse EP routes may prefer narrower N tiles or fewer persistent blocks, without changing FP32 reduction or NVFP4 math.
Unmeasured; require fresh-input local confidence before a judge checkhoursokc-moe-expertsI-20261027-ef65
parkedmuse_moe_w4a16Instantiate narrow NVFP4 Marlin tiles
Existing launch sweep found automatic geometry near best and N64 unavailable for NVFP4. Source has an N64,K128,128-thread instantiation for MXFP4, while the generic Marlin template also accepts NVFP4/E4M3 scales. Test an isolated explicit NVFP4 instantiation; validate scale addressing and FP32 reductions before timing. This changes compiled geometry without shared SGLang edits.
Unmeasured; narrower tiles may improve sparse expert utilizationhoursokc-moe-expertsI-20261027-08a1
parkedmuse_moe_w4a16Compose validated small-batch expert launch changes
Explicit N64 tile is correct but generally slower for M>=8; M1 down GEMM screening shows about6% whole-path gain, potentially additive with gate geometry and previously parked M16 rowtile16. Re-test selected combinations on40 fresh routes and nearby sizes to distinguish useful aggregate improvement from control outliers before a judge check.
Unmeasured combined gain; seek more than1.5% aggregate improvementhoursokc-moe-expertsI-20261027-5cb8
parkedmuse_moe_w4a16Deepen NVFP4 Marlin asynchronous pipeline
Reducing pipeline depth4 to3 slowed every shape about5%;depth2 slowed much more at fixed geometry/residency. Test explicit depths6/8 with unchanged tile and block count to examine latency hiding. Bound actual requested shared memory conservatively using the four-stage cache estimate scaled by depth, retaining the same FP32 reduction and weight layout. Counter access is blocked, so causality comes from controlled graph comparisons.
Unmeasured; deeper staging may improve overlaphoursokc-moe-expertsI-20261027-8d02
parkedmuse_moe_w4a16Compose six-stage expert pipeline with tiny-batch geometry
Guarded fixed-geometry six-stage sweep passes all fresh-input cases and preserves3resident blocks perSM; both GEMMs improve about4% atM8/16 and1% atlarger shapes. Combine six-stage K64/N128 geometry for M>4 with previously validated tiny-batch wide-gate/narrow-down geometry, preserving row tiles, metadata and FP32 reductions. Eight stages regress.
Unmeasured combined gain; direct fresh-input comparison before judgehoursokc-moe-expertsI-20261027-5f1f
parkedmuse_moe_w4a16Retune expert residency for six-stage pipeline
Six-stage kernels improved over four stages at a fixed three blocks per SM, while previous launch geometry tuning used four stages. Test two and four resident blocks independently for gate/up and down at fixed six-stage K64/N128/128-thread geometry. Enforce shared-memory and actual occupancy guards, retain tiny-batch path and shared row schedule, and compare changing-input cold-L2 graphs against the validated six-stage file.
Unmeasured; examine interaction between deeper staging and residencyhoursokc-moe-expertsI-20261027-7532
parkedservingFuse speculative write-window translation
Upstream 2e9b2cacda fuses virtual slot gathering and FULL/SWA translation, but its DSpark caller differs from pinned DFlash.
Remove metadata launches if reachable from DFlash; must measurehoursokc-serving-strategyI-20261027-f272

Decisions

WhenScopeDecisionRationaleBy
Oct 27 12:08 CTmuse_bf16_projReplace the failed Kimi dense-projections owner with a Codex agentKimi ended a turn with no response or error at about 17:00 UTC (app_agent_failed). Its 1.118x kernel and ideas are committed; the replacement (gpt-6.1-sol, xhigh) received a handoff with files, jobs, dead ends and open ideas. The Kimi thread is retired in the roster.orchestrator
Oct 27 11:39 CTallDeliver scheduled orchestration passes through a macOS LaunchAgent instead of Ribbon Field schedules or agent-coord job timersRibbon Field schedule_prompt requests stayed queued (stored as queued_schedule, which the running app does not process) and were cancelled to avoid duplicates; five agent-coord sleep-job timers lost their supervisors within 40 minutes and fired early. The LaunchAgent dev.rquinonez.okc-campaign-scheduler runs scripts/campaign_scheduler.py every 5 minutes and sends each due event (51 events: daily 23:10 and 00:10 UTC through Nov 5, kernel close, serving kickoff, serving dailies, serving close) once as an agent-coord message to the orchestrator; state in ~/.local/state/okc-campaign-scheduler. Requires the Mac to be awake.orchestrator
Oct 27 11:12 CTinfrastructureUse the pod's local container disk for work and the GPU lock; never the RunPod network volumeAt 16:05 UTC on 2026-10-27 the /workspace volume of pod pezfd1k0c76rk8 (MooseFS, mfs#us-co-1.runpod.net) stalled: writes and directory listings hang and three benchmark processes are stuck in uninterruptible I/O. Local disk and SSH work. Lock moved to /var/lock/okc_gpu.lock and work to /root/work/<agent>. Future pods: no persistent network volume on hot paths.orchestrator
Oct 27 11:00 CTallDaily orchestration cadence through the kernel close and a serving-track kickoff, scheduled in Ribbon FieldUnused scored runs expire at 00:00 UTC and idle agents need a nudge each judge day. Scheduled self-prompts: 23:10 UTC pre-reset sweep (Oct 27 to Nov 4), 00:10 UTC new-day pass (Oct 28 to Nov 5), Nov 5 04:30 UTC kernel-close push, Nov 5 20:10 UTC serving kickoff (which schedules the serving-phase dailies). The Mac must be awake and Ribbon Field running.orchestrator
Oct 27 10:51 CTinfrastructureRent one RunPod secure H200 (CUDA 13.2 host) now for harness validation and shared developmentThe judge image needs CUDA 13 (driver >= 580). RunPod showed H200 SXM CUDA-13 stock only on secure cloud at 5.29 USD/h; community 3.59 USD/h hosts were CUDA 12.9 only. H100 had zero stock on the first read. A provider comparison is assigned to decide the long-run machine.orchestrator
Oct 27 10:51 CTallOne owning agent per kernel problem; only the owner spends that problem's judge runsDaily limits are per problem (10 scored, 20 checks); single ownership prevents two agents exhausting one budget and keeps the ledger attributable.orchestrator
Oct 27 10:51 CTallKeep kernel sources and serving patches in a private companion repository until each track closesThe repository is public by request; 1328 teams with agent fleets can read GitHub. Infrastructure, ledger and status stay public; code that would hand a competitor our speedup stays private until Nov 5 (kernels) and Nov 20 (serving), then is published.orchestrator

Recent judge runs

WhenProblemModeResultSpeedupHypothesisAgentJob
Oct 27 13:25 CTmuse_decode_layersubmitpass1.381xPromote exact check932fe81888db478e at1.3900966x CI[1.369245,1.410948], lower above scored1.335054;13nearby/threshold cases and12cold-L2 shapes pass locally.okc-decode-layerb5aaec708c9145bb
Oct 27 13:22 CTmuse_decode_layercheckpass1.390xCompose selective FFN with six-stage non-EP experts, tiny-token geometry and filtered local routes. Thirteen nearby/threshold graphs and all12cold-L2 scored cases pass locally with promising decoder gains.okc-decode-layer932fe81888db478e
Oct 27 13:13 CTmuse_decode_layercheckpass1.341xRetain locally faster small-token FFN projections and singleton unpooler; use stockcombined FP32 elsewhere to remove regressions. Ten nearby graphs and four cold-L2 cases pass; quantify incremental whole-decoder gain.okc-decode-layere7970d768e064b51
Oct 27 13:13 CTmuse_bf16_projsubmitpass1.219xPromote checked singleton GEMV and deterministic final-CTA pooler completion; check0dbd345400b74d5a lower bound1.198316 exceeds scored1.197021, full175 nearby cases passokc-dense-projections7e1e3ca4eae241d9
Oct 27 13:11 CTmuse_bf16_projcheckpass1.217xFuse fixed-order split-K completion into the final pooler CTA and retain validated singleton GEMV; full175 nearby cases pass and paired pooler gains1.7-4.2percentokc-dense-projections0dbd345400b74d5a
Oct 27 13:07 CTmuse_moe_w4a16checkpass1.371xCompose six-stage Marlin expert pipelines with validated tiny-batch launch geometryokc-moe-experts53a911aac7a04df8
Oct 27 12:58 CTmuse_decode_layersubmitpass1.335xPromote check60de8e30065544b3 at1.328504x CI[1.308576,1.348432], whose lower bound exceeds scored1.280506; exact-source ten nearby graphs and four cold-L2 cases pass.okc-decode-layer68bb4291c12646f3
Oct 27 12:54 CTmuse_decode_layercheckpass1.329xCombine both KV heads in128-byte cache loads and deepen the long-context pipeline, preserving FP32partials and BF16before norm; ten nearby graphs and four cold-L2 cases pass locally.okc-decode-layer60de8e30065544b3
Oct 27 12:49 CTmuse_attention_decodesubmitpass1.748xBlock-diagonal [16, 128] single-load 128-byte coalesced decode kernel + 5-stage software pipelining + BF16 split buffers (verified 1.749x in check 6e1d80f8da7447ab)okc-attention-decode27abf122aa954f53
Oct 27 12:53 CTmuse_bf16_projcheckpass1.207xSelective original-precision GEMV at m1 QKV/unpooler/router; direct30-replay wins19/45/42%, full175-case scored+nearby correctness passed, unchanged O/pooler/other batches.okc-dense-projections4347e30b97864255
Oct 27 12:41 CTmuse_attention_decodecheckpass1.749xBlock-diagonal [16, 128] single-load 128-byte coalesced decode kernel + 5-stage software pipelining + BF16 split buffersokc-attention-decode6e1d80f8da7447ab
Oct 27 12:40 CTmuse_moe_w4a16checkpass1.364xTune small-batch Marlin launch geometry and instantiate narrow NVFP4 down-projection tileokc-moe-expertsa5c5014655cd4071
Oct 27 12:36 CTmuse_decode_layersubmitpass1.281xScore split-KV decoder after passing check CI lower bound exceeds scored best and exact-file nearby reference validationokc-decode-layerd4dfbb2c4aef412e
Oct 27 12:32 CTmuse_bf16_projcheckpass1.194xRestrict single-pass BN32 BK512 pooler to singleton/aligned16-token tiles and original-FP32 router to m1; directpaired gains6-7% pooler and11% router,175 scored+nearby correctness passed without large cliffs.okc-dense-projections8f444f50d9074e88
Oct 27 12:30 CTmuse_decode_layercheckpass1.276xIntegrate split-KV FP8 attention with BF16-rounded combine and fused norm gate; retain exact preprocessing and scored dense/router/expert pathsokc-decode-layer558c47e25d1d4408
Oct 27 12:25 CTmuse_attention_decodesubmitpass1.383xSubmit verified 1.40x coalesced split-KV Flash-Decoding + fused prologue + fused combine-norm-gate epilogue to get on the boardokc-attention-decode698a5ce2a9ea4b7f
Oct 27 12:27 CTmuse_bf16_projsubmitpass1.197xCheck f24562501f164303 passed1.17864x CI[1.16096,1.19632], lower bound above scored1.1177; unchanged source passed all175 local scored+nearby cases with no severe cliffs.okc-dense-projectionsef5f707514cf4ef3
Oct 27 12:25 CTmuse_bf16_projcheckpass1.179xParallel reduction plus generic tile-utilization cuBLAS dispatch retains scored gains while removing nearby cliffs;175 local scored+nearby cases passed, nearby1.1556x, worst0.878x.okc-dense-projectionsf24562501f164303
Oct 27 12:19 CTmuse_attention_decodecheckpass1.399xCoalesced vectorized FP8 split-KV Flash-Decoding with exact 2048 sliding-window bounds, fused prologue, and single-pass fused combine-norm-gate epilogueokc-attention-decode6acd36cb62154b0d
Oct 27 12:12 CTmuse_decode_layercheckpass1.237xFilter and pad existing local expert routes in one CTA at tokens<=32 and use non-EP Marlin; retain current router and original medium scheduleokc-decode-layer75afe45528d44114
Oct 27 12:13 CTmuse_bf16_projcheckpass1.194xParallel tree reduction removes serialized split-K router loads, while cuBLAS dispatch recovers larger qkv/O/unpooler shapes; full local scored suite passed at 1.1954x.okc-dense-projections11ef2ffcc2484412
Oct 27 12:02 CTmuse_moe_w4a16submitpass1.358xScore fused small-batch route preparation after confident check and nearby validationokc-moe-experts8393702059ca48c4
Oct 27 12:00 CTmuse_decode_layercheckpass1.237xUse packed BF16 pooler/unpooler only at small batches; preserve stock QKV/O and current scored attention/router path to avoid routing driftokc-decode-layer55f0e68d8fb44abc
Oct 27 11:55 CTmuse_moe_w4a16checkpass1.349xFuse small-batch expert mapping, nonlocal filtering, sorting and Marlin padding in one CTAokc-moe-experts3558f029424a409a
Oct 27 11:53 CTmuse_decode_layercheckfail–Integrate validated packed split-K BF16 qkv, O, unpooler and small-token pooler; retain current router, combined pooler-router and M16 expertsokc-decode-layer335a36cc48644883
Oct 27 11:49 CTmuse_bf16_projsubmitpass1.118xBoard entry: packed split-K, judge check 9d3fc93a 1.12x CI[1.103,1.137]okc-dense-projections9d41ecbeddc4471f
Oct 27 11:47 CTmuse_bf16_projcheckpass1.120xk-major packed weight tiles + per-shape split-K configs; router hi+lo at all m; local geomean ~1.14xokc-dense-projections9d3fc93a472246e9
Oct 27 11:42 CTmuse_decode_layercheckpass1.235xFilter nonlocal expert routes before non-EP Marlin scheduling at tokens>=64, preserving combined pooler-router and small-token pathsokc-decode-layer0d67368d3f5b42e4
Oct 27 11:38 CTmuse_moe_w4a16submitpass1.315xScore filtered local-only Marlin for M>=64 after confident check and exact-file nearby validationokc-moe-experts463b8db9cef04df3
Oct 27 11:31 CTmuse_prefill_layersubmitpass1.034xAdopt filtered local-expert alignment after passing check 296645f9d3cb40c0 with CI lower bound 1.01749 and all five nearby-shape reference checksokc-prefill-layera31ef34796454e30
Oct 27 11:32 CTmuse_decode_layersubmitpass1.228xScore concatenated pooler-router projection plus M16 expert tiles after ten nearby canonical graph/reference checks; check lower bound1.22062 exceeds scored1.20033okc-decode-layer1555e5739a3343b0
Oct 27 11:28 CTmuse_moe_w4a16checkpass1.314xFilter nonlocal routes and use local-only Marlin scheduling for M>=64 while preserving small batchesokc-moe-experts905a7ab1b74e4c84
Oct 27 11:27 CTmuse_prefill_layercheckpass1.033xExclude nonlocal routes from expert alignment after fused prefill boundaries and long full-attention cache gatheringokc-prefill-layer296645f9d3cb40c0
Oct 27 11:31 CTmuse_decode_layersubmitpass1.200xRecover completed M16 scored result after local Agent Coord supervisor exitedokc-decode-layer2acc91eadaed4546
Oct 27 11:28 CTmuse_decode_layercheckpass1.239xCombine validated M16 expert tiles with concatenated pooler-router BF16 GEMM and fused normalization-routing epilogue after ten nearby FP32-reference graph tests passokc-decode-layerfcf8a6d7049f4f0a
Oct 27 11:28 CTmuse_decode_layersubmitfail–Recover completed M16 scored job 2acc91eadaed4546 after the local supervisor exited; identical file reuses its resultokc-decode-layerc1ebb807bfe742c2
Oct 27 11:22 CTmuse_decode_layercheckpass1.170xLosslessly concatenate pooler and router hi-lo weights into one BF16 GEMM with FP32 output and fuse expert normalization routing and rank mappingokc-decode-layer76e361402018449a
Oct 27 11:21 CTmuse_decode_layercheckpass1.213xIntegrate expert owner M16 tiles for token counts at least 64 to reduce large-verify Marlin latencyokc-decode-layer8cc823437b384958
Oct 27 11:21 CTmuse_moe_w4a16submitpass1.286xScore Marlin M16 large-batch tiles after confident check and nearby correctness passokc-moe-experts7abd4026af8c4e8c
Oct 27 11:17 CTmuse_prefill_layercheckpass1.013xCombine boundary fusion with contiguous FP8 scratch only for long full-attention blocksokc-prefill-layer4e832c3cc25a43ef
Oct 27 11:18 CTmuse_moe_w4a16checkpass1.279xUse Marlin M16 tiles for tokens>=64 after local graph sweep showed large batch gainsokc-moe-experts47a3ba9ac45f4537
Oct 27 11:17 CTmuse_decode_layersubmitpass1.158xScore full-precision small-token router plus fused decoder boundaries after six nearby CUDA graph correctness passesokc-decode-layer85e649964e524d1b
Oct 27 11:12 CTmuse_prefill_layercheckpass1.012xCombine prefill-sized attention preprocessing and expert route-normalization fusions with a wider hi-lo router projectionokc-prefill-layerd285b1c1b9ff472c
Oct 27 11:11 CTmuse_prefill_layercheckpass1.005xOne wider GEMM computes the BF16 hi and lo router projections with the same FP32 output sumokc-prefill-layer7c040b9b86604204
Oct 27 11:11 CTmuse_moe_w4a16checkpass0.902xUse split-K direct NVFP4 GEMV for M<=4 while retaining fused Marlin for larger batchesokc-moe-expertsb73543230a8840f4
Oct 27 11:09 CTmuse_decode_layercheckpass1.156xReplace small-token FP32 router GEMM with direct FP32 per-expert reductions while preserving original FP32 inputs and weightsokc-decode-layer4d1aa5b56f724027
Oct 27 11:06 CTmuse_prefill_layercheckpass1.004xGather complete visible KV spans byte-exactly for contiguous FA3 on full and sliding attentionokc-prefill-layer484f86bc7ead4e54
Oct 27 11:06 CTmuse_decode_layersubmitpass1.104xScore integrated fused preprocessing, local route mapping, and small-batch post-expert normalization after six nearby graph correctness passesokc-decode-layer64c3295f016c445a
Oct 27 11:06 CTmuse_moe_w4a16checkpass0.356xReplace Marlin with grouped BF16 tensor-core GEMM and in-kernel NVFP4 dequantizationokc-moe-expertsc97f4764e0b94363
Oct 27 11:05 CTmuse_decode_layercheckpass1.124xFuse pre-expert RMS normalization and global-to-local route mapping at the pooler-to-expert boundaryokc-decode-layer39dcce7bba1b4eab

Notes

WhenScopeNoteBy
Oct 27 13:38 CTmuse_bf16_projDecoder owner independently validated final-CTA integration:448 bit-exact BF16 pooler freshgraphcalls atT1-8 and320 FP32 combined calls atT12-16 with maxrelativeL2 5.37e-8; all counters reset, eachvariant also passed10nearbydecodercases. Paired BF16 latency approximately unchanged; FP32combined pairedratios0.9677-0.9997, no benefit. No judge run spent; do not promote this decoder integration from standalone projection gains. Projectionbest1.219499x unchanged, decoderowner reports scored1.381116x.okc-dense-projections
Oct 27 13:34 CTmuse_bf16_projCurrent bounded projection search lines are exhausted without another judge-worthy candidate: compression remained veryslow after eliminating spills, occupancy retuning failed, generic completion refinement passed175nearby but didnot establish aggregate gain. Keep scored7e1e3ca4eae241d9=1.21949868x;13checks7scored remain, reserve3checks until22UTC. Stop60second timer loop per AGENTS no-good-idea exception; resume on a viable new design or scheduled campaign pass. No GPU job remains pending for this owner.okc-dense-projections
Oct 27 13:24 CTmuse_bf16_projCompiler-resource probe found heavy pressure in rejected lossless BF16 projection: pooler m1/128/256 each uses255registers and reports130/146/152spills, versus raw56/115/124registers and0spills. Compressed sharedmemory139264/147456/147456bytes. This is concrete evidence for a targeted lower-lifetime/smaller-tile probe, not proof of a speedup or attribution of the entire0.2056x regression. No judge budget used.okc-dense-projections
Oct 27 13:12 CTmuse_moe_w4a16Six-stage composition check53a911aac7a04df8 passes1.371225x CI[1.350657,1.391794], not promotion-qualified vs scored1.358391. M256 ratio1.3958 vs scored1.4574 merits seed-transfer review; previous small-only checka5c501 also had lowerM2561.4094 despite unchanged timed GPU path there. Current judge case seeds differ, so cross-run shape differences cannot alone establish a kernel regression. Direct same-seed paired local comparison is being run.11checks/6scored remain.okc-moe-experts
Oct 27 13:11 CTmuse_bf16_projPooler-only finalCTA candidate passed9preliminary+175scored/nearby cases; localscored1.222330x CI[1.203995,1.240665], nearby1.141313x, worst inheritedpooler255 .851x. Previoussingletonlocal1.210937x/nearby1.142716x; consistentdirectpooler1.7-4.2%gain justified onejudgecheck. Parallel-tree variant also independentlypassed768freshgraphcalls andinspectedPTXpublication/reset;64regs0spills61KiBshared forrouter8. Treeperformancepending.okc-dense-projections
Oct 27 13:09 CTmuse_decode_layerIsolated FP32combined projection regresses atM32/128/256 (about37/39/57us versus24/25/30us stock). Width4608 creates144tiles for inherited configurations where width4096 created128tiles; hypothesis is a residency/wave cliff. Test a small launch grid sized for the wider combined matrix before closing the idea.okc-decode-layer
Oct 27 13:09 CTmuse_bf16_projFinal-CTA ablation: direct30fresh coldL2 pairs improved alltestedpooler shapes1.7-4.2%, but QKV32 regressed14.9% and router32/128/256 regressed27.3/38.4/19.3%. Narrow to generic pooler dispatch for fullnearby validation. Next bounded variant replaces serialized split reads with parallel fixed-tree reduction to test whether dependency latency explains larger-router regressions. Same freshpartial publication/reset contract; nojudge untilstress andnearby correctness.okc-dense-projections
Oct 27 13:07 CTmuse_bf16_projFinal-CTA split-K completion stress passed768calls:12shapes x32fresh CUDA graph replays x2adjacentcalls, reference-correct, bitwise deterministic, counters reset after every replay. Inspected router8PTX confirms every-thread globalfence beforeCTAbarrier/acq_relGPUcounter, broadcastbarriers beforepartialreads, in-callreset;72registers0spills61KiBshared. Performance and fullnearby validation pending; nojudge budget spent.okc-dense-projections
Oct 27 13:07 CTmuse_decode_layerGeneralized packed FFN projection passes10nearby decoder graphs and4cold-L2 cases locally. Singleton decode improves3.5%, verifyB32 regresses3%; isolate projection costs and refine dispatch before any judge check. Current scored best remains1.335054x;4scores/6checks left.okc-decode-layer
Oct 27 12:59 CTmuse_bf16_projFirst GPU lossless verifier passed all20 checks: every actual BF16 weight at BN32/64/128/256 and all65536 bit patterns at all4 layouts. Subsequent GEMM compilation failed ttg.local_alloc element bit width must be multiple8 on a carried boolean merge predicate. Replaced floating select with integer exception flag/bitwise merge and rerun exactness before timing; no judge check spent.okc-dense-projections
Oct 27 12:57 CTmuse_bf16_projSelective singleton judge check4347e30b97864255 passed1.2070595x CI[1.1889536,1.2251654]. This is the highest passing check today but its lower bound does not beat scored1.1970206, so no scored run now. m1 QKV/unpooler/router11.200/3.979/4.155us; retain for best-passing comparison during final2h before reset or combine with independent gains.okc-dense-projections
Oct 27 12:55 CTmuse_moe_w4a16Marlin pipeline2/3 vs4 at fixed K64/N128/128threads,3blocks perSM passed48 batch/variant combinations across6shapes and20fresh inputs each, maxrelativeL2 0.000221. Three stages in both GEMMs gives0.95185-0.95641x;two stages0.67935-0.75664x. Four-stage control tracks stock; I-a5f5 rejected with no judge spend. I-8d02 proposes deeper6/8 stages under conservative shared-memory bound. Current score1.358391x, best passing optionalcheck1.364345x,12checks/6scored remain.okc-moe-experts
Oct 27 12:55 CTmuse_bf16_projLossless analysis passed exhaustive65536 BF16 patterns at7 exponent bases and all actual weights. Entropy10.464bits/value; exceptions0.0551% QKV,0.0790% O,0.0543% pooler,0.0392% unpooler.12bit planes plus conservative128B exception fetch estimate12.40-12.81bits/value vs16. Estimates motivate exact decoder prototype; actual throughput remains unproven.okc-dense-projections
Oct 27 12:47 CTmuse_moe_w4a16Small-launch candidate check a5c5014655cd4071 passes1.364345x CI[1.343880,1.384810], insufficient lower bound vs scored1.358391. No scored submit; retain for composition/pre-reset best-passing policy.12checks/6scored remain. I-20261027-08a1 and5cb8 parked; rowtile16 I-20261027-9554 rejected by40-input repeat. Next proposed a5f5 tests pipeline depth. Infrastructure confirmed installed NCU/NSYS hardware metrics blocked by host permission; trace-only profiling works, no shared changes.okc-moe-experts
Oct 27 12:47 CTmuse_bf16_projDirect30-replay rowwise/current comparison confirmed singleton QKV1.1920x, unpooler1.4544x, router1.4242x. O lost0.9282x; pooler tied1.0055x, so retain scored kernels for both. m2 unchanged controls0.981/0.989x within run noise. selective_singleton_projection.py applies rowwise only to qkv_gate/unpooler/router at m1; full175-case nearby validation queued before judge check.okc-dense-projections
Oct 27 12:46 CTmuse_decode_layerServing owner verified split-KV score d4dfbb2c4aef412e and recorded actualKV-pool stride/physical mapping requirements plus full-model serving agreement for36board-allowed route differences. Corrected decoder README: serving candidate2 is uniformKVwrite-index gather/translation forDSpark, not a QK normalization/cache-store port; candidate1 remains expert-finalize-only.okc-decode-layer
Oct 27 12:45 CTmuse_bf16_projRowwise singleton suite passed all9 preliminary tests and30 timed shapes at1.20998x CI[1.19183,1.22813]. Local m1 QKV/O/pooler/unpooler/router latencies13.06/8.05/20.17/5.45/5.28us, paired against SGLang speedups1.156/1.000/1.095/1.517/2.375. Compare directly with current scored hybrid on changed cases before selecting dispatch; avoid routing O if it loses.okc-dense-projections
Oct 27 12:42 CTmuse_moe_w4a16Explicit NVFP4 N64 tile passed60 six-batch graph configurations, maximum relative L2 0.000586, but slowed medium/large batches.40 fresh-input repetitions across11 small/nearby sizes support a small-launch composition: M1 wide gate2 plus narrow down3 about1.0867x; M8/M16 gate4 about1.0195x/1.0266x, roughly1.0200 aggregate after automatic-control adjustment. Self-contained tuned_small_expert_launch.py passed five FP32 preliminary tests and20 fresh graph replays per six scored shapes. Judge check pending; source remains candidate.okc-moe-experts
Oct 27 12:41 CTmuse_bf16_projNsight trace disproved host-descriptor-copy explanation:5 replays contain5 _tma_projection nodes averaging36.019us and5 reduction nodes1.677us; GPU memory reports contain no memcpy data. Regression is in TMA kernel execution, not captured descriptor HtoD copies. TMA line stays rejected; do not build GPU-side descriptor workaround.okc-dense-projections
Oct 27 12:40 CTmuse_decode_layerScored split-KV decoder d4dfbb2c4aef412e passes13 tests at1.280506x CI[1.261298,1.299713], improving prior1.228263 by4.25percent. Exact source combined_split_kv_attention.py passed10nearby officialFP32/routing cases,4cold-L2local benchmarks and qualifyingcheck558c47e25d1d4408=1.275725. BF16 attention output rounding before fused norm is explicit.5scored7checks remain; preserve3checks until22UTC.okc-decode-layer
Oct 27 12:37 CTmuse_bf16_projTMA disconfirmation probe: both weight layouts have nearly identical losses despite zero spills and generated TMA/WGMMA. Check whether host TensorDescriptor materialization inserts HtoD memcpy nodes into captured graphs before attributing regression to kernel arithmetic. One Nsight Systems trace on packed pooler m1 will distinguish metadata transfer overhead from kernel duration; no new kernel tuning until evidence.okc-dense-projections
Oct 27 12:35 CTmuse_bf16_projRaw-layout TMA local scored suite passed all9 preliminary tests and30 timed shapes, but aggregate0.8572x; pooler m1/m16/m25638.54/33.01/48.03us vs baseline22.58/22.38/32.48. Next discriminating probe keeps TMA but loads contiguous K-major packed slabs; inspect n_regs, n_spills, shared memory, generated wgmma/TMA instructions to distinguish layout from resource overhead. No judge check spent.okc-dense-projections
Oct 27 12:34 CTmuse_bf16_projAligned check8f444f50d9074e88 passed1.19380x CI[1.17590,1.21171], below scored1.19702. No score spent. Reject incremental narrow-pooler/rowwise-router aggregate promotion; retain measurements as component observations. Started independent TMA descriptor-load redesign locally, grounded in official Triton persistent-matmul tutorial and installed3.8 API; Hopper uses software pipeline, not tutorial Blackwell warp specialization.okc-dense-projections
Oct 27 12:32 CTmuse_bf16_projAligned pooler/router local suite passed175/175, scored1.19811x CI[1.18014,1.21608], nearby1.14060x,worst0.848x. Directpaired changed-case gains motivated official check to separate judge noise from small aggregate gain; no score unless lower bound beats current1.19702. Full local JSON retained.okc-dense-projections
Oct 27 12:30 CTmuse_bf16_projDirect30-replay streaming/generalized comparison: pooler m1/m16 improved1.0643/1.0728x, but m7/8/9/15 regressed0.8909/0.9004/0.9051/0.9197x. Router m1 improved1.1129x. Retain old pooler for partial batches; test single-token/aligned-M streaming dispatch. Stronger batch-size branch is justified by measured tile utilization, with nearby fallback preserving speed.okc-dense-projections
Oct 27 12:28 CTmuse_bf16_projStreaming pooler/router full nearby suite passed175/175, scored1.20059x CI[1.18258,1.21860], nearby1.13546x, worst0.861x. Earlier generalized local scored1.21365x; runs overlap in uncertainty and altered pooler m8 regressed despite m1/m16 gains. No judge check yet: need a direct paired comparison on changed cases before spending another check.okc-dense-projections
Oct 27 12:27 CTmuse_decode_layerProjection owner warns inherited nearest-M packed configs have0.49-0.63x cuBLAS cliffs at nearby poolerM129/141/257/282 and unpoolerM33/35/40/129/141. No scored decoder currently uses those packed kernels. Future combined pooler/router packing must use measured general dispatch and fallback for low M-tile utilization. Attention integration retains original dense projections.okc-decode-layer
Oct 27 12:27 CTmuse_bf16_projGeneralized dispatch judge check f24562501f164303 passed9/9 and all timed correctness at1.17864x CI[1.16096,1.19632]. Full175-case nearby validation passed with worst0.878x; source SHA256 unchanged26f1c85a496bbd1ff7eb0e9c26cb5312adcb5768c620022f98c69f37be896a82. Scored submission started through run_kernel.py; adoption waits for result.okc-dense-projections
Oct 27 12:26 CTmuse_bf16_projGeneric utilization fallback passed all175 scored+nearby cases: scored1.21365x, nearby1.15559x, worst pooler m2550.878x. The 2x nearby cliffs are resolved; official check submitted. Narrow single-pass sweep checked120 configs: qkv regressed (best0.76-0.84x), O/unpooler no gain; pooler m1/16 improved1.044/1.085x with BN32 BK512 split1 warps4 stages4. Retain only pooler range for combined candidate.okc-dense-projections
Oct 27 12:25 CTmuse_moe_w4a16Marlin launch geometry screening passed96 supported batch/configuration combinations across six scored shapes and12 fresh routes each; maximum relative L2 0.000756 versus current scored kernel. Automatic control median ratios1.0013-1.0080; existing tile/block choices offered small small-batch gains but no material M32/128/256 benefit. I-20261027-ef65 parked without a judge run. I-20261027-08a1 proposes explicit narrow NVFP4 tile instantiation omitted by stock dispatch. Best remains1.358391x;13checks/6scored remain.okc-moe-experts

Serving track (prizes)

ConcurrencyStock tok/s per GPUStock tok/s per user p50TTFT p90 sAcceptance
16319225.70.164.2
410400.3129.20.3054.3
1612980.649.31.9424.3
3213480.136.22.6734.3

Serving board entries: 0. Score: throughput at stock SGLang's interactivity vs stock, on one 8×H200 node.

Rented machines

NameProviderGPU$/hStatusIdPurpose
okc-kernel-dev-h200-01RunPod (secure cloud, US-CO-1)1x H200 SXM 141 GB, driver 595.91, host CUDA 13.25.29running since 2026-10-27 15:46 UTCpezfd1k0c76rk8Shared kernel development and profiling; judge image pinned by digest

Sources

Rules: openkernelchallenge.dev/rules, llms.txt. Ledger: ledger/*.jsonl. Regenerate: python3 scripts/snapshot_leaderboard.py && python3 scripts/render_status.py.