Complete Guide to Liger Kernel: Efficient LLM Training with Triton Kernels

Introduction

Training and fine-tuning Large Language Models (LLMs) demands enormous computational resources, particularly GPU VRAM. For many machine learning practitioners, VRAM limitations represent the primary bottleneck — whether performing full fine-tuning or parameter-efficient approaches like LoRA. This is where Liger Kernel comes in as an exceptionally practical solution.

Liger Kernel is a collection of Triton kernels developed by LinkedIn's engineering team, specifically designed to accelerate LLM training while significantly reducing VRAM consumption. The name "Liger" stands for LinkedIn GPU Efficient Runtimes. By replacing standard PyTorch operations with optimized Triton kernels, Liger Kernel can deliver up to 60% memory savings and 20% throughput improvement across popular model architectures like LLaMA, Mistral, Gemma, and Qwen.

In this tutorial, we will explore how to use Liger Kernel across various scenarios — from simple one-line integration with HuggingFace Transformers to individual kernel usage in custom training loops. This tutorial is aimed at ML practitioners who want to maximize their hardware efficiency without sacrificing training quality.

Why Liger Kernel?

There are several compelling reasons to consider Liger Kernel:

Available Kernels

Liger Kernel provides optimized Triton implementations for the following operations:

• RMSNorm: Root Mean Square Layer Normalization
• RoPE (Rotary Positional Embedding): Rotary positional encoding
• SwiGLU: The SwiGLU activation function used in LLaMA and similar models
• CrossEntropyLoss: Memory-efficient cross-entropy loss function
• FusedLinearCrossEntropy: Fusion of the final linear layer with cross-entropy loss
• GeGLU and ReGLU: Other gated linear unit variants
• Layer Norm: Standard layer normalization
• GroupNorm: Group normalization

Installation

Prerequisites

Before installing Liger Kernel, ensure your environment meets these requirements:

• Python 3.8 or newer
• PyTorch 2.1.2 or newer
• CUDA toolkit compatible with your PyTorch version
• NVIDIA GPU with compute capability 7.0 or higher (Volta, Turing, Ampere, Hopper, or Ada Lovelace)

Installation via pip

The easiest way to install Liger Kernel:

pip install liger-kernel

Installation from Source

To use the latest version directly from the repository:

git clone https://github.com/linkedin/Liger-Kernel.git
cd Liger-Kernel
pip install -e .

Installation with Additional Dependencies

For development and testing:

pip install liger-kernel[dev]

Verifying Installation

After installation, verify that Liger Kernel is properly installed:

import ligerkernel
print(f"Liger Kernel version: {ligerkernel.version}")

Verify Triton is available
import triton
print(f"Triton version: {triton.version}")

Quick test
from ligerkernel.transformers import LigerRMSNorm

import torch

norm = LigerRMSNorm(hiddensize=4096).cuda()
x = torch.randn(2, 128, 4096, device="cuda")
output = norm(x)
print(f"RMSNorm output shape: {output.shape}")
print("Liger Kernel successfully installed and working!")

Basic Usage

Integration with HuggingFace Transformers (One Line of Code)

The easiest way to use Liger Kernel is through the monkey-patching approach. With this method, you only need to add one line of code before loading your model, and Liger Kernel will automatically replace the model's built-in operations with optimized Triton implementations.

import torch
from transformers import AutoModelForCausalLM, AutoTokenizer, TrainingArguments
from trl import SFTTrainer
from datasets import loaddataset


This one line activates Liger Kernel!
from ligerkernel.transformers import applyligerkerneltollama
applyligerkerneltollama()

Now load the model as usual — kernels are automatically patched
modelname = "meta-llama/Llama-3.1-8B"

model = AutoModelForCausalLM.frompretrained(
    modelname,

    torchdtype=torch.bfloat16,
    devicemap="auto",

)
tokenizer = AutoTokenizer.frompretrained(modelname)

tokenizer.padtoken = tokenizer.eostoken


Prepare dataset
dataset = loaddataset("tatsu-lab/alpaca", split="train[:1000]")

def formatprompt(example):

    if example["input"]:
        return f"### Instruction:\n{example['instruction']}\n\n### Input:\n{example['input']}\n\n### Response:\n{example['output']}"
    return f"### Instruction:\n{example['instruction']}\n\n### Response:\n{example['output']}"

dataset = dataset.map(lambda x: {"text": formatprompt(x)})

