From 6177f27a9e16c4b081a5a27d3e387474810f2cda Mon Sep 17 00:00:00 2001
From: Justin Hu <181588904+justinhh4@users.noreply.github.com>
Date: Sun, 5 Jul 2026 20:49:13 +0000
Subject: [PATCH 1/5] [CuteDSL] CE backward memory parity + benchmark tooling
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

Follow-up to #1279 addressing the review follow-ups (memory + benchmark scripts).

Memory parity fix
-----------------
The CuTe DSL CE backward, for the scalar-grad (reduction mean/sum) path, computed
`_input * grad_output`, allocating a *second* BT×V buffer. That doubled peak memory
vs the Triton CE, which scales the saved gradient IN PLACE via a raw element-wise
kernel (no autograd version bump). Measured on B200 (full fwd+bwd, BT=8192,
V=128256): cutedsl 4008 MB vs triton 2004 MB in bf16.

Fix: mirror Triton exactly — scale in place with `element_mul_kernel`. A raw Triton
kernel (not `_input *= grad_output`) is used so the in-place write doesn't bump the
autograd version counter and trip backward-through-backward. Peak memory now equals
Triton across the whole vocab/BT sweep. CE parity suite: 158 passed.

Benchmark tooling (reviewer ask: "add benchmarking scripts for cutedsl, like cutile")
-------------------------------------------------------------------------------------
- benchmark/scripts/run_cutedsl_compare.py: Triton-vs-CuTe-DSL compare driver
  (mirrors run_cutile_compare.py), runs benchmark_<kernel>.py twice under
  LIGER_KERNEL_IMPL and tags the providers (liger_triton / liger_cutedsl) into
  benchmark/data/all_benchmark_data_cutedsl.csv. Supports cross_entropy.
- benchmark/README.md: document the compare drivers (CuTile + CuTe-DSL).
- benchmark/data/all_benchmark_data_cutedsl.csv: sample CE speed+memory dataset
  (liger_triton / liger_cutedsl / torch), generated by the new driver on B200.

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
---
 benchmark/README.md                           |  37 +++-
 benchmark/data/all_benchmark_data_cutedsl.csv | 166 ++++++++++++++++++
 benchmark/scripts/run_cutedsl_compare.py      |  71 ++++++++
 .../ops/cutedsl/ops/cross_entropy.py          |  25 ++-
 4 files changed, 295 insertions(+), 4 deletions(-)
 create mode 100644 benchmark/data/all_benchmark_data_cutedsl.csv
 create mode 100644 benchmark/scripts/run_cutedsl_compare.py

diff --git a/benchmark/README.md b/benchmark/README.md
index 9597d94e3..0ba6f9f9a 100644
--- a/benchmark/README.md
+++ b/benchmark/README.md
@@ -148,4 +148,39 @@ The benchmarking system is designed to provide a **consistent, low-boilerplate w
 
 4. View results
    - Generated plots will be saved in `benchmark/visualizations/`
