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, 19:45 UTC
forward fix PR 14245, restore build on ROCm 6.2 (#14709)
Signed-off-by: Jeff Daily <jeff.daily@amd.com>
Jeff Daily committed
Mar 13, 2025 at 20:40 UTC
2a602b055a180c982126e8a438d188325fdb01a5
1 file changed
+12
csrc/quantization/fp8/amd/quant_utils.cuh
+12
@@ -19,12 +19,24 @@ __device__ __forceinline__ fp8_type cvt_c10(float const r) {
19
return {};
20
}
21
22
+// __hip_fp8_e4m3 only exists starting in ROCm 6.3. The macro
23
+// HIP_FP8_TYPE_OCP comes from the hip_fp8.h header and also makes
24
+// its first appearance in ROCm 6.3. Since VLLM_DISPATCH_FP8_TYPES
25
+// on ROCm instantiates both OCP and FNUZ kernels, we need to replace
26
+// the new HW cvt with something reasonable that doesn't rely on the
27
+// ROCm 6.3 feature. This allows compiling on ROCm 6.2 or newer.
28
template <>
29
__device__ __forceinline__ c10::Float8_e4m3fn cvt_c10(float const r) {
30
+ #if HIP_FP8_TYPE_OCP
31
return c10::Float8_e4m3fn(
32
__hip_cvt_float_to_fp8(r, __hip_fp8_e4m3::__default_saturation,
33
__hip_fp8_e4m3::__default_interpret),
34
c10::Float8_e4m3fn::from_bits());
35
+ #else
36
+ // Cast implemented by pytorch. Uses bit manipulation instead of HW cvt.
37
+ // HW cvt above is faster when it is available (ROCm 6.3 or newer).
38
+ return static_cast<c10::Float8_e4m3fn>(r);
39
+ #endif
40
}
41
42
template <>