Training arguments
trainingargs = TrainingArguments(

    outputdir="./output",
    numtrainepochs=1,
    perdevicetrainbatchsize=4,
    gradientaccumulationsteps=4,
    learningrate=2e-5,

    bf16=True,
    loggingsteps=10,
    savestrategy="epoch",

    reportto="none",
)

Create trainer and start training
trainer = SFTTrainer(
    model=model,
    args=trainingargs,

    traindataset=dataset,
    tokenizer=tokenizer,
    datasettextfield="text",
    maxseqlength=2048,
)

trainer.train()

Monkey-Patch Functions for Different Architectures

Liger Kernel provides architecture-specific monkey-patch functions for each supported model:

# For LLaMA models (LLaMA 2, LLaMA 3, LLaMA 3.1, etc.)
from ligerkernel.transformers import applyligerkerneltollama

applyligerkerneltollama()


For Mistral models
from ligerkernel.transformers import applyligerkerneltomistral
applyligerkerneltomistral()

For Gemma and Gemma 2
from ligerkernel.transformers import applyligerkerneltogemma

applyligerkerneltogemma()


from ligerkernel.transformers import applyligerkerneltogemma2
applyligerkerneltogemma2()

For Qwen 2 and Qwen 2.5
from ligerkernel.transformers import applyligerkerneltoqwen2

applyligerkerneltoqwen2()


For Phi-3
from ligerkernel.transformers import applyligerkerneltophi3
applyligerkerneltophi3()

Controlling Monkey-Patch Parameters

Each monkey-patch function accepts optional parameters to control which kernels are activated:

from ligerkernel.transformers import applyligerkerneltollama

applyligerkerneltollama(

    rope=True,           # Enable Triton RoPE (default: True)
    crossentropy=True,  # Enable Triton CrossEntropy (default: True)
    fusedlinearcrossentropy=True,  # Enable FusedLinearCrossEntropy (default: True)

    rmsnorm=True,       # Enable Triton RMSNorm (default: True)
    swiglu=True,         # Enable Triton SwiGLU (default: True)
)

Example: only enable RMSNorm and SwiGLU
applyligerkerneltollama(
    rope=False,
    crossentropy=False,

    fusedlinearcrossentropy=False,
    rmsnorm=True,

    swiglu=True,
)

from transformers import TrainingArguments

trainingargs = TrainingArguments(
    outputdir="./output",

    useligerkernel=True,  # Automatically enable Liger Kernel

    # ... other parameters
)

import torch
from ligerkernel.transformers import LigerRMSNorm

Replace nn.RMSNorm or custom RMSNorm implementations
rmsnorm = LigerRMSNorm(hiddensize=4096, eps=1e-6).cuda()

Forward pass
hiddenstates = torch.randn(2, 512, 4096, device="cuda", dtype=torch.bfloat16)

output = rmsnorm(hiddenstates)

print(f"Input shape: {hiddenstates.shape}, Output shape: {output.shape}")

LigerSwiGLUMLP

SwiGLU is an activation function combining Swish and Gated Linear Unit, used in LLaMA, Mistral, and similar models. Liger's implementation fuses the gate and up projection operations for efficiency.

import torch
from ligerkernel.transformers import LigerSwiGLUMLP


Create SwiGLU MLP layer
swiglumlp = LigerSwiGLUMLP(
    hiddensize=4096,

    intermediatesize=11008,
).cuda().to(torch.bfloat16)

Forward pass
hiddenstates = torch.randn(2, 512, 4096, device="cuda", dtype=torch.bfloat16)

output = swiglumlp(hiddenstates)

print(f"SwiGLU MLP output shape: {output.shape}")

import torch
from ligerkernel.transformers.rope import ligerrotaryposemb


Simulate query and key tensors
batchsize, numheads, seqlen, headdim = 2, 32, 512, 128

q = torch.randn(batchsize, numheads, seqlen, headdim, device="cuda", dtype=torch.bfloat16)

k = torch.randn(batchsize, numheads, seqlen, headdim, device="cuda", dtype=torch.bfloat16)


Create cos and sin embeddings
cos = torch.randn(1, 1, seqlen, headdim, device="cuda", dtype=torch.bfloat16)

sin = torch.randn(1, 1, seqlen, headdim, device="cuda", dtype=torch.bfloat16)


Apply rotary embedding using Liger
qembed, kembed = ligerrotaryposemb(q, k, cos, sin)
print(f"Q embedded shape: {qembed.shape}")

print(f"K embedded shape: {kembed.shape}")

LigerCrossEntropyLoss

Liger's optimized cross-entropy loss uses chunked computation to drastically reduce memory footprint, especially for large vocabularies.