-   - Filenames include the sweep mode when specified (e.g. `geglu_speed_full_model_config.png`)
\ No newline at end of file
+   - Filenames include the sweep mode when specified (e.g. `geglu_speed_full_model_config.png`)
+
+### Comparing kernel backends (Triton vs CuTile / CuTe-DSL)
+
+Some kernels have an alternative backend (a different DSL on the same CUDA
+device) that is opt-in via the `LIGER_KERNEL_IMPL` environment variable:
+
+* **CuTile** (`LIGER_KERNEL_IMPL=cutile`) — for `cross_entropy`, `fused_linear_jsd`,
+  `geglu`, `jsd`, `layer_norm`.
+* **CuTe-DSL** (`LIGER_KERNEL_IMPL=cutedsl`) — for `cross_entropy`.
+
+To benchmark a kernel's Triton and alternative backend **side by side in one
+CSV**, use the corresponding compare driver. It runs the standard
+`benchmark_<kernel>.py` twice — once on Triton, once on the alternative — and
+tags the `liger` provider for each run (`liger_triton` / `liger_cutile` /
+`liger_cutedsl`) so both land in a dedicated CSV without colliding.
+
+```bash
+cd benchmark/scripts
+
+# Triton vs CuTile (writes data/all_benchmark_data_cutile.csv)
+python run_cutile_compare.py --kernel cross_entropy [--model llama_3_8b] [--overwrite]
+
+# Triton vs CuTe-DSL (writes data/all_benchmark_data_cutedsl.csv)
+python run_cutedsl_compare.py --kernel cross_entropy [--overwrite]
+```
+
+Any extra args (`--model`, `--sweep-mode`, `--bt`, `--overwrite`) are forwarded
+to the underlying benchmark script. Plot the merged CSV by pointing the
+visualizer at it with `--data-file`:
+
+```bash
+python ../benchmarks_visualizer.py \
+    --kernel-name cross_entropy --metric-name speed \
+    --data-file data/all_benchmark_data_cutedsl.csv
+```
\ No newline at end of file
diff --git a/benchmark/data/all_benchmark_data_cutedsl.csv b/benchmark/data/all_benchmark_data_cutedsl.csv
new file mode 100644
index 000000000..4d21c1022
--- /dev/null
+++ b/benchmark/data/all_benchmark_data_cutedsl.csv
@@ -0,0 +1,166 @@
+kernel_name,kernel_provider,kernel_operation_mode,metric_name,metric_unit,x_name,x_label,x_value,y_value_50,y_value_20,y_value_80,extra_benchmark_config_str,gpu_name,timestamp,liger_version
+cross_entropy,liger_triton,forward,speed,ms,model_config,model configuration,llama_2_7b,0.23815999925136566,0.23594879806041716,0.2401663988828659,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:19,0.8.0
+cross_entropy,liger_triton,forward,speed,ms,model_config,model configuration,llama_3_8b,0.5447679758071899,0.54197758436203,0.546617591381073,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:19,0.8.0
+cross_entropy,liger_triton,forward,speed,ms,model_config,model configuration,qwen2.5_7b,0.6508639752864838,0.6478015899658203,0.6537535786628723,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:19,0.8.0
+cross_entropy,liger_triton,forward,speed,ms,model_config,model configuration,qwen2.5_14b,0.6493279933929443,0.6451200246810913,0.6543359756469727,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:19,0.8.0
+cross_entropy,liger_triton,forward,speed,ms,model_config,model configuration,qwen2.5_72b,0.6482400000095367,0.6461439728736877,0.6533248066902161,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:19,0.8.0
+cross_entropy,liger_triton,forward,speed,ms,model_config,model configuration,deepseek_v2_lite,0.49457600712776184,0.492576003074646,0.495142388343811,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:19,0.8.0
+cross_entropy,liger_triton,forward,speed,ms,model_config,model configuration,deepseek_v3,0.5417280197143555,0.5398848056793213,0.5435456275939942,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:19,0.8.0
+cross_entropy,torch,forward,speed,ms,model_config,model configuration,llama_2_7b,0.15859200060367584,0.15772800147533417,0.158720001578331,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:28,0.8.0
+cross_entropy,torch,forward,speed,ms,model_config,model configuration,llama_3_8b,0.6932479739189148,0.6932160258293152,0.6944000124931335,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:28,0.8.0
+cross_entropy,torch,forward,speed,ms,model_config,model configuration,qwen2.5_7b,0.8027520179748535,0.8017920255661011,0.8028159737586975,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:28,0.8.0
+cross_entropy,torch,forward,speed,ms,model_config,model configuration,qwen2.5_14b,0.8028799891471863,0.8028159737586975,0.8038079738616943,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:28,0.8.0
+cross_entropy,torch,forward,speed,ms,model_config,model configuration,qwen2.5_72b,0.8028159737586975,0.8018559813499451,0.8038399815559387,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:28,0.8.0
+cross_entropy,torch,forward,speed,ms,model_config,model configuration,deepseek_v2_lite,0.5535359978675842,0.5529599785804749,0.5540159940719604,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:28,0.8.0
+cross_entropy,torch,forward,speed,ms,model_config,model configuration,deepseek_v3,0.6963199973106384,0.6952959895133972,0.6973312020301818,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:28,0.8.0
+cross_entropy,liger_triton,backward,speed,ms,model_config,model configuration,llama_2_7b,0.21299199759960175,0.21217280328273774,0.21398399770259857,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:20,0.8.0
+cross_entropy,liger_triton,backward,speed,ms,model_config,model configuration,llama_3_8b,0.6993920207023621,0.6973440051078796,0.7014399766921997,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:20,0.8.0
+cross_entropy,liger_triton,backward,speed,ms,model_config,model configuration,qwen2.5_7b,0.8202880024909973,0.8202175855636596,0.8222656011581421,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:20,0.8.0
+cross_entropy,liger_triton,backward,speed,ms,model_config,model configuration,qwen2.5_14b,0.8202239871025085,0.8179327726364136,0.8212928056716919,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:20,0.8.0
+cross_entropy,liger_triton,backward,speed,ms,model_config,model configuration,qwen2.5_72b,0.8203840255737305,0.8198463916778564,0.8216128110885621,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:20,0.8.0
+cross_entropy,liger_triton,backward,speed,ms,model_config,model configuration,deepseek_v2_lite,0.5662560164928436,0.5647871971130372,0.567296028137207,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:20,0.8.0
+cross_entropy,liger_triton,backward,speed,ms,model_config,model configuration,deepseek_v3,0.7004479765892029,0.6994240283966064,0.7014719843864441,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:20,0.8.0
+cross_entropy,torch,backward,speed,ms,model_config,model configuration,llama_2_7b,0.22631999850273132,0.2253119945526123,0.22732800245285034,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:29,0.8.0
+cross_entropy,torch,backward,speed,ms,model_config,model configuration,llama_3_8b,0.8253440260887146,0.8243520259857178,0.8255168080329895,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:29,0.8.0
+cross_entropy,torch,backward,speed,ms,model_config,model configuration,qwen2.5_7b,0.970304012298584,0.969760000705719,0.9716799855232239,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:29,0.8.0
+cross_entropy,torch,backward,speed,ms,model_config,model configuration,qwen2.5_14b,0.9702720046043396,0.9697279930114746,0.9708160161972046,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:29,0.8.0
+cross_entropy,torch,backward,speed,ms,model_config,model configuration,qwen2.5_72b,0.9712960124015808,0.969760000705719,0.9728320240974426,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:29,0.8.0
+cross_entropy,torch,backward,speed,ms,model_config,model configuration,deepseek_v2_lite,0.6615999937057495,0.6605439782142639,0.6629823923110962,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:29,0.8.0
+cross_entropy,torch,backward,speed,ms,model_config,model configuration,deepseek_v3,0.8254079818725586,0.8243648171424866,0.8272319912910462,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:29,0.8.0
+cross_entropy,liger_triton,full,speed,ms,model_config,model configuration,llama_2_7b,0.48102399706840515,0.4790463924407959,0.48521599769592283,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:20,0.8.0
+cross_entropy,liger_triton,full,speed,ms,model_config,model configuration,llama_3_8b,1.273855984210968,1.2728639841079712,1.274880051612854,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:20,0.8.0
+cross_entropy,liger_triton,full,speed,ms,model_config,model configuration,qwen2.5_7b,1.5001920461654663,1.4999615907669068,1.502623987197876,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:20,0.8.0
+cross_entropy,liger_triton,full,speed,ms,model_config,model configuration,qwen2.5_14b,1.505344033241272,1.5040831804275514,1.5067455768585205,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:20,0.8.0
+cross_entropy,liger_triton,full,speed,ms,model_config,model configuration,qwen2.5_72b,1.5012160539627075,1.4982975721359253,1.502854347229004,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:20,0.8.0
+cross_entropy,liger_triton,full,speed,ms,model_config,model configuration,deepseek_v2_lite,1.094208002090454,1.0889536142349243,1.0966271877288818,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:20,0.8.0
+cross_entropy,liger_triton,full,speed,ms,model_config,model configuration,deepseek_v3,1.275920033454895,1.2738560438156128,1.2789759635925293,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:20,0.8.0
+cross_entropy,torch,full,speed,ms,model_config,model configuration,llama_2_7b,0.3845280110836029,0.38393598794937134,0.3859967887401581,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:30,0.8.0
+cross_entropy,torch,full,speed,ms,model_config,model configuration,llama_3_8b,1.5231679677963257,1.5225919961929322,1.524076771736145,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:30,0.8.0
+cross_entropy,torch,full,speed,ms,model_config,model configuration,qwen2.5_7b,1.7715519666671753,1.7705663681030273,1.7725375652313233,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:30,0.8.0
+cross_entropy,torch,full,speed,ms,model_config,model configuration,qwen2.5_14b,1.7704959511756897,1.7704511642456056,1.771347165107727,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:30,0.8.0
+cross_entropy,torch,full,speed,ms,model_config,model configuration,qwen2.5_72b,1.7715519666671753,1.77054078578949,1.77248637676239,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:30,0.8.0
+cross_entropy,torch,full,speed,ms,model_config,model configuration,deepseek_v2_lite,1.215455949306488,1.21452796459198,1.2154879570007324,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:30,0.8.0
+cross_entropy,torch,full,speed,ms,model_config,model configuration,deepseek_v3,1.5231519937515259,1.5219584226608276,1.524127984046936,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:30,0.8.0
+cross_entropy,liger_triton,no-grad-forward,speed,ms,model_config,model configuration,llama_2_7b,0.23555199801921844,0.23449599742889404,0.23862400650978088,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:21,0.8.0
+cross_entropy,liger_triton,no-grad-forward,speed,ms,model_config,model configuration,llama_3_8b,0.5416960120201111,0.5409152150154114,0.5437440276145935,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:21,0.8.0
+cross_entropy,liger_triton,no-grad-forward,speed,ms,model_config,model configuration,qwen2.5_7b,0.6476959884166718,0.6461631774902343,0.6500927925109863,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:21,0.8.0
+cross_entropy,liger_triton,no-grad-forward,speed,ms,model_config,model configuration,qwen2.5_14b,0.649728000164032,0.6467711806297303,0.6533439755439758,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:21,0.8.0
+cross_entropy,liger_triton,no-grad-forward,speed,ms,model_config,model configuration,qwen2.5_72b,0.6472159922122955,0.6457535982131959,0.6523008227348328,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:21,0.8.0
+cross_entropy,liger_triton,no-grad-forward,speed,ms,model_config,model configuration,deepseek_v2_lite,0.493584007024765,0.4915199875831604,0.4952191889286041,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:21,0.8.0
+cross_entropy,liger_triton,no-grad-forward,speed,ms,model_config,model configuration,deepseek_v3,0.5447679758071899,0.54170241355896,0.5455935835838318,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:21,0.8.0
+cross_entropy,torch,no-grad-forward,speed,ms,model_config,model configuration,llama_2_7b,0.15859200060367584,0.15769599378108978,0.158720001578331,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:30,0.8.0
+cross_entropy,torch,no-grad-forward,speed,ms,model_config,model configuration,llama_3_8b,0.6943359971046448,0.6937151908874511,0.6952959895133972,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:30,0.8.0
+cross_entropy,torch,no-grad-forward,speed,ms,model_config,model configuration,qwen2.5_7b,0.8037440180778503,0.8017920255661011,0.8040000200271606,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:30,0.8.0
+cross_entropy,torch,no-grad-forward,speed,ms,model_config,model configuration,qwen2.5_14b,0.8030400276184082,0.8019840121269226,0.8038399815559387,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:30,0.8.0
+cross_entropy,torch,no-grad-forward,speed,ms,model_config,model configuration,qwen2.5_72b,0.8029760122299194,0.8028479814529419,0.8038399815559387,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:30,0.8.0
+cross_entropy,torch,no-grad-forward,speed,ms,model_config,model configuration,deepseek_v2_lite,0.5519839823246002,0.551904022693634,0.5529599785804749,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:30,0.8.0
+cross_entropy,torch,no-grad-forward,speed,ms,model_config,model configuration,deepseek_v3,0.6963199973106384,0.6957247972488403,0.6973440051078796,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:30,0.8.0
+cross_entropy,liger_triton,full,memory,MB,model_config,model configuration,llama_2_7b,500.04150390625,500.04150390625,500.04150390625,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:21,0.8.0
+cross_entropy,liger_triton,full,memory,MB,model_config,model configuration,llama_3_8b,2004.04150390625,2004.04150390625,2004.04150390625,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:21,0.8.0
+cross_entropy,liger_triton,full,memory,MB,model_config,model configuration,qwen2.5_7b,2376.04150390625,2376.04150390625,2376.04150390625,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:21,0.8.0
+cross_entropy,liger_triton,full,memory,MB,model_config,model configuration,qwen2.5_14b,2376.04150390625,2376.04150390625,2376.04150390625,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:21,0.8.0
+cross_entropy,liger_triton,full,memory,MB,model_config,model configuration,qwen2.5_72b,2376.04150390625,2376.04150390625,2376.04150390625,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:21,0.8.0
+cross_entropy,liger_triton,full,memory,MB,model_config,model configuration,deepseek_v2_lite,1600.04150390625,1600.04150390625,1600.04150390625,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:21,0.8.0
+cross_entropy,liger_triton,full,memory,MB,model_config,model configuration,deepseek_v3,2020.04150390625,2020.04150390625,2020.04150390625,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:21,0.8.0
+cross_entropy,torch,full,memory,MB,model_config,model configuration,llama_2_7b,1250.01708984375,1250.01708984375,1250.01708984375,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:30,0.8.0
+cross_entropy,torch,full,memory,MB,model_config,model configuration,llama_3_8b,5010.01708984375,5010.01708984375,5010.01708984375,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:30,0.8.0
+cross_entropy,torch,full,memory,MB,model_config,model configuration,qwen2.5_7b,5940.01708984375,5940.01708984375,5940.01708984375,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:30,0.8.0
+cross_entropy,torch,full,memory,MB,model_config,model configuration,qwen2.5_14b,5940.01708984375,5940.01708984375,5940.01708984375,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:30,0.8.0
+cross_entropy,torch,full,memory,MB,model_config,model configuration,qwen2.5_72b,5940.01708984375,5940.01708984375,5940.01708984375,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:30,0.8.0
+cross_entropy,torch,full,memory,MB,model_config,model configuration,deepseek_v2_lite,4000.01708984375,4000.01708984375,4000.01708984375,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:30,0.8.0
+cross_entropy,torch,full,memory,MB,model_config,model configuration,deepseek_v3,5050.01708984375,5050.01708984375,5050.01708984375,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:30,0.8.0
+cross_entropy,liger_cutedsl,forward,speed,ms,model_config,model configuration,llama_2_7b,0.24065600335597992,0.2375359982252121,0.2457599937915802,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:28,0.8.0
+cross_entropy,liger_cutedsl,forward,speed,ms,model_config,model configuration,llama_3_8b,0.5181760191917419,0.5144767999649048,0.52019202709198,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:28,0.8.0
+cross_entropy,liger_cutedsl,forward,speed,ms,model_config,model configuration,qwen2.5_7b,0.6172800064086914,0.6116415858268738,0.6214655876159668,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:28,0.8.0
+cross_entropy,liger_cutedsl,forward,speed,ms,model_config,model configuration,qwen2.5_14b,0.6225919723510742,0.6182784080505371,0.6246655941009521,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:28,0.8.0
+cross_entropy,liger_cutedsl,forward,speed,ms,model_config,model configuration,qwen2.5_72b,0.6195520162582397,0.6141696214675904,0.6240511894226074,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:28,0.8.0
+cross_entropy,liger_cutedsl,forward,speed,ms,model_config,model configuration,deepseek_v2_lite,0.43110400438308716,0.43007999658584595,0.43222400546073914,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:28,0.8.0
+cross_entropy,liger_cutedsl,forward,speed,ms,model_config,model configuration,deepseek_v3,0.5192480087280273,0.516262412071228,0.5202560067176819,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:28,0.8.0
+cross_entropy,liger_cutedsl,backward,speed,ms,model_config,model configuration,llama_2_7b,0.2170880064368248,0.217056006193161,0.22038400173187256,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:29,0.8.0
+cross_entropy,liger_cutedsl,backward,speed,ms,model_config,model configuration,llama_3_8b,0.7014399766921997,0.7014080286026001,0.7024319767951965,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:29,0.8.0
+cross_entropy,liger_cutedsl,backward,speed,ms,model_config,model configuration,qwen2.5_7b,0.8232960104942322,0.8212800025939941,0.8257855772972107,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:29,0.8.0
+cross_entropy,liger_cutedsl,backward,speed,ms,model_config,model configuration,qwen2.5_14b,0.8212800025939941,0.8202880024909973,0.822329616546631,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:29,0.8.0
+cross_entropy,liger_cutedsl,backward,speed,ms,model_config,model configuration,qwen2.5_72b,0.8241919875144958,0.8225088119506836,0.8243520259857178,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:29,0.8.0
+cross_entropy,liger_cutedsl,backward,speed,ms,model_config,model configuration,deepseek_v2_lite,0.567296028137207,0.5662400126457214,0.5692928075790405,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:29,0.8.0
+cross_entropy,liger_cutedsl,backward,speed,ms,model_config,model configuration,deepseek_v3,0.7034879922866821,0.7014719843864441,0.7036479711532593,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:29,0.8.0
+cross_entropy,liger_cutedsl,full,speed,ms,model_config,model configuration,llama_2_7b,0.47839999198913574,0.4753279983997345,0.4821376025676727,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:29,0.8.0
+cross_entropy,liger_cutedsl,full,speed,ms,model_config,model configuration,llama_3_8b,1.2410880327224731,1.2384703636169434,1.2444287538528442,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:29,0.8.0
+cross_entropy,liger_cutedsl,full,speed,ms,model_config,model configuration,qwen2.5_7b,1.474560022354126,1.474560022354126,1.4776639938354492,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:29,0.8.0
+cross_entropy,liger_cutedsl,full,speed,ms,model_config,model configuration,qwen2.5_14b,1.4683359861373901,1.4632960557937622,1.4727040529251099,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:29,0.8.0
+cross_entropy,liger_cutedsl,full,speed,ms,model_config,model configuration,qwen2.5_72b,1.4653120040893555,1.4618111848831177,1.467404770851135,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:29,0.8.0
+cross_entropy,liger_cutedsl,full,speed,ms,model_config,model configuration,deepseek_v2_lite,1.028223991394043,1.0264383792877196,1.0317760467529298,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:29,0.8.0
+cross_entropy,liger_cutedsl,full,speed,ms,model_config,model configuration,deepseek_v3,1.2492480278015137,1.2443328380584717,1.252031970024109,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:29,0.8.0
+cross_entropy,liger_cutedsl,no-grad-forward,speed,ms,model_config,model configuration,llama_2_7b,0.23553600162267685,0.23449599742889404,0.238060799241066,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:30,0.8.0
+cross_entropy,liger_cutedsl,no-grad-forward,speed,ms,model_config,model configuration,llama_3_8b,0.5160959959030151,0.515020775794983,0.5177152037620545,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:30,0.8.0
+cross_entropy,liger_cutedsl,no-grad-forward,speed,ms,model_config,model configuration,qwen2.5_7b,0.6217600107192993,0.6203264117240905,0.6260223865509034,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:30,0.8.0
+cross_entropy,liger_cutedsl,no-grad-forward,speed,ms,model_config,model configuration,qwen2.5_14b,0.6224640011787415,0.6078847885131836,0.6260223865509034,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:30,0.8.0
+cross_entropy,liger_cutedsl,no-grad-forward,speed,ms,model_config,model configuration,qwen2.5_72b,0.6184960007667542,0.6127616167068481,0.6207103967666626,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:30,0.8.0
+cross_entropy,liger_cutedsl,no-grad-forward,speed,ms,model_config,model configuration,deepseek_v2_lite,0.4301759898662567,0.42905598878860474,0.4341759979724884,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:30,0.8.0
+cross_entropy,liger_cutedsl,no-grad-forward,speed,ms,model_config,model configuration,deepseek_v3,0.5190880000591278,0.5171200037002563,0.527135980129242,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:30,0.8.0
+cross_entropy,liger_cutedsl,full,memory,MB,model_config,model configuration,llama_2_7b,500.04931640625,500.04931640625,500.04931640625,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:30,0.8.0
+cross_entropy,liger_cutedsl,full,memory,MB,model_config,model configuration,llama_3_8b,2004.04931640625,2004.04931640625,2004.04931640625,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:30,0.8.0
+cross_entropy,liger_cutedsl,full,memory,MB,model_config,model configuration,qwen2.5_7b,2376.04931640625,2376.04931640625,2376.04931640625,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:30,0.8.0
+cross_entropy,liger_cutedsl,full,memory,MB,model_config,model configuration,qwen2.5_14b,2376.04931640625,2376.04931640625,2376.04931640625,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:30,0.8.0
+cross_entropy,liger_cutedsl,full,memory,MB,model_config,model configuration,qwen2.5_72b,2376.04931640625,2376.04931640625,2376.04931640625,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:30,0.8.0
+cross_entropy,liger_cutedsl,full,memory,MB,model_config,model configuration,deepseek_v2_lite,1600.04931640625,1600.04931640625,1600.04931640625,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:30,0.8.0
+cross_entropy,liger_cutedsl,full,memory,MB,model_config,model configuration,deepseek_v3,2020.04931640625,2020.04931640625,2020.04931640625,"{""bsz"": 1, ""seq_len"": 2048}",NVIDIA B200,2026-07-05 21:00:30,0.8.0
+cross_entropy,liger_triton,forward,speed,ms,BT,B * T,1024,0.32872000336647034,0.32688639163970945,0.3326016008853912,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:34,0.8.0
+cross_entropy,liger_triton,forward,speed,ms,BT,B * T,2048,0.5417280197143555,0.5398784041404724,0.5445951819419861,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:34,0.8.0
+cross_entropy,liger_triton,forward,speed,ms,BT,B * T,4096,0.9615359902381897,0.9595199823379517,0.9623744010925293,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:34,0.8.0
+cross_entropy,liger_triton,forward,speed,ms,BT,B * T,8192,1.7879040241241455,1.7844288110733033,1.7879743814468383,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:34,0.8.0
+cross_entropy,torch,forward,speed,ms,BT,B * T,1024,0.3779039978981018,0.37781120538711543,0.3788927972316742,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:34,0.8.0
+cross_entropy,torch,forward,speed,ms,BT,B * T,2048,0.6953920125961304,0.6952767848968506,0.6973952174186706,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:34,0.8.0
+cross_entropy,torch,forward,speed,ms,BT,B * T,4096,1.3506560325622559,1.3488383769989014,1.3514880418777466,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:34,0.8.0
+cross_entropy,torch,forward,speed,ms,BT,B * T,8192,2.6470720767974854,2.647052764892578,2.64768648147583,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:34,0.8.0
+cross_entropy,liger_triton,backward,speed,ms,BT,B * T,1024,0.37268801033496857,0.3718079924583435,0.37276801466941833,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:35,0.8.0
+cross_entropy,liger_triton,backward,speed,ms,BT,B * T,2048,0.6901760101318359,0.689087986946106,0.6912000179290771,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:35,0.8.0
+cross_entropy,liger_triton,backward,speed,ms,BT,B * T,4096,1.333296000957489,1.3322559595108032,1.335263967514038,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:35,0.8.0
+cross_entropy,liger_triton,backward,speed,ms,BT,B * T,8192,2.6091840267181396,2.607936048507691,2.6097984313964844,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:35,0.8.0
+cross_entropy,torch,backward,speed,ms,BT,B * T,1024,0.4577600061893463,0.4571136057376861,0.45777921080589296,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:35,0.8.0
+cross_entropy,torch,backward,speed,ms,BT,B * T,2048,0.8261759877204895,0.8255359768867492,0.8270655989646911,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:35,0.8.0
+cross_entropy,torch,backward,speed,ms,BT,B * T,4096,1.623039960861206,1.6218111515045166,1.6230783700942992,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:35,0.8.0
+cross_entropy,torch,backward,speed,ms,BT,B * T,8192,3.167167901992798,3.167167901992798,3.167167901992798,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:35,0.8.0
+cross_entropy,liger_triton,full,speed,ms,BT,B * T,1024,0.7117120027542114,0.7094399929046631,0.7137280106544495,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:35,0.8.0
+cross_entropy,liger_triton,full,speed,ms,BT,B * T,2048,1.24617600440979,1.2425408363342285,1.2462079524993896,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:35,0.8.0
+cross_entropy,liger_triton,full,speed,ms,BT,B * T,4096,2.300960063934326,2.300326490402222,2.308275270462036,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:35,0.8.0
+cross_entropy,liger_triton,full,speed,ms,BT,B * T,8192,4.44209623336792,4.440243434906006,4.443949031829834,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:35,0.8.0
+cross_entropy,torch,full,speed,ms,BT,B * T,1024,0.8367679715156555,0.8365951895713806,0.8385791897773742,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:35,0.8.0
+cross_entropy,torch,full,speed,ms,BT,B * T,2048,1.5242239832878113,1.5233216047286988,1.5247168064117433,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:35,0.8.0
+cross_entropy,torch,full,speed,ms,BT,B * T,4096,2.9778239727020264,2.977190399169922,2.978457546234131,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:35,0.8.0
+cross_entropy,torch,full,speed,ms,BT,B * T,8192,5.815328121185303,5.815328121185303,5.815328121185303,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:35,0.8.0
+cross_entropy,liger_triton,no-grad-forward,speed,ms,BT,B * T,1024,0.3307519853115082,0.3270911931991577,0.332832008600235,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:35,0.8.0
+cross_entropy,liger_triton,no-grad-forward,speed,ms,BT,B * T,2048,0.5396479964256287,0.5386303901672362,0.5419072031974792,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:35,0.8.0
+cross_entropy,liger_triton,no-grad-forward,speed,ms,BT,B * T,4096,0.9599519968032837,0.9582911849021911,0.9619711995124817,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:35,0.8.0
+cross_entropy,liger_triton,no-grad-forward,speed,ms,BT,B * T,8192,1.783903956413269,1.7815488338470458,1.7850687980651856,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:35,0.8.0
+cross_entropy,torch,no-grad-forward,speed,ms,BT,B * T,1024,0.37785598635673523,0.37781120538711543,0.37895039916038514,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:35,0.8.0
+cross_entropy,torch,no-grad-forward,speed,ms,BT,B * T,2048,0.6954240202903748,0.6952959895133972,0.6973440051078796,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:35,0.8.0
+cross_entropy,torch,no-grad-forward,speed,ms,BT,B * T,4096,1.3518719673156738,1.350988817214966,1.352729606628418,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:35,0.8.0
+cross_entropy,torch,no-grad-forward,speed,ms,BT,B * T,8192,2.642944097518921,2.6428096294403076,2.646015977859497,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:35,0.8.0
+cross_entropy,liger_triton,full,memory,MB,BT,B * T,1024,1003.02099609375,1003.02099609375,1003.02099609375,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:36,0.8.0
+cross_entropy,liger_triton,full,memory,MB,BT,B * T,2048,2004.04150390625,2004.04150390625,2004.04150390625,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:36,0.8.0
+cross_entropy,liger_triton,full,memory,MB,BT,B * T,4096,4008.08251953125,4008.08251953125,4008.08251953125,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:36,0.8.0
+cross_entropy,liger_triton,full,memory,MB,BT,B * T,8192,8016.16455078125,8016.16455078125,8016.16455078125,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:36,0.8.0
+cross_entropy,torch,full,memory,MB,BT,B * T,1024,2506.00927734375,2506.00927734375,2506.00927734375,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:36,0.8.0
+cross_entropy,torch,full,memory,MB,BT,B * T,2048,5010.01708984375,5010.01708984375,5010.01708984375,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:36,0.8.0
+cross_entropy,torch,full,memory,MB,BT,B * T,4096,10020.033203125,10020.033203125,10020.033203125,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:36,0.8.0
+cross_entropy,torch,full,memory,MB,BT,B * T,8192,20040.064453125,20040.064453125,20040.064453125,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:36,0.8.0
+cross_entropy,liger_cutedsl,forward,speed,ms,BT,B * T,1024,0.3113119900226593,0.309254401922226,0.3132736086845398,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:41,0.8.0
+cross_entropy,liger_cutedsl,forward,speed,ms,BT,B * T,2048,0.5152319967746735,0.5144831895828247,0.5171200037002563,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:41,0.8.0
+cross_entropy,liger_cutedsl,forward,speed,ms,BT,B * T,4096,0.9226239919662476,0.9172287940979005,0.925056004524231,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:41,0.8.0
+cross_entropy,liger_cutedsl,forward,speed,ms,BT,B * T,8192,1.747167944908142,1.7424511671066283,1.7489344120025634,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:41,0.8.0
+cross_entropy,liger_cutedsl,backward,speed,ms,BT,B * T,1024,0.3850240111351013,0.3836351990699768,0.3880959928035736,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:41,0.8.0
+cross_entropy,liger_cutedsl,backward,speed,ms,BT,B * T,2048,0.7014080286026001,0.7003840208053589,0.7024639844894409,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:41,0.8.0
+cross_entropy,liger_cutedsl,backward,speed,ms,BT,B * T,4096,1.3424000144004822,1.3414080142974854,1.3434560298919678,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:41,0.8.0
+cross_entropy,liger_cutedsl,backward,speed,ms,BT,B * T,8192,2.6193599700927734,2.6181696414947506,2.6199935436248776,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:41,0.8.0
+cross_entropy,liger_cutedsl,full,speed,ms,BT,B * T,1024,0.7210559844970703,0.7186880111694336,0.7229440212249756,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:42,0.8.0
+cross_entropy,liger_cutedsl,full,speed,ms,BT,B * T,2048,1.2433279752731323,1.2404608249664306,1.253324770927429,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:42,0.8.0
+cross_entropy,liger_cutedsl,full,speed,ms,BT,B * T,4096,2.287839889526367,2.2870719909667967,2.2889535427093506,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:42,0.8.0
+cross_entropy,liger_cutedsl,full,speed,ms,BT,B * T,8192,4.428784132003784,4.423820972442627,4.433747291564941,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:42,0.8.0
+cross_entropy,liger_cutedsl,no-grad-forward,speed,ms,BT,B * T,1024,0.30934399366378784,0.30822399258613586,0.3141183972358703,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:42,0.8.0
+cross_entropy,liger_cutedsl,no-grad-forward,speed,ms,BT,B * T,2048,0.5155520141124725,0.5130879878997803,0.5171712160110473,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:42,0.8.0
+cross_entropy,liger_cutedsl,no-grad-forward,speed,ms,BT,B * T,4096,0.9210560023784637,0.9193472266197205,0.9225024223327637,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:42,0.8.0
+cross_entropy,liger_cutedsl,no-grad-forward,speed,ms,BT,B * T,8192,1.7428480386734009,1.7378752470016479,1.7455807447433473,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:42,0.8.0
+cross_entropy,liger_cutedsl,full,memory,MB,BT,B * T,1024,1003.02490234375,1003.02490234375,1003.02490234375,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:42,0.8.0
+cross_entropy,liger_cutedsl,full,memory,MB,BT,B * T,2048,2004.04931640625,2004.04931640625,2004.04931640625,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:42,0.8.0
+cross_entropy,liger_cutedsl,full,memory,MB,BT,B * T,4096,4008.09814453125,4008.09814453125,4008.09814453125,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:42,0.8.0
+cross_entropy,liger_cutedsl,full,memory,MB,BT,B * T,8192,8016.19580078125,8016.19580078125,8016.19580078125,"{""vocab_size"": 128256, ""bsz"": 1, ""seq_len"": 8192}",NVIDIA B200,2026-07-05 21:00:42,0.8.0
diff --git a/benchmark/scripts/run_cutedsl_compare.py b/benchmark/scripts/run_cutedsl_compare.py
new file mode 100644
index 000000000..c0b358d12
--- /dev/null
+++ b/benchmark/scripts/run_cutedsl_compare.py
@@ -0,0 +1,71 @@
+"""Run a kernel benchmark twice (Triton + CuTe-DSL) and merge results into one CSV.
+
+Workflow:
+    python scripts/run_cutedsl_compare.py --kernel cross_entropy [benchmark args...]
+
+This driver spawns the per-kernel benchmark script in two subprocesses with
+different env vars, so all series (liger_triton / liger_cutedsl / torch) land in
+`benchmark/data/all_benchmark_data_cutedsl.csv` under distinct `kernel_provider`
+values, ready for direct plotting via:
+
+    python ../benchmarks_visualizer.py \
+        --kernel-name <name> --metric-name speed \
+        --data-file data/all_benchmark_data_cutedsl.csv
+
+CuTe-DSL (``cutedsl``) is the CUDA-only Python DSL shipped with NVIDIA CUTLASS
+(``import cutlass.cute``), targeting Hopper (SM90) / Blackwell (SM100/SM110). It
+is opt-in via ``LIGER_KERNEL_IMPL=cutedsl``; the benchmark script's ``liger``
+provider then dispatches to the CuTe-DSL kernel with no script changes needed.
+"""
+
+import argparse
+import os
+import subprocess
+import sys
+
+CUTEDSL_ENABLED_KERNELS = [
+    "cross_entropy",
+]
+
+
+def main():
+    parser = argparse.ArgumentParser(
+        description="Compare Triton vs CuTe-DSL Liger kernels in one CSV.",
+        # Unknown args are forwarded to the underlying benchmark script.
+    )
+    parser.add_argument(
+        "--kernel",
+        required=True,
+        choices=CUTEDSL_ENABLED_KERNELS,
+        help="Kernel to compare. Must have a CuTe-DSL (cutedsl) backend.",
+    )
+    args, passthrough = parser.parse_known_args()
+
+    script_dir = os.path.dirname(os.path.abspath(__file__))
+    bench_script = os.path.join(script_dir, f"benchmark_{args.kernel}.py")
+    if not os.path.isfile(bench_script):
+        print(f"error: benchmark script not found: {bench_script}", file=sys.stderr)
+        sys.exit(1)
+
+    # Both runs target the same _cutedsl.csv; provider_tag disambiguates the
+    # "liger" rows so they don't overwrite each other on the dedup key.
+    runs = [
+        ("triton baseline", {"LIGER_KERNEL_IMPL": "", "LIGER_BENCH_PROVIDER_TAG": "liger_triton"}),
+        ("cutedsl", {"LIGER_KERNEL_IMPL": "cutedsl", "LIGER_BENCH_PROVIDER_TAG": "liger_cutedsl"}),
+    ]
+
+    for label, run_env in runs:
+        print(f"\n========== {args.kernel}: {label} ==========\n", flush=True)
+        env = {**os.environ, "LIGER_BENCH_TARGET": "cutedsl", **run_env}
+        result = subprocess.run(
+            [sys.executable, bench_script, *passthrough],
+            env=env,
+            cwd=script_dir,
+        )
+        if result.returncode != 0:
+            print(f"error: {label} run failed with exit code {result.returncode}", file=sys.stderr)
+            sys.exit(result.returncode)
+
+
+if __name__ == "__main__":
+    main()
diff --git a/src/liger_kernel/ops/cutedsl/ops/cross_entropy.py b/src/liger_kernel/ops/cutedsl/ops/cross_entropy.py
index 807659dcf..d77a9da5e 100644
--- a/src/liger_kernel/ops/cutedsl/ops/cross_entropy.py
+++ b/src/liger_kernel/ops/cutedsl/ops/cross_entropy.py
@@ -7,6 +7,7 @@
 import cutlass.cute as cute
 import cutlass.utils
 import torch
