94057c3d3e
PR Test (NPU) / check-changes (push) Has been cancelled
PR Test (NPU) / pr-gate (push) Has been cancelled
PR Test (NPU) / set-image-config (push) Has been cancelled
PR Test (NPU) / stage-b-test-1-npu-a2 (0) (push) Has been cancelled
PR Test (NPU) / stage-b-test-1-npu-a2 (1) (push) Has been cancelled
PR Test (NPU) / stage-b-test-2-npu-a2 (0) (push) Has been cancelled
PR Test (NPU) / stage-b-test-2-npu-a2 (1) (push) Has been cancelled
PR Test (NPU) / stage-b-test-4-npu-a3 (push) Has been cancelled
PR Test (NPU) / stage-b-test-16-npu-a3 (push) Has been cancelled
PR Test (NPU) / multimodal-gen-test-1-npu-a3 (push) Has been cancelled
PR Test (NPU) / multimodal-gen-test-2-npu-a3 (push) Has been cancelled
PR Test (Arm64) / pr-gate (push) Has been cancelled
PR Test (Arm64) / check-changes (push) Has been cancelled
PR Test (Arm64) / build-test (push) Has been cancelled
PR Test (sgl-router) / gate (push) Has been cancelled
PR Test (sgl-router) / tier-1 — lint (push) Has been cancelled
PR Test (sgl-router) / tier-2 — build + test (push) Has been cancelled
PR Test (sgl-router) / tier-3 — docker (placeholder) (push) Has been cancelled
PR Test (sgl-router) / tier-3 — k8s integration (push) Has been cancelled
PR Test (sgl-router) / tier-3 — e2e (push) Has been cancelled
PR Test (sgl-router) / finish (push) Has been cancelled
PR Test (NPU) / single-node-poc (map[name:qwen3_6_27b_w8a8_1p_in64k_out1k_50ms runner:linux-aarch64-a3-2 test_case:test/registered/ascend/performance/qwen3_6_27b/test_npu_qwen3_6_27b_w8a8_1p_in64k_out1k_50ms.py test_type:perf]) (push) Has been cancelled
PR Test (NPU) / pr-test-npu-finish (push) Has been cancelled
PR Test (Xeon) / pr-gate (push) Has been cancelled
PR Test (Xeon) / check-changes (push) Has been cancelled
PR Test (Xeon) / build-test (, xeon-gnr, base-b-test-cpu) (push) Has been cancelled
PR Test (XPU) / check-changes (push) Has been cancelled
PR Test (XPU) / pr-gate (push) Has been cancelled
PR Test (XPU) / stage-a-test-1-gpu-xpu (push) Has been cancelled
PR Test (XPU) / wait-for-stage-a (push) Has been cancelled
PR Test (XPU) / stage-b-test-1-gpu-xpu (push) Has been cancelled
PR Test (XPU) / finish (push) Has been cancelled
CI Model Inventory / build-inventory (push) Has been cancelled
Lint / lint (push) Has been cancelled
PR Benchmark (SMG Components) / Benchmark Compilation Check (push) Has been cancelled
PR Benchmark (SMG Components) / Benchmark - Manual Policy (push) Has been cancelled
PR Benchmark (SMG Components) / Benchmark - Request Processing (push) Has been cancelled
PR Benchmark (SMG Components) / Benchmark Summary (push) Has been cancelled
PR Test (SMG) / build-wheel (push) Has been cancelled
Release SGLang Model Gateway to PyPI / build on windows (x86_64 - auto) (push) Has been cancelled
Release SGLang Model Gateway to PyPI / build on macos (x86_64 - auto) (push) Has been cancelled
PR Test (SMG) / python-unit-tests (push) Has been cancelled
PR Test (SMG) / unit-tests (push) Has been cancelled
PR Test (SMG) / benchmarks (push) Has been cancelled
PR Test (SMG) / chat-completions (push) Has been cancelled
PR Test (SMG) / chat-completions-4gpu (push) Has been cancelled
PR Test (SMG) / e2e (push) Has been cancelled
PR Test (SMG) / docker-build-test (push) Has been cancelled
PR Test (SMG) / k8s-integration (push) Has been cancelled
PR Test (SMG) / finish (push) Has been cancelled
PR Test (SMG) / summarize-benchmarks (push) Has been cancelled
Release SGLang Model Gateway Docker Image / publish (push) Has been cancelled
Release SGLang Model Gateway to PyPI / build on macos (aarch64 - auto) (push) Has been cancelled
Release SGLang Model Gateway to PyPI / build on linux (aarch64 - auto) (push) Has been cancelled
Release SGLang Model Gateway to PyPI / build on linux (x86_64 - auto) (push) Has been cancelled
Release SGLang Model Gateway to PyPI / build on linux (aarch64 - musllinux_1_1) (push) Has been cancelled
Release SGLang Model Gateway to PyPI / build on linux (x86_64 - musllinux_1_1) (push) Has been cancelled
Release SGLang Model Gateway to PyPI / Build SDist (push) Has been cancelled
Release SGLang Model Gateway to PyPI / Upload to PyPI (push) Has been cancelled
Release SGLang Kernels / build-cu129-matrix (aarch64, 12.9, 3.10, arm-kernel-build-node) (push) Has been cancelled
Release SGLang Kernels / build-cu129-matrix (x86_64, 12.9, 3.10, x64-kernel-build-node) (push) Has been cancelled
Release SGLang Kernels / release-cu129 (push) Has been cancelled
Release SGLang Kernels / build-cu130-matrix (aarch64, 13.0, 3.10, arm-kernel-build-node) (push) Has been cancelled
Release SGLang Kernels / build-cu130-matrix (x86_64, 13.0, 3.10, x64-kernel-build-node) (push) Has been cancelled
Release SGLang Kernels / release-cu130 (push) Has been cancelled
Release SGLang Kernels / build-rocm-matrix (3.10, 700) (push) Has been cancelled
Release SGLang Kernels / build-rocm-matrix (3.10, 720) (push) Has been cancelled
Release SGLang Kernels / release-rocm700 (push) Has been cancelled
Release SGLang Kernels / release-rocm720 (push) Has been cancelled
Release SGLang Kernels / build-musa43 (43, 3.10) (push) Has been cancelled
Release SGLang Kernels / release-musa43 (push) Has been cancelled
96 lines
2.8 KiB
Python
96 lines
2.8 KiB
Python
import argparse
|
|
|
|
import torch
|
|
import triton
|
|
from vllm._custom_ops import scaled_int8_quant as vllm_scaled_int8_quant
|
|
|
|
from sglang.benchmark.bench_utils import run_bench
|
|
from sglang.srt.layers.quantization.int8_kernel import per_token_quant_int8
|
|
|
|
|
|
@torch.compile(backend="inductor")
|
|
def torch_int8_quant(x):
|
|
int8_max = torch.iinfo(torch.int8).max
|
|
|
|
abs_max = x.abs().max(dim=-1, keepdim=True).values
|
|
scales = abs_max.to(torch.float32) / float(int8_max)
|
|
|
|
q_x = (x / scales).round().to(torch.int8)
|
|
|
|
return q_x, scales
|
|
|
|
|
|
def _test_accuracy_once(M, K, input_dtype, device):
|
|
x = torch.randn(M, K, dtype=input_dtype, device=device) * 5000
|
|
out, scales, _ = vllm_scaled_int8_quant(x, symmetric=True)
|
|
out1, scales1 = per_token_quant_int8(x)
|
|
out2, scales2 = torch_int8_quant(x)
|
|
torch.testing.assert_close(out, out2, atol=1, rtol=0)
|
|
torch.testing.assert_close(out, out1, atol=1, rtol=0)
|
|
torch.testing.assert_close(scales, scales2)
|
|
torch.testing.assert_close(scales1, scales2)
|
|
print(f"M: {M}, K: {K}, type: {input_dtype} OK")
|
|
|
|
|
|
def test_accuracy():
|
|
Ms = [1, 13, 128, 1024, 2048, 4096]
|
|
Ks = [512, 1024, 2048, 8192]
|
|
input_dtypes = [torch.float16, torch.bfloat16]
|
|
for M in Ms:
|
|
for K in Ks:
|
|
for input_dtype in input_dtypes:
|
|
_test_accuracy_once(M, K, input_dtype, "cuda")
|
|
|
|
|
|
@triton.testing.perf_report(
|
|
triton.testing.Benchmark(
|
|
x_names=["batch_size"],
|
|
x_vals=[1, 16, 32, 64, 128, 256, 512, 1024, 2048],
|
|
x_log=False,
|
|
line_arg="provider",
|
|
line_vals=["vllm op", "triton", "torch.compile"],
|
|
line_names=["vllm op", "triton", "torch.compile"],
|
|
styles=[("blue", "-"), ("orange", "-"), ("red", "-")],
|
|
ylabel="ms",
|
|
plot_name="int8 per token quant",
|
|
args={},
|
|
)
|
|
)
|
|
def benchmark(batch_size, provider):
|
|
M, K = batch_size, 16384
|
|
x = torch.randn(M, K, dtype=torch.float16, device="cuda") * 1000
|
|
|
|
quantiles = (0.5, 0.2, 0.8)
|
|
if provider == "vllm op":
|
|
ms, min_ms, max_ms = run_bench(
|
|
lambda: vllm_scaled_int8_quant(x, symmetric=True),
|
|
quantiles=quantiles,
|
|
)
|
|
if provider == "triton":
|
|
ms, min_ms, max_ms = run_bench(
|
|
lambda: per_token_quant_int8(x),
|
|
quantiles=quantiles,
|
|
)
|
|
if provider == "torch.compile":
|
|
ms, min_ms, max_ms = run_bench(
|
|
lambda: torch_int8_quant(x),
|
|
quantiles=quantiles,
|
|
)
|
|
|
|
return ms, min_ms, max_ms
|
|
|
|
|
|
if __name__ == "__main__":
|
|
parser = argparse.ArgumentParser()
|
|
parser.add_argument(
|
|
"--save_path",
|
|
type=str,
|
|
default="./bench_int8_quant_res",
|
|
help="Path to save int8 quant benchmark results",
|
|
)
|
|
args = parser.parse_args()
|
|
|
|
test_accuracy()
|
|
|
|
benchmark.run(print_data=True, show_plots=True, save_path=args.save_path)
|