benchmark_distill_jsd_loss.py

import os
import sys

import torch
import triton

from utils import QUANTILES
from utils import SingleBenchmarkRunInput
from utils import SingleBenchmarkRunOutput
from utils import _test_memory
from utils import parse_benchmark_script_args
from utils import run_benchmarks

from liger_kernel.chunked_loss.jsd_loss import LigerFusedLinearJSDFunction
from liger_kernel.utils import get_total_gpu_memory
from liger_kernel.utils import infer_device

device = infer_device()

sys.path.insert(0, os.path.abspath(os.path.join(os.path.dirname(__file__), "../..")))


class TorchJSDLoss(torch.nn.Module):
    def __init__(
        self,
        H: int,
        V: int,
        dtype: torch.dtype,
        weight_hard_loss: float = 0.5,
        weight_soft_loss: float = 0.5,
        ignore_index: int = -100,
        temperature: float = 1.0,
        bias: bool = False,
    ):
        from test.chunked_loss.test_jsd_loss import HFJSDLoss

        super().__init__()
        self.student_lin = torch.nn.Linear(in_features=H // 2, out_features=V, bias=bias, dtype=dtype)
        self.teacher_lin = torch.nn.Linear(in_features=H, out_features=V, bias=bias, dtype=dtype)
        self.jsd_loss = HFJSDLoss(
            ignore_index=ignore_index,
            weight_hard_loss=weight_hard_loss,
            weight_soft_loss=weight_soft_loss,
            temperature=temperature,
        ).get_batch_loss_metrics

    def forward(self, student, teacher, target):
        return self.jsd_loss(
            student,
            self.student_lin.weight,
            teacher,
            self.teacher_lin.weight,
            target,
        )


class LigerJSDLoss(torch.nn.Module):
    def __init__(
        self,
        H: int,
        V: int,
        dtype: torch.dtype,
        weight_hard_loss: float = 0.5,
        weight_soft_loss: float = 0.5,
        ignore_index: int = -100,
        temperature: float = 1.0,
        bias: bool = False,
    ):
        super().__init__()
        self.student_lin = torch.nn.Linear(in_features=H // 2, out_features=V, bias=bias, dtype=dtype)
        self.teacher_lin = torch.nn.Linear(in_features=H, out_features=V, bias=bias, dtype=dtype)
        self.weight_hard_loss = weight_hard_loss
        self.weight_soft_loss = weight_soft_loss
        self.ignore_index = ignore_index
        self.temperature = temperature
        self.jsd_loss = LigerFusedLinearJSDFunction.apply

    def forward(self, student, teacher, target):
        return self.jsd_loss(
            student,
            self.student_lin.weight,
            teacher,
            self.teacher_lin.weight,
            target,
            self.student_lin.bias,
            self.teacher_lin.bias,
            self.weight_hard_loss,
            self.weight_soft_loss,
        )


def bench_memory_jsd_loss(input: SingleBenchmarkRunInput) -> SingleBenchmarkRunOutput:
    BT = input.x
    H = input.extra_benchmark_config["H"]
    V = input.extra_benchmark_config["V"]
    dtype = input.extra_benchmark_config["dtype"]
    bias = input.extra_benchmark_config["bias"]
    weight_hard_loss = input.extra_benchmark_config["weight_hard_loss"]
    weight_soft_loss = input.extra_benchmark_config["weight_soft_loss"]
    ignore_index = input.extra_benchmark_config["ignore_index"]
    provider = input.kernel_provider

    torch_jsd_loss = TorchJSDLoss(
        H=H,
        V=V,
        dtype=dtype,
        ignore_index=ignore_index,
        bias=bias,
        weight_hard_loss=weight_hard_loss,
        weight_soft_loss=weight_soft_loss,
    ).to(device)
    liger_jsd_loss = LigerJSDLoss(
        H=H,
        V=V,
        dtype=dtype,
        ignore_index=ignore_index,
        bias=bias,
        weight_hard_loss=weight_hard_loss,
        weight_soft_loss=weight_soft_loss,
    ).to(device)

    _tensor = torch.rand(BT, H // 2, device=device, dtype=dtype)
    student_input1 = _tensor.detach().clone().requires_grad_(True)
    student_input2 = _tensor.detach().clone().requires_grad_(True)

    teacher_input = torch.rand(BT, H, device=device, dtype=dtype)

    target = torch.randint(0, V, (BT,), device=device, dtype=torch.long)

    def fwd():
        if provider == "liger":
            return liger_jsd_loss(student_input1, teacher_input, target)
        elif provider == "torch":
            return torch_jsd_loss(student_input2, teacher_input, target)

    def full():
        y = fwd()
        y.backward()

    mem_50, mem_20, mem_80 = _test_memory(full, _iter=10, quantiles=QUANTILES)
    return SingleBenchmarkRunOutput(
        y_20=mem_20,
        y_50=mem_50,
        y_80=mem_80,
    )


def bench_speed_jsd_loss(input: SingleBenchmarkRunInput) -> SingleBenchmarkRunOutput:
    BT = input.x
    H = input.extra_benchmark_config["H"]
    V = input.extra_benchmark_config["V"]
    dtype = input.extra_benchmark_config["dtype"]
    bias = input.extra_benchmark_config["bias"]
    weight_hard_loss = input.extra_benchmark_config["weight_hard_loss"]
    weight_soft_loss = input.extra_benchmark_config["weight_soft_loss"]
    ignore_index = input.extra_benchmark_config["ignore_index"]
    provider = input.kernel_provider
    mode = input.kernel_operation_mode

    torch_jsd_loss = TorchJSDLoss(
        H=H,
        V=V,
        dtype=dtype,
        ignore_index=ignore_index,
        bias=bias,
        weight_hard_loss=weight_hard_loss,
        weight_soft_loss=weight_soft_loss,
    ).to(device)
    liger_jsd_loss = LigerJSDLoss(
        H=H,
        V=V,
        dtype=dtype,
        ignore_index=ignore_index,
        bias=bias,
        weight_hard_loss=weight_hard_loss,
        weight_soft_loss=weight_soft_loss,
    ).to(device)

    _tensor = torch.rand(BT, H // 2, device=device, dtype=dtype)
    student_input1 = _tensor.detach().clone().requires_grad_(True)
    student_input2 = _tensor.detach().clone().requires_grad_(True)

    teacher_input = torch.rand(BT, H, device=device, dtype=dtype)

    target = torch.randint(0, V, (BT,), device=device, dtype=torch.long)

    def fwd():
        if provider == "liger":
            return liger_jsd_loss(student_input1, teacher_input, target)
        elif provider == "torch":
            return torch_jsd_loss(student_input2, teacher_input, target)

    if mode == "forward":
        ms_50, ms_20, ms_80 = triton.testing.do_bench(
            fwd,
            rep=100,
            quantiles=QUANTILES,
        )
    elif mode == "backward":
        y = fwd()
        ms_50, ms_20, ms_80 = triton.testing.do_bench(
            lambda: y.backward(retain_graph=True),
            grad_to_none=[student_input1, student_input2],
            rep=100,
            quantiles=QUANTILES,
        )
    elif mode == "full":

        def full():
            y = fwd()
            y.backward()

        ms_50, ms_20, ms_80 = triton.testing.do_bench(
            full,
            rep=100,
            quantiles=QUANTILES,
        )

    return SingleBenchmarkRunOutput(
        y_20=ms_20,
        y_50=ms_50,
        y_80=ms_80,
    )


if __name__ == "__main__":
    args = parse_benchmark_script_args()
    gpu_memory_gbs = get_total_gpu_memory()
    # We know that the full test will require 69GBs for vocab size 2^13 and 39GBs for vocab size 2^12 on torch
    if gpu_memory_gbs >= 69:
        x_max = 13
    elif gpu_memory_gbs >= 39:
        x_max = 12
    else:
        x_max = 11

    common_configs = {
        "kernel_name": "distill_jsd_loss",
        "x_name": "BT",
        "x_label": "B x T",
        "x_values": [2**i for i in range(10, x_max + 1)],
        "kernel_providers": ["liger", "torch"],
        "extra_benchmark_configs": [
            {
                "H": 4096,
                "V": 128256,
                "mode": "forward",
                "dtype": torch.bfloat16,
                "bias": False,
                "weight_hard_loss": 0.5,
                "weight_soft_loss": 0.5,
                "ignore_index": -100,
            }
        ],
        "overwrite": args.overwrite,
    }

    run_benchmarks(
        bench_test_fn=bench_speed_jsd_loss,
        kernel_operation_modes=["forward", "backward", "full"],
        metric_name="speed",
        metric_unit="ms",
        **common_configs,
    )

    run_benchmarks(
        bench_test_fn=bench_memory_jsd_loss,
        kernel_operation_modes=["full"],
        metric_name="memory",
        metric_unit="MB",
        **common_configs,
    )