+import triton
 
 from cutlass import Float32
 from cutlass import Int32
@@ -17,6 +18,10 @@
 from cutlass.cutlass_dsl import dsl_user_op
 
 from liger_kernel.ops.cutedsl.ops.utils import to_cute_tensor
+from liger_kernel.ops.utils import element_mul_kernel
+
+# Matches the Triton CE backward cap (NVIDIA-only path, so the NPU 2048 cap is irrelevant).
+_MAX_FUSED_SIZE = 65536 // 2
 
 # log2(e) and ln(2): the online-softmax math is done in base-2 (hardware ex2.approx)
 # then converted, exactly mirroring the Triton kernel for numerical parity.
@@ -827,9 +832,23 @@ def cross_entropy_backward(_input, grad_output):
     # reduction="none": per-row upstream grad.
     if grad_output.ndim > 0:
         return _input * grad_output.unsqueeze(dim=1)
-    # reduction in {mean, sum}: scalar upstream grad. Fresh tensor (not in-place)
-    # to avoid the autograd anomalies the Triton path uses a kernel to dodge.
-    return _input * grad_output
+    # reduction in {mean, sum}: scalar upstream grad. Scale the saved gradient IN PLACE so we
+    # never materialize a second BT×V buffer (Triton-parity peak memory: 1x logits, not 2x).
+    # A raw Triton element-wise kernel is used instead of `_input *= grad_output` because an
+    # in-place torch mul on the tensor returned from forward bumps its autograd version counter
+    # and trips backward-through-backward anomalies; the raw kernel writes through the pointer
+    # without that bookkeeping — exactly how the Triton CE backward dodges the same issue.
+    BT, V = _input.shape
+    BLOCK_SIZE = min(_MAX_FUSED_SIZE, triton.next_power_of_2(V))
+    element_mul_kernel[(BT,)](
+        _input,
+        _input.stride(-2),
+        grad_output,
+        V,
+        BLOCK_SIZE=BLOCK_SIZE,
+        num_warps=32,
+    )
+    return _input
 
 
 class LigerCrossEntropyFunction(torch.autograd.Function):

