From de8993f82af07e7fb9025ae6e506f2db14851266 Mon Sep 17 00:00:00 2001 From: alay2shah Date: Fri, 31 Jul 2026 17:13:48 +0000 Subject: [PATCH] feat: add LFM2 model family support --- Makefile | 2 + README.md | 5 + benchmark/data/all_benchmark_data.csv | 336 ++++++++++++++++++ .../scripts/benchmark_lfm2_moe_router.py | 116 ++++++ .../scripts/benchmark_lfm2_short_conv.py | 103 ++++++ benchmark/scripts/benchmark_model_configs.py | 29 ++ docs/High-Level-APIs.md | 21 ++ src/liger_kernel/ops/__init__.py | 2 + src/liger_kernel/ops/lfm2_moe_router.py | 170 +++++++++ src/liger_kernel/ops/lfm2_short_conv.py | 264 ++++++++++++++ src/liger_kernel/transformers/__init__.py | 13 + .../transformers/lfm2_moe_router.py | 13 + .../transformers/lfm2_short_conv.py | 27 ++ .../transformers/model/lfm2_vl.py | 92 +++++ .../transformers/model/output_classes.py | 15 + src/liger_kernel/transformers/monkey_patch.py | 227 ++++++++++++ src/liger_kernel/transformers/swiglu.py | 75 ++++ test/convergence/bf16/test_lfm2_models.py | 25 ++ test/convergence/fp32/test_lfm2_models.py | 25 ++ test/convergence/lfm2_utils.py | 167 +++++++++ test/transformers/test_lfm2.py | 253 +++++++++++++ test/transformers/test_lfm2_moe_dispatch.py | 100 ++++++ test/transformers/test_lfm2_moe_router.py | 55 +++ test/transformers/test_lfm2_short_conv.py | 102 ++++++ 24 files changed, 2237 insertions(+) create mode 100644 benchmark/scripts/benchmark_lfm2_moe_router.py create mode 100644 benchmark/scripts/benchmark_lfm2_short_conv.py create mode 100644 src/liger_kernel/ops/lfm2_moe_router.py create mode 100644 src/liger_kernel/ops/lfm2_short_conv.py create mode 100644 src/liger_kernel/transformers/lfm2_moe_router.py create mode 100644 src/liger_kernel/transformers/lfm2_short_conv.py create mode 100644 src/liger_kernel/transformers/model/lfm2_vl.py create mode 100644 test/convergence/bf16/test_lfm2_models.py create mode 100644 test/convergence/fp32/test_lfm2_models.py create mode 100644 test/convergence/lfm2_utils.py create mode 100644 test/transformers/test_lfm2.py create mode 100644 test/transformers/test_lfm2_moe_dispatch.py create mode 100644 test/transformers/test_lfm2_moe_router.py create mode 100644 test/transformers/test_lfm2_short_conv.py diff --git a/Makefile b/Makefile index 713e3771a..7826dac37 100644 --- a/Makefile +++ b/Makefile @@ -38,6 +38,8 @@ checkstyle: # We have to explicitly set HF_DATASETS_OFFLINE=1, or dataset will silently try to send metrics and timeout (80s) https://github.com/huggingface/datasets/blob/37a603679f451826cfafd8aae00738b01dcb9d58/src/datasets/load.py#L286 test-convergence: HF_DATASETS_OFFLINE=1 python -m pytest --disable-warnings \ + test/convergence/fp32/test_lfm2_models.py \ + test/convergence/bf16/test_lfm2_models.py \ test/convergence/fp32/test_mini_models.py \ test/convergence/fp32/test_mini_models_multimodal.py \ test/convergence/fp32/test_mini_models_with_logits.py \ diff --git a/README.md b/README.md index fecc68518..78ffe16e4 100644 --- a/README.md +++ b/README.md @@ -297,6 +297,9 @@ loss.backward() | Llama4 (Text) & (Multimodal) | `liger_kernel.transformers.apply_liger_kernel_to_llama4` | RMSNorm, LayerNorm, GeGLU, CrossEntropyLoss, FusedLinearCrossEntropy | | LLaMA 2 & 3 | `liger_kernel.transformers.apply_liger_kernel_to_llama` | RoPE, RMSNorm, SwiGLU, CrossEntropyLoss, FusedLinearCrossEntropy | | LLaMA 3.2-Vision | `liger_kernel.transformers.apply_liger_kernel_to_mllama` | RoPE, RMSNorm, SwiGLU, CrossEntropyLoss, FusedLinearCrossEntropy | +| LFM2 | `liger_kernel.transformers.apply_liger_kernel_to_lfm2` | RoPE, RMSNorm, SwiGLU, ShortConv, CrossEntropyLoss, FusedLinearCrossEntropy | +| LFM2MoE | `liger_kernel.transformers.apply_liger_kernel_to_lfm2_moe` | RoPE, RMSNorm, SwiGLU, ShortConv, FusedMoE, MoERouter, CrossEntropyLoss, FusedLinearCrossEntropy | +| LFM2VL | `liger_kernel.transformers.apply_liger_kernel_to_lfm2_vl` | SigLIP2 LayerNorm, RoPE, RMSNorm, SwiGLU, ShortConv, CrossEntropyLoss, FusedLinearCrossEntropy | | Ministral | `liger_kernel.transformers.apply_liger_kernel_to_ministral` | RoPE, RMSNorm, SwiGLU, CrossEntropyLoss, FusedLinearCrossEntropy | | Mistral | `liger_kernel.transformers.apply_liger_kernel_to_mistral` | RoPE, RMSNorm, SwiGLU, CrossEntropyLoss, FusedLinearCrossEntropy | | Mixtral | `liger_kernel.transformers.apply_liger_kernel_to_mixtral` | RoPE, RMSNorm, SwiGLU, CrossEntropyLoss, FusedLinearCrossEntropy | @@ -346,6 +349,8 @@ loss.backward() | CrossEntropy | `liger_kernel.transformers.LigerCrossEntropyLoss` | | Fused Linear CrossEntropy | `liger_kernel.transformers.LigerFusedLinearCrossEntropyLoss`| | Multi Token Attention | `liger_kernel.transformers.LigerMultiTokenAttention` | +| LFM2 Short Convolution | `liger_kernel.ops.LigerLfm2ShortConvFunction` | +| LFM2 MoE Router | `liger_kernel.ops.LigerLfm2MoeRouterFunction` | | Softmax | `liger_kernel.transformers.LigerSoftmax` | | Sparsemax | `liger_kernel.transformers.LigerSparsemax` | | mHC (Hyper-Connections) | `liger_kernel.transformers.LigerMHC` | diff --git a/benchmark/data/all_benchmark_data.csv b/benchmark/data/all_benchmark_data.csv index 84e73583c..8d1074d14 100644 --- a/benchmark/data/all_benchmark_data.csv +++ b/benchmark/data/all_benchmark_data.csv @@ -2351,3 +2351,339 @@ megatron_cross_entropy,megatron-unfused,full,memory,MB,V,vocab size,16384,1024.1 megatron_cross_entropy,megatron-unfused,full,memory,MB,V,vocab size,32768,2048.16552734375,2048.16552734375,2048.16552734375,"{""S"": 2048, ""B"": 4, ""TP"": 1}",NVIDIA H100 80GB HBM3,2026-06-15 22:13:03,0.8.0 megatron_cross_entropy,megatron-unfused,full,memory,MB,V,vocab size,65536,4096.16552734375,4096.16552734375,4096.16552734375,"{""S"": 2048, ""B"": 4, ""TP"": 1}",NVIDIA H100 80GB HBM3,2026-06-15 22:13:03,0.8.0 megatron_cross_entropy,megatron-unfused,full,memory,MB,V,vocab size,131072,8192.166015625,8192.166015625,8192.166015625,"{""S"": 2048, ""B"": 4, ""TP"": 1}",NVIDIA H100 80GB HBM3,2026-06-15 22:13:03,0.8.0 +lfm2_moe_router,huggingface,forward,speed,ms,T,total tokens,1024,0.053979501128196716,0.05314320027828216,0.05571899935603142,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:15,0.8.1 +lfm2_moe_router,huggingface,forward,speed,ms,T,total tokens,2048,0.060139000415802,0.05831899866461754,0.06150299981236458,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:15,0.8.1 +lfm2_moe_router,huggingface,forward,speed,ms,T,total tokens,4096,0.06907899677753448,0.0682390034198761,0.06971900165081024,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:15,0.8.1 +lfm2_moe_router,huggingface,forward,speed,ms,T,total tokens,8192,0.08307799696922302,0.08195800334215164,0.08463899791240692,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:15,0.8.1 +lfm2_moe_router,huggingface,forward,speed,ms,T,total tokens,16384,0.1083579994738102,0.10691040158271789,0.10989400148391723,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:15,0.8.1 +lfm2_moe_router,huggingface,forward,speed,ms,T,total tokens,32768,0.1688780039548874,0.16687700152397156,0.17072499692440032,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:15,0.8.1 +lfm2_moe_router,huggingface,forward,speed,ms,T,total tokens,65536,0.27623599767684937,0.27567601203918457,0.2779960036277771,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:15,0.8.1 +lfm2_moe_router,liger,forward,speed,ms,T,total tokens,1024,0.00803999975323677,0.007599399983882904,0.008223400078713895,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:16,0.8.1 +lfm2_moe_router,liger,forward,speed,ms,T,total tokens,2048,0.010080000385642052,0.009840000420808792,0.010400000028312206,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:16,0.8.1 +lfm2_moe_router,liger,forward,speed,ms,T,total tokens,4096,0.014960000291466713,0.014879000373184681,0.015127799659967424,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:16,0.8.1 +lfm2_moe_router,liger,forward,speed,ms,T,total tokens,8192,0.024560000747442245,0.02443999983370304,0.02476000040769577,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:16,0.8.1 +lfm2_moe_router,liger,forward,speed,ms,T,total tokens,16384,0.044139498844742775,0.04405559822916985,0.04427900165319443,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:16,0.8.1 +lfm2_moe_router,liger,forward,speed,ms,T,total tokens,32768,0.08317850157618523,0.08307900279760361,0.08331900089979172,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:16,0.8.1 +lfm2_moe_router,liger,forward,speed,ms,T,total tokens,65536,0.16099700331687927,0.1609169989824295,0.16109400391578677,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:16,0.8.1 +lfm2_moe_router,huggingface,backward,speed,ms,T,total tokens,1024,0.03919899836182594,0.03838300108909607,0.04095980077981949,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:16,0.8.1 +lfm2_moe_router,huggingface,backward,speed,ms,T,total tokens,2048,0.03967899829149246,0.03935899958014488,0.040430998802185064,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:16,0.8.1 +lfm2_moe_router,huggingface,backward,speed,ms,T,total tokens,4096,0.04001949913799763,0.039719000458717346,0.040663999319076535,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:16,0.8.1 +lfm2_moe_router,huggingface,backward,speed,ms,T,total tokens,8192,0.04135899990797043,0.04095999896526337,0.04215899854898453,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:16,0.8.1 +lfm2_moe_router,huggingface,backward,speed,ms,T,total tokens,16384,0.042559001594781876,0.04222339913249015,0.04293539896607399,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:16,0.8.1 +lfm2_moe_router,huggingface,backward,speed,ms,T,total tokens,32768,0.044078998267650604,0.043775601685047155,0.04699940010905266,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:16,0.8.1 +lfm2_moe_router,huggingface,backward,speed,ms,T,total tokens,65536,0.04875899851322174,0.048239000141620636,0.049254998564720154,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:16,0.8.1 +lfm2_moe_router,liger,backward,speed,ms,T,total tokens,1024,0.009560000151395798,0.009320000186562538,0.009720000438392162,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:16,0.8.1 +lfm2_moe_router,liger,backward,speed,ms,T,total tokens,2048,0.009879999794065952,0.009704000316560268,0.010119999758899212,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:16,0.8.1 +lfm2_moe_router,liger,backward,speed,ms,T,total tokens,4096,0.01104000024497509,0.01083999965339899,0.01131999958306551,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:16,0.8.1 +lfm2_moe_router,liger,backward,speed,ms,T,total tokens,8192,0.013240000233054161,0.013000000268220901,0.01348000019788742,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:16,0.8.1 +lfm2_moe_router,liger,backward,speed,ms,T,total tokens,16384,0.0175199992954731,0.017358999699354172,0.017880000174045563,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:16,0.8.1 +lfm2_moe_router,liger,backward,speed,ms,T,total tokens,32768,0.026079000905156136,0.025999000295996666,0.026159999892115593,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:16,0.8.1 +lfm2_moe_router,liger,backward,speed,ms,T,total tokens,65536,0.04292000085115433,0.04284000024199486,0.04335999861359596,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:16,0.8.1 +lfm2_moe_router,huggingface,full,speed,ms,T,total tokens,1024,0.09571900218725204,0.0934229999780655,0.10327040255069733,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:16,0.8.1 +lfm2_moe_router,huggingface,full,speed,ms,T,total tokens,2048,0.09723900258541107,0.09541460126638412,0.10228640139102936,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:16,0.8.1 +lfm2_moe_router,huggingface,full,speed,ms,T,total tokens,4096,0.10699799656867981,0.10590999871492385,0.11059020310640336,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:16,0.8.1 +lfm2_moe_router,huggingface,full,speed,ms,T,total tokens,8192,0.12139800190925598,0.12075799703598022,0.12247899919748306,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:16,0.8.1 +lfm2_moe_router,huggingface,full,speed,ms,T,total tokens,16384,0.14973700046539307,0.1477411985397339,0.15191720128059388,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:16,0.8.1 +lfm2_moe_router,huggingface,full,speed,ms,T,total tokens,32768,0.21223650127649307,0.21073299646377563,0.21379700005054475,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:16,0.8.1 +lfm2_moe_router,huggingface,full,speed,ms,T,total tokens,65536,0.32883499562740326,0.32811498641967773,0.3311150074005127,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:16,0.8.1 +lfm2_moe_router,liger,full,speed,ms,T,total tokens,1024,0.0386389996856451,0.037535198777914044,0.04092000052332878,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,liger,full,speed,ms,T,total tokens,2048,0.03901950083673,0.0369190014898777,0.041078999638557434,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,liger,full,speed,ms,T,total tokens,4096,0.03811950050294399,0.036559998989105225,0.03995899856090546,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,liger,full,speed,ms,T,total tokens,8192,0.038359999656677246,0.03736000135540962,0.041159000247716904,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,liger,full,speed,ms,T,total tokens,16384,0.05995900183916092,0.05979900062084198,0.06019899994134903,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,liger,full,speed,ms,T,total tokens,32768,0.11087799817323685,0.11059859693050386,0.11107800155878067,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,liger,full,speed,ms,T,total tokens,65536,0.2078769952058792,0.20748439729213713,0.2081963986158371,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,huggingface,forward,memory,MB,T,total tokens,1024,0.30712890625,0.30712890625,0.30712890625,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,huggingface,forward,memory,MB,T,total tokens,2048,0.61376953125,0.61376953125,0.61376953125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,huggingface,forward,memory,MB,T,total tokens,4096,1.22705078125,1.22705078125,1.22705078125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,huggingface,forward,memory,MB,T,total tokens,8192,2.45361328125,2.45361328125,2.45361328125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,huggingface,forward,memory,MB,T,total tokens,16384,4.90673828125,4.90673828125,4.90673828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,huggingface,forward,memory,MB,T,total tokens,32768,9.81298828125,9.81298828125,9.81298828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,huggingface,forward,memory,MB,T,total tokens,65536,19.62548828125,19.62548828125,19.62548828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,liger,forward,memory,MB,T,total tokens,1024,0.08642578125,0.08642578125,0.08642578125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,liger,forward,memory,MB,T,total tokens,2048,0.17236328125,0.17236328125,0.17236328125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,liger,forward,memory,MB,T,total tokens,4096,0.34423828125,0.34423828125,0.34423828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,liger,forward,memory,MB,T,total tokens,8192,0.68798828125,0.68798828125,0.68798828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,liger,forward,memory,MB,T,total tokens,16384,1.37548828125,1.37548828125,1.37548828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,liger,forward,memory,MB,T,total tokens,32768,2.75048828125,2.75048828125,2.75048828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,liger,forward,memory,MB,T,total tokens,65536,5.50048828125,5.50048828125,5.50048828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,huggingface,backward,memory,MB,T,total tokens,1024,0.36962890625,0.36962890625,0.36962890625,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,huggingface,backward,memory,MB,T,total tokens,2048,0.73876953125,0.73876953125,0.73876953125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,huggingface,backward,memory,MB,T,total tokens,4096,1.47705078125,1.47705078125,1.47705078125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,huggingface,backward,memory,MB,T,total tokens,8192,2.95361328125,2.95361328125,2.95361328125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,huggingface,backward,memory,MB,T,total tokens,16384,5.90673828125,5.90673828125,5.90673828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,huggingface,backward,memory,MB,T,total tokens,32768,11.81298828125,11.81298828125,11.81298828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,huggingface,backward,memory,MB,T,total tokens,65536,23.62548828125,23.62548828125,23.62548828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,liger,backward,memory,MB,T,total tokens,1024,0.23486328125,0.23486328125,0.23486328125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,liger,backward,memory,MB,T,total tokens,2048,0.46923828125,0.46923828125,0.46923828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,liger,backward,memory,MB,T,total tokens,4096,0.93798828125,0.93798828125,0.93798828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,liger,backward,memory,MB,T,total tokens,8192,1.87548828125,1.87548828125,1.87548828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,liger,backward,memory,MB,T,total tokens,16384,3.75048828125,3.75048828125,3.75048828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,liger,backward,memory,MB,T,total tokens,32768,7.50048828125,7.50048828125,7.50048828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,liger,backward,memory,MB,T,total tokens,65536,15.00048828125,15.00048828125,15.00048828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,huggingface,full,memory,MB,T,total tokens,1024,0.36962890625,0.36962890625,0.36962890625,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,huggingface,full,memory,MB,T,total tokens,2048,0.73876953125,0.73876953125,0.73876953125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,huggingface,full,memory,MB,T,total tokens,4096,1.47705078125,1.47705078125,1.47705078125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,huggingface,full,memory,MB,T,total tokens,8192,2.95361328125,2.95361328125,2.95361328125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,huggingface,full,memory,MB,T,total tokens,16384,5.90673828125,5.90673828125,5.90673828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,huggingface,full,memory,MB,T,total tokens,32768,11.81298828125,11.81298828125,11.81298828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,huggingface,full,memory,MB,T,total tokens,65536,23.62548828125,23.62548828125,23.62548828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,liger,full,memory,MB,T,total tokens,1024,0.23486328125,0.23486328125,0.23486328125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,liger,full,memory,MB,T,total tokens,2048,0.46923828125,0.46923828125,0.46923828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,liger,full,memory,MB,T,total tokens,4096,0.93798828125,0.93798828125,0.93798828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,liger,full,memory,MB,T,total tokens,8192,1.87548828125,1.87548828125,1.87548828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,liger,full,memory,MB,T,total tokens,16384,3.75048828125,3.75048828125,3.75048828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,liger,full,memory,MB,T,total tokens,32768,7.50048828125,7.50048828125,7.50048828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_moe_router,liger,full,memory,MB,T,total tokens,65536,15.00048828125,15.00048828125,15.00048828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:13:17,0.8.1 +lfm2_short_conv,huggingface,forward,speed,ms,T,total tokens,1024,0.06215900182723999,0.061183000355958944,0.0631429985165596,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:36,0.8.1 +lfm2_short_conv,huggingface,forward,speed,ms,T,total tokens,2048,0.0918779969215393,0.09130199849605561,0.09283900260925293,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:36,0.8.1 +lfm2_short_conv,huggingface,forward,speed,ms,T,total tokens,4096,0.17417750507593155,0.17319700121879578,0.17527799308300018,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:36,0.8.1 +lfm2_short_conv,huggingface,forward,speed,ms,T,total tokens,8192,0.3145560026168823,0.3133392035961151,0.3158598005771637,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:36,0.8.1 +lfm2_short_conv,huggingface,forward,speed,ms,T,total tokens,16384,0.6063510179519653,0.6035110116004945,0.6084389805793762,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:36,0.8.1 +lfm2_short_conv,huggingface,forward,speed,ms,T,total tokens,32768,1.5736589431762695,1.5710920095443726,1.591739797592163,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:36,0.8.1 +lfm2_short_conv,huggingface,forward,speed,ms,T,total tokens,65536,4.47641396522522,4.46810998916626,4.484717941284179,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:36,0.8.1 +lfm2_short_conv,liger,forward,speed,ms,T,total tokens,1024,0.01023900043219328,0.009624000079929829,0.010479999706149101,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:36,0.8.1 +lfm2_short_conv,liger,forward,speed,ms,T,total tokens,2048,0.014039999805390835,0.013880000449717045,0.014279999770224094,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:36,0.8.1 +lfm2_short_conv,liger,forward,speed,ms,T,total tokens,4096,0.024399999529123306,0.024080000817775726,0.02507999911904335,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:36,0.8.1 +lfm2_short_conv,liger,forward,speed,ms,T,total tokens,8192,0.05279900133609772,0.05218299925327301,0.054463401436805725,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:36,0.8.1 +lfm2_short_conv,liger,forward,speed,ms,T,total tokens,16384,0.10221900045871735,0.10031079798936844,0.10519079864025116,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:36,0.8.1 +lfm2_short_conv,liger,forward,speed,ms,T,total tokens,32768,0.18731699883937836,0.1867497980594635,0.19035700261592864,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:36,0.8.1 +lfm2_short_conv,liger,forward,speed,ms,T,total tokens,65536,0.3651140034198761,0.36428300142288206,0.36769100427627566,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:36,0.8.1 +lfm2_short_conv,huggingface,backward,speed,ms,T,total tokens,1024,0.22743700444698334,0.22743700444698334,0.22743700444698334,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:37,0.8.1 +lfm2_short_conv,huggingface,backward,speed,ms,T,total tokens,2048,0.3773549944162369,0.3758019983768463,0.378338211774826,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:37,0.8.1 +lfm2_short_conv,huggingface,backward,speed,ms,T,total tokens,4096,0.6660299897193909,0.6653100252151489,0.6675981998443603,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:37,0.8.1 +lfm2_short_conv,huggingface,backward,speed,ms,T,total tokens,8192,1.2021019458770752,1.1990219831466675,1.2048539638519287,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:37,0.8.1 +lfm2_short_conv,huggingface,backward,speed,ms,T,total tokens,16384,2.5534420013427734,2.5483539581298826,2.5573305606842043,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:37,0.8.1 +lfm2_short_conv,huggingface,backward,speed,ms,T,total tokens,32768,6.396505832672119,6.396505832672119,6.396505832672119,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:37,0.8.1 +lfm2_short_conv,huggingface,backward,speed,ms,T,total tokens,65536,15.992445945739746,15.992445945739746,15.992445945739746,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:37,0.8.1 +lfm2_short_conv,liger,backward,speed,ms,T,total tokens,1024,0.1125979982316494,0.11194239854812622,0.1136139988899231,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:37,0.8.1 +lfm2_short_conv,liger,backward,speed,ms,T,total tokens,2048,0.2088169977068901,0.20767700672149658,0.21063700318336487,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:37,0.8.1 +lfm2_short_conv,liger,backward,speed,ms,T,total tokens,4096,0.4244140088558197,0.42103400230407717,0.4277619898319245,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:37,0.8.1 +lfm2_short_conv,liger,backward,speed,ms,T,total tokens,8192,0.870107501745224,0.8681949973106384,0.8727152109146119,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:37,0.8.1 +lfm2_short_conv,liger,backward,speed,ms,T,total tokens,16384,1.774533987045288,1.768086051940918,1.7844299554824827,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:37,0.8.1 +lfm2_short_conv,liger,backward,speed,ms,T,total tokens,32768,3.562108039855957,3.5594439983367923,3.5647720813751222,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:37,0.8.1 +lfm2_short_conv,liger,backward,speed,ms,T,total tokens,65536,7.216012954711914,7.216012954711914,7.216012954711914,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:37,0.8.1 +lfm2_short_conv,huggingface,full,speed,ms,T,total tokens,1024,0.2812359929084778,0.2805880010128021,0.2824519991874695,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,huggingface,full,speed,ms,T,total tokens,2048,0.4641930013895035,0.46206499338150026,0.4650890052318573,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,huggingface,full,speed,ms,T,total tokens,4096,0.8196679949760437,0.8157079815864563,0.8214679956436157,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,huggingface,full,speed,ms,T,total tokens,8192,1.532657504081726,1.5320969820022583,1.5331779718399048,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,huggingface,full,speed,ms,T,total tokens,16384,3.1960530281066895,3.184016990661621,3.208089065551758,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,huggingface,full,speed,ms,T,total tokens,32768,7.932443141937256,7.932443141937256,7.932443141937256,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,huggingface,full,speed,ms,T,total tokens,65536,20.603696823120117,20.603696823120117,20.603696823120117,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,liger,full,speed,ms,T,total tokens,1024,0.12961800396442413,0.12866200208663942,0.13386999964714053,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,liger,full,speed,ms,T,total tokens,2048,0.2303764969110489,0.22950820326805116,0.2317721962928772,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,liger,full,speed,ms,T,total tokens,4096,0.46323299407958984,0.46224119067192077,0.46820100545883175,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,liger,full,speed,ms,T,total tokens,8192,0.9439460039138794,0.9423784136772155,0.950626015663147,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,liger,full,speed,ms,T,total tokens,16384,1.9137319922447205,1.9018359899520876,1.9199316024780273,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,liger,full,speed,ms,T,total tokens,32768,3.8448840379714966,3.838200044631958,3.8515680313110354,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,liger,full,speed,ms,T,total tokens,65536,8.056200981140137,8.056200981140137,8.056200981140137,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,huggingface,forward,memory,MB,T,total tokens,1024,24.01953125,24.01953125,24.01953125,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,huggingface,forward,memory,MB,T,total tokens,2048,48.01953125,48.01953125,48.01953125,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,huggingface,forward,memory,MB,T,total tokens,4096,96.01953125,96.01953125,96.01953125,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,huggingface,forward,memory,MB,T,total tokens,8192,192.01953125,192.01953125,192.01953125,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,huggingface,forward,memory,MB,T,total tokens,16384,384.01953125,384.01953125,384.01953125,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,huggingface,forward,memory,MB,T,total tokens,32768,768.01953125,768.01953125,768.01953125,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,huggingface,forward,memory,MB,T,total tokens,65536,1536.01953125,1536.01953125,1536.01953125,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,liger,forward,memory,MB,T,total tokens,1024,16.01171875,16.01171875,16.01171875,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,liger,forward,memory,MB,T,total tokens,2048,32.01171875,32.01171875,32.01171875,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,liger,forward,memory,MB,T,total tokens,4096,64.01171875,64.01171875,64.01171875,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,liger,forward,memory,MB,T,total tokens,8192,128.01171875,128.01171875,128.01171875,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,liger,forward,memory,MB,T,total tokens,16384,256.01171875,256.01171875,256.01171875,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,liger,forward,memory,MB,T,total tokens,32768,512.01171875,512.01171875,512.01171875,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,liger,forward,memory,MB,T,total tokens,65536,1024.01171875,1024.01171875,1024.01171875,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,huggingface,backward,memory,MB,T,total tokens,1024,64.2578125,64.2578125,64.2578125,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,huggingface,backward,memory,MB,T,total tokens,2048,128.2578125,128.2578125,128.2578125,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,huggingface,backward,memory,MB,T,total tokens,4096,256.2578125,256.2578125,256.2578125,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,huggingface,backward,memory,MB,T,total tokens,8192,512.2578125,512.2578125,512.2578125,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,huggingface,backward,memory,MB,T,total tokens,16384,1024.2578125,1024.2578125,1024.2578125,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,huggingface,backward,memory,MB,T,total tokens,32768,2048.2578125,2048.2578125,2048.2578125,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,huggingface,backward,memory,MB,T,total tokens,65536,4096.2578125,4096.2578125,4096.2578125,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,liger,backward,memory,MB,T,total tokens,1024,44.15234375,44.15234375,44.15234375,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,liger,backward,memory,MB,T,total tokens,2048,88.24609375,88.24609375,88.24609375,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,liger,backward,memory,MB,T,total tokens,4096,176.43359375,176.43359375,176.43359375,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,liger,backward,memory,MB,T,total tokens,8192,352.80859375,352.80859375,352.80859375,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,liger,backward,memory,MB,T,total tokens,16384,705.55859375,705.55859375,705.55859375,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,liger,backward,memory,MB,T,total tokens,32768,1411.05859375,1411.05859375,1411.05859375,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,liger,backward,memory,MB,T,total tokens,65536,2822.05859375,2822.05859375,2822.05859375,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,huggingface,full,memory,MB,T,total tokens,1024,64.2578125,64.2578125,64.2578125,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,huggingface,full,memory,MB,T,total tokens,2048,128.2578125,128.2578125,128.2578125,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,huggingface,full,memory,MB,T,total tokens,4096,256.2578125,256.2578125,256.2578125,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,huggingface,full,memory,MB,T,total tokens,8192,512.2578125,512.2578125,512.2578125,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,huggingface,full,memory,MB,T,total tokens,16384,1024.2578125,1024.2578125,1024.2578125,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,huggingface,full,memory,MB,T,total tokens,32768,2048.2578125,2048.2578125,2048.2578125,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,huggingface,full,memory,MB,T,total tokens,65536,4096.2578125,4096.2578125,4096.2578125,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,liger,full,memory,MB,T,total tokens,1024,44.15234375,44.15234375,44.15234375,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,liger,full,memory,MB,T,total tokens,2048,88.24609375,88.24609375,88.24609375,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,liger,full,memory,MB,T,total tokens,4096,176.43359375,176.43359375,176.43359375,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,liger,full,memory,MB,T,total tokens,8192,352.80859375,352.80859375,352.80859375,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,liger,full,memory,MB,T,total tokens,16384,705.55859375,705.55859375,705.55859375,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,liger,full,memory,MB,T,total tokens,32768,1411.05859375,1411.05859375,1411.05859375,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,liger,full,memory,MB,T,total tokens,65536,2822.05859375,2822.05859375,2822.05859375,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",AMD Instinct MI325X,2026-07-23 18:12:38,0.8.1 +lfm2_short_conv,huggingface,forward,speed,ms,T,total tokens,1024,0.15972799807786942,0.15697279572486877,0.1644544005393982,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:09,0.8.1 +lfm2_short_conv,huggingface,forward,speed,ms,T,total tokens,2048,0.13254399597644806,0.1315584033727646,0.13606399297714233,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:09,0.8.1 +lfm2_short_conv,huggingface,forward,speed,ms,T,total tokens,4096,0.24983999878168106,0.24830719828605652,0.2504703938961029,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:09,0.8.1 +lfm2_short_conv,huggingface,forward,speed,ms,T,total tokens,8192,0.4848319888114929,0.4836800038814545,0.48528001308441165,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:09,0.8.1 +lfm2_short_conv,huggingface,forward,speed,ms,T,total tokens,16384,0.9461439847946167,0.9450240135192871,0.9467008113861084,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:09,0.8.1 +lfm2_short_conv,huggingface,forward,speed,ms,T,total tokens,32768,1.892416000366211,1.8912063837051392,1.8925567626953126,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:09,0.8.1 +lfm2_short_conv,huggingface,forward,speed,ms,T,total tokens,65536,3.787503957748413,3.786899137496948,3.788108777999878,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:09,0.8.1 +lfm2_short_conv,liger,forward,speed,ms,T,total tokens,1024,0.14800000190734863,0.12143999934196473,0.17244800329208376,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:13,0.8.1 +lfm2_short_conv,liger,forward,speed,ms,T,total tokens,2048,0.09483199939131737,0.09008640050888062,0.11547520011663437,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:13,0.8.1 +lfm2_short_conv,liger,forward,speed,ms,T,total tokens,4096,0.04203199967741966,0.04073600098490715,0.04380160123109818,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:13,0.8.1 +lfm2_short_conv,liger,forward,speed,ms,T,total tokens,8192,0.05825600028038025,0.05737600103020668,0.059193599969148636,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:13,0.8.1 +lfm2_short_conv,liger,forward,speed,ms,T,total tokens,16384,0.10412800312042236,0.10362879931926727,0.10463999956846237,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:13,0.8.1 +lfm2_short_conv,liger,forward,speed,ms,T,total tokens,32768,0.19734400510787964,0.19701120257377625,0.19779199361801147,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:13,0.8.1 +lfm2_short_conv,liger,forward,speed,ms,T,total tokens,65536,0.382176011800766,0.3816959857940674,0.38284799456596375,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:13,0.8.1 +lfm2_short_conv,huggingface,backward,speed,ms,T,total tokens,1024,1.5607680082321167,1.5607680082321167,1.5607680082321167,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:14,0.8.1 +lfm2_short_conv,huggingface,backward,speed,ms,T,total tokens,2048,0.4657920002937317,0.44010881185531614,1.6382976055145264,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:14,0.8.1 +lfm2_short_conv,huggingface,backward,speed,ms,T,total tokens,4096,0.672080010175705,0.6701568126678468,0.6821439862251282,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:14,0.8.1 +lfm2_short_conv,huggingface,backward,speed,ms,T,total tokens,8192,1.2852800488471985,1.284608006477356,1.287392020225525,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:14,0.8.1 +lfm2_short_conv,huggingface,backward,speed,ms,T,total tokens,16384,2.524224042892456,2.5196735858917236,2.524992084503174,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:14,0.8.1 +lfm2_short_conv,huggingface,backward,speed,ms,T,total tokens,32768,5.169151782989502,5.169151782989502,5.169151782989502,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:14,0.8.1 +lfm2_short_conv,huggingface,backward,speed,ms,T,total tokens,65536,10.381407737731934,10.381407737731934,10.381407737731934,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:14,0.8.1 +lfm2_short_conv,liger,backward,speed,ms,T,total tokens,1024,1.8704479932785034,1.7834688425064087,1.9295231580734256,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:18,0.8.1 +lfm2_short_conv,liger,backward,speed,ms,T,total tokens,2048,1.891584038734436,1.87708158493042,1.936966371536255,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:18,0.8.1 +lfm2_short_conv,liger,backward,speed,ms,T,total tokens,4096,1.9956159591674805,1.9661824464797975,2.1324031352996826,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:18,0.8.1 +lfm2_short_conv,liger,backward,speed,ms,T,total tokens,8192,2.087552070617676,2.025977611541748,2.124409627914429,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:18,0.8.1 +lfm2_short_conv,liger,backward,speed,ms,T,total tokens,16384,1.6147359609603882,1.6130943775177002,1.61496958732605,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:18,0.8.1 +lfm2_short_conv,liger,backward,speed,ms,T,total tokens,32768,2.937359929084778,2.9365247249603272,2.938195133209229,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:18,0.8.1 +lfm2_short_conv,liger,backward,speed,ms,T,total tokens,65536,5.839903831481934,5.839903831481934,5.839903831481934,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:18,0.8.1 +lfm2_short_conv,huggingface,full,speed,ms,T,total tokens,1024,2.374127984046936,2.35118727684021,2.3915390968322754,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:18,0.8.1 +lfm2_short_conv,huggingface,full,speed,ms,T,total tokens,2048,2.3436319828033447,2.3258623600006105,2.356473684310913,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:18,0.8.1 +lfm2_short_conv,huggingface,full,speed,ms,T,total tokens,4096,2.3510559797286987,2.3312000751495363,2.3841855049133303,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:18,0.8.1 +lfm2_short_conv,huggingface,full,speed,ms,T,total tokens,8192,2.440608024597168,2.386067247390747,2.452281618118286,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:18,0.8.1 +lfm2_short_conv,huggingface,full,speed,ms,T,total tokens,16384,3.5435519218444824,3.5418623447418214,3.5452414989471435,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:18,0.8.1 +lfm2_short_conv,huggingface,full,speed,ms,T,total tokens,32768,7.148928165435791,7.148928165435791,7.148928165435791,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:18,0.8.1 +lfm2_short_conv,huggingface,full,speed,ms,T,total tokens,65536,14.624192237854004,14.624192237854004,14.624192237854004,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:18,0.8.1 +lfm2_short_conv,liger,full,speed,ms,T,total tokens,1024,2.1708641052246094,2.1500224590301515,2.202969694137573,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:18,0.8.1 +lfm2_short_conv,liger,full,speed,ms,T,total tokens,2048,2.362048029899597,2.348377561569214,2.3652480125427244,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:18,0.8.1 +lfm2_short_conv,liger,full,speed,ms,T,total tokens,4096,2.50545597076416,2.496895933151245,2.5324096202850344,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:18,0.8.1 +lfm2_short_conv,liger,full,speed,ms,T,total tokens,8192,2.5763840675354004,2.5432640075683595,2.6552192687988283,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:18,0.8.1 +lfm2_short_conv,liger,full,speed,ms,T,total tokens,16384,2.582751989364624,2.5270720481872555,2.6026432514190674,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:18,0.8.1 +lfm2_short_conv,liger,full,speed,ms,T,total tokens,32768,3.2891520261764526,3.2883840560913087,3.289919996261597,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:18,0.8.1 +lfm2_short_conv,liger,full,speed,ms,T,total tokens,65536,6.52236795425415,6.52236795425415,6.52236795425415,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:18,0.8.1 +lfm2_short_conv,huggingface,forward,memory,MB,T,total tokens,1024,24.01953125,24.01953125,24.01953125,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:18,0.8.1 +lfm2_short_conv,huggingface,forward,memory,MB,T,total tokens,2048,48.01953125,48.01953125,48.01953125,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:18,0.8.1 +lfm2_short_conv,huggingface,forward,memory,MB,T,total tokens,4096,96.01953125,96.01953125,96.01953125,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:18,0.8.1 +lfm2_short_conv,huggingface,forward,memory,MB,T,total tokens,8192,192.01953125,192.01953125,192.01953125,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:18,0.8.1 +lfm2_short_conv,huggingface,forward,memory,MB,T,total tokens,16384,384.01953125,384.01953125,384.01953125,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:18,0.8.1 +lfm2_short_conv,huggingface,forward,memory,MB,T,total tokens,32768,768.01953125,768.01953125,768.01953125,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:18,0.8.1 +lfm2_short_conv,huggingface,forward,memory,MB,T,total tokens,65536,1536.01953125,1536.01953125,1536.01953125,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:18,0.8.1 +lfm2_short_conv,liger,forward,memory,MB,T,total tokens,1024,16.01171875,16.01171875,16.01171875,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:19,0.8.1 +lfm2_short_conv,liger,forward,memory,MB,T,total tokens,2048,32.01171875,32.01171875,32.01171875,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:19,0.8.1 +lfm2_short_conv,liger,forward,memory,MB,T,total tokens,4096,64.01171875,64.01171875,64.01171875,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:19,0.8.1 +lfm2_short_conv,liger,forward,memory,MB,T,total tokens,8192,128.01171875,128.01171875,128.01171875,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:19,0.8.1 +lfm2_short_conv,liger,forward,memory,MB,T,total tokens,16384,256.01171875,256.01171875,256.01171875,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:19,0.8.1 +lfm2_short_conv,liger,forward,memory,MB,T,total tokens,32768,512.01171875,512.01171875,512.01171875,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:19,0.8.1 +lfm2_short_conv,liger,forward,memory,MB,T,total tokens,65536,1024.01171875,1024.01171875,1024.01171875,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:19,0.8.1 +lfm2_short_conv,huggingface,backward,memory,MB,T,total tokens,1024,64.0390625,64.0390625,64.0390625,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:19,0.8.1 +lfm2_short_conv,huggingface,backward,memory,MB,T,total tokens,2048,128.03125,128.03125,128.03125,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:19,0.8.1 +lfm2_short_conv,huggingface,backward,memory,MB,T,total tokens,4096,256.0390625,256.0390625,256.0390625,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:19,0.8.1 +lfm2_short_conv,huggingface,backward,memory,MB,T,total tokens,8192,512.0390625,512.0390625,512.0390625,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:19,0.8.1 +lfm2_short_conv,huggingface,backward,memory,MB,T,total tokens,16384,1024.0390625,1024.0390625,1024.0390625,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:19,0.8.1 +lfm2_short_conv,huggingface,backward,memory,MB,T,total tokens,32768,2048.03125,2048.03125,2048.03125,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:19,0.8.1 +lfm2_short_conv,huggingface,backward,memory,MB,T,total tokens,65536,4096.03125,4096.03125,4096.03125,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:19,0.8.1 +lfm2_short_conv,liger,backward,memory,MB,T,total tokens,1024,44.15234375,44.15234375,44.15234375,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:19,0.8.1 +lfm2_short_conv,liger,backward,memory,MB,T,total tokens,2048,88.24609375,88.24609375,88.24609375,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:19,0.8.1 +lfm2_short_conv,liger,backward,memory,MB,T,total tokens,4096,176.43359375,176.43359375,176.43359375,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:19,0.8.1 +lfm2_short_conv,liger,backward,memory,MB,T,total tokens,8192,352.80859375,352.80859375,352.80859375,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:19,0.8.1 +lfm2_short_conv,liger,backward,memory,MB,T,total tokens,16384,707.05859375,707.05859375,707.05859375,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:19,0.8.1 +lfm2_short_conv,liger,backward,memory,MB,T,total tokens,32768,1420.05859375,1420.05859375,1420.05859375,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:19,0.8.1 +lfm2_short_conv,liger,backward,memory,MB,T,total tokens,65536,2840.05859375,2840.05859375,2840.05859375,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:19,0.8.1 +lfm2_short_conv,huggingface,full,memory,MB,T,total tokens,1024,64.03125,64.03125,64.03125,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:19,0.8.1 +lfm2_short_conv,huggingface,full,memory,MB,T,total tokens,2048,128.03125,128.03125,128.03125,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:19,0.8.1 +lfm2_short_conv,huggingface,full,memory,MB,T,total tokens,4096,256.046875,256.046875,256.046875,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:19,0.8.1 +lfm2_short_conv,huggingface,full,memory,MB,T,total tokens,8192,512.0390625,512.0390625,512.0390625,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:19,0.8.1 +lfm2_short_conv,huggingface,full,memory,MB,T,total tokens,16384,1024.0390625,1024.0390625,1024.0390625,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:19,0.8.1 +lfm2_short_conv,huggingface,full,memory,MB,T,total tokens,32768,2048.03125,2048.03125,2048.03125,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:19,0.8.1 +lfm2_short_conv,huggingface,full,memory,MB,T,total tokens,65536,4096.03125,4096.03125,4096.03125,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:19,0.8.1 +lfm2_short_conv,liger,full,memory,MB,T,total tokens,1024,44.15234375,44.15234375,44.15234375,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:19,0.8.1 +lfm2_short_conv,liger,full,memory,MB,T,total tokens,2048,88.24609375,88.24609375,88.24609375,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:19,0.8.1 +lfm2_short_conv,liger,full,memory,MB,T,total tokens,4096,176.43359375,176.43359375,176.43359375,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:19,0.8.1 +lfm2_short_conv,liger,full,memory,MB,T,total tokens,8192,352.80859375,352.80859375,352.80859375,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:19,0.8.1 +lfm2_short_conv,liger,full,memory,MB,T,total tokens,16384,707.05859375,707.05859375,707.05859375,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:19,0.8.1 +lfm2_short_conv,liger,full,memory,MB,T,total tokens,32768,1420.05859375,1420.05859375,1420.05859375,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:19,0.8.1 +lfm2_short_conv,liger,full,memory,MB,T,total tokens,65536,2840.05859375,2840.05859375,2840.05859375,"{""hidden_size"": 2048, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""kernel_size"": 3, ""bias_enabled"": false, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:19,0.8.1 +lfm2_moe_router,huggingface,forward,speed,ms,T,total tokens,1024,0.04508800059556961,0.04473600164055824,0.04572800174355507,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:38,0.8.1 +lfm2_moe_router,huggingface,forward,speed,ms,T,total tokens,2048,0.04678399860858917,0.04623999819159508,0.04751360043883324,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:38,0.8.1 +lfm2_moe_router,huggingface,forward,speed,ms,T,total tokens,4096,0.05246400088071823,0.05185920000076294,0.053036801517009735,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:38,0.8.1 +lfm2_moe_router,huggingface,forward,speed,ms,T,total tokens,8192,0.06351999938488007,0.06286720186471939,0.06424959748983383,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:38,0.8.1 +lfm2_moe_router,huggingface,forward,speed,ms,T,total tokens,16384,0.08323200047016144,0.08226560056209564,0.0839167982339859,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:38,0.8.1 +lfm2_moe_router,huggingface,forward,speed,ms,T,total tokens,32768,0.12169600278139114,0.12083840072154999,0.12266240268945694,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:38,0.8.1 +lfm2_moe_router,huggingface,forward,speed,ms,T,total tokens,65536,0.23199999332427979,0.23071999847888947,0.23292799293994904,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:38,0.8.1 +lfm2_moe_router,liger,forward,speed,ms,T,total tokens,1024,0.14033599942922592,0.1365056037902832,0.14862720370292665,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:39,0.8.1 +lfm2_moe_router,liger,forward,speed,ms,T,total tokens,2048,0.11617599800229073,0.09512320160865784,0.12019199877977371,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:39,0.8.1 +lfm2_moe_router,liger,forward,speed,ms,T,total tokens,4096,0.046592000871896744,0.04426239877939224,0.05018239840865135,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:39,0.8.1 +lfm2_moe_router,liger,forward,speed,ms,T,total tokens,8192,0.017520000226795673,0.017343999817967415,0.020416000857949257,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:39,0.8.1 +lfm2_moe_router,liger,forward,speed,ms,T,total tokens,16384,0.028511999174952507,0.028384000062942505,0.028780799359083176,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:39,0.8.1 +lfm2_moe_router,liger,forward,speed,ms,T,total tokens,32768,0.050912000238895416,0.05075199902057648,0.051072001457214355,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:39,0.8.1 +lfm2_moe_router,liger,forward,speed,ms,T,total tokens,65536,0.09529600292444229,0.09513600170612335,0.09548799693584442,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:39,0.8.1 +lfm2_moe_router,huggingface,backward,speed,ms,T,total tokens,1024,0.3251520097255707,0.321011209487915,0.33043840527534485,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:40,0.8.1 +lfm2_moe_router,huggingface,backward,speed,ms,T,total tokens,2048,0.32444798946380615,0.3222911894321442,0.32783361077308654,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:40,0.8.1 +lfm2_moe_router,huggingface,backward,speed,ms,T,total tokens,4096,0.33983999490737915,0.32618879675865176,0.881920003890992,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:40,0.8.1 +lfm2_moe_router,huggingface,backward,speed,ms,T,total tokens,8192,0.17097599804401398,0.16739200353622435,0.18328959643840792,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:40,0.8.1 +lfm2_moe_router,huggingface,backward,speed,ms,T,total tokens,16384,0.2987839877605438,0.29191678762435913,0.3213375985622406,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:40,0.8.1 +lfm2_moe_router,huggingface,backward,speed,ms,T,total tokens,32768,0.32819201052188873,0.3226560056209564,0.4341823995113376,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:40,0.8.1 +lfm2_moe_router,huggingface,backward,speed,ms,T,total tokens,65536,0.36476799845695496,0.3285312056541443,0.6833407759666443,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:40,0.8.1 +lfm2_moe_router,liger,backward,speed,ms,T,total tokens,1024,0.23825599998235703,0.23617279827594756,0.24745599925518036,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:41,0.8.1 +lfm2_moe_router,liger,backward,speed,ms,T,total tokens,2048,0.22387200593948364,0.22039039731025697,0.2831423997879029,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:41,0.8.1 +lfm2_moe_router,liger,backward,speed,ms,T,total tokens,4096,1.4731199741363525,0.900767993927002,1.514355182647705,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:41,0.8.1 +lfm2_moe_router,liger,backward,speed,ms,T,total tokens,8192,1.5130879878997803,1.3292352199554442,1.5443584203720093,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:41,0.8.1 +lfm2_moe_router,liger,backward,speed,ms,T,total tokens,16384,1.4589279890060425,1.3670079708099365,1.4958720207214355,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:41,0.8.1 +lfm2_moe_router,liger,backward,speed,ms,T,total tokens,32768,1.5173759460449219,1.3729535818099976,1.5512319564819337,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:41,0.8.1 +lfm2_moe_router,liger,backward,speed,ms,T,total tokens,65536,1.484000027179718,1.4716479778289795,1.5045119524002075,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:41,0.8.1 +lfm2_moe_router,huggingface,full,speed,ms,T,total tokens,1024,2.3496479988098145,2.321094417572022,2.383296012878418,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:41,0.8.1 +lfm2_moe_router,huggingface,full,speed,ms,T,total tokens,2048,2.353392004966736,2.153926372528076,2.3577088356018066,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:41,0.8.1 +lfm2_moe_router,huggingface,full,speed,ms,T,total tokens,4096,2.336416006088257,2.314604711532593,2.3909951210021974,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:41,0.8.1 +lfm2_moe_router,huggingface,full,speed,ms,T,total tokens,8192,2.3352320194244385,2.3170816898345947,2.3488255977630614,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:41,0.8.1 +lfm2_moe_router,huggingface,full,speed,ms,T,total tokens,16384,2.3631999492645264,2.335302400588989,2.3779776096343994,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:41,0.8.1 +lfm2_moe_router,huggingface,full,speed,ms,T,total tokens,32768,2.3592480421066284,2.3513407707214355,2.3790464878082274,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:41,0.8.1 +lfm2_moe_router,huggingface,full,speed,ms,T,total tokens,65536,2.3815040588378906,2.3666048526763914,2.4175743579864504,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:41,0.8.1 +lfm2_moe_router,liger,full,speed,ms,T,total tokens,1024,2.0937600135803223,2.0891263484954834,2.141926383972168,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,liger,full,speed,ms,T,total tokens,2048,1.980735957622528,1.656607985496521,2.1323071479797364,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,liger,full,speed,ms,T,total tokens,4096,2.1509439945220947,2.138432025909424,2.1622400283813477,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,liger,full,speed,ms,T,total tokens,8192,2.113152027130127,2.1073280811309814,2.13372163772583,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,liger,full,speed,ms,T,total tokens,16384,2.1474560499191284,2.1093376636505123,2.1852031707763673,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,liger,full,speed,ms,T,total tokens,32768,1.5060959458351135,1.1497535705566406,1.7032063722610475,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,liger,full,speed,ms,T,total tokens,65536,2.144416093826294,2.1296064853668213,2.161004734039307,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,huggingface,forward,memory,MB,T,total tokens,1024,0.30712890625,0.30712890625,0.30712890625,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,huggingface,forward,memory,MB,T,total tokens,2048,0.61376953125,0.61376953125,0.61376953125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,huggingface,forward,memory,MB,T,total tokens,4096,1.22705078125,1.22705078125,1.22705078125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,huggingface,forward,memory,MB,T,total tokens,8192,2.45361328125,2.45361328125,2.45361328125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,huggingface,forward,memory,MB,T,total tokens,16384,4.90673828125,4.90673828125,4.90673828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,huggingface,forward,memory,MB,T,total tokens,32768,9.81298828125,9.81298828125,9.81298828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,huggingface,forward,memory,MB,T,total tokens,65536,19.62548828125,19.62548828125,19.62548828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,liger,forward,memory,MB,T,total tokens,1024,0.08642578125,0.08642578125,0.08642578125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,liger,forward,memory,MB,T,total tokens,2048,0.17236328125,0.17236328125,0.17236328125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,liger,forward,memory,MB,T,total tokens,4096,0.34423828125,0.34423828125,0.34423828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,liger,forward,memory,MB,T,total tokens,8192,0.68798828125,0.68798828125,0.68798828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,liger,forward,memory,MB,T,total tokens,16384,1.37548828125,1.37548828125,1.37548828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,liger,forward,memory,MB,T,total tokens,32768,2.75048828125,2.75048828125,2.75048828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,liger,forward,memory,MB,T,total tokens,65536,5.50048828125,5.50048828125,5.50048828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,huggingface,backward,memory,MB,T,total tokens,1024,0.36962890625,0.36962890625,0.36962890625,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,huggingface,backward,memory,MB,T,total tokens,2048,0.73876953125,0.73876953125,0.73876953125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,huggingface,backward,memory,MB,T,total tokens,4096,1.47705078125,1.47705078125,1.47705078125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,huggingface,backward,memory,MB,T,total tokens,8192,2.95361328125,2.95361328125,2.95361328125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,huggingface,backward,memory,MB,T,total tokens,16384,5.90673828125,5.90673828125,5.90673828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,huggingface,backward,memory,MB,T,total tokens,32768,11.81298828125,11.81298828125,11.81298828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,huggingface,backward,memory,MB,T,total tokens,65536,23.62548828125,23.62548828125,23.62548828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,liger,backward,memory,MB,T,total tokens,1024,0.23486328125,0.23486328125,0.23486328125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,liger,backward,memory,MB,T,total tokens,2048,0.46923828125,0.46923828125,0.46923828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,liger,backward,memory,MB,T,total tokens,4096,0.93798828125,0.93798828125,0.93798828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,liger,backward,memory,MB,T,total tokens,8192,1.87548828125,1.87548828125,1.87548828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,liger,backward,memory,MB,T,total tokens,16384,3.75048828125,3.75048828125,3.75048828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,liger,backward,memory,MB,T,total tokens,32768,7.50048828125,7.50048828125,7.50048828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,liger,backward,memory,MB,T,total tokens,65536,15.00048828125,15.00048828125,15.00048828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,huggingface,full,memory,MB,T,total tokens,1024,0.36962890625,0.36962890625,0.36962890625,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,huggingface,full,memory,MB,T,total tokens,2048,0.73876953125,0.73876953125,0.73876953125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,huggingface,full,memory,MB,T,total tokens,4096,1.47705078125,1.47705078125,1.47705078125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,huggingface,full,memory,MB,T,total tokens,8192,2.95361328125,2.95361328125,2.95361328125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,huggingface,full,memory,MB,T,total tokens,16384,5.90673828125,5.90673828125,5.90673828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,huggingface,full,memory,MB,T,total tokens,32768,11.81298828125,11.81298828125,11.81298828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,huggingface,full,memory,MB,T,total tokens,65536,23.62548828125,23.62548828125,23.62548828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,liger,full,memory,MB,T,total tokens,1024,0.23486328125,0.23486328125,0.23486328125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,liger,full,memory,MB,T,total tokens,2048,0.46923828125,0.46923828125,0.46923828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,liger,full,memory,MB,T,total tokens,4096,0.93798828125,0.93798828125,0.93798828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,liger,full,memory,MB,T,total tokens,8192,1.87548828125,1.87548828125,1.87548828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,liger,full,memory,MB,T,total tokens,16384,3.75048828125,3.75048828125,3.75048828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,liger,full,memory,MB,T,total tokens,32768,7.50048828125,7.50048828125,7.50048828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 +lfm2_moe_router,liger,full,memory,MB,T,total tokens,65536,15.00048828125,15.00048828125,15.00048828125,"{""num_experts"": 32, ""topk"": 4, ""dtype"": ""torch.bfloat16"", ""bsz"": 1, ""seq_len"": 65536}",NVIDIA H100 80GB HBM3,2026-07-31 16:41:42,0.8.1 diff --git a/benchmark/scripts/benchmark_lfm2_moe_router.py b/benchmark/scripts/benchmark_lfm2_moe_router.py new file mode 100644 index 000000000..5220acb47 --- /dev/null +++ b/benchmark/scripts/benchmark_lfm2_moe_router.py @@ -0,0 +1,116 @@ +import torch +import torch.nn as nn + +from benchmark_model_configs import MODEL_REGISTRY +from benchmark_model_configs import build_model_config_sweep +from benchmark_model_configs import build_token_length_sweep +from benchmark_model_configs import get_benchmark_model_config +from test.transformers.test_lfm2_moe_router import _reference +from utils import SingleBenchmarkRunInput +from utils import build_memory_bench_fn +from utils import build_speed_bench_fn +from utils import parse_benchmark_script_args +from utils import run_benchmarks + +from liger_kernel.ops import LigerLfm2MoeRouterFunction +from liger_kernel.utils import infer_device + +device = infer_device() + + +class _Router(nn.Module): + def __init__(self, num_experts, top_k, dtype, use_liger): + super().__init__() + self.register_buffer("expert_bias", torch.zeros(num_experts, device=device, dtype=torch.float32)) + self.top_k = top_k + self.use_liger = use_liger + self.dtype = dtype + + def forward(self, router_logits): + if self.use_liger: + return LigerLfm2MoeRouterFunction.apply( + router_logits, + self.expert_bias, + self.top_k, + True, + 1.0, + )[1] + return _reference(router_logits, self.expert_bias, self.top_k, True, 1.0)[1] + + +def setup_lfm2_moe_router(input: SingleBenchmarkRunInput): + cfg = input.extra_benchmark_config + if isinstance(input.x, str): + model_cfg = MODEL_REGISTRY[input.x] + num_tokens = cfg["bsz"] * cfg["seq_len"] + num_experts = model_cfg.num_experts + top_k = model_cfg.topk + dtype = model_cfg.dtype + else: + num_tokens = cfg["bsz"] * input.x + num_experts = cfg["num_experts"] + top_k = cfg["topk"] + dtype = cfg["dtype"] + + if num_experts is None or top_k is None: + raise ValueError("LFM2 MoE router benchmarks require an MoE model configuration") + + router_logits = torch.randn( + num_tokens, + num_experts, + device=device, + dtype=dtype, + requires_grad=True, + ) + if input.kernel_provider == "liger": + layer = _Router(num_experts, top_k, dtype, use_liger=True) + elif input.kernel_provider == "huggingface": + layer = _Router(num_experts, top_k, dtype, use_liger=False) + else: + raise ValueError(f"Invalid provider: {input.kernel_provider} for LFM2 MoE router") + return router_logits, layer + + +if __name__ == "__main__": + args = parse_benchmark_script_args() + + if args.sweep_mode == "model_config": + moe_configs = [config for config in MODEL_REGISTRY.values() if config.is_moe] + common_configs = build_model_config_sweep( + kernel_name="lfm2_moe_router", + all_model_configs=moe_configs, + setup_fn=setup_lfm2_moe_router, + model_keys=["num_experts", "topk", "dtype"], + probe_provider="huggingface", + extra_configs={"bsz": 1}, + probe_dim="T", + bt=args.bt, + overwrite=args.overwrite, + ) + else: + model = get_benchmark_model_config(args.model or "lfm2_moe_8b_a1b") + common_configs = build_token_length_sweep( + kernel_name="lfm2_moe_router", + probe_x=1024, + model=model, + setup_fn=setup_lfm2_moe_router, + model_keys=["num_experts", "topk", "dtype"], + extra_configs={"bsz": 1}, + scale_dim="T", + x_label="total tokens", + probe_provider="huggingface", + overwrite=args.overwrite, + ) + + common_configs["kernel_providers"] = ["huggingface", "liger"] + for metric_name, metric_unit, bench_fn in ( + ("speed", "ms", build_speed_bench_fn(setup_lfm2_moe_router)), + ("memory", "MB", build_memory_bench_fn(setup_lfm2_moe_router)), + ): + run_benchmarks( + bench_test_fn=bench_fn, + kernel_operation_modes=["forward", "backward", "full"], + metric_name=metric_name, + metric_unit=metric_unit, + **common_configs, + ) diff --git a/benchmark/scripts/benchmark_lfm2_short_conv.py b/benchmark/scripts/benchmark_lfm2_short_conv.py new file mode 100644 index 000000000..5dbfdcf2d --- /dev/null +++ b/benchmark/scripts/benchmark_lfm2_short_conv.py @@ -0,0 +1,103 @@ +import torch +import torch.nn as nn + +from benchmark_model_configs import MODEL_REGISTRY +from benchmark_model_configs import build_model_config_sweep +from benchmark_model_configs import build_token_length_sweep +from benchmark_model_configs import get_benchmark_model_config +from test.transformers.test_lfm2_short_conv import _reference +from utils import SingleBenchmarkRunInput +from utils import build_memory_bench_fn +from utils import build_speed_bench_fn +from utils import parse_benchmark_script_args +from utils import run_benchmarks + +from liger_kernel.ops import LigerLfm2ShortConvFunction +from liger_kernel.utils import infer_device + +device = infer_device() + + +class _ShortConv(nn.Module): + def __init__(self, hidden_size, kernel_size, dtype, use_liger, bias_enabled): + super().__init__() + self.weight = nn.Parameter(torch.randn(hidden_size, 1, kernel_size, device=device, dtype=dtype) * 0.02) + self.bias = nn.Parameter(torch.zeros(hidden_size, device=device, dtype=dtype)) if bias_enabled else None + self.use_liger = use_liger + + def forward(self, bcx): + if self.use_liger: + return LigerLfm2ShortConvFunction.apply(bcx, self.weight, self.bias) + return _reference(bcx, self.weight, self.bias) + + +def setup_lfm2_short_conv(input: SingleBenchmarkRunInput): + cfg = input.extra_benchmark_config + if isinstance(input.x, str): + model_cfg = MODEL_REGISTRY[input.x] + seq_len = cfg["seq_len"] + hidden_size = model_cfg.hidden_size + dtype = model_cfg.dtype + else: + seq_len = input.x + hidden_size = cfg["hidden_size"] + dtype = cfg["dtype"] + + bcx = torch.randn( + cfg["bsz"], + seq_len, + 3 * hidden_size, + device=device, + dtype=dtype, + requires_grad=True, + ) + if input.kernel_provider == "liger": + layer = _ShortConv(hidden_size, cfg["kernel_size"], dtype, use_liger=True, bias_enabled=cfg["bias_enabled"]) + elif input.kernel_provider == "huggingface": + layer = _ShortConv(hidden_size, cfg["kernel_size"], dtype, use_liger=False, bias_enabled=cfg["bias_enabled"]) + else: + raise ValueError(f"Invalid provider: {input.kernel_provider} for LFM2 short convolution") + return bcx, layer + + +if __name__ == "__main__": + args = parse_benchmark_script_args() + + if args.sweep_mode == "model_config": + common_configs = build_model_config_sweep( + kernel_name="lfm2_short_conv", + setup_fn=setup_lfm2_short_conv, + model_keys=["hidden_size", "dtype"], + probe_provider="huggingface", + extra_configs={"bsz": 1, "kernel_size": 3, "bias_enabled": False}, + probe_dim="T", + bt=args.bt, + overwrite=args.overwrite, + ) + else: + model = get_benchmark_model_config(args.model or "lfm2_1.2b") + common_configs = build_token_length_sweep( + kernel_name="lfm2_short_conv", + probe_x=1024, + model=model, + setup_fn=setup_lfm2_short_conv, + model_keys=["hidden_size", "dtype"], + extra_configs={"bsz": 1, "kernel_size": 3, "bias_enabled": False}, + scale_dim="T", + x_label="total tokens", + probe_provider="huggingface", + overwrite=args.overwrite, + ) + + common_configs["kernel_providers"] = ["huggingface", "liger"] + for metric_name, metric_unit, bench_fn in ( + ("speed", "ms", build_speed_bench_fn(setup_lfm2_short_conv)), + ("memory", "MB", build_memory_bench_fn(setup_lfm2_short_conv)), + ): + run_benchmarks( + bench_test_fn=bench_fn, + kernel_operation_modes=["forward", "backward", "full"], + metric_name=metric_name, + metric_unit=metric_unit, + **common_configs, + ) diff --git a/benchmark/scripts/benchmark_model_configs.py b/benchmark/scripts/benchmark_model_configs.py index 192495c96..4d8447033 100644 --- a/benchmark/scripts/benchmark_model_configs.py +++ b/benchmark/scripts/benchmark_model_configs.py @@ -223,6 +223,33 @@ class MoEModelConfig: topk=8, ) +LFM2_1_2B = ModelConfig( + name="lfm2_1.2b", + hidden_size=2048, + intermediate_size=12288, + vocab_size=65536, + num_attention_heads=32, + num_key_value_heads=8, + head_dim=64, + hidden_act="silu", + max_position_embeddings=128000, +) + +LFM2_MOE_8B_A1B = ModelConfig( + name="lfm2_moe_8b_a1b", + hidden_size=2048, + intermediate_size=7168, + vocab_size=65536, + num_attention_heads=32, + num_key_value_heads=8, + head_dim=64, + hidden_act="silu", + max_position_embeddings=128000, + num_experts=32, + topk=4, + moe_intermediate_size=1792, +) + MODEL_REGISTRY: Dict[str, ModelConfig] = { "llama_2_7b": LLAMA_2_7B, "llama_3_8b": LLAMA_3_8B, @@ -231,6 +258,8 @@ class MoEModelConfig: "qwen2.5_72b": QWEN_2_5_72B, "deepseek_v2_lite": DEEPSEEK_V2_LITE, "deepseek_v3": DEEPSEEK_V3, + "lfm2_1.2b": LFM2_1_2B, + "lfm2_moe_8b_a1b": LFM2_MOE_8B_A1B, } DEFAULT_MODEL_CONFIG = LLAMA_3_8B diff --git a/docs/High-Level-APIs.md b/docs/High-Level-APIs.md index 6bbe008a9..789d68c29 100644 --- a/docs/High-Level-APIs.md +++ b/docs/High-Level-APIs.md @@ -28,6 +28,9 @@ You can also use the Patching APIs to use the kernels for a specific model archi |-------------|--------------------------------------------------------------|-------------------------------------------------------------------------| | LLaMA 2 & 3 | `liger_kernel.transformers.apply_liger_kernel_to_llama` | RoPE, RMSNorm, SwiGLU, CrossEntropyLoss, FusedLinearCrossEntropy | | LLaMA 3.2-Vision | `liger_kernel.transformers.apply_liger_kernel_to_mllama` | RoPE, RMSNorm, SwiGLU, CrossEntropyLoss, FusedLinearCrossEntropy | +| LFM2 | `liger_kernel.transformers.apply_liger_kernel_to_lfm2` | RoPE, RMSNorm, SwiGLU, ShortConv, CrossEntropyLoss, FusedLinearCrossEntropy | +| LFM2MoE | `liger_kernel.transformers.apply_liger_kernel_to_lfm2_moe` | RoPE, RMSNorm, SwiGLU, ShortConv, FusedMoE, MoERouter, CrossEntropyLoss, FusedLinearCrossEntropy | +| LFM2VL | `liger_kernel.transformers.apply_liger_kernel_to_lfm2_vl` | SigLIP2 LayerNorm, RoPE, RMSNorm, SwiGLU, ShortConv, CrossEntropyLoss, FusedLinearCrossEntropy | | Mistral | `liger_kernel.transformers.apply_liger_kernel_to_mistral` | RoPE, RMSNorm, SwiGLU, CrossEntropyLoss, FusedLinearCrossEntropy | | Mixtral | `liger_kernel.transformers.apply_liger_kernel_to_mixtral` | RoPE, RMSNorm, SwiGLU, CrossEntropyLoss, FusedLinearCrossEntropy | | Gemma1 | `liger_kernel.transformers.apply_liger_kernel_to_gemma` | RoPE, RMSNorm, GeGLU, CrossEntropyLoss, FusedLinearCrossEntropy | @@ -62,6 +65,24 @@ You can also use the Patching APIs to use the kernels for a specific model archi show_docstring: true show_signature: true +::: liger_kernel.transformers.apply_liger_kernel_to_lfm2 + options: + extra: + show_docstring: true + show_signature: true + +::: liger_kernel.transformers.apply_liger_kernel_to_lfm2_moe + options: + extra: + show_docstring: true + show_signature: true + +::: liger_kernel.transformers.apply_liger_kernel_to_lfm2_vl + options: + extra: + show_docstring: true + show_signature: true + ::: liger_kernel.transformers.apply_liger_kernel_to_gemma options: extra: diff --git a/src/liger_kernel/ops/__init__.py b/src/liger_kernel/ops/__init__.py index f6676c07e..32182ed5f 100644 --- a/src/liger_kernel/ops/__init__.py +++ b/src/liger_kernel/ops/__init__.py @@ -65,6 +65,8 @@ from liger_kernel.ops.layer_norm import LigerLayerNormFunction # noqa: F401 from liger_kernel.ops.layer_norm import layer_norm_backward # noqa: F401 from liger_kernel.ops.layer_norm import layer_norm_forward # noqa: F401 +from liger_kernel.ops.lfm2_moe_router import LigerLfm2MoeRouterFunction # noqa: F401 +from liger_kernel.ops.lfm2_short_conv import LigerLfm2ShortConvFunction # noqa: F401 from liger_kernel.ops.llama4_rope import LigerLlama4RopeFunction # noqa: F401 from liger_kernel.ops.mhc import LigerMHCCoeffsFunction # noqa: F401 from liger_kernel.ops.mhc import LigerMHCPostResFunction # noqa: F401 diff --git a/src/liger_kernel/ops/lfm2_moe_router.py b/src/liger_kernel/ops/lfm2_moe_router.py new file mode 100644 index 000000000..df28dbd20 --- /dev/null +++ b/src/liger_kernel/ops/lfm2_moe_router.py @@ -0,0 +1,170 @@ +import torch +import triton +import triton.language as tl + + +@triton.jit +def _lfm2_moe_router_forward( + router_logits, + expert_bias, + selected_experts, + routing_weights, + n_tokens, + n_experts, + stride_token, + routed_scaling_factor: tl.constexpr, + norm_topk_prob: tl.constexpr, + has_expert_bias: tl.constexpr, + TOP_K: tl.constexpr, + BLOCK_K: tl.constexpr, + BLOCK_EXPERTS: tl.constexpr, +): + token = tl.program_id(0) + expert_offsets = tl.arange(0, BLOCK_EXPERTS) + expert_mask = expert_offsets < n_experts + logits = tl.load(router_logits + token * stride_token + expert_offsets, mask=expert_mask, other=-float("inf")) + # Match torch.sigmoid's output dtype before top-k; BF16 rounding can + # otherwise change expert selection for nearly equal router scores. + probabilities = tl.sigmoid(logits.to(tl.float32)).to(logits.dtype) + scores = probabilities + if has_expert_bias: + scores = probabilities.to(tl.float32) + tl.load(expert_bias + expert_offsets, mask=expert_mask, other=0.0).to( + tl.float32 + ) + + k_offsets = tl.arange(0, BLOCK_K) + topk_probabilities = tl.zeros((BLOCK_K,), dtype=tl.float32) + topk_indices = tl.zeros((BLOCK_K,), dtype=tl.int32) + for k in range(TOP_K): + expert = tl.argmax(scores, axis=0, tie_break_left=True) + probability = tl.sum(tl.where(expert_offsets == expert, probabilities, 0.0), axis=0) + topk_probabilities = tl.where(k_offsets == k, probability, topk_probabilities) + topk_indices = tl.where(k_offsets == k, expert, topk_indices) + scores = tl.where(expert_offsets == expert, -float("inf"), scores) + + if norm_topk_prob: + denominator = tl.sum(tl.where(k_offsets < TOP_K, topk_probabilities, 0.0), axis=0) + 1e-6 + topk_probabilities /= denominator + topk_probabilities *= routed_scaling_factor + + output_mask = k_offsets < TOP_K + output_offsets = token * TOP_K + k_offsets + tl.store(selected_experts + output_offsets, topk_indices, mask=output_mask) + tl.store(routing_weights + output_offsets, topk_probabilities, mask=output_mask) + + +@triton.jit +def _lfm2_moe_router_backward( + grad_routing_weights, + router_logits, + selected_experts, + grad_router_logits, + n_tokens, + n_experts, + stride_token, + routed_scaling_factor: tl.constexpr, + norm_topk_prob: tl.constexpr, + TOP_K: tl.constexpr, + BLOCK_K: tl.constexpr, + BLOCK_EXPERTS: tl.constexpr, +): + token = tl.program_id(0) + expert_offsets = tl.arange(0, BLOCK_EXPERTS) + expert_mask = expert_offsets < n_experts + k_offsets = tl.arange(0, BLOCK_K) + mask = k_offsets < TOP_K + offsets = token * TOP_K + k_offsets + experts = tl.load(selected_experts + offsets, mask=mask, other=0).to(tl.int32) + grad_weights = tl.load(grad_routing_weights + offsets, mask=mask, other=0.0).to(tl.float32) + probabilities = tl.sigmoid( + tl.load(router_logits + token * stride_token + experts, mask=mask, other=0.0).to(tl.float32) + ) + + if norm_topk_prob: + denominator = tl.sum(probabilities, axis=0) + 1e-6 + weighted_grad_sum = tl.sum(grad_weights * probabilities, axis=0) + grad_probabilities = routed_scaling_factor * ( + grad_weights / denominator - weighted_grad_sum / (denominator * denominator) + ) + else: + grad_probabilities = routed_scaling_factor * grad_weights + + grad_logits = grad_probabilities * probabilities * (1.0 - probabilities) + dense_grad_logits = tl.sum( + tl.where( + expert_offsets[:, None] == experts[None, :], + grad_logits[None, :], + 0.0, + ), + axis=1, + ) + tl.store( + grad_router_logits + token * stride_token + expert_offsets, + dense_grad_logits, + mask=expert_mask, + ) + + +class LigerLfm2MoeRouterFunction(torch.autograd.Function): + """Fused sigmoid, biased top-k selection, and routing-weight normalization for LFM2-MoE.""" + + @staticmethod + def forward(ctx, router_logits, expert_bias, top_k, norm_topk_prob, routed_scaling_factor): + if router_logits.ndim != 2: + raise ValueError("router_logits must have shape [tokens, experts]") + n_tokens, n_experts = router_logits.shape + if not 0 < top_k <= n_experts: + raise ValueError("top_k must be between one and the number of experts") + + selected_experts = torch.empty((n_tokens, top_k), dtype=torch.int32, device=router_logits.device) + routing_weights = torch.empty((n_tokens, top_k), dtype=router_logits.dtype, device=router_logits.device) + saved_bias = ( + expert_bias if expert_bias is not None else torch.empty(0, dtype=torch.float32, device=router_logits.device) + ) + block_k = triton.next_power_of_2(top_k) + _lfm2_moe_router_forward[(n_tokens,)]( + router_logits, + saved_bias, + selected_experts, + routing_weights, + n_tokens, + n_experts, + router_logits.stride(0), + routed_scaling_factor=routed_scaling_factor, + norm_topk_prob=norm_topk_prob, + has_expert_bias=expert_bias is not None, + TOP_K=top_k, + BLOCK_K=block_k, + BLOCK_EXPERTS=triton.next_power_of_2(n_experts), + ) + ctx.save_for_backward(router_logits, selected_experts) + ctx.n_experts = n_experts + ctx.stride_token = router_logits.stride(0) + ctx.top_k = top_k + ctx.norm_topk_prob = norm_topk_prob + ctx.routed_scaling_factor = routed_scaling_factor + ctx.mark_non_differentiable(selected_experts) + return selected_experts, routing_weights + + @staticmethod + def backward(ctx, grad_selected_experts, grad_routing_weights): + router_logits, selected_experts = ctx.saved_tensors + n_tokens = selected_experts.shape[0] + grad_router_logits = torch.empty( + (n_tokens, ctx.n_experts), dtype=router_logits.dtype, device=router_logits.device + ) + _lfm2_moe_router_backward[(n_tokens,)]( + grad_routing_weights, + router_logits, + selected_experts, + grad_router_logits, + n_tokens, + ctx.n_experts, + ctx.stride_token, + routed_scaling_factor=ctx.routed_scaling_factor, + norm_topk_prob=ctx.norm_topk_prob, + TOP_K=ctx.top_k, + BLOCK_K=triton.next_power_of_2(ctx.top_k), + BLOCK_EXPERTS=triton.next_power_of_2(ctx.n_experts), + ) + return grad_router_logits, None, None, None, None diff --git a/src/liger_kernel/ops/lfm2_short_conv.py b/src/liger_kernel/ops/lfm2_short_conv.py new file mode 100644 index 000000000..bdffe96b2 --- /dev/null +++ b/src/liger_kernel/ops/lfm2_short_conv.py @@ -0,0 +1,264 @@ +import torch +import triton +import triton.language as tl + +from liger_kernel.utils import infer_device_arch + + +def _short_conv_weight_backward_config(batch_tokens): + if infer_device_arch() == "hopper": + if batch_tokens >= 32768: + return 64, 2, 2 + if batch_tokens >= 16384: + return 128, 4, 2 + return 256, None, None + + +@triton.jit +def _short_conv_forward( + bcx, + weight, + bias, + output, + n_elements, + seq_len, + hidden_size, + stride_b, + stride_t, + stride_h, + stride_wh, + stride_wk, + K: tl.constexpr, + HAS_BIAS: tl.constexpr, + BLOCK: tl.constexpr, +): + offsets = tl.program_id(0) * BLOCK + tl.arange(0, BLOCK) + mask = offsets < n_elements + hidden = offsets % hidden_size + tokens = offsets // hidden_size + time = tokens % seq_len + batch = tokens // seq_len + + conv = tl.zeros((BLOCK,), dtype=tl.float32) + for lag in range(K): + source_time = time - lag + source_mask = mask & (source_time >= 0) + source = batch * stride_b + source_time * stride_t + hidden * stride_h + gate_b = tl.load(bcx + source, mask=source_mask, other=0.0).to(tl.float32) + value = tl.load(bcx + source + 2 * hidden_size * stride_h, mask=source_mask, other=0.0).to(tl.float32) + kernel = tl.load(weight + hidden * stride_wh + (K - 1 - lag) * stride_wk, mask=mask, other=0.0).to(tl.float32) + conv += gate_b * value * kernel + + if HAS_BIAS: + conv += tl.load(bias + hidden, mask=mask, other=0.0).to(tl.float32) + + current = batch * stride_b + time * stride_t + hidden * stride_h + gate_c = tl.load(bcx + current + hidden_size * stride_h, mask=mask, other=0.0).to(tl.float32) + tl.store(output + offsets, gate_c * conv, mask=mask) + + +@triton.jit +def _short_conv_input_backward( + grad_output, + bcx, + weight, + bias, + grad_bcx, + n_elements, + seq_len, + hidden_size, + stride_b, + stride_t, + stride_h, + stride_wh, + stride_wk, + K: tl.constexpr, + HAS_BIAS: tl.constexpr, + BLOCK: tl.constexpr, +): + offsets = tl.program_id(0) * BLOCK + tl.arange(0, BLOCK) + mask = offsets < n_elements + hidden = offsets % hidden_size + tokens = offsets // hidden_size + time = tokens % seq_len + batch = tokens // seq_len + current = batch * stride_b + time * stride_t + hidden * stride_h + + conv = tl.zeros((BLOCK,), dtype=tl.float32) + grad_product = tl.zeros((BLOCK,), dtype=tl.float32) + for lag in range(K): + source_time = time - lag + source_mask = mask & (source_time >= 0) + source = batch * stride_b + source_time * stride_t + hidden * stride_h + source_b = tl.load(bcx + source, mask=source_mask, other=0.0).to(tl.float32) + source_x = tl.load(bcx + source + 2 * hidden_size * stride_h, mask=source_mask, other=0.0).to(tl.float32) + kernel = tl.load(weight + hidden * stride_wh + (K - 1 - lag) * stride_wk, mask=mask, other=0.0).to(tl.float32) + conv += source_b * source_x * kernel + + output_time = time + lag + output_mask = mask & (output_time < seq_len) + output_offset = (batch * seq_len + output_time) * hidden_size + hidden + output_base = batch * stride_b + output_time * stride_t + hidden * stride_h + grad_y = tl.load(grad_output + output_offset, mask=output_mask, other=0.0).to(tl.float32) + output_c = tl.load(bcx + output_base + hidden_size * stride_h, mask=output_mask, other=0.0).to(tl.float32) + grad_product += grad_y * output_c * kernel + + if HAS_BIAS: + conv += tl.load(bias + hidden, mask=mask, other=0.0).to(tl.float32) + + current_b = tl.load(bcx + current, mask=mask, other=0.0).to(tl.float32) + current_x = tl.load(bcx + current + 2 * hidden_size * stride_h, mask=mask, other=0.0).to(tl.float32) + grad_y = tl.load(grad_output + offsets, mask=mask, other=0.0).to(tl.float32) + + grad_base = tokens * 3 * hidden_size + hidden + tl.store(grad_bcx + grad_base, grad_product * current_x, mask=mask) + tl.store(grad_bcx + grad_base + hidden_size, grad_y * conv, mask=mask) + tl.store(grad_bcx + grad_base + 2 * hidden_size, grad_product * current_b, mask=mask) + + +@triton.jit +def _short_conv_weight_backward( + grad_output, + bcx, + weight_partials, + bias_partials, + seq_len, + hidden_size, + batch_tokens, + n_chunks, + stride_b, + stride_t, + stride_h, + K: tl.constexpr, + HAS_BIAS: tl.constexpr, + BLOCK: tl.constexpr, +): + hidden = tl.program_id(0) + chunk = tl.program_id(1) + batch_token = chunk * BLOCK + tl.arange(0, BLOCK) + mask = batch_token < batch_tokens + batch = batch_token // seq_len + output_time = batch_token % seq_len + output_offset = batch_token * hidden_size + hidden + output_base = batch * stride_b + output_time * stride_t + hidden * stride_h + grad_y = tl.load(grad_output + output_offset, mask=mask, other=0.0).to(tl.float32) + gate_c = tl.load(bcx + output_base + hidden_size * stride_h, mask=mask, other=0.0).to(tl.float32) + grad_conv = grad_y * gate_c + + # One program handles all taps for a hidden channel, reusing grad_y and + # gate_c and reducing the launch grid by K. + for kernel_idx in range(K): + source_time = output_time + kernel_idx - (K - 1) + source_mask = mask & (source_time >= 0) + source = batch * stride_b + source_time * stride_t + hidden * stride_h + gate_b = tl.load(bcx + source, mask=source_mask, other=0.0).to(tl.float32) + value = tl.load(bcx + source + 2 * hidden_size * stride_h, mask=source_mask, other=0.0).to(tl.float32) + partial_offset = (hidden * K + kernel_idx) * n_chunks + chunk + tl.store(weight_partials + partial_offset, tl.sum(grad_conv * gate_b * value, axis=0)) + + if HAS_BIAS: + tl.store(bias_partials + hidden * n_chunks + chunk, tl.sum(grad_conv, axis=0)) + + +class LigerLfm2ShortConvFunction(torch.autograd.Function): + """Fuses the B*x gate, causal depthwise convolution, and C gate.""" + + @staticmethod + def forward(ctx, bcx, weight, bias=None): + if bcx.ndim != 3 or weight.ndim != 3 or weight.shape[1] != 1: + raise ValueError("expected bcx [batch, time, 3*hidden] and weight [hidden, 1, kernel]") + batch_size, seq_len, three_hidden = bcx.shape + hidden_size = weight.shape[0] + if three_hidden != 3 * hidden_size: + raise ValueError("bcx final dimension must equal three times the convolution hidden size") + + output = torch.empty((batch_size, seq_len, hidden_size), dtype=bcx.dtype, device=bcx.device) + n_elements = output.numel() + _short_conv_forward[(triton.cdiv(n_elements, 512),)]( + bcx, + weight, + bias, + output, + n_elements, + seq_len, + hidden_size, + bcx.stride(0), + bcx.stride(1), + bcx.stride(2), + weight.stride(0), + weight.stride(2), + K=weight.shape[2], + HAS_BIAS=bias is not None, + BLOCK=512, + ) + saved_bias = bias if bias is not None else torch.empty(0, dtype=bcx.dtype, device=bcx.device) + ctx.save_for_backward(bcx, weight, saved_bias) + ctx.has_bias = bias is not None + return output + + @staticmethod + def backward(ctx, grad_output): + bcx, weight, bias = ctx.saved_tensors + grad_output = grad_output.contiguous() + batch_size, seq_len, three_hidden = bcx.shape + hidden_size = three_hidden // 3 + kernel_size = weight.shape[2] + n_elements = batch_size * seq_len * hidden_size + + grad_bcx = torch.empty_like(bcx) + _short_conv_input_backward[(triton.cdiv(n_elements, 512),)]( + grad_output, + bcx, + weight, + bias, + grad_bcx, + n_elements, + seq_len, + hidden_size, + bcx.stride(0), + bcx.stride(1), + bcx.stride(2), + weight.stride(0), + weight.stride(2), + K=kernel_size, + HAS_BIAS=ctx.has_bias, + BLOCK=512, + ) + + batch_tokens = batch_size * seq_len + # Hopper benefits from finer-grained long-sequence weight reductions: + # the additional partials expose more parallel work and reduce the + # dominant backward latency. Keep every other architecture on the + # existing 256-token reduction to preserve ROCm and portable behavior. + weight_block, weight_warps, weight_stages = _short_conv_weight_backward_config(batch_tokens) + use_hopper_weight_config = weight_warps is not None + n_chunks = triton.cdiv(batch_tokens, weight_block) + partial_shape = (hidden_size * kernel_size, n_chunks) + weight_partials = torch.empty(partial_shape, dtype=torch.float32, device=bcx.device) + bias_partials = ( + torch.empty((hidden_size, n_chunks), dtype=torch.float32, device=bcx.device) + if ctx.has_bias + else weight_partials + ) + _short_conv_weight_backward[(hidden_size, n_chunks)]( + grad_output, + bcx, + weight_partials, + bias_partials, + seq_len, + hidden_size, + batch_tokens, + n_chunks, + bcx.stride(0), + bcx.stride(1), + bcx.stride(2), + K=kernel_size, + HAS_BIAS=ctx.has_bias, + BLOCK=weight_block, + **({"num_warps": weight_warps, "num_stages": weight_stages} if use_hopper_weight_config else {}), + ) + grad_weight = weight_partials.sum(1).reshape(hidden_size, 1, kernel_size).to(weight.dtype) + grad_bias = None + if ctx.has_bias: + grad_bias = bias_partials.sum(1).to(bias.dtype) + return grad_bcx, grad_weight, grad_bias diff --git a/src/liger_kernel/transformers/__init__.py b/src/liger_kernel/transformers/__init__.py index 26bdef91b..d7b06ece0 100644 --- a/src/liger_kernel/transformers/__init__.py +++ b/src/liger_kernel/transformers/__init__.py @@ -26,6 +26,8 @@ from liger_kernel.transformers.swiglu import LigerBlockSparseTop2MLP # noqa: F401 from liger_kernel.transformers.swiglu import LigerExperts # noqa: F401 from liger_kernel.transformers.swiglu import LigerFalconH1SwiGLUMLP # noqa: F401 +from liger_kernel.transformers.swiglu import LigerLfm2MoeExperts # noqa: F401 +from liger_kernel.transformers.swiglu import LigerLfm2SwiGLUMLP # noqa: F401 from liger_kernel.transformers.swiglu import LigerPhi3SwiGLUMLP # noqa: F401 from liger_kernel.transformers.swiglu import LigerQwen3MoeSwiGLUMLP # noqa: F401 from liger_kernel.transformers.swiglu import LigerSwiGLUMLP # noqa: F401 @@ -56,6 +58,9 @@ from liger_kernel.transformers.monkey_patch import apply_liger_kernel_to_hunyuan_v1_dense # noqa: F401 from liger_kernel.transformers.monkey_patch import apply_liger_kernel_to_hunyuan_v1_moe # noqa: F401 from liger_kernel.transformers.monkey_patch import apply_liger_kernel_to_internvl # noqa: F401 + from liger_kernel.transformers.monkey_patch import apply_liger_kernel_to_lfm2 # noqa: F401 + from liger_kernel.transformers.monkey_patch import apply_liger_kernel_to_lfm2_moe # noqa: F401 + from liger_kernel.transformers.monkey_patch import apply_liger_kernel_to_lfm2_vl # noqa: F401 from liger_kernel.transformers.monkey_patch import apply_liger_kernel_to_llama # noqa: F401 from liger_kernel.transformers.monkey_patch import apply_liger_kernel_to_llama4 # noqa: F401 from liger_kernel.transformers.monkey_patch import apply_liger_kernel_to_llava # noqa: F401 @@ -131,6 +136,9 @@ def __getattr__(name: str): "apply_liger_kernel_to_gpt_oss", "apply_liger_kernel_to_granite", "apply_liger_kernel_to_internvl", + "apply_liger_kernel_to_lfm2", + "apply_liger_kernel_to_lfm2_moe", + "apply_liger_kernel_to_lfm2_vl", "apply_liger_kernel_to_llama", "apply_liger_kernel_to_llava", "apply_liger_kernel_to_llama4", @@ -188,6 +196,8 @@ def __getattr__(name: str): "liger_llama4_text_rotary_pos_emb", "liger_llama4_vision_rotary_pos_emb", "LigerBlockSparseTop2MLP", + "LigerLfm2MoeExperts", + "LigerLfm2SwiGLUMLP", "LigerFalconH1SwiGLUMLP", "LigerPhi3SwiGLUMLP", "LigerQwen3MoeSwiGLUMLP", @@ -222,6 +232,9 @@ def __getattr__(name: str): "apply_liger_kernel_to_gpt_oss", "apply_liger_kernel_to_granite", "apply_liger_kernel_to_internvl", + "apply_liger_kernel_to_lfm2", + "apply_liger_kernel_to_lfm2_moe", + "apply_liger_kernel_to_lfm2_vl", "apply_liger_kernel_to_llama", "apply_liger_kernel_to_llava", "apply_liger_kernel_to_llama4", diff --git a/src/liger_kernel/transformers/lfm2_moe_router.py b/src/liger_kernel/transformers/lfm2_moe_router.py new file mode 100644 index 000000000..264df5eb7 --- /dev/null +++ b/src/liger_kernel/transformers/lfm2_moe_router.py @@ -0,0 +1,13 @@ +from liger_kernel.ops import LigerLfm2MoeRouterFunction + + +def liger_lfm2_moe_route_tokens_to_experts(self, router_logits): + """Route LFM2-MoE tokens without materializing full sigmoid routing weights.""" + expert_bias = self.expert_bias if self.use_expert_bias else None + return LigerLfm2MoeRouterFunction.apply( + router_logits, + expert_bias, + self.top_k, + self.norm_topk_prob, + self.routed_scaling_factor, + ) diff --git a/src/liger_kernel/transformers/lfm2_short_conv.py b/src/liger_kernel/transformers/lfm2_short_conv.py new file mode 100644 index 000000000..65136a652 --- /dev/null +++ b/src/liger_kernel/transformers/lfm2_short_conv.py @@ -0,0 +1,27 @@ +import torch + +from liger_kernel.ops import LigerLfm2ShortConvFunction + + +def liger_lfm2_short_conv_forward( + self, + hidden_states: torch.Tensor, + past_key_values=None, + cache_position=None, + attention_mask=None, +): + """Fused full-sequence training forward for LFM2 short convolution.""" + if past_key_values is not None: + return self.slow_forward( + hidden_states, + past_key_values=past_key_values, + cache_position=cache_position, + attention_mask=attention_mask, + ) + + if attention_mask is not None and attention_mask.shape[1] > 1 and attention_mask.shape[0] > 1: + hidden_states = (hidden_states * attention_mask[:, :, None]).to(hidden_states.dtype) + + bcx = self.in_proj(hidden_states) + hidden_states = LigerLfm2ShortConvFunction.apply(bcx, self.conv.weight, self.conv.bias) + return self.out_proj(hidden_states) diff --git a/src/liger_kernel/transformers/model/lfm2_vl.py b/src/liger_kernel/transformers/model/lfm2_vl.py new file mode 100644 index 000000000..47a694125 --- /dev/null +++ b/src/liger_kernel/transformers/model/lfm2_vl.py @@ -0,0 +1,92 @@ +from typing import Optional +from typing import Tuple +from typing import Union + +import torch + +from transformers.cache_utils import Cache +from transformers.utils import can_return_tuple + +from liger_kernel.transformers.model.loss_utils import LigerForCausalLMLoss +from liger_kernel.transformers.model.loss_utils import unpack_cross_entropy_result +from liger_kernel.transformers.model.output_classes import LigerLfm2VlCausalLMOutputWithPast + + +@can_return_tuple +def lce_forward( + self, + input_ids: Optional[torch.LongTensor] = None, + pixel_values: Optional[torch.FloatTensor] = None, + spatial_shapes: Optional[torch.Tensor] = None, + pixel_attention_mask: Optional[torch.Tensor] = None, + attention_mask: Optional[torch.Tensor] = None, + position_ids: Optional[torch.LongTensor] = None, + past_key_values: Optional[Cache] = None, + inputs_embeds: Optional[torch.FloatTensor] = None, + labels: Optional[torch.LongTensor] = None, + use_cache: Optional[bool] = None, + cache_position: Optional[torch.LongTensor] = None, + logits_to_keep: Union[int, torch.Tensor] = 0, + skip_logits: Optional[bool] = None, + **kwargs, +) -> Union[Tuple, LigerLfm2VlCausalLMOutputWithPast]: + """LFM2-VL forward with fused linear cross entropy during training.""" + outputs = self.model( + input_ids=input_ids, + pixel_values=pixel_values, + spatial_shapes=spatial_shapes, + pixel_attention_mask=pixel_attention_mask, + attention_mask=attention_mask, + position_ids=position_ids, + past_key_values=past_key_values, + inputs_embeds=inputs_embeds, + use_cache=use_cache, + cache_position=cache_position, + **kwargs, + ) + + hidden_states = outputs[0] + slice_indices = slice(-logits_to_keep, None) if isinstance(logits_to_keep, int) else logits_to_keep + kept_hidden_states = hidden_states[:, slice_indices, :] + shift_labels = kwargs.pop("shift_labels", None) + + if skip_logits and labels is None and shift_labels is None: + raise ValueError("skip_logits is True, but labels and shift_labels are None") + if skip_logits is None: + skip_logits = self.training and (labels is not None or shift_labels is not None) + + logits = None + loss = None + token_accuracy = None + predicted_tokens = None + if skip_logits: + result = LigerForCausalLMLoss( + hidden_states=kept_hidden_states, + lm_head_weight=self.lm_head.weight, + labels=labels, + shift_labels=shift_labels, + hidden_size=self.config.text_config.hidden_size, + **kwargs, + ) + loss, _, token_accuracy, predicted_tokens = unpack_cross_entropy_result(result) + else: + logits = self.lm_head(kept_hidden_states) + if labels is not None or shift_labels is not None: + loss = self.loss_function( + logits=logits, + labels=labels, + shift_labels=shift_labels, + vocab_size=self.config.text_config.vocab_size, + **kwargs, + ) + + return LigerLfm2VlCausalLMOutputWithPast( + loss=loss, + logits=logits, + past_key_values=outputs.past_key_values, + hidden_states=outputs.hidden_states, + attentions=outputs.attentions, + image_hidden_states=outputs.image_hidden_states, + token_accuracy=token_accuracy, + predicted_tokens=predicted_tokens, + ) diff --git a/src/liger_kernel/transformers/model/output_classes.py b/src/liger_kernel/transformers/model/output_classes.py index c2ef4a7c7..16b302c46 100644 --- a/src/liger_kernel/transformers/model/output_classes.py +++ b/src/liger_kernel/transformers/model/output_classes.py @@ -43,6 +43,13 @@ except Exception: _LlavaCausalLMOutputWithPast = None +try: + from transformers.models.lfm2_vl.modeling_lfm2_vl import ( + Lfm2VlCausalLMOutputWithPast as _Lfm2VlCausalLMOutputWithPast, + ) +except Exception: + _Lfm2VlCausalLMOutputWithPast = None + try: from transformers.models.paligemma.modeling_paligemma import ( PaliGemmaCausalLMOutputWithPast as _PaliGemmaCausalLMOutputWithPast, @@ -137,6 +144,14 @@ class LigerLlavaCausalLMOutputWithPast(_LlavaCausalLMOutputWithPast): predicted_tokens: Optional[torch.LongTensor] = None +if _Lfm2VlCausalLMOutputWithPast is not None: + + @dataclass + class LigerLfm2VlCausalLMOutputWithPast(_Lfm2VlCausalLMOutputWithPast): + token_accuracy: Optional[torch.FloatTensor] = None + predicted_tokens: Optional[torch.LongTensor] = None + + if _InternVLCausalLMOutputWithPast is not None: @dataclass diff --git a/src/liger_kernel/transformers/monkey_patch.py b/src/liger_kernel/transformers/monkey_patch.py index 4d33d7e41..f8520e75d 100755 --- a/src/liger_kernel/transformers/monkey_patch.py +++ b/src/liger_kernel/transformers/monkey_patch.py @@ -11,11 +11,14 @@ from packaging import version from transformers import PreTrainedModel +from liger_kernel.ops.utils import is_hip from liger_kernel.transformers.cross_entropy import LigerCrossEntropyLoss from liger_kernel.transformers.functional import liger_cross_entropy from liger_kernel.transformers.geglu import LigerGEGLUMLP from liger_kernel.transformers.geglu import LigerGEGLUMLPForGemma4 from liger_kernel.transformers.layer_norm import LigerLayerNorm +from liger_kernel.transformers.lfm2_moe_router import liger_lfm2_moe_route_tokens_to_experts +from liger_kernel.transformers.lfm2_short_conv import liger_lfm2_short_conv_forward from liger_kernel.transformers.model.falcon_h1 import lce_forward as falcon_h1_lce_forward from liger_kernel.transformers.model.gemma import lce_forward as gemma_lce_forward from liger_kernel.transformers.model.gemma2 import lce_forward as gemma2_lce_forward @@ -36,6 +39,8 @@ from liger_kernel.transformers.rope import liger_rotary_pos_emb_vision from liger_kernel.transformers.swiglu import LigerBlockSparseTop2MLP from liger_kernel.transformers.swiglu import LigerExperts +from liger_kernel.transformers.swiglu import LigerLfm2MoeExperts +from liger_kernel.transformers.swiglu import LigerLfm2SwiGLUMLP from liger_kernel.transformers.swiglu import LigerPhi3SwiGLUMLP from liger_kernel.transformers.swiglu import LigerSwiGLUMLP @@ -134,6 +139,11 @@ def _patch_layer_norm_module(module, eps=1e-6): def _patch_swiglu_module(module, liger_module): + if liger_module is LigerLfm2MoeExperts: + module.has_gate = True + module.has_bias = False + module.is_transposed = False + _bind_method_to_module(module, "_apply_gate", LigerLfm2MoeExperts._apply_gate) _bind_method_to_module(module, "forward", liger_module.forward) _bind_method_to_module(module, "_get_name", lambda self: liger_module.__name__) @@ -3530,6 +3540,220 @@ def __init__(self, hidden_size, eps=1e-6, **kwargs): _patch_rms_norm_module(decoder_layer.self_attn.k_norm, in_place=False) +def _patch_lfm2_base_model( + base_model, + rms_norm: bool, + swiglu: bool, + short_conv: bool, + fused_moe: bool = True, + fused_moe_router: bool = True, +) -> None: + if rms_norm: + _patch_rms_norm_module(base_model.embedding_norm) + + for decoder_layer in base_model.layers: + feed_forward = decoder_layer.feed_forward + if hasattr(feed_forward, "experts"): + if swiglu and fused_moe: + _patch_swiglu_module(feed_forward.experts, LigerLfm2MoeExperts) + if fused_moe_router: + _bind_method_to_module(feed_forward, "route_tokens_to_experts", liger_lfm2_moe_route_tokens_to_experts) + elif swiglu: + _patch_swiglu_module(feed_forward, LigerLfm2SwiGLUMLP) + + if short_conv and hasattr(decoder_layer, "conv"): + _bind_method_to_module(decoder_layer.conv, "forward", liger_lfm2_short_conv_forward) + + if rms_norm: + _patch_rms_norm_module(decoder_layer.operator_norm) + _patch_rms_norm_module(decoder_layer.ffn_norm) + if hasattr(decoder_layer, "self_attn"): + _patch_rms_norm_module(decoder_layer.self_attn.q_layernorm) + _patch_rms_norm_module(decoder_layer.self_attn.k_layernorm) + + +def apply_liger_kernel_to_lfm2( + rope: bool = True, + cross_entropy: bool = False, + fused_linear_cross_entropy: Optional[bool] = None, + rms_norm: bool = True, + swiglu: bool = True, + short_conv: bool = True, + model: PreTrainedModel = None, +) -> None: + """Apply Liger kernels to Hugging Face LFM2 models. + + Fused linear cross entropy defaults to enabled on ROCm and disabled on + CUDA, where native compiled chunked loss is faster. Pass an explicit bool + to override the backend default. + """ + if fused_linear_cross_entropy is None: + fused_linear_cross_entropy = is_hip() + assert not (cross_entropy and fused_linear_cross_entropy), ( + "cross_entropy and fused_linear_cross_entropy cannot both be True." + ) + + from transformers.models.lfm2 import modeling_lfm2 + from transformers.models.lfm2.modeling_lfm2 import Lfm2Model + + if rope: + modeling_lfm2.apply_rotary_pos_emb = liger_rotary_pos_emb + if rms_norm: + modeling_lfm2.Lfm2RMSNorm = LigerRMSNorm + if swiglu: + modeling_lfm2.Lfm2MLP = LigerLfm2SwiGLUMLP + if short_conv: + modeling_lfm2.Lfm2ShortConv.forward = liger_lfm2_short_conv_forward + if cross_entropy: + from transformers.loss.loss_utils import nn + + nn.functional.cross_entropy = liger_cross_entropy + if fused_linear_cross_entropy: + if model is not None: + model.forward = MethodType(qwen2_lce_forward, model) + else: + modeling_lfm2.Lfm2ForCausalLM.forward = qwen2_lce_forward + + if model is not None: + base_model: Lfm2Model = getattr(model, model.base_model_prefix, model) + _patch_lfm2_base_model(base_model, rms_norm=rms_norm, swiglu=swiglu, short_conv=short_conv) + + +def apply_liger_kernel_to_lfm2_moe( + rope: bool = True, + cross_entropy: bool = False, + fused_linear_cross_entropy: Optional[bool] = None, + rms_norm: bool = True, + swiglu: bool = True, + fused_moe: bool = True, + fused_moe_router: bool = True, + short_conv: bool = True, + model: PreTrainedModel = None, +) -> None: + """Apply Liger kernels to Hugging Face LFM2-MoE models.""" + if fused_linear_cross_entropy is None: + fused_linear_cross_entropy = is_hip() + assert not (cross_entropy and fused_linear_cross_entropy), ( + "cross_entropy and fused_linear_cross_entropy cannot both be True." + ) + + from transformers.models.lfm2_moe import modeling_lfm2_moe + from transformers.models.lfm2_moe.modeling_lfm2_moe import Lfm2MoeModel + + if rope: + modeling_lfm2_moe.apply_rotary_pos_emb = liger_rotary_pos_emb + if rms_norm: + modeling_lfm2_moe.Lfm2MoeRMSNorm = LigerRMSNorm + if swiglu: + modeling_lfm2_moe.Lfm2MoeMLP = LigerLfm2SwiGLUMLP + if fused_moe: + modeling_lfm2_moe.Lfm2MoeExperts = LigerLfm2MoeExperts + if fused_moe_router: + modeling_lfm2_moe.Lfm2MoeSparseMoeBlock.route_tokens_to_experts = liger_lfm2_moe_route_tokens_to_experts + if short_conv: + modeling_lfm2_moe.Lfm2MoeShortConv.forward = liger_lfm2_short_conv_forward + if cross_entropy: + from transformers.loss.loss_utils import nn + + nn.functional.cross_entropy = liger_cross_entropy + if fused_linear_cross_entropy: + if model is not None: + model.forward = MethodType(qwen2_lce_forward, model) + else: + modeling_lfm2_moe.Lfm2MoeForCausalLM.forward = qwen2_lce_forward + + if model is not None: + base_model: Lfm2MoeModel = getattr(model, model.base_model_prefix, model) + _patch_lfm2_base_model( + base_model, + rms_norm=rms_norm, + swiglu=swiglu, + short_conv=short_conv, + fused_moe=fused_moe, + fused_moe_router=fused_moe_router, + ) + + +def apply_liger_kernel_to_lfm2_vl( + rope: bool = True, + cross_entropy: bool = False, + fused_linear_cross_entropy: Optional[bool] = None, + layer_norm: Optional[bool] = None, + rms_norm: bool = True, + swiglu: bool = True, + short_conv: bool = True, + model: PreTrainedModel = None, +) -> None: + """Apply Liger kernels to LFM2-VL's LFM2 decoder and SigLIP2 tower. + + LayerNorm defaults to disabled because PyTorch's implementation is faster + for the SigLIP2 shapes on both CUDA and ROCm. Pass ``layer_norm=True`` to + opt in explicitly. + """ + if fused_linear_cross_entropy is None: + fused_linear_cross_entropy = is_hip() + assert not (cross_entropy and fused_linear_cross_entropy), ( + "cross_entropy and fused_linear_cross_entropy cannot both be True." + ) + if layer_norm is None: + layer_norm = False + + from transformers.models.lfm2_vl import modeling_lfm2_vl + from transformers.models.lfm2_vl.modeling_lfm2_vl import Lfm2VlForConditionalGeneration + from transformers.models.siglip2 import modeling_siglip2 + from transformers.models.siglip2.modeling_siglip2 import Siglip2VisionModel + + from liger_kernel.transformers.model.lfm2_vl import lce_forward as lfm2_vl_lce_forward + + if layer_norm and model is None: + modeling_siglip2.nn.LayerNorm = LigerLayerNorm + + apply_liger_kernel_to_lfm2( + rope=rope, + cross_entropy=False, + fused_linear_cross_entropy=False, + rms_norm=rms_norm, + swiglu=swiglu, + short_conv=short_conv, + ) + + if cross_entropy: + from transformers.loss.loss_utils import nn + + nn.functional.cross_entropy = liger_cross_entropy + if fused_linear_cross_entropy: + if model is not None: + model.forward = MethodType(lfm2_vl_lce_forward, model) + else: + modeling_lfm2_vl.Lfm2VlForConditionalGeneration.forward = lfm2_vl_lce_forward + + if model is not None: + if not isinstance(model, Lfm2VlForConditionalGeneration): + raise TypeError("model must be an Lfm2VlForConditionalGeneration instance") + + apply_liger_kernel_to_lfm2( + rope=rope, + cross_entropy=False, + fused_linear_cross_entropy=False, + rms_norm=rms_norm, + swiglu=swiglu, + model=model.model.language_model, + short_conv=short_conv, + ) + + if layer_norm: + projector_norm = model.model.multi_modal_projector.layer_norm + if projector_norm is not None: + _patch_layer_norm_module(projector_norm) + + vision_tower: Siglip2VisionModel = model.model.vision_tower + vision_model = getattr(vision_tower, "vision_model", vision_tower) + _patch_layer_norm_module(vision_model.post_layernorm) + for encoder_layer in vision_model.encoder.layers: + _patch_layer_norm_module(encoder_layer.layer_norm1) + _patch_layer_norm_module(encoder_layer.layer_norm2) + + # Model type corresponds to the keys defined in transformers/models/auto/modeling_auto.py MODEL_TYPE_TO_APPLY_LIGER_FN = { "deepseek_v4": apply_liger_kernel_to_deepseek_v4, @@ -3544,6 +3768,9 @@ def __init__(self, hidden_size, eps=1e-6, **kwargs): "glm4v_moe": apply_liger_kernel_to_glm4v_moe, "gpt_oss": apply_liger_kernel_to_gpt_oss, "internvl": apply_liger_kernel_to_internvl, + "lfm2": apply_liger_kernel_to_lfm2, + "lfm2_moe": apply_liger_kernel_to_lfm2_moe, + "lfm2_vl": apply_liger_kernel_to_lfm2_vl, "llama": apply_liger_kernel_to_llama, "llama4_text": apply_liger_kernel_to_llama4, "llama4": apply_liger_kernel_to_llama4, diff --git a/src/liger_kernel/transformers/swiglu.py b/src/liger_kernel/transformers/swiglu.py index 4836478e5..5a8cb07d0 100644 --- a/src/liger_kernel/transformers/swiglu.py +++ b/src/liger_kernel/transformers/swiglu.py @@ -77,6 +77,81 @@ def forward(self, hidden_states, top_k_index, top_k_weights): return out.view(orig_shape) +class LigerLfm2SwiGLUMLP(nn.Module): + """LFM2 SwiGLU MLP using the fused SiLU-multiply kernel. + + LFM2 names its projections w1, w3, and w2 and computes + w2(silu(w1(x)) * w3(x)). Its configuration also adjusts the dense + intermediate size before module construction, so that calculation must be + preserved when the class is patched before model initialization. + """ + + def __init__(self, config, intermediate_size=None): + super().__init__() + if intermediate_size is None: + intermediate_size = config.intermediate_size + if getattr(config, "block_auto_adjust_ff_dim", False): + intermediate_size = int(2 * intermediate_size / 3) + if config.block_ffn_dim_multiplier is not None: + intermediate_size = int(config.block_ffn_dim_multiplier * intermediate_size) + intermediate_size = config.block_multiple_of * ( + (intermediate_size + config.block_multiple_of - 1) // config.block_multiple_of + ) + + self.w1 = nn.Linear(config.hidden_size, intermediate_size, bias=False) + self.w3 = nn.Linear(config.hidden_size, intermediate_size, bias=False) + self.w2 = nn.Linear(intermediate_size, config.hidden_size, bias=False) + + def forward(self, x): + return self.w2(LigerSiLUMulFunction.apply(self.w1(x), self.w3(x))) + + +# MI325X and H100 sweeps of the 8B-A1B and 24B-A2B shapes put the +# large-workload crossover near 256 routed rows per expert. Below it, +# preserve the portable fused Triton path to avoid small-batch regressions. +_GROUPED_MM_MIN_ROWS_PER_EXPERT = 256 + + +class LigerLfm2MoeExperts(LigerExperts): + """LFM2-MoE experts with workload-aware grouped-MM dispatch.""" + + def __init__(self, config): + nn.Module.__init__(self) + self.num_experts = config.num_experts + self.hidden_dim = config.hidden_size + self.intermediate_dim = config.moe_intermediate_size + self.has_gate = True + self.has_bias = False + self.is_transposed = False + self.act_fn = torch.nn.functional.silu + self.gate_up_proj = nn.Parameter(torch.empty(self.num_experts, 2 * self.intermediate_dim, self.hidden_dim)) + self.down_proj = nn.Parameter(torch.empty(self.num_experts, self.hidden_dim, self.intermediate_dim)) + + def _apply_gate(self, gate_up_out): + gate, up = gate_up_out.chunk(2, dim=-1) + return self.act_fn(gate) * up + + def forward(self, hidden_states, top_k_index, top_k_weights): + grouped_mm_available = hasattr(torch.nn.functional, "grouped_mm") or hasattr(torch, "_grouped_mm") + if grouped_mm_available: + tokens = hidden_states.numel() // self.hidden_dim + top_k = top_k_index.shape[-1] + enough_work_per_expert = tokens * top_k >= self.num_experts * _GROUPED_MM_MIN_ROWS_PER_EXPERT + if enough_work_per_expert: + try: + from transformers.integrations.moe import grouped_mm_experts_forward + except ImportError: + pass + else: + orig_shape = hidden_states.shape + x = hidden_states.view(-1, self.hidden_dim) + out = grouped_mm_experts_forward( + self, x, top_k_index.view(x.shape[0], -1), top_k_weights.view(x.shape[0], -1) + ) + return out.view(orig_shape) + return LigerExperts.forward(self, hidden_states, top_k_index, top_k_weights) + + class LigerPhi3SwiGLUMLP(nn.Module): """ Patch Phi3MLP to use LigerSiLUMulFunction diff --git a/test/convergence/bf16/test_lfm2_models.py b/test/convergence/bf16/test_lfm2_models.py new file mode 100644 index 000000000..a27438b0d --- /dev/null +++ b/test/convergence/bf16/test_lfm2_models.py @@ -0,0 +1,25 @@ +import pytest +import torch + +from test.convergence.lfm2_utils import run_lfm2_convergence +from test.utils import require_deterministic +from test.utils import supports_bfloat16 + + +@pytest.mark.skipif(not supports_bfloat16(), reason="bfloat16 not supported") +@require_deterministic +@pytest.mark.parametrize("model_kind", ["lfm2", "lfm2_moe", "lfm2_vl"]) +def test_lfm2_models_converge_bf16(model_kind): + expected_losses, expected_parameters = run_lfm2_convergence(model_kind, torch.bfloat16, with_liger=False) + actual_losses, actual_parameters = run_lfm2_convergence(model_kind, torch.bfloat16, with_liger=True) + + torch.testing.assert_close(actual_losses, expected_losses, atol=5e-2, rtol=5e-2) + assert actual_parameters.keys() == expected_parameters.keys() + for name in actual_parameters: + torch.testing.assert_close( + actual_parameters[name], + expected_parameters[name], + atol=2e-2, + rtol=2e-2, + msg=lambda message, parameter=name: f"{parameter}: {message}", + ) diff --git a/test/convergence/fp32/test_lfm2_models.py b/test/convergence/fp32/test_lfm2_models.py new file mode 100644 index 000000000..8441e2960 --- /dev/null +++ b/test/convergence/fp32/test_lfm2_models.py @@ -0,0 +1,25 @@ +import pytest +import torch + +from test.convergence.lfm2_utils import run_lfm2_convergence +from test.utils import require_deterministic + + +@require_deterministic +@pytest.mark.parametrize("model_kind", ["lfm2", "lfm2_moe", "lfm2_vl"]) +def test_lfm2_models_converge_fp32(model_kind): + expected_losses, expected_parameters = run_lfm2_convergence(model_kind, torch.float32, with_liger=False) + actual_losses, actual_parameters = run_lfm2_convergence(model_kind, torch.float32, with_liger=True) + + atol = 5e-4 if model_kind == "lfm2_moe" else 2e-5 + rtol = 5e-4 if model_kind == "lfm2_moe" else 2e-4 + torch.testing.assert_close(actual_losses, expected_losses, atol=atol, rtol=rtol) + assert actual_parameters.keys() == expected_parameters.keys() + for name in actual_parameters: + torch.testing.assert_close( + actual_parameters[name], + expected_parameters[name], + atol=atol, + rtol=rtol, + msg=lambda message, parameter=name: f"{parameter}: {message}", + ) diff --git a/test/convergence/lfm2_utils.py b/test/convergence/lfm2_utils.py new file mode 100644 index 000000000..bfac58df7 --- /dev/null +++ b/test/convergence/lfm2_utils.py @@ -0,0 +1,167 @@ +import gc +import importlib + +import torch +import torch.nn as nn + +from liger_kernel.transformers import apply_liger_kernel_to_lfm2 +from liger_kernel.transformers import apply_liger_kernel_to_lfm2_moe +from liger_kernel.transformers import apply_liger_kernel_to_lfm2_vl +from liger_kernel.utils import infer_device + +device = infer_device() +_ORIGINAL_LAYER_NORM = nn.LayerNorm + + +def _text_config(): + from transformers.models.lfm2.configuration_lfm2 import Lfm2Config + + return Lfm2Config( + vocab_size=128, + hidden_size=64, + intermediate_size=128, + num_hidden_layers=3, + num_attention_heads=4, + num_key_value_heads=2, + max_position_embeddings=128, + block_multiple_of=8, + block_auto_adjust_ff_dim=False, + layer_types=["conv", "full_attention", "conv"], + rope_parameters={"rope_type": "default", "rope_theta": 10000.0}, + conv_bias=True, + use_cache=False, + attn_implementation="sdpa", + ) + + +def _moe_config(): + from transformers.models.lfm2_moe.configuration_lfm2_moe import Lfm2MoeConfig + + return Lfm2MoeConfig( + vocab_size=128, + hidden_size=64, + intermediate_size=128, + moe_intermediate_size=32, + num_hidden_layers=3, + num_attention_heads=4, + num_key_value_heads=2, + max_position_embeddings=128, + num_dense_layers=1, + num_experts=8, + num_experts_per_tok=2, + use_expert_bias=True, + norm_topk_prob=True, + routed_scaling_factor=1.3, + layer_types=["conv", "full_attention", "conv"], + rope_parameters={"rope_type": "default", "rope_theta": 10000.0}, + conv_bias=True, + use_cache=False, + attn_implementation="sdpa", + ) + + +def _vl_config(): + from transformers.models.lfm2_vl.configuration_lfm2_vl import Lfm2VlConfig + + return Lfm2VlConfig( + text_config=_text_config().to_dict(), + vision_config={ + "model_type": "siglip2_vision_model", + "hidden_size": 64, + "intermediate_size": 128, + "num_hidden_layers": 2, + "num_attention_heads": 4, + "num_channels": 3, + "num_patches": 16, + "patch_size": 2, + "hidden_act": "gelu_pytorch_tanh", + "vision_use_head": False, + "attn_implementation": "sdpa", + }, + image_token_id=127, + projector_hidden_size=64, + downsample_factor=2, + ) + + +def _reset_modules(model_kind): + if model_kind == "lfm2": + from transformers.models.lfm2 import modeling_lfm2 + + importlib.reload(modeling_lfm2) + return modeling_lfm2.Lfm2ForCausalLM + if model_kind == "lfm2_moe": + from transformers.models.lfm2_moe import modeling_lfm2_moe + + importlib.reload(modeling_lfm2_moe) + return modeling_lfm2_moe.Lfm2MoeForCausalLM + if model_kind == "lfm2_vl": + from transformers.models.lfm2 import modeling_lfm2 + from transformers.models.lfm2_vl import modeling_lfm2_vl + from transformers.models.siglip2 import modeling_siglip2 + + nn.LayerNorm = _ORIGINAL_LAYER_NORM + importlib.reload(modeling_lfm2) + importlib.reload(modeling_siglip2) + importlib.reload(modeling_lfm2_vl) + return modeling_lfm2_vl.Lfm2VlForConditionalGeneration + raise ValueError(f"unknown model kind: {model_kind}") + + +def _config(model_kind): + if model_kind == "lfm2": + return _text_config() + if model_kind == "lfm2_moe": + return _moe_config() + return _vl_config() + + +def _batch(model_kind, dtype): + input_ids = (torch.arange(32, device=device).reshape(2, 16) % 120) + 3 + attention_mask = torch.ones_like(input_ids) + labels = input_ids.clone() + batch = {"input_ids": input_ids, "attention_mask": attention_mask, "labels": labels} + if model_kind == "lfm2_vl": + input_ids = input_ids[:1] + input_ids[:, :4] = 127 + labels = input_ids.clone() + labels[:, :4] = -100 + batch = { + "input_ids": input_ids, + "attention_mask": torch.ones_like(input_ids), + "labels": labels, + "pixel_values": torch.linspace(-1, 1, 16 * 12, device=device, dtype=dtype).reshape(1, 16, 12), + "spatial_shapes": torch.tensor([[4, 4]], device=device), + "pixel_attention_mask": torch.ones(1, 16, dtype=torch.bool, device=device), + } + return batch + + +def run_lfm2_convergence(model_kind, dtype, with_liger, num_steps=2): + model_class = _reset_modules(model_kind) + if with_liger: + { + "lfm2": apply_liger_kernel_to_lfm2, + "lfm2_moe": apply_liger_kernel_to_lfm2_moe, + "lfm2_vl": apply_liger_kernel_to_lfm2_vl, + }[model_kind]() + + torch.manual_seed(42) + model = model_class(_config(model_kind)).to(device=device, dtype=dtype) + model.train() + optimizer = torch.optim.AdamW(model.parameters(), lr=1e-4) + batch = _batch(model_kind, dtype) + losses = [] + for _ in range(num_steps): + optimizer.zero_grad(set_to_none=True) + loss = model(**batch).loss + loss.backward() + optimizer.step() + losses.append(loss.detach().float().cpu()) + + parameters = {name: parameter.detach().float().cpu() for name, parameter in model.named_parameters()} + del model, optimizer + gc.collect() + torch.cuda.empty_cache() + _reset_modules(model_kind) + return torch.stack(losses), parameters diff --git a/test/transformers/test_lfm2.py b/test/transformers/test_lfm2.py new file mode 100644 index 000000000..96afd6230 --- /dev/null +++ b/test/transformers/test_lfm2.py @@ -0,0 +1,253 @@ +import inspect + +import pytest +import torch + +from liger_kernel.ops.utils import is_hip +from liger_kernel.transformers.auto_model import AutoLigerKernelForCausalLM +from liger_kernel.transformers.lfm2_moe_router import liger_lfm2_moe_route_tokens_to_experts +from liger_kernel.transformers.lfm2_short_conv import liger_lfm2_short_conv_forward +from liger_kernel.transformers.model.qwen2 import lce_forward as lfm2_lce_forward +from liger_kernel.transformers.monkey_patch import _apply_liger_kernel_to_instance +from liger_kernel.transformers.rms_norm import LigerRMSNorm +from liger_kernel.transformers.swiglu import LigerLfm2MoeExperts +from liger_kernel.transformers.swiglu import LigerLfm2SwiGLUMLP +from liger_kernel.utils import infer_device + + +def _has_module(name): + try: + __import__(name) + return True + except ImportError: + return False + + +HAS_LFM2 = _has_module("transformers.models.lfm2.modeling_lfm2") +HAS_LFM2_MOE = _has_module("transformers.models.lfm2_moe.modeling_lfm2_moe") +HAS_LFM2_VL = _has_module("transformers.models.lfm2_vl.modeling_lfm2_vl") +device = infer_device() + + +def _lfm2_config(**overrides): + from transformers.models.lfm2.configuration_lfm2 import Lfm2Config + + config = { + "vocab_size": 128, + "hidden_size": 32, + "intermediate_size": 64, + "num_hidden_layers": 2, + "num_attention_heads": 4, + "num_key_value_heads": 2, + "max_position_embeddings": 128, + "block_multiple_of": 8, + "block_auto_adjust_ff_dim": False, + "layer_types": ["conv", "full_attention"], + "rope_parameters": {"rope_type": "default", "rope_theta": 10000.0}, + "use_cache": False, + } + config.update(overrides) + return Lfm2Config(**config) + + +@pytest.mark.skipif(not HAS_LFM2, reason="lfm2 module not available") +def test_apply_liger_kernel_to_lfm2_instance(): + from transformers.models.lfm2.modeling_lfm2 import Lfm2ForCausalLM + + model = Lfm2ForCausalLM(_lfm2_config()) + original_forward = inspect.getsource(model.forward) + _apply_liger_kernel_to_instance(model) + + expected_forward = inspect.getsource(lfm2_lce_forward) if is_hip() else original_forward + assert inspect.getsource(model.forward) == expected_forward + assert inspect.getsource(model.model.embedding_norm.forward) == inspect.getsource(LigerRMSNorm.forward) + for layer in model.model.layers: + assert inspect.getsource(layer.feed_forward.forward) == inspect.getsource(LigerLfm2SwiGLUMLP.forward) + assert inspect.getsource(layer.operator_norm.forward) == inspect.getsource(LigerRMSNorm.forward) + assert inspect.getsource(layer.ffn_norm.forward) == inspect.getsource(LigerRMSNorm.forward) + if hasattr(layer, "conv"): + assert inspect.getsource(layer.conv.forward) == inspect.getsource(liger_lfm2_short_conv_forward) + if hasattr(layer, "self_attn"): + assert inspect.getsource(layer.self_attn.q_layernorm.forward) == inspect.getsource(LigerRMSNorm.forward) + assert inspect.getsource(layer.self_attn.k_layernorm.forward) == inspect.getsource(LigerRMSNorm.forward) + + +@pytest.mark.skipif(not HAS_LFM2_MOE, reason="lfm2_moe module not available") +def test_apply_liger_kernel_to_lfm2_moe_instance(): + from transformers.models.lfm2_moe.configuration_lfm2_moe import Lfm2MoeConfig + from transformers.models.lfm2_moe.modeling_lfm2_moe import Lfm2MoeForCausalLM + + config = Lfm2MoeConfig( + vocab_size=128, + hidden_size=32, + intermediate_size=64, + moe_intermediate_size=16, + num_hidden_layers=2, + num_attention_heads=4, + num_key_value_heads=2, + num_dense_layers=1, + num_experts=4, + num_experts_per_tok=2, + layer_types=["conv", "full_attention"], + rope_parameters={"rope_type": "default", "rope_theta": 10000.0}, + use_cache=False, + ) + model = Lfm2MoeForCausalLM(config) + original_forward = inspect.getsource(model.forward) + _apply_liger_kernel_to_instance(model) + + expected_forward = inspect.getsource(lfm2_lce_forward) if is_hip() else original_forward + assert inspect.getsource(model.forward) == expected_forward + assert inspect.getsource(model.model.embedding_norm.forward) == inspect.getsource(LigerRMSNorm.forward) + dense_layer, sparse_layer = model.model.layers + assert inspect.getsource(dense_layer.feed_forward.forward) == inspect.getsource(LigerLfm2SwiGLUMLP.forward) + assert inspect.getsource(sparse_layer.feed_forward.experts.forward) == inspect.getsource( + LigerLfm2MoeExperts.forward + ) + assert sparse_layer.feed_forward.experts.has_gate + assert not sparse_layer.feed_forward.experts.has_bias + assert not sparse_layer.feed_forward.experts.is_transposed + assert inspect.getsource(sparse_layer.feed_forward.experts._apply_gate) == inspect.getsource( + LigerLfm2MoeExperts._apply_gate + ) + assert inspect.getsource(sparse_layer.feed_forward.route_tokens_to_experts) == inspect.getsource( + liger_lfm2_moe_route_tokens_to_experts + ) + + assert inspect.getsource(dense_layer.conv.forward) == inspect.getsource(liger_lfm2_short_conv_forward) + + +@pytest.mark.skipif(not HAS_LFM2_VL, reason="lfm2_vl module not available") +@pytest.mark.parametrize( + ("layer_norm", "expected_liger_layer_norm"), + [ + (None, False), + (True, True), + ], +) +def test_apply_liger_kernel_to_lfm2_vl_instance(layer_norm, expected_liger_layer_norm): + from transformers.models.lfm2_vl.configuration_lfm2_vl import Lfm2VlConfig + from transformers.models.lfm2_vl.modeling_lfm2_vl import Lfm2VlForConditionalGeneration + + from liger_kernel.transformers import monkey_patch + from liger_kernel.transformers.model.lfm2_vl import lce_forward as lfm2_vl_lce_forward + + text_config = _lfm2_config().to_dict() + vision_config = { + "model_type": "siglip2_vision_model", + "hidden_size": 32, + "intermediate_size": 64, + "num_hidden_layers": 2, + "num_attention_heads": 4, + "num_channels": 3, + "num_patches": 16, + "patch_size": 2, + "vision_use_head": False, + } + config = Lfm2VlConfig( + text_config=text_config, + vision_config=vision_config, + image_token_id=127, + projector_hidden_size=32, + downsample_factor=2, + ) + model = Lfm2VlForConditionalGeneration(config) + original_forward = inspect.getsource(model.forward) + monkey_patch.apply_liger_kernel_to_lfm2_vl(model=model, layer_norm=layer_norm) + + expected_forward = inspect.getsource(lfm2_vl_lce_forward) if is_hip() else original_forward + assert inspect.getsource(model.forward) == expected_forward + language_model = model.model.language_model + assert inspect.getsource(language_model.embedding_norm.forward) == inspect.getsource(LigerRMSNorm.forward) + for layer in language_model.layers: + assert inspect.getsource(layer.feed_forward.forward) == inspect.getsource(LigerLfm2SwiGLUMLP.forward) + + vision_model = getattr(model.model.vision_tower, "vision_model", model.model.vision_tower) + expected_layer_norm_name = "LigerLayerNorm" if expected_liger_layer_norm else "LayerNorm" + assert vision_model.post_layernorm._get_name() == expected_layer_norm_name + for layer in vision_model.encoder.layers: + assert layer.layer_norm1._get_name() == expected_layer_norm_name + assert layer.layer_norm2._get_name() == expected_layer_norm_name + assert model.model.multi_modal_projector.layer_norm._get_name() == expected_layer_norm_name + + +@pytest.mark.skipif(not HAS_LFM2, reason="lfm2 module not available") +@pytest.mark.parametrize( + ("hip", "fused_linear_cross_entropy", "expect_liger"), + [ + (False, None, False), + (True, None, True), + (False, True, True), + (True, False, False), + ], +) +def test_lfm2_fused_linear_cross_entropy_backend_default(monkeypatch, hip, fused_linear_cross_entropy, expect_liger): + from transformers.models.lfm2.modeling_lfm2 import Lfm2ForCausalLM + + from liger_kernel.transformers import monkey_patch + + model = Lfm2ForCausalLM(_lfm2_config()) + original_forward = inspect.getsource(model.forward) + monkeypatch.setattr(monkey_patch, "is_hip", lambda: hip) + monkey_patch.apply_liger_kernel_to_lfm2( + model=model, + fused_linear_cross_entropy=fused_linear_cross_entropy, + ) + + expected_forward = inspect.getsource(lfm2_lce_forward) if expect_liger else original_forward + assert inspect.getsource(model.forward) == expected_forward + + +@pytest.mark.skipif(not HAS_LFM2 or device == "cpu", reason="requires LFM2 and an accelerator") +def test_lfm2_explicit_fused_linear_cross_entropy_forward_backward(): + from transformers.models.lfm2.modeling_lfm2 import Lfm2ForCausalLM + + from liger_kernel.transformers import monkey_patch + + model = Lfm2ForCausalLM(_lfm2_config()).to(device) + monkey_patch.apply_liger_kernel_to_lfm2(model=model, fused_linear_cross_entropy=True) + input_ids = torch.randint(0, model.config.vocab_size, (2, 16), device=device) + + loss = model(input_ids=input_ids, labels=input_ids).loss + loss.backward() + + assert torch.isfinite(loss) + assert model.lm_head.weight.grad is not None + assert torch.isfinite(model.lm_head.weight.grad).all() + + +@pytest.mark.skipif(not HAS_LFM2, reason="lfm2 module not available") +def test_auto_liger_kernel_for_lfm2_from_config(): + model = AutoLigerKernelForCausalLM.from_config(_lfm2_config()) + + assert isinstance(model.model.layers[0].feed_forward, LigerLfm2SwiGLUMLP) + assert inspect.getsource(model.model.layers[0].conv.forward) == inspect.getsource(liger_lfm2_short_conv_forward) + + +@pytest.mark.skipif(not HAS_LFM2_MOE, reason="lfm2_moe module not available") +def test_auto_liger_kernel_for_lfm2_moe_from_config(): + from transformers.models.lfm2_moe.configuration_lfm2_moe import Lfm2MoeConfig + + config = Lfm2MoeConfig( + vocab_size=128, + hidden_size=32, + intermediate_size=64, + moe_intermediate_size=16, + num_hidden_layers=2, + num_attention_heads=4, + num_key_value_heads=2, + num_dense_layers=1, + num_experts=4, + num_experts_per_tok=2, + layer_types=["conv", "full_attention"], + rope_parameters={"rope_type": "default", "rope_theta": 10000.0}, + use_cache=False, + ) + model = AutoLigerKernelForCausalLM.from_config(config) + + dense_layer, sparse_layer = model.model.layers + assert isinstance(dense_layer.feed_forward, LigerLfm2SwiGLUMLP) + assert isinstance(sparse_layer.feed_forward.experts, LigerLfm2MoeExperts) + assert inspect.getsource(sparse_layer.feed_forward.route_tokens_to_experts) == inspect.getsource( + liger_lfm2_moe_route_tokens_to_experts + ) diff --git a/test/transformers/test_lfm2_moe_dispatch.py b/test/transformers/test_lfm2_moe_dispatch.py new file mode 100644 index 000000000..5adf47b75 --- /dev/null +++ b/test/transformers/test_lfm2_moe_dispatch.py @@ -0,0 +1,100 @@ +from types import SimpleNamespace + +import pytest +import torch + +from liger_kernel.transformers.monkey_patch import _patch_swiglu_module +from liger_kernel.transformers.swiglu import LigerExperts +from liger_kernel.transformers.swiglu import LigerLfm2MoeExperts + + +@pytest.mark.parametrize( + ("tokens", "expected"), + [ + (256, "liger"), + (512, "grouped_mm"), + ], +) +def test_lfm2_moe_shape_dispatch(monkeypatch, tokens, expected): + from transformers.integrations import moe + + config = SimpleNamespace( + hidden_size=8, + moe_intermediate_size=4, + num_experts=4, + ) + from transformers.models.lfm2_moe.modeling_lfm2_moe import Lfm2MoeExperts + + experts = Lfm2MoeExperts(config) + _patch_swiglu_module(experts, LigerLfm2MoeExperts) + hidden_states = torch.randn(tokens, config.hidden_size) + top_k_index = torch.zeros(tokens, 2, dtype=torch.long) + top_k_weights = torch.full((tokens, 2), 0.5) + calls = [] + + monkeypatch.setattr(torch.nn.functional, "grouped_mm", object(), raising=False) + + def fake_grouped_mm(module, hidden, indices, weights): + assert module.has_gate + assert not module.has_bias + assert not module.is_transposed + torch.testing.assert_close( + module._apply_gate(torch.ones(tokens, 8)), + torch.nn.functional.silu(torch.ones(tokens, 4)), + ) + calls.append("grouped_mm") + return hidden + + def fake_liger(module, hidden, indices, weights): + calls.append("liger") + return hidden + + monkeypatch.setattr(moe, "grouped_mm_experts_forward", fake_grouped_mm) + monkeypatch.setattr(LigerExperts, "forward", fake_liger) + + output = experts(hidden_states, top_k_index, top_k_weights) + + assert calls == [expected] + torch.testing.assert_close(output, hidden_states) + + +@pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA or ROCm required") +def test_lfm2_moe_grouped_mm_dispatch_forward_backward(): + from transformers.models.lfm2_moe.modeling_lfm2_moe import Lfm2MoeExperts + + grouped_mm_available = hasattr(torch.nn.functional, "grouped_mm") or hasattr(torch, "_grouped_mm") + if not grouped_mm_available: + pytest.skip("PyTorch grouped MM is unavailable") + + config = SimpleNamespace( + hidden_size=64, + moe_intermediate_size=32, + num_experts=4, + _experts_implementation=None, + ) + reference = Lfm2MoeExperts(config).to(device="cuda", dtype=torch.bfloat16) + actual = Lfm2MoeExperts(config).to(device="cuda", dtype=torch.bfloat16) + with torch.no_grad(): + reference.gate_up_proj.normal_(std=0.02) + reference.down_proj.normal_(std=0.02) + actual.load_state_dict(reference.state_dict()) + _patch_swiglu_module(actual, LigerLfm2MoeExperts) + torch.manual_seed(42) + hidden_reference = torch.randn(512, config.hidden_size, device="cuda", dtype=torch.bfloat16, requires_grad=True) + hidden_actual = hidden_reference.detach().clone().requires_grad_(True) + indices = torch.randint(0, config.num_experts, (512, 2), device="cuda") + weights_reference = ( + torch.softmax(torch.randn(512, 2, device="cuda"), dim=-1).to(torch.bfloat16).requires_grad_(True) + ) + weights_actual = weights_reference.detach().clone().requires_grad_(True) + + output_reference = reference(hidden_reference, indices, weights_reference) + output_actual = actual(hidden_actual, indices, weights_actual) + torch.testing.assert_close(output_actual, output_reference, atol=0.08, rtol=0.02) + + output_reference.float().sum().backward() + output_actual.float().sum().backward() + torch.testing.assert_close(hidden_actual.grad, hidden_reference.grad, atol=0.08, rtol=0.02) + torch.testing.assert_close(weights_actual.grad, weights_reference.grad, atol=0.08, rtol=0.02) + for actual_parameter, reference_parameter in zip(actual.parameters(), reference.parameters(), strict=True): + torch.testing.assert_close(actual_parameter.grad, reference_parameter.grad, atol=0.5, rtol=0.05) diff --git a/test/transformers/test_lfm2_moe_router.py b/test/transformers/test_lfm2_moe_router.py new file mode 100644 index 000000000..84f4e0347 --- /dev/null +++ b/test/transformers/test_lfm2_moe_router.py @@ -0,0 +1,55 @@ +import pytest +import torch + +from liger_kernel.ops import LigerLfm2MoeRouterFunction +from liger_kernel.utils import infer_device + +device = infer_device() + + +def _reference(router_logits, expert_bias, top_k, norm_topk_prob, routed_scaling_factor): + probabilities = router_logits.sigmoid() + scores = probabilities if expert_bias is None else probabilities + expert_bias + selected_experts = torch.topk(scores, k=top_k, dim=-1).indices + routing_weights = torch.gather(probabilities, 1, selected_experts) + if norm_topk_prob: + routing_weights = routing_weights / (routing_weights.sum(dim=-1, keepdim=True) + 1e-6) + return selected_experts, routing_weights * routed_scaling_factor + + +@pytest.mark.parametrize("dtype", [torch.float32, torch.bfloat16]) +@pytest.mark.parametrize("use_expert_bias", [False, True]) +@pytest.mark.parametrize("norm_topk_prob", [False, True]) +@pytest.mark.parametrize(("num_experts", "top_k"), [(8, 2), (48, 4)]) +def test_lfm2_moe_router_forward_backward(dtype, use_expert_bias, norm_topk_prob, num_experts, top_k): + torch.manual_seed(42) + router_logits = torch.randn(37, num_experts, device=device, dtype=dtype) * 0.5 + expert_bias = torch.randn(num_experts, device=device, dtype=torch.float32) * 0.01 if use_expert_bias else None + grad = torch.randn(37, top_k, device=device, dtype=dtype) + routed_scaling_factor = 1.7 + + logits_ref = router_logits.detach().clone().requires_grad_(True) + indices_ref, weights_ref = _reference( + logits_ref, + expert_bias, + top_k, + norm_topk_prob, + routed_scaling_factor, + ) + weights_ref.backward(grad) + + logits_liger = router_logits.detach().clone().requires_grad_(True) + indices_liger, weights_liger = LigerLfm2MoeRouterFunction.apply( + logits_liger, + expert_bias, + top_k, + norm_topk_prob, + routed_scaling_factor, + ) + weights_liger.backward(grad) + + torch.testing.assert_close(indices_liger.long(), indices_ref) + atol = 1e-6 if dtype == torch.float32 else 2e-2 + rtol = 1e-6 if dtype == torch.float32 else 2e-2 + torch.testing.assert_close(weights_liger, weights_ref, atol=atol, rtol=rtol) + torch.testing.assert_close(logits_liger.grad, logits_ref.grad, atol=atol, rtol=rtol) diff --git a/test/transformers/test_lfm2_short_conv.py b/test/transformers/test_lfm2_short_conv.py new file mode 100644 index 000000000..9c2033be6 --- /dev/null +++ b/test/transformers/test_lfm2_short_conv.py @@ -0,0 +1,102 @@ +import pytest +import torch +import torch.nn.functional as F + +import liger_kernel.ops.lfm2_short_conv as short_conv_ops + +from liger_kernel.ops import LigerLfm2ShortConvFunction +from liger_kernel.utils import infer_device + +device = infer_device() + + +@pytest.mark.parametrize( + ("arch", "batch_tokens", "expected"), + [ + ("hopper", 16383, (256, None, None)), + ("hopper", 16384, (128, 4, 2)), + ("hopper", 32767, (128, 4, 2)), + ("hopper", 32768, (64, 2, 2)), + ("cdna3", 65536, (256, None, None)), + ("ampere_ada", 65536, (256, None, None)), + ("blackwell", 65536, (256, None, None)), + ], +) +def test_lfm2_short_conv_weight_backward_dispatch(monkeypatch, arch, batch_tokens, expected): + monkeypatch.setattr(short_conv_ops, "infer_device_arch", lambda: arch) + assert short_conv_ops._short_conv_weight_backward_config(batch_tokens) == expected + + +def _reference(bcx, weight, bias): + gate_b, gate_c, value = bcx.chunk(3, dim=-1) + product = (gate_b * value).transpose(1, 2) + conv = F.conv1d(product, weight, bias=bias, padding=weight.shape[-1] - 1, groups=weight.shape[0]) + conv = conv[..., : bcx.shape[1]].transpose(1, 2) + return gate_c * conv + + +@pytest.mark.parametrize("dtype", [torch.float32, torch.bfloat16]) +@pytest.mark.parametrize("bias_enabled", [False, True]) +@pytest.mark.parametrize("shape", [(2, 17, 32, 3), (1, 128, 64, 4)]) +def test_lfm2_short_conv_forward_backward(dtype, bias_enabled, shape): + batch, seq_len, hidden_size, kernel_size = shape + torch.manual_seed(42) + bcx = torch.randn(batch, seq_len, 3 * hidden_size, device=device, dtype=dtype) + weight = torch.randn(hidden_size, 1, kernel_size, device=device, dtype=dtype) * 0.02 + bias = torch.randn(hidden_size, device=device, dtype=dtype) * 0.02 if bias_enabled else None + grad = torch.randn(batch, seq_len, hidden_size, device=device, dtype=dtype) + + bcx_ref = bcx.detach().clone().requires_grad_(True) + weight_ref = weight.detach().clone().requires_grad_(True) + bias_ref = bias.detach().clone().requires_grad_(True) if bias_enabled else None + output_ref = _reference(bcx_ref, weight_ref, bias_ref) + output_ref.backward(grad) + + bcx_liger = bcx.detach().clone().requires_grad_(True) + weight_liger = weight.detach().clone().requires_grad_(True) + bias_liger = bias.detach().clone().requires_grad_(True) if bias_enabled else None + output_liger = LigerLfm2ShortConvFunction.apply(bcx_liger, weight_liger, bias_liger) + output_liger.backward(grad) + + atol = 1e-5 if dtype == torch.float32 else 5e-2 + rtol = 1e-5 if dtype == torch.float32 else 5e-2 + torch.testing.assert_close(output_liger, output_ref, atol=atol, rtol=rtol) + torch.testing.assert_close(bcx_liger.grad, bcx_ref.grad, atol=atol, rtol=rtol) + torch.testing.assert_close(weight_liger.grad, weight_ref.grad, atol=atol, rtol=rtol) + if bias_enabled: + torch.testing.assert_close(bias_liger.grad, bias_ref.grad, atol=atol, rtol=rtol) + + +def test_lfm2_short_conv_long_sequence_backward(monkeypatch): + """Exercise the Hopper-only long-sequence weight-reduction configuration.""" + torch.manual_seed(123) + shape = (1, 16384, 32, 3) + batch, seq_len, hidden_size, kernel_size = shape + bcx = torch.randn(batch, seq_len, 3 * hidden_size, device=device, dtype=torch.bfloat16) + weight = torch.randn(hidden_size, 1, kernel_size, device=device, dtype=torch.bfloat16) * 0.02 + grad = torch.randn(batch, seq_len, hidden_size, device=device, dtype=torch.bfloat16) + + bcx_ref = bcx.detach().clone().requires_grad_(True) + weight_ref = weight.detach().clone().requires_grad_(True) + output_ref = _reference(bcx_ref, weight_ref, None) + output_ref.backward(grad) + + monkeypatch.setattr(short_conv_ops, "infer_device_arch", lambda: "ampere_ada") + bcx_portable = bcx.detach().clone().requires_grad_(True) + weight_portable = weight.detach().clone().requires_grad_(True) + output_portable = LigerLfm2ShortConvFunction.apply(bcx_portable, weight_portable, None) + output_portable.backward(grad) + + monkeypatch.setattr(short_conv_ops, "infer_device_arch", lambda: "hopper") + bcx_liger = bcx.detach().clone().requires_grad_(True) + weight_liger = weight.detach().clone().requires_grad_(True) + output_liger = LigerLfm2ShortConvFunction.apply(bcx_liger, weight_liger, None) + output_liger.backward(grad) + + torch.testing.assert_close(output_liger, output_ref, atol=5e-2, rtol=5e-2) + torch.testing.assert_close(bcx_liger.grad, bcx_ref.grad, atol=5e-2, rtol=5e-2) + # Long BF16 reductions have the same rounding error on the portable and + # Hopper paths. Require the tuned launch to match the portable result + # exactly, then compare that common result with a reduction-aware tolerance. + torch.testing.assert_close(weight_liger.grad, weight_portable.grad, atol=0, rtol=0) + torch.testing.assert_close(weight_liger.grad, weight_ref.grad, atol=5e-1, rtol=8e-2)