import torch
from ligerkernel.transformers import LigerCrossEntropyLoss


Create loss function
lossfn = LigerCrossEntropyLoss()

Simulate logits and labels
vocabsize = 128256  # LLaMA 3 vocabulary size

batchsize, seqlen = 2, 512


logits = torch.randn(batchsize  seqlen, vocabsize, device="cuda", dtype=torch.bfloat16)

labels = torch.randint(0, vocabsize, (batchsize  seqlen,), device="cuda")


Compute loss
loss = lossfn(logits, labels)
print(f"Cross-entropy loss: {loss.item():.4f}")

LigerFusedLinearCrossEntropyLoss

This is Liger Kernel's most powerful kernel. It fuses the final linear layer (lmhead) with cross-entropy loss, eliminating the need to store the full logits tensor in memory. For models with large vocabularies (e.g., 128K tokens), this provides massive memory savings.

import torch
from ligerkernel.transformers import LigerFusedLinearCrossEntropyLoss

Create fused loss function
fusedlossfn = LigerFusedLinearCrossEntropyLoss()

Simulate hidden states and weight matrix
hiddensize = 4096

vocabsize = 128256
batchseqlen = 1024  # batchsize  seqlen


hiddenstates = torch.randn(batchseqlen, hiddensize, device="cuda", dtype=torch.bfloat16)
weight = torch.randn(vocabsize, hiddensize, device="cuda", dtype=torch.bfloat16)
labels = torch.randint(0, vocabsize, (batchseqlen,), device="cuda")


Compute loss without materializing the full logits tensor!
loss = fusedlossfn(hiddenstates, weight, labels)
print(f"Fused linear cross-entropy loss: {loss.item():.4f}")

import torch
from transformers import AutoModelForCausalLM, AutoTokenizer
from ligerkernel.transformers import applyligerkerneltollama


Activate Liger Kernel before loading the model
applyligerkerneltollama()


Load model and tokenizer
modelname = "meta-llama/Llama-3.1-8B"
model = AutoModelForCausalLM.frompretrained(

    modelname,
    torchdtype=torch.bfloat16,

    devicemap="auto",
)
tokenizer = AutoTokenizer.frompretrained(modelname)
tokenizer.padtoken = tokenizer.eostoken

Setup optimizer
optimizer = torch.optim.AdamW(model.parameters(), lr=2e-5, weightdecay=0.01)


Training loop
model.train()
numepochs = 3
gradientaccumulationsteps = 4

for epoch in range(numepochs):

    totalloss = 0
    optimizer.zerograd()


    # In practice, use a DataLoader with your dataset
    for step in range(100):
        inputids = torch.randint(0, 32000, (2, 512), device="cuda")
        attentionmask = torch.oneslike(inputids)

        labels = inputids.clone()

        # Forward pass - Liger Kernel is used automatically
        outputs = model(
            inputids=inputids,
            attentionmask=attentionmask,
            labels=labels,
        )
        loss = outputs.loss / gradientaccumulationsteps
        totalloss += loss.item()


        # Backward pass
        loss.backward()

        # Update weights every N steps
        if (step + 1) % gradientaccumulationsteps == 0:

            torch.nn.utils.clipgradnorm(model.parameters(), maxnorm=1.0)

            optimizer.step()
            optimizer.zerograd()

            if (step + 1) % 20 == 0:
                avgloss = totalloss / (step + 1)
                print(f"Epoch {epoch+1}, Step {step+1}, Loss: {avgloss:.4f}")


    print(f"Epoch {epoch+1} complete. Average loss: {totalloss / 100:.4f}")

import torch
from transformers import AutoModelForCausalLM, AutoTokenizer, TrainingArguments
from peft import LoraConfig, getpeftmodel, TaskType
from trl import SFTTrainer
from datasets import loaddataset

from ligerkernel.transformers import applyligerkerneltollama

Activate Liger Kernel
applyligerkerneltollama()

Load base model
model = AutoModelForCausalLM.frompretrained(

    "meta-llama/Llama-3.1-8B",
    torchdtype=torch.bfloat16,
    devicemap="auto",

)
tokenizer = AutoTokenizer.frompretrained("meta-llama/Llama-3.1-8B")
tokenizer.padtoken = tokenizer.eostoken

LoRA configuration
loraconfig = LoraConfig(

    tasktype=TaskType.CAUSALLM,

    r=16,
    loraalpha=32,
    loradropout=0.05,

    targetmodules=["qproj", "kproj", "vproj", "oproj",
                     "gateproj", "upproj", "downproj"],

)

