@tduncan / Nvidia-Optimization-Stack-VLLM Mirror · read-only / commits / a31614e386
Fork 0

A high-throughput and memory-efficient inference and serving engine for LLMs

Mirrored from https://github.com/tduncan-lum/Nvidia-Optimization-Stack---VLLM.git

Synced Nov 15, 22:15 UTC

[ROCm][Quantization][Kernel] Use FP8 FNUZ when OCP flag is 0 or undefined (#13851)

Signed-off-by: Hollow Man <hollowman@opensuse.org>

ℍ𝕠𝕝𝕝𝕠𝕨 𝕄𝕒𝕟 committed Feb 27, 2025 at 04:39 UTC a31614e386edf86ea949dd7214b13f82f734b05c
1 file changed +4 -4
csrc/quantization/fp8/amd/quant_utils.cuh
+4 -4
@@ -24,12 +24,12 @@ __inline__ __device__ Tout scaled_vec_conversion(const Tin& x,
24 return x;
25 }
26
27 - #if HIP_FP8_TYPE_FNUZ
28 -using fp8_type = __hip_fp8_e4m3_fnuz;
29 -using fp8x2_type = __hip_fp8x2_e4m3_fnuz;
30 - #elif HIP_FP8_TYPE_OCP
27 + #if HIP_FP8_TYPE_OCP
28 using fp8_type = __hip_fp8_e4m3;
29 using fp8x2_type = __hip_fp8x2_e4m3;
30 + #else
31 +using fp8_type = __hip_fp8_e4m3_fnuz;
32 +using fp8x2_type = __hip_fp8x2_e4m3_fnuz;
33 #endif
34
35 // fp8 -> half