From d2c3d0c7c95dbaea4e36a0740bab53490489431f Mon Sep 17 00:00:00 2001
From: Justin Hu <181588904+justinhh4@users.noreply.github.com>
Date: Sun, 5 Jul 2026 22:45:08 +0000
Subject: [PATCH 2/5] [CuteDSL] CE: arch-aware num_warps (fix Hopper
 regression)

The CuTe DSL CE forward baked 8 warps/CTA for any 2-byte dtype (a
Blackwell-only tuning). On Hopper (sm_90) that underfills the SMs and
loses ~0.90x to the 32-warp Triton forward. Mirror the Triton CE
convention exactly via infer_device_arch()/is_hip():
  - Blackwell (sm_100+): bf16/fp16 -> 8, fp32 -> 32
  - Hopper (sm_90) and earlier: 32 for all dtypes
  - AMD (ROCm): 16

num_warps is baked into the kernel and is part of the compile-cache key,
so this is a pure launch-config change with no numerics impact. On H100
bf16 goes from ~0.89-0.91x to parity-to-a-win on the vocab sweep
(0.91-1.11x); memory is unchanged (exact Triton parity).

Update the white-box test_num_warps_matches_dtype_convention to assert
the arch-aware value instead of the fixed 8/8/32.

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
---
 .../ops/cutedsl/ops/cross_entropy.py          | 18 +++++++++++----
 .../test_cutedsl_cross_entropy.py             | 22 ++++++++++++++-----
 2 files changed, 30 insertions(+), 10 deletions(-)