Apply LoRA
model = getpeftmodel(model, loraconfig)
model.printtrainableparameters()

Training arguments with Liger Kernel active
trainingargs = TrainingArguments(

    outputdir="./lora-output",
    numtrainepochs=1,
    perdevicetrainbatchsize=8,  # Larger batch size thanks to Liger!
    gradientaccumulationsteps=2,
    learningrate=2e-4,

    bf16=True,
    loggingsteps=10,
    savestrategy="epoch",

    reportto="none",
)

Prepare dataset
dataset = loaddataset("tatsu-lab/alpaca", split="train[:5000]")

dataset = dataset.map(lambda x: {
    "text": f"### Instruction:\n{x['instruction']}\n\n### Response:\n{x['output']}"
})

trainer = SFTTrainer(
    model=model,
    args=trainingargs,
    traindataset=dataset,

    tokenizer=tokenizer,
    datasettextfield="text",

    maxseqlength=2048,

)

trainer.train()

# axolotlconfig.yml
basemodel: meta-llama/Llama-3.1-8B

modeltype: LlamaForCausalLM

Enable Liger Kernel - just this one line!
plugins:
ligerkernel


loadin8bit: false

loadin4bit: true


adapter: qlora
lorar: 16
loraalpha: 32

loradropout: 0.05
loratargetmodules:
qproj

kproj
vproj

oproj
gateproj

upproj
downproj


datasetpreparedpath: lastrunprepared

datasets:
path: tatsu-lab/alpaca
    type: alpaca


sequencelen: 4096
microbatchsize: 4
gradientaccumulationsteps: 4
numepochs: 3

learningrate: 2e-4
bf16: true

accelerate launch -m axolotl.cli.train axolotlconfig.yml

llamafactory-cli train \
    --modelnameorpath meta-llama/Llama-3.1-8B \
    --stage sft \
    --dotrain \

    --finetuningtype lora \
    --lorarank 16 \

    --loraalpha 32 \
    --loratarget qproj,kproj,vproj,oproj,gateproj,upproj,downproj \
    --dataset alpacaen \

    --template llama3 \
    --outputdir ./output \
    --perdevicetrainbatchsize 4 \
    --gradientaccumulationsteps 4 \
    --learningrate 2e-4 \

    --numtrainepochs 3 \

    --bf16 \
    --enableligerkernel true

# trainingconfig.yaml
modelnameorpath: meta-llama/Llama-3.1-8B

stage: sft
dotrain: true
finetuningtype: lora

lorarank: 16
loraalpha: 32

dataset: alpacaen
template: llama3
outputdir: ./output

perdevicetrainbatchsize: 4

gradientaccumulationsteps: 4

learningrate: 2.0e-4
numtrainepochs: 3
bf16: true
enableligerkernel: true

from trl import SFTTrainer, SFTConfig

sftconfig = SFTConfig(

    outputdir="./output",
    useligerkernel=True,
    perdevicetrainbatchsize=4,
    numtrainepochs=3,
    learningrate=2e-5,

    bf16=True,
)

import torch
import time
from transformers import AutoModelForCausalLM, AutoTokenizer

def benchmarktraining(modelname, useliger=False, numsteps=50):

    """Benchmark training throughput and VRAM usage."""
    torch.cuda.resetpeakmemorystats()
    torch.cuda.emptycache()


    if useliger:
        from ligerkernel.transformers import applyligerkerneltollama

        applyligerkerneltollama()


    model = AutoModelForCausalLM.frompretrained(
        modelname,

        torchdtype=torch.bfloat16,
        devicemap="cuda",

    )
    tokenizer = AutoTokenizer.frompretrained(modelname)


    optimizer = torch.optim.AdamW(model.parameters(), lr=2e-5)
    model.train()

    # Warmup
    for  in range(3):
        inputids = torch.randint(0, 32000, (2, 512), device="cuda")

        outputs = model(inputids=inputids, labels=inputids)
        outputs.loss.backward()
        optimizer.step()
        optimizer.zerograd()


    torch.cuda.resetpeakmemorystats()

    # Benchmark
    starttime = time.time()

    totaltokens = 0

    for step in range(numsteps):

        inputids = torch.randint(0, 32000, (2, 512), device="cuda")
        outputs = model(inputids=inputids, labels=inputids)

        outputs.loss.backward()
        optimizer.step()
        optimizer.zerograd()
        totaltokens += 2  512

    elapsed = time.time() - starttime
    peakvram = torch.cuda.maxmemoryallocated() / 10243


    print(f"{'With Liger' if useliger else 'Without Liger'}:")
    print(f"  Peak VRAM: {peakvram:.2f} GB")

    print(f"  Throughput: {totaltokens / elapsed:.0f} tokens/sec")
    print(f"  Total time: {elapsed:.1f} sec")
    print()

    del model, optimizer
    torch.cuda.emptycache()


