-
Notifications
You must be signed in to change notification settings - Fork 252
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
Add JSD Loss for Distillation (#425)
## Summary > [!CAUTION] > This PR depends on #417. Do not merge until #417 (later #432) is merged. This is a pure torch compiled, chunked fused linear JSD Loss, aiming for knowledge distillation. #### Jensen-Shannon Divergence Loss This PR implements Jensen-Shannon Divergence (JSD) loss as the soft learning objective in a distillation setting (teacher & student). This component can be replaced with other losses (e.g., KL divergence) as `distillation_loss_fn`. JSD is defined as the average of the KL divergences between each distribution and the mean distribution: ```math \text{JSD}(P || Q) = \frac{1}{2} \text{KL}(P || M) + \frac{1}{2} \text{KL}(Q || M), \quad \text{where } M = \frac{1}{2}(P + Q) ``` Here, `P`and `Q` are the two probability distributions, and `M` is their average. ## Testing Done Below figures are benchmark results with different `chunk_size`, which also significantly affects performance. #### Hint: User can tune their `chunk_size` as suggested by the liger [paper](https://arxiv.org/pdf/2306.13649) for the moment: ```math 2^{\lceil \log_2 \lceil \frac{BT}{V/H} \rceil \rceil} ``` #### Memory 1. `chunk_size` = 1  2. `chunk_size` = 1024  #### Speed (Elapsed Time) 1. `chunk_size` = 1  2. `chunk_size` = 1024  - Hardware Type: NVIDIA H100 80GB HBM3 (SXM5) - [X] run `make test` to ensure correctness - [X] run `make checkstyle` to ensure code style - [ ] run `make test-convergence` to ensure convergence --------- Signed-off-by: Austin Liu <[email protected]>
- Loading branch information
1 parent
b80bf95
commit aa2d23d
Showing
8 changed files
with
778 additions
and
6 deletions.
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,261 @@ | ||
| 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 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.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() | ||
|
|
||
| common_configs = { | ||
| "kernel_name": "distill_jsd_loss", | ||
| "x_name": "BT", | ||
| "x_label": "B x T", | ||
| "x_values": [2**i for i in range(10, 14)], | ||
| "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", "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, | ||
| ) |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -1,5 +1,6 @@ | ||
| from liger_kernel.chunked_loss.cpo_loss import LigerFusedLinearCPOLoss # noqa: F401 | ||
| from liger_kernel.chunked_loss.dpo_loss import LigerFusedLinearDPOLoss # noqa: F401 | ||
| from liger_kernel.chunked_loss.jsd_loss import LigerFusedLinearJSDLoss # noqa: F401 | ||
| from liger_kernel.chunked_loss.kto_loss import LigerFusedLinearKTOLoss # noqa: F401 | ||
| from liger_kernel.chunked_loss.orpo_loss import LigerFusedLinearORPOLoss # noqa: F401 | ||
| from liger_kernel.chunked_loss.simpo_loss import LigerFusedLinearSimPOLoss # noqa: F401 |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -1,11 +1,13 @@ | ||
| from liger_kernel.chunked_loss.cpo_loss import LigerFusedLinearCPOFunction | ||
| from liger_kernel.chunked_loss.dpo_loss import LigerFusedLinearDPOFunction | ||
| from liger_kernel.chunked_loss.jsd_loss import LigerFusedLinearJSDFunction | ||
| from liger_kernel.chunked_loss.kto_loss import LigerFusedLinearKTOFunction | ||
| from liger_kernel.chunked_loss.orpo_loss import LigerFusedLinearORPOFunction | ||
| from liger_kernel.chunked_loss.simpo_loss import LigerFusedLinearSimPOFunction | ||
|
|
||
| liger_fused_linear_orpo = LigerFusedLinearORPOFunction.apply | ||
| liger_fused_linear_dpo = LigerFusedLinearDPOFunction.apply | ||
| liger_fused_linear_jsd = LigerFusedLinearJSDFunction.apply | ||
| liger_fused_linear_cpo = LigerFusedLinearCPOFunction.apply | ||
| liger_fused_linear_simpo = LigerFusedLinearSimPOFunction.apply | ||
| liger_fused_linear_kto = LigerFusedLinearKTOFunction.apply |
Oops, something went wrong.