diff --git a/src/liger_kernel/ops/cutedsl/ops/cross_entropy.py b/src/liger_kernel/ops/cutedsl/ops/cross_entropy.py
index d77a9da5e..d6228e201 100644
--- a/src/liger_kernel/ops/cutedsl/ops/cross_entropy.py
+++ b/src/liger_kernel/ops/cutedsl/ops/cross_entropy.py
@@ -19,6 +19,8 @@
 
 from liger_kernel.ops.cutedsl.ops.utils import to_cute_tensor
 from liger_kernel.ops.utils import element_mul_kernel
+from liger_kernel.ops.utils import is_hip
+from liger_kernel.utils import infer_device_arch
 
 # Matches the Triton CE backward cap (NVIDIA-only path, so the NPU 2048 cap is irrelevant).
 _MAX_FUSED_SIZE = 65536 // 2
@@ -647,10 +649,18 @@ def _launch_ce_fwd(
     pt_ct = to_cute_tensor(pred_tok_out, assumed_align=8) if return_predicted_tokens else y_ct
     # weight is a fp32 (V,) vector when present (caller upcasts); dummy reuses int64 `y`.
     w_ct = to_cute_tensor(weight, assumed_align=4) if has_weight else y_ct
-    # warps/CTA: mirror Triton's Blackwell CE convention — 2-byte dtypes (bf16/fp16) are
-    # instruction-issue-bound -> 8 warps; fp32 is bandwidth-bound -> 32 warps. Baked into the
-    # kernel, so it's part of the compile key.
-    num_warps = 8 if x.element_size() == 2 else 32
+    # warps/CTA: mirror the Triton CE convention exactly (arch- and dtype-dependent):
+    #   Blackwell (B200, sm_100+) bf16/fp16 -> 8 (instruction-issue-bound); fp32 -> 32
+    #   Hopper (H100, sm_90) and earlier    -> 32 for all dtypes (bandwidth-bound)
+    #   AMD (ROCm)                          -> 16
+    # On Hopper the 8-warp bf16 kernel underfills the SMs and loses to the 32-warp Triton
+    # forward, so we gate the 8-warp choice on Blackwell only (matches ops/cross_entropy.py).
+    # Baked into the kernel, so it's part of the compile key.
+    if is_hip():
+        num_warps = 16
+    else:
+        is_blackwell = infer_device_arch().startswith("blackwell")
+        num_warps = 8 if (x.element_size() == 2 and is_blackwell) else 32
     key = (
         x.dtype,
         y.dtype,
diff --git a/test/transformers/test_cutedsl_cross_entropy.py b/test/transformers/test_cutedsl_cross_entropy.py
index 1463fc7d3..b712b7358 100644
--- a/test/transformers/test_cutedsl_cross_entropy.py
+++ b/test/transformers/test_cutedsl_cross_entropy.py
@@ -359,17 +359,27 @@ def test_ce_noncontiguous_input_matches_triton(dtype):
 
 # =============================================================================
 # B. Warp-count selection (white-box on the compile cache): the streaming kernel bakes
-#    num_warps per dtype, mirroring the Triton CE Blackwell convention — 8 warps for 2-byte
-#    dtypes (instruction-issue-bound), 32 for fp32 (bandwidth-bound). num_warps is the last
-#    element of the compile-cache key.
+#    num_warps per dtype AND arch, mirroring the Triton CE convention exactly —
+#    Blackwell (sm_100+) bf16/fp16 -> 8 (instruction-issue-bound), fp32 -> 32; Hopper
+#    (sm_90) and earlier -> 32 for all dtypes (bandwidth-bound); AMD (ROCm) -> 16.
+#    num_warps is the last element of the compile-cache key.
 # =============================================================================
 @cuda_required
 @pytest.mark.parametrize(
-    "dtype, expected_warps",
-    [(torch.bfloat16, 8), (torch.float16, 8), (torch.float32, 32)],
+    "dtype",
+    [torch.bfloat16, torch.float16, torch.float32],
     ids=["bf16", "fp16", "fp32"],
 )
-def test_num_warps_matches_dtype_convention(dtype, expected_warps):
+def test_num_warps_matches_dtype_convention(dtype):
+    from liger_kernel.ops.utils import is_hip
+    from liger_kernel.utils import infer_device_arch
+
+    if is_hip():
+        expected_warps = 16
+    else:
+        is_blackwell = infer_device_arch().startswith("blackwell")
+        expected_warps = 8 if (dtype in (torch.bfloat16, torch.float16) and is_blackwell) else 32
+
     mod = _cutedsl_ce_module()
     base = torch.randn(64, 4096, device="cuda", dtype=torch.float32)
     target = torch.randint(0, 4096, (64,), device="cuda", dtype=torch.long)

From 4314af1d6cb93c6f69082d8691f66ea5df7460b3 Mon Sep 17 00:00:00 2001
From: Justin Hu <181588904+justinhh4@users.noreply.github.com>
Date: Wed, 22 Jul 2026 16:39:27 +0000
Subject: [PATCH 3/5] Replace Triton CE scale with CuTe DSL

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
---
 .../ops/cutedsl/ops/cross_entropy.py          | 59 +++++++++++++------
 .../test_cutedsl_cross_entropy.py             |  2 +-
 2 files changed, 41 insertions(+), 20 deletions(-)

diff --git a/src/liger_kernel/ops/cutedsl/ops/cross_entropy.py b/src/liger_kernel/ops/cutedsl/ops/cross_entropy.py
index d6228e201..50fff6eb6 100644
--- a/src/liger_kernel/ops/cutedsl/ops/cross_entropy.py
+++ b/src/liger_kernel/ops/cutedsl/ops/cross_entropy.py
@@ -7,7 +7,6 @@
 import cutlass.cute as cute
 import cutlass.utils
 import torch
-import triton
 
 from cutlass import Float32
 from cutlass import Int32
@@ -18,13 +17,9 @@
 from cutlass.cutlass_dsl import dsl_user_op
 
 from liger_kernel.ops.cutedsl.ops.utils import to_cute_tensor
-from liger_kernel.ops.utils import element_mul_kernel
 from liger_kernel.ops.utils import is_hip
 from liger_kernel.utils import infer_device_arch
 
-# Matches the Triton CE backward cap (NVIDIA-only path, so the NPU 2048 cap is irrelevant).
-_MAX_FUSED_SIZE = 65536 // 2
-
 # log2(e) and ln(2): the online-softmax math is done in base-2 (hardware ex2.approx)
 # then converted, exactly mirroring the Triton kernel for numerical parity.
 LOG2_E = 1.4426950408889634
@@ -77,6 +72,7 @@ def fmax(a, b, c=None, *, loc=None, ip=None) -> Float32:
 # kernel bakes. V/BT are dynamic so one compile serves all shapes. REQUIRED: without it the
 # @cute.jit host fn recompiles on every call (~30 ms that dwarfs the kernel).
 _compile_cache = {}
+_scale_compile_cache = {}
 
 # Per-call host overhead is constant (~25 us): it doesn't scale with BT/V, so it dominates small
 # shapes and vanishes at scale. Cache the CUstream wrapper keyed on torch's raw stream handle so
@@ -149,6 +145,42 @@ def _advance(idx, n: cutlass.Constexpr):
     return idx + 1 if idx < n - 1 else 0
 
 
+# =============================================================================
+# Backward scale kernel
+# =============================================================================
+@cute.kernel
+def _scale_in_place_kernel(mX: cute.Tensor, mScale: cute.Tensor):
+    tid, _, _ = cute.arch.thread_idx()
+    row, _, _ = cute.arch.block_idx()
+    V = mX.shape[1]
+    scale = mScale[0].to(Float32)
+
+    for i in cutlass.range(0, cute.ceil_div(V, 256)):
+        col = tid + i * 256
+        if col < V:
+            value = mX[row, col].to(Float32) * scale
+            mX[row, col] = value.to(mX.element_type)
+
+
+@cute.jit
+def _scale_in_place_host(mX: cute.Tensor, mScale: cute.Tensor, stream: cuda.CUstream = None):
+    _scale_in_place_kernel(mX, mScale).launch(
+        grid=[mX.shape[0], 1, 1],
+        block=[256, 1, 1],
+        stream=stream,
+    )
+
+
+def _scale_in_place(x, scale):
+    x_ct = to_cute_tensor(x)
+    scale_ct = to_cute_tensor(scale.reshape(1), assumed_align=2)
+    stream = _cute_stream()
+    key = (x.dtype, scale.dtype)
+    if key not in _scale_compile_cache:
+        _scale_compile_cache[key] = cute.compile(_scale_in_place_host, x_ct, scale_ct, stream)
+    _scale_compile_cache[key](x_ct, scale_ct, stream)
+
+
 # =============================================================================
 # Device kernel
 # =============================================================================
@@ -844,20 +876,9 @@ def cross_entropy_backward(_input, grad_output):
         return _input * grad_output.unsqueeze(dim=1)
     # reduction in {mean, sum}: scalar upstream grad. Scale the saved gradient IN PLACE so we
     # never materialize a second BT×V buffer (Triton-parity peak memory: 1x logits, not 2x).
-    # A raw Triton element-wise kernel is used instead of `_input *= grad_output` because an
-    # in-place torch mul on the tensor returned from forward bumps its autograd version counter
-    # and trips backward-through-backward anomalies; the raw kernel writes through the pointer
-    # without that bookkeeping — exactly how the Triton CE backward dodges the same issue.
-    BT, V = _input.shape
-    BLOCK_SIZE = min(_MAX_FUSED_SIZE, triton.next_power_of_2(V))
-    element_mul_kernel[(BT,)](
-        _input,
-        _input.stride(-2),
-        grad_output,
-        V,
-        BLOCK_SIZE=BLOCK_SIZE,
-        num_warps=32,
-    )
+    # A raw CuTe DSL kernel is used instead of `_input *= grad_output` because the torch op
+    # bumps the forward output's autograd version counter and breaks repeated backward.
+    _scale_in_place(_input, grad_output)
     return _input
 
 
diff --git a/test/transformers/test_cutedsl_cross_entropy.py b/test/transformers/test_cutedsl_cross_entropy.py
index b712b7358..dd340cf81 100644
--- a/test/transformers/test_cutedsl_cross_entropy.py
+++ b/test/transformers/test_cutedsl_cross_entropy.py
@@ -325,7 +325,7 @@ def test_ce_forward_only_matches_triton(dtype):
 @pytest.mark.parametrize("dtype", _DTYPES, ids=_DTYPE_IDS)
 @pytest.mark.parametrize("reduction", ["mean", "sum"])
 def test_ce_not_last_layer_grad_matches_triton(reduction, dtype):
-    """grad_output != 1.0 (scalar): exercises the `_input * grad_output` backward branch."""
+    """grad_output != 1.0 (scalar): exercises the in-place CuTe DSL scale kernel."""
     set_seed()
     BT, V = 256, 4096
     base = torch.randn(BT, V, device="cuda", dtype=torch.float32)

From e61efa9151b695a6483ed2f04c78a0e9c1623ce2 Mon Sep 17 00:00:00 2001
From: Justin Hu <181588904+justinhh4@users.noreply.github.com>
Date: Wed, 22 Jul 2026 16:46:15 +0000
Subject: [PATCH 4/5] Vectorize CuTe CE gradient scaling

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
---
 .../ops/cutedsl/ops/cross_entropy.py          | 25 +++++++++++++------
 1 file changed, 18 insertions(+), 7 deletions(-)

diff --git a/src/liger_kernel/ops/cutedsl/ops/cross_entropy.py b/src/liger_kernel/ops/cutedsl/ops/cross_entropy.py
index 50fff6eb6..d337e3f73 100644
--- a/src/liger_kernel/ops/cutedsl/ops/cross_entropy.py
+++ b/src/liger_kernel/ops/cutedsl/ops/cross_entropy.py
@@ -152,21 +152,32 @@ def _advance(idx, n: cutlass.Constexpr):
 def _scale_in_place_kernel(mX: cute.Tensor, mScale: cute.Tensor):
     tid, _, _ = cute.arch.thread_idx()
     row, _, _ = cute.arch.block_idx()
-    V = mX.shape[1]
     scale = mScale[0].to(Float32)
 
-    for i in cutlass.range(0, cute.ceil_div(V, 256)):
-        col = tid + i * 256
-        if col < V:
-            value = mX[row, col].to(Float32) * scale
-            mX[row, col] = value.to(mX.element_type)
+    gX = mX[row, None]
+    V = gX.shape[0]
+    gX = cute.make_tensor(
+        cute.make_ptr(mX.element_type, gX.iterator.toint(), cute.AddressSpace.gmem, assumed_align=16),
+        cute.make_layout((V,)),
+    )
+    VEC = const_expr(128 // gX.element_type.width)
+    gXv = cute.tiled_divide(gX, (VEC,))
+    num_vec = V // VEC
+    x_frag = cute.make_rmem_tensor((VEC,), gX.element_type)
+
+    for i in cutlass.range(0, cute.ceil_div(num_vec, 1024)):
+        vec_idx = tid + i * 1024
+        if vec_idx < num_vec:
+            cute.autovec_copy(gXv[None, vec_idx], x_frag)
+            x_frag.store((x_frag.load().to(Float32) * scale).to(gX.element_type))
+            cute.autovec_copy(x_frag, gXv[None, vec_idx])
 
 
 @cute.jit
 def _scale_in_place_host(mX: cute.Tensor, mScale: cute.Tensor, stream: cuda.CUstream = None):
     _scale_in_place_kernel(mX, mScale).launch(
         grid=[mX.shape[0], 1, 1],
-        block=[256, 1, 1],
+        block=[1024, 1, 1],
         stream=stream,
     )
 

From fe0a6cc4801b9a8b87ed67ae936ba472e0e984ff Mon Sep 17 00:00:00 2001
From: Justin Hu <181588904+justinhh4@users.noreply.github.com>
Date: Wed, 22 Jul 2026 17:00:36 +0000
Subject: [PATCH 5/5] Mirror generic CuTile benchmark driver

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
---
 benchmark/scripts/run_cutedsl_compare.py | 55 +++++++++---------------
 1 file changed, 20 insertions(+), 35 deletions(-)

diff --git a/benchmark/scripts/run_cutedsl_compare.py b/benchmark/scripts/run_cutedsl_compare.py
index 5a0d02969..6814b9543 100644
--- a/benchmark/scripts/run_cutedsl_compare.py
+++ b/benchmark/scripts/run_cutedsl_compare.py
@@ -1,23 +1,16 @@
-"""Run a kernel benchmark with Triton and CuTe DSL backends into one CSV.
+"""Run a kernel benchmark twice (Triton + CuTe DSL) and merge results into one CSV.
 
-Usage (from ``benchmark/scripts``)::
+Workflow:
+    python scripts/run_cutedsl_compare.py --kernel cross_entropy [benchmark args...]
 
-    python run_cutedsl_compare.py --kernel cross_entropy [benchmark args...]
-    python run_cutedsl_compare.py --kernel rms_norm --source h100 [benchmark args...]
+This driver spawns the per-kernel benchmark script in two subprocesses with
+different env vars, so all three series (liger_triton / liger_cutedsl /
+huggingface or torch) land in `benchmark/data/all_benchmark_data_cutedsl.csv`
+under distinct `kernel_provider` values, ready for direct plotting via:
 
-Backend selection is process-global and happens at import time, so this driver
-spawns the underlying benchmark script twice:
-
-  * baseline: ``LIGER_KERNEL_IMPL`` unset
-  * CuTe DSL: ``LIGER_KERNEL_IMPL=cutedsl``
-
-``LIGER_BENCH_PROVIDER_TAG`` keeps the ``liger_triton`` and ``liger_cutedsl``
-rows distinct. Runs are sequential because the CSV writer performs an unlocked
-read-modify-write.
-
-For backward compatibility, cross-entropy without ``--source`` writes
-``all_benchmark_data_cutedsl.csv``. Other runs use a kernel-specific target,
-such as ``all_benchmark_data_cutedsl_rms_norm_h100.csv``.
+    python ../benchmarks_visualizer.py \
+        --kernel-name <name> --metric-name speed \
+        --data-file data/all_benchmark_data_cutedsl.csv
 """
 
 import argparse
@@ -25,25 +18,23 @@
 import subprocess
 import sys
 
-CUTEDSL_ENABLED_KERNELS = ["cross_entropy", "rms_norm"]
+CUTEDSL_ENABLED_KERNELS = [
+    "cross_entropy",
+    "rms_norm",
+]
 
 
 def main():
     parser = argparse.ArgumentParser(
-        description="Compare Triton and CuTe DSL Liger kernels in one CSV.",
+        description="Compare Triton vs CuTe DSL Liger kernels in one CSV.",
         # Unknown args are forwarded to the underlying benchmark script.
     )
     parser.add_argument(
         "--kernel",
-        default="rms_norm",
+        required=True,
         choices=CUTEDSL_ENABLED_KERNELS,
         help="Kernel to compare. Must have a CuTe DSL backend.",
     )
-    parser.add_argument(
-        "--source",
-        default="",
-        help="Optional label (for example, 'h100' or 'b200') appended to the CSV target name.",
-    )
     args, passthrough = parser.parse_known_args()
 
     script_dir = os.path.dirname(os.path.abspath(__file__))
@@ -52,20 +43,16 @@ def main():
         print(f"error: benchmark script not found: {bench_script}", file=sys.stderr)
         sys.exit(1)
 
-    source = args.source.strip().lower()
-    if args.kernel == "cross_entropy" and not source:
-        target = "cutedsl"
-    else:
-        target = f"cutedsl_{args.kernel}" + (f"_{source}" if source else "")
-
+    # Both runs target the same _cutedsl.csv; provider_tag disambiguates the
+    # "liger" rows so they don't overwrite each other on the dedup key.
     runs = [
         ("triton baseline", {"LIGER_KERNEL_IMPL": "", "LIGER_BENCH_PROVIDER_TAG": "liger_triton"}),
         ("cutedsl", {"LIGER_KERNEL_IMPL": "cutedsl", "LIGER_BENCH_PROVIDER_TAG": "liger_cutedsl"}),
     ]
 
     for label, run_env in runs:
-        print(f"\n========== {args.kernel} [{source or 'default'}]: {label} ==========\n", flush=True)
-        env = {**os.environ, "LIGER_BENCH_TARGET": target, **run_env}
+        print(f"\n========== {args.kernel}: {label} ==========\n", flush=True)
+        env = {**os.environ, "LIGER_BENCH_TARGET": "cutedsl", **run_env}
         result = subprocess.run(
             [sys.executable, bench_script, *passthrough],
             env=env,
@@ -75,8 +62,6 @@ def main():
             print(f"error: {label} run failed with exit code {result.returncode}", file=sys.stderr)
             sys.exit(result.returncode)
 
-    print(f"\nWrote merged results to benchmark/data/all_benchmark_data_{target}.csv", flush=True)
-
 
 if __name__ == "__main__":
     main()