| 1 | # SPDX-License-Identifier: Apache-2.0 |
| 2 | # SPDX-FileCopyrightText: Copyright contributors to the vLLM project |
| 3 | |
| 4 | import torch |
| 5 | |
| 6 | from vllm.model_executor.layers.fusion.quant_activation import QuantizedActivation |
| 7 | from vllm.model_executor.layers.linear import LinearBase |
| 8 | from vllm.model_executor.layers.quantization.utils.quant_utils import ( |
| 9 | get_fp8_min_max, |
| 10 | kFp8StaticTensorSym, |
| 11 | ) |
| 12 | from vllm.platforms import current_platform |
| 13 | from vllm.triton_utils import tl, triton |
| 14 | |
| 15 | _FP8_MIN, _FP8_MAX = get_fp8_min_max() |
| 16 | |
| 17 | |
| 18 | @triton.jit |
| 19 | def _relu_squared_static_fp8_quant_kernel( |
| 20 | x_ptr, |
| 21 | scale_ptr, |
| 22 | output_ptr, |
| 23 | num_elements, |
| 24 | FP8_MIN: tl.constexpr, |
| 25 | FP8_MAX: tl.constexpr, |
| 26 | BLOCK_SIZE: tl.constexpr, |
| 27 | ): |
| 28 | offsets = tl.program_id(0) * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE) |
| 29 | mask = offsets < num_elements |
| 30 | x = tl.load(x_ptr + offsets, mask=mask, other=0.0).to(tl.float32) |
| 31 | input_is_nan = x != x |
| 32 | relu = tl.maximum(x, 0.0) |
| 33 | # Preserve the standalone ReLU2 kernel's BF16 output boundary. |
| 34 | activated = (relu * relu).to(tl.bfloat16).to(tl.float32) |
| 35 | # Match O2's Inductor-compiled QuantFP8 reciprocal-then-multiply. |
| 36 | quantized = activated * (1.0 / tl.load(scale_ptr)) |
| 37 | quantized = tl.maximum(tl.minimum(quantized, FP8_MAX), FP8_MIN) |
| 38 | # Match the O2 native chain by preserving NaN through the FP8 conversion. |
| 39 | quantized = tl.where(input_is_nan, x, quantized) |
| 40 | tl.store(output_ptr + offsets, quantized, mask=mask) |
| 41 | |
| 42 | |
| 43 | def relu_squared_static_fp8_quant( |
| 44 | x: torch.Tensor, linear: LinearBase |
| 45 | ) -> QuantizedActivation: |
| 46 | """BF16 ReLU2 followed by static per-tensor FP8 quantization.""" |
| 47 | assert x.dtype == torch.bfloat16 |
| 48 | assert x.is_contiguous() |
| 49 | scale = linear.input_scale |
| 50 | assert scale.dtype == torch.float32 |
| 51 | assert x.device == scale.device |
| 52 | assert scale.numel() == 1 |
| 53 | |
| 54 | output = torch.empty_like(x, dtype=current_platform.fp8_dtype()) |
| 55 | if x.numel() != 0: |
| 56 | block_size = min(triton.next_power_of_2(x.shape[-1]), 2048) |
| 57 | num_warps = min(max(block_size // 256, 1), 4) |
| 58 | grid = lambda meta: (triton.cdiv(x.numel(), meta["BLOCK_SIZE"]),) |
| 59 | _relu_squared_static_fp8_quant_kernel[grid]( |
| 60 | x, |
| 61 | scale, |
| 62 | output, |
| 63 | x.numel(), |
| 64 | FP8_MIN=_FP8_MIN, |
| 65 | FP8_MAX=_FP8_MAX, |
| 66 | BLOCK_SIZE=block_size, |
| 67 | num_warps=num_warps, |
| 68 | ) |
| 69 | |
| 70 | return QuantizedActivation( |
| 71 | data=output, |
| 72 | scale=scale, |
| 73 | orig_dtype=x.dtype, |
| 74 | orig_shape=x.shape, |
| 75 | quant_key=kFp8StaticTensorSym, |
| 76 | ) |