Run benchmark
modelname = "meta-llama/Llama-3.1-8B"
benchmarktraining(modelname, useliger=False)

benchmarktraining(modelname, useliger=True)

applyligerkerneltollama(
    fusedlinearcrossentropy=True,  # This is the most important one!

    rmsnorm=True,
    swiglu=True,
    rope=True,
)

• Increasing batch size: Larger batch sizes generally result in more stable training and faster convergence.
• Extending sequence length: Longer context helps models learn long-range dependencies.
• Avoiding gradient checkpointing: If VRAM is sufficient, disable gradient checkpointing for maximum throughput.

trainingargs = TrainingArguments(
    perdevicetrainbatchsize=8,    # Increased from 4

    gradientcheckpointing=False,      # Disable if VRAM is sufficient
    maxseqlength=4096,               # Longer sequence length
)

from transformers import TrainingArguments, BitsAndBytesConfig

Combination: Liger Kernel + QLoRA + Flash Attention
bnbconfig = BitsAndBytesConfig(

    loadin4bit=True,

    bnb4bitquanttype="nf4",
    bnb4bitcomputedtype=torch.bfloat16,

    bnb4bitusedoublequant=True,

)

model = AutoModelForCausalLM.frompretrained(
    modelname,

    quantizationconfig=bnbconfig,

    attnimplementation="flashattention2",  # Flash Attention
    devicemap="auto",

)

trainingargs = TrainingArguments(
    useligerkernel=True,            # Liger Kernel
    perdevicetrainbatchsize=16,   # Large batch size!
    bf16=True,
    gradientcheckpointing=True,      # For very large models

    gradientcheckpointingkwargs={"usereentrant": False},
)

import torch
from transformers import AutoModelForCausalLM

Load model WITHOUT Liger
modelvanilla = AutoModelForCausalLM.frompretrained(
    modelname, torchdtype=torch.float32, devicemap="cuda"

)

Load model WITH Liger
from ligerkernel.transformers import applyligerkerneltollama
applyligerkerneltollama()
modelliger = AutoModelForCausalLM.frompretrained(
    modelname, torchdtype=torch.float32, devicemap="cuda"

)

Compare outputs
inputids = torch.randint(0, 32000, (1, 128), device="cuda")
with torch.nograd():

    outvanilla = modelvanilla(inputids).logits
    outliger = modelliger(inputids).logits


diff = (outvanilla - outliger).abs().max()

print(f"Max absolute difference: {diff.item()}")
Should be very small (< 1e-5 for float32)

import torch

def printgpumemory():

    """Print current GPU memory usage."""
    allocated = torch.cuda.memoryallocated() / 10243
    reserved = torch.cuda.memoryreserved() / 10243

    peak = torch.cuda.maxmemoryallocated() / 10243

    print(f"GPU Memory - Allocated: {allocated:.2f} GB, "
          f"Reserved: {reserved:.2f} GB, Peak: {peak:.2f} GB")

Call at various points in your training loop
printgpumemory()

# Check installed versions
pip show liger-kernel torch transformers triton

Update if needed
pip install --upgrade liger-kernel

# If using torch.compile, test compatibility
model = torch.compile(model)
If errors occur, try disabling specific kernels
applyligerkerneltollama(

    fusedlinearcrossentropy=False,  # Sometimes conflicts with compile
)

Common Troubleshooting

Error: Triton Not Found

# Install Triton explicitly
pip install triton>=2.1.0

Error: CUDA Version Mismatch

Ensure your CUDA toolkit is compatible:

# Check CUDA version
nvcc --version
python -c "import torch; print(torch.version.cuda)"

Kernel Crash or NaN Loss

If you experience NaN loss after enabling Liger Kernel:

# Debug: enable kernels one at a time
applyligerkerneltollama(
    rope=False,
    crossentropy=False,

    fusedlinearcrossentropy=False,
    rmsnorm=True,   # Try one at a time

    swiglu=False,
)

By adopting Liger Kernel in your training workflow, you can focus on what truly matters — data quality, model architecture, and hyperparameter tuning — without worrying about hardware limitations. Happy training!

For further questions about LLM training optimization, feel free to reach out at ruby@rubythalib.ai.