sycl: fix UE4M3 parsing (#25608)
The NVFP4 quantization format stores a scaling factor for every group of 16 weights, packed into a single UE4M3 byte. The SYCL GPU code was converting these scale values using the E4M3 path, but that's *signed*, and these are unsigned values.
This commit is contained in:
@@ -1022,9 +1022,20 @@ static T block_reduce(T val, T * shared_vals, int block_size_template) {
|
|||||||
}
|
}
|
||||||
|
|
||||||
static __dpct_inline__ float ggml_sycl_ue4m3_to_fp32(uint8_t x) {
|
static __dpct_inline__ float ggml_sycl_ue4m3_to_fp32(uint8_t x) {
|
||||||
const uint32_t bits = x * (x != 0x7F && x != 0xFF);
|
// UE4M3 is unsigned: 4 exp bits (bias 7), 3 mantissa bits, no sign, no NaN.
|
||||||
const __nv_fp8_e4m3 xf = *reinterpret_cast<const __nv_fp8_e4m3 *>(&bits);
|
// exp == 0xF is a valid exponent (256-448 range), not NaN.
|
||||||
return static_cast<float>(xf) / 2;
|
if (x == 0 || x == 0x7F) {
|
||||||
|
return 0.0f;
|
||||||
|
}
|
||||||
|
const int exp = (x >> 3) & 0xF;
|
||||||
|
const int man = x & 0x7;
|
||||||
|
float raw;
|
||||||
|
if (exp == 0) {
|
||||||
|
raw = man * (1.0f / 8.0f) * sycl::pow(2.0f, -6.0f);
|
||||||
|
} else {
|
||||||
|
raw = (1.0f + man / 8.0f) * sycl::pow(2.0f, (float) exp - 7.0f);
|
||||||
|
}
|
||||||
|
return raw * 0.5f;
|
||||||
}
|
}
|
||||||
|
|
||||||
#endif // GGML_SYCL_COMMON_HPP
|
#endif // GGML_SYCL_COMMON_HPP
|
||||||
|
|||||||
Reference in New Issue
Block a user