mirror of
https://github.com/linkedin/Liger-Kernel.git
synced 2026-10-12 02:17:47 +00:00
## Summary Adds an optional **CuTe DSL** (`nvidia-cutlass-dsl`) backend implementing Liger cross-entropy (forward + backward) as a single fused kernel, signature-compatible with the Triton `LigerCrossEntropyFunction`. Feature parity: z-loss, softcap, label smoothing, class weights, `ignore_index`, token-accuracy, predicted-tokens, reductions `mean`/`sum`/`none`, across bf16/fp16/fp32. The fused-linear-CE (FLCE) variant will follow in a later PR. ### Kernel design / optimizations (B200) - Online-softmax forward pass with a **cp.async multi-stage smem pipeline** (128-bit vectorized). - **FMA-folded `exp2` arguments** — `(x - m)·log2e` → `fma(x, log2e, -m·log2e)` with the scalar hoisted, cutting one issue slot per element on the instruction-bound pass. - **Power-of-2 ring-buffer stage advance** (`(i+1) & (n-1)`) — one fewer ALU op per tile. - **Reduced per-call host overhead** (constant ~25–40 µs, dominates small shapes): cached `CUstream`, batched the 3 target-validation D2H syncs into 1, and reused DLPack handles for unused optional outputs. ### Correctness `test/transformers/test_cutedsl_cross_entropy.py` — **158 passed, 1 skipped**, full parity vs the Triton kernel across dtype × reduction × features × grad (and vs `torch` where applicable). ## Benchmarks — CuTe DSL vs Triton (NVIDIA B200) Median of 40 iters, warmup 20. **speedup = Triton ÷ CuteDSL** (>1.00 → CuteDSL faster). Reference config: BT=8192, V=128256. `full` = fwd+bwd. ### bf16 **full (fwd+bwd) — vocab sweep (BT=8192)** | vocab | cutedsl (ms) | triton (ms) | speedup | |------:|-------------:|------------:|--------:| | 32000 | 0.4024 | 0.4519 | 1.123× | | 50304 | 0.5616 | 0.6703 | 1.194× | | 102400 | 0.9643 | 1.1170 | 1.158× | | 128256 | 1.1595 | 1.1855 | 1.022× | | 152064 | 1.3336 | 1.4445 | 1.083× | | 201088 | 1.7063 | 1.9767 | 1.158× | | 262144 | 2.2164 | 2.3127 | 1.043× | **full — BT sweep (V=128256)** | BT | cutedsl (ms) | triton (ms) | speedup | |---:|-------------:|------------:|--------:| | 1024 | 0.3364 | 0.3348 | 0.995× | | 2048 | 0.4554 | 0.4570 | 1.003× | | 4096 | 0.6902 | 0.7008 | 1.015× | | 8192 | 1.1565 | 1.1825 | 1.023× | | 16384 | 2.0945 | 2.1952 | 1.048× | | 32768 | 4.0096 | 4.1110 | 1.025× | | 65536 | 7.7274 | 7.9275 | 1.026× | ### fp32 **full (fwd+bwd) — vocab sweep (BT=8192)** | vocab | cutedsl (ms) | triton (ms) | speedup | |------:|-------------:|------------:|--------:| | 32000 | 0.7084 | 0.6213 | 0.877× | | 50304 | 0.9226 | 0.9231 | 1.001× | | 102400 | 1.4332 | 1.6221 | 1.132× | | 128256 | 1.8311 | 1.7849 | 0.975× | | 152064 | 2.2882 | 2.2613 | 0.988× | | 201088 | 3.0892 | 3.1684 | 1.026× | | 262144 | 3.9836 | 3.9752 | 0.998× | **full — BT sweep (V=128256)** | BT | cutedsl (ms) | triton (ms) | speedup | |---:|-------------:|------------:|--------:| | 1024 | 0.3967 | 0.3883 | 0.979× | | 2048 | 0.5990 | 0.5944 | 0.992× | | 4096 | 1.0071 | 0.9909 | 0.984× | | 8192 | 1.8288 | 1.7867 | 0.977× | | 16384 | 3.5309 | 3.4171 | 0.968× | | 32768 | 6.8136 | 6.5726 | 0.965× | | 65536 | 13.6012 | 13.0380 | 0.959× | ### Notes - CuteDSL wins/ties Triton across bf16 (both modes) and fp32 inference; the remaining gap is **fp32 `full`** (~0.96–0.98×, a steady-state per-element gap, not launch overhead) plus the small-vocab V=32k fp32 corner — targeted for a follow-up. - Requires `pip install 'liger-kernel[cutedsl]'` (adds `nvidia-cutlass-dsl`). NVIDIA GPUs only. --------- Co-authored-by: justinhh4 <[email protected]> Co-authored-by: Copilot <[email protected]> Co-authored-by: Justin Hu <[email protected]> Co-authored-by: Vaibhav Jindal <[email protected]>
145 lines
3.8 KiB
Python
145 lines
3.8 KiB
Python
# setup.py
|
|
|
|
import subprocess
|
|
|
|
from typing import Literal
|
|
|
|
from setuptools import setup
|
|
|
|
|
|
def get_default_dependencies():
|
|
"""Determine the appropriate dependencies based on detected hardware."""
|
|
platform = get_platform()
|
|
|
|
if platform in ["cuda", "cpu"]:
|
|
return [
|
|
"torch>=2.1.2",
|
|
"triton>=2.3.1",
|
|
]
|
|
elif platform == "rocm":
|
|
return [
|
|
"triton>=3.0.0",
|
|
]
|
|
elif platform == "xpu":
|
|
return [
|
|
"torch>=2.6.0",
|
|
]
|
|
elif platform == "npu":
|
|
return ["torch==2.7.1", "torch_npu==2.7.1", "triton-ascend==3.2.1"]
|
|
|
|
|
|
def get_optional_dependencies():
|
|
"""Get optional dependency groups."""
|
|
cutile_deps = [
|
|
"cuda-tile",
|
|
]
|
|
cutile_tileiras_deps = [
|
|
"cuda-tile[tileiras]",
|
|
]
|
|
cutedsl_deps = [
|
|
"nvidia-cutlass-dsl",
|
|
]
|
|
dev_deps = [
|
|
"transformers>=4.52.0",
|
|
"matplotlib>=3.7.2",
|
|
"ruff>=0.12.0",
|
|
"pytest>=7.1.2",
|
|
"pytest-xdist",
|
|
"pytest-cov",
|
|
"pytest-asyncio",
|
|
"pytest-rerunfailures",
|
|
"datasets>=2.19.2",
|
|
"seaborn",
|
|
"mkdocs-material",
|
|
"torchvision>=0.20",
|
|
"prek>=0.2.28",
|
|
]
|
|
return {
|
|
"cutile": cutile_deps,
|
|
"cutile-tileiras": cutile_tileiras_deps,
|
|
"cutedsl": cutedsl_deps,
|
|
"dev": dev_deps,
|
|
}
|
|
|
|
|
|
def is_xpu_available():
|
|
"""
|
|
Check if Intel XPU is available.
|
|
xpu-smi is often missing right now.
|
|
"""
|
|
try:
|
|
subprocess.run(["xpu-smi"], check=True)
|
|
return True
|
|
except (subprocess.SubprocessError, FileNotFoundError):
|
|
pass
|
|
|
|
try:
|
|
result = subprocess.run("sycl-ls", check=True, capture_output=True, shell=True)
|
|
if "level_zero:gpu" in result.stdout.decode():
|
|
return True
|
|
except (subprocess.SubprocessError, FileNotFoundError):
|
|
pass
|
|
|
|
return False
|
|
|
|
|
|
def is_ascend_available() -> bool:
|
|
"""Best-effort Ascend detection.
|
|
|
|
Checks for common Ascend environment variables and a possible `npu-smi`
|
|
utility if present.
|
|
"""
|
|
try:
|
|
subprocess.run(["npu-smi", "info"], check=True)
|
|
return True
|
|
except (subprocess.SubprocessError, FileNotFoundError):
|
|
pass
|
|
return False
|
|
|
|
|
|
def get_platform() -> Literal["cuda", "rocm", "cpu", "xpu", "npu"]:
|
|
"""
|
|
Detect whether the system has NVIDIA or AMD GPU without torch dependency.
|
|
"""
|
|
# Try nvidia-smi first
|
|
try:
|
|
subprocess.run(["nvidia-smi"], check=True)
|
|
print("NVIDIA GPU detected")
|
|
return "cuda"
|
|
except (subprocess.SubprocessError, FileNotFoundError):
|
|
# If nvidia-smi fails, check for ROCm
|
|
try:
|
|
subprocess.run(["rocm-smi"], check=True)
|
|
print("ROCm GPU detected")
|
|
return "rocm"
|
|
except (subprocess.SubprocessError, FileNotFoundError):
|
|
if is_xpu_available():
|
|
print("Intel GPU detected")
|
|
return "xpu"
|
|
elif is_ascend_available():
|
|
print("Ascend NPU detected")
|
|
return "npu"
|
|
else:
|
|
print("No GPU detected")
|
|
return "cpu"
|
|
|
|
|
|
setup(
|
|
name="liger_kernel",
|
|
package_dir={"": "src"},
|
|
packages=["liger_kernel"],
|
|
install_requires=get_default_dependencies(),
|
|
extras_require=get_optional_dependencies(),
|
|
classifiers=[
|
|
"Development Status :: 5 - Production/Stable",
|
|
"Intended Audience :: Developers",
|
|
"Intended Audience :: Education",
|
|
"Intended Audience :: Science/Research",
|
|
"Programming Language :: Python :: 3",
|
|
"Topic :: Scientific/Engineering :: Artificial Intelligence",
|
|
"Topic :: Software Development :: Libraries :: Python Modules",
|
|
"License :: OSI Approved :: BSD-2-Clause Software License",
|
|
"Operating System :: OS Independent",
|
|
],
|
|
)
|