-
Notifications
You must be signed in to change notification settings - Fork 243
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
Signed-off-by: Austin Liu <[email protected]>
- Loading branch information
1 parent
46298ac
commit 7799ef1
Showing
2 changed files
with
294 additions
and
0 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,270 @@ | ||
import os | ||
import sys | ||
|
||
import torch | ||
import triton | ||
from utils import ( | ||
QUANTILES, | ||
SingleBenchmarkRunInput, | ||
SingleBenchmarkRunOutput, | ||
_test_memory, | ||
parse_benchmark_script_args, | ||
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 | ||
) |