Replies: 16 comments 51 replies
-
|
https://github.com/IntelLabs/FP8-Emulation-Toolkit Look a good starting point I find some day ago... I'll have a look. |
Beta Was this translation helpful? Give feedback.
-
|
I can't find any code in there about quantizing / dequantizing fp8. I need something with clarity. |
Beta Was this translation helpful? Give feedback.
-
|
I would even accept a hash map generated from the CUDA API. |
Beta Was this translation helpful? Give feedback.
-
|
There is... Don't "waste" your time with that now, I will study that and hopefuly i'll can give you something simple. 🤞 |
Beta Was this translation helpful? Give feedback.
-
|
What a grim reminder that just because something's open source, doesn't mean it's open source. |
Beta Was this translation helpful? Give feedback.
-
|
https://arxiv.org/pdf/2209.05433v2 I'll have to read it more carefully, but look like we need to apply a scale factor to. Possibly is common for all weight or per weight. |
Beta Was this translation helpful? Give feedback.
-
// optimized :
static float from_fp8(uint8_t fp8) {
union {
float f;
uint32_t i;
} u;
uint32_t t = fp8;
int exp = ((fp8 >> 3) & 0x0F) - 7;
u.i = (t&0x80)<<24; // signe
u.i |= (exp+127)<<23; // exponent
u.i |= (t&0x7)<<(20) ; // mantissa: bit 2-0 -> 22-20
return u.f;
}
[edit]: it miss subnormaldon't have time to test it now... I'll do it tomorrow (with more optimisation / correction?) |
Beta Was this translation helpful? Give feedback.
-
|
Thanks! I've pushed an fp8 branch with our latest FP8 code. llamafile/llama.cpp/ggml-impl.h Lines 123 to 166 in 42fa422 You can build and use the branch like this: You can get that gguf file as follows: |
Beta Was this translation helpful? Give feedback.
-
|
In order to make this work, we need an arithmetic expression for encoding Nemo's precious subnormals in |
Beta Was this translation helpful? Give feedback.
-
|
Nice work! 👍 |
Beta Was this translation helpful? Give feedback.
-
|
for nemo llamafile when extract the GGUF I get: not 131072 (128k) as the max ? |
Beta Was this translation helpful? Give feedback.
-
static float from_fp8(uint8_t fp8) {
union {
float f;
uint32_t i;
} u;
const uint32_t t = fp8;
if ((fp8 & 0x7F) == 0x7F) return (fp8==0xFF)?(-NAN):(+NAN);
int exp = ((fp8 >> 3) & 0x0F) - 7;
if (fp8 & 0x78) {
u.i = (t & 7) << 20; // mantissa: bit 2-0 -> 22-20
} else if (fp8 & 0x04) { // denormalise:
u.i = (t & 3) << 21;
exp = -7;
} else if (fp8 & 0x02) {
u.i = (t & 1) << 22;
exp = -8;
} else if (fp8 & 0x01) {
u.i = 0;
exp = -9;
} else {
u.i = 0;
exp = -127;
}
u.i |= (exp + 127) << 23; // exponent
u.i |= (t & 128) << 24; // sign
return u.f;
}with more arithmetic expression, not sure it is best. At least the calculation is good (or at least the same as the current one). |
Beta Was this translation helpful? Give feedback.
-
|
Thanks for spotting the pattern. You've got the magic eye. Those else clauses can be made into the expression: static float from_fp8(uint8_t fp8) {
union {
float f;
uint32_t i;
} u;
const uint32_t t = fp8;
if ((fp8 & 0x7F) == 0x7F)
return (fp8 == 0xFF) ? (-NAN) : (+NAN);
int exp = ((fp8 >> 3) & 0x0F) - 7;
if (fp8 & 0x78) {
u.i = (t & 7) << 20; // mantissa: bit 2-0 -> 22-20
} else if (!(fp8 & 7)) {
u.i = 0;
exp = -127;
} else {
int lg2mant = bsr(fp8 & 7);
u.i = ((t & 3) << (23 - lg2mant)) & 0x007fffff;
exp = -9 + lg2mant;
}
u.i |= (exp + 127) << 23; // exponent
u.i |= (t & 128) << 24; // sign
return u.f;
}But no vector ISA really has a floor(log2(x)) instruction. So, by applying the select trick: We can decode subnormals as follows: static float from_fp8(uint8_t fp8) {
union {
float f;
uint32_t i;
} u;
const uint32_t t = fp8;
if ((fp8 & 0x7F) == 0x7F)
return (fp8 == 0xFF) ? (-NAN) : (+NAN);
int exp = ((fp8 >> 3) & 0x0F) - 7;
if (fp8 & 0x78) {
u.i = (t & 7) << 20; // mantissa: bit 2-0 -> 22-20
} else if (!(fp8 & 7)) {
u.i = 0;
exp = -127;
} else {
uint32_t mant = fp8 & 7;
uint32_t p1 = (mant >> 2) - 1;
uint32_t p2 = ((mant >> 1) & 1) - 1;
int lg2mant = 2 & ~p1;
lg2mant |= 1 & ~p2 & p1;
u.i = ((t & 3) << (23 - lg2mant)) & 0x007fffff;
exp = -9 + lg2mant;
}
u.i |= (exp + 127) << 23; // exponent
u.i |= (t & 128) << 24; // sign
return u.f;
} |
Beta Was this translation helpful? Give feedback.
-
|
Nice. We can decode 16 fp8 numbers in 24 cycles on znver4 using this.
|
Beta Was this translation helpful? Give feedback.
-
|
I love the "timeline view"... how you get it? |
Beta Was this translation helpful? Give feedback.
-
|
I do a few statistical calculations on weights (mean, standard deviation, min max)... |
Beta Was this translation helpful? Give feedback.
-
|
I wrote an avx512 vectorized fp8 loader. On my znver4 Threadripper I managed to get prompt processing tokens per second up from 10 tok/sec to 56 tok/sec on my workstation for Mistral Nemo. See the commit to my fp8 branch here: c10a65c On the other hand, using BF16 gives me 385.92 tok/sec for prompt processing. The K quants probably go just as fast, and token generation goes even faster. So unless someone can come up with a significantly better faster way of dequantizing FP8 then I really don't think it has anything to offer us until native CPU hardware support becomes available. Except it is worth supporting, because NVIDIA has hardware support for it. However I don't own an NVIDIA graphics card that has FP8 yet. If someone is willing to devote the eng resources to implementing CUDA support, then I will merge that and ensure it at least works on CPU too. Although realistically, anyone who chooses FP8 will only be interested in doing it on GPU. There's also a question of what you'd do for AMD GPUs. As I mentioned, as long as it can be made to work I'm happy. Even if it's only primarily good for NVIDIA GPU owners. That's my judgement for now. If anyone wants to volunteer, then send me pull requests to my FP8 branch and we'll keep working on it. |
Beta Was this translation helpful? Give feedback.
-
static float from_fp8_4(uint8_t fp8) {
union {
float f;
uint32_t i;
} u;
const uint32_t t = fp8;
if ((fp8 & 0x7F) == 0x7F) return (fp8==0xFF)?(-NAN):(+NAN);
int exp = ((fp8 >> 3) & 0x0F) - 7;
if (fp8 & 0x78) {
u.i = (t & 7) << 20; // mantissa: bit 2-0 -> 22-20
u.i |= (exp + 127) << 23; // exponent
} else { // denormalise:
float m = (float)(t & 0x7);
u.f = m/ (1<<9);
}
u.i |= (t & 128) << 24; // sign
return u.f;
}I test this one... and it work. and __m256h _mm256_cvtepi16_ph (when available) can do the float cast vecterized. https://www.intel.com/content/www/us/en/docs/intrinsics-guide/index.html#text=_mm256_cvt |
Beta Was this translation helpful? Give feedback.
-
static float from_fp8_5(uint8_t fp8) {
union {
float f;
uint32_t i;
} u;
const uint32_t t = fp8;
if ((fp8 & 0x7F) == 0x7F) return (fp8==0xFF)?(-NAN):(+NAN);
int exp = ((fp8 >> 3) & 0x0F) - 7;
if (fp8 & 0x78) {
u.i = (t & 7) << 20; // mantissa: bit 2-0 -> 22-20
u.i |= (exp + 127) << 23; // exponent
} else { // denormalise:
u.i = (t & 7) << (20-6);
u.i |= (0 + 127) << 23; // exponent
u.f -= 1;
}
u.i |= (t & 128) << 24; // sign
return u.f;
}This one works too. |
Beta Was this translation helpful? Give feedback.
-
static float from_fp8_6(uint8_t fp8) {
union {
float f;
uint32_t i;
} u;
const uint32_t t = fp8;
if ((fp8 & 0x7F) == 0x7F) return (fp8==0xFF)?(-NAN):(+NAN);
u.i = (t & 7) << 20; // mantissa: bit 2-0 -> 22-20
if (fp8 & 0x78) {
const int exp = ((fp8 >> 3) & 0x0F) - 7;
u.i |= (exp + 127) << 23; // exponent
} else { // denormalise:
u.i |= (-6 + 127) << 23; // exponent
u.f -= 1.0/64; // 2⁻⁶
}
u.i |= (t & 128) << 24; // sign
return u.f;
}this forme can be use with fp16 I think (for ARM?) |
Beta Was this translation helpful? Give feedback.
-
static float from_fp8_7(uint8_t fp8) {
union {
float f;
uint32_t i;
} u;
const uint32_t t = fp8;
if ((fp8 & 0x7F) == 0x7F) return (fp8==0xFF)?(-NAN):(+NAN);
auto denorm = fp8 & 0x78;
int exp = -6;
u.i = (t & 7) << 20; // mantissa: bit 2-0 -> 22-20
if (denorm) { exp = ((fp8 >> 3) & 0x0F) - 7; }
u.i |= (exp + 127) << 23; // exponent
if (!denorm) { u.f -= 1.0/64; } // 2⁻⁶
u.i |= (t & 128) << 24; // sign
return u.f;
}This form may be simple to vectorise. I need to update my fp8 branch and try to adapte it with that. |
Beta Was this translation helpful? Give feedback.
-
|
Brilliant idea on the floating point subtraction. Latest code update is 52d042f
static inline float
llamafile_fp8_e4m3_to_fp32(ggml_fp8_t f)
{
union
{
ggml_fp8_t f;
unsigned char i;
} in = { f };
union
{
float f;
unsigned i;
} u;
unsigned x = in.i;
if (x & 127) {
if (x & 120) {
u.i = (x & 7) << 20;
u.i |= (((x >> 3) & 15) + 120) << 23;
} else {
u.i = (x & 7) << 14;
u.i |= 127 << 23;
u.f -= 1;
}
} else {
u.i = 0;
}
u.i |= (x & 128) << 24;
return u.f;
}I'm going to try your BF16 and F16 transcoding idea. Those would certainly help on znver4 and armv8.2. |
Beta Was this translation helpful? Give feedback.
-
|
OK some bench with this fp8 branch:
Now using this forme of computing fp8 > fp16: static float from_fp8_8(uint8_t fp8) {
union {
float f;
uint32_t i;
} u;
const uint32_t t = fp8;
// if ((fp8 & 0x7F) == 0x7F) return (fp8==0xFF)?(-NAN):(+NAN); // not needed here there is no NAN after quantisation.
auto exp_8 = t & 0x78;
auto exp_32 = (exp_8+(120<<3))<<20;
u.i = (t & 7) << 20; // mantissa: bit 2-0 -> 22-20
u.i |= exp_8 ? exp_32 : (-6 + 127) << 23; // exponent
if (!exp_8) { u.f -= 1.0/64; } // 2⁻⁶
u.i |= (t & 0x80) << 24; // sign
return u.f;
}I get this vectorised implemetation: #if defined(__AVX512F__) && defined(__AVX512VL__) && defined(__AVX512BW__)
#include <immintrin.h>
static __m512 llamafile_from_fp8_e4m3_avx512(__m128i fp8_vec) {
// extract componants:
__m128i expo_8 = _mm_and_si128(fp8_vec, _mm_set1_epi8(0x78));
__m128i mant_8 = _mm_and_si128(fp8_vec, _mm_set1_epi8(0x07));
__m128i sign_8 = _mm_and_si128(fp8_vec, _mm_set1_epi8(0x80));
// denorm mask
__mmask16 is_denorm = _mm_cmpeq_epi8_mask(expo_8, _mm_set1_epi8(0));
// convert to 32 bits
__m512i expo_32 = _mm512_cvtepu8_epi32(expo_8);
__m512i mant_32 = _mm512_cvtepu8_epi32(mant_8);
__m512i sign_32 = _mm512_cvtepu8_epi32(sign_8);
// shift
expo_32 = _mm512_slli_epi32(_mm512_add_epi32(expo_32,_mm512_set1_epi32(120<<3)), 20);
mant_32 = _mm512_slli_epi32(mant_32, 20);
sign_32 = _mm512_slli_epi32(sign_32, 24);
// correction denorm expo:
expo_32 = _mm512_mask_blend_epi32(is_denorm, expo_32, _mm512_set1_epi32((-6 + 127) << 23));
// merge mantissa+exponent
__m512 result = _mm512_castsi512_ps(_mm512_or_si512(expo_32,mant_32));
// correction denorm mantissa:
result = _mm512_mask_add_ps(result, is_denorm, result, _mm512_set1_ps(-1.0/64));
// add sign
return _mm512_castsi512_ps(_mm512_or_si512(sign_32,_mm512_castps_si512(result)));
}
#endifnow some bench:
Now Token generation is good. and prompt processing is not that bad if we think that we do not use bf16. pretty sure we can get a x2 if we convert fp8 to bf16 and use it for compute as is done with bf16 quant. I may study closely this branch and see if I can add fp8/bf16 in tinyblas. It can be good exercise. Or update my test on blas_bf16 (https://github.com/Djip007/llama.cpp/tree/poc/bf16) but may take some time with all re-factoring that have be done on llama.cpp 😉. may be faster to rewrite it on top of fp8 branch... |
Beta Was this translation helpful? Give feedback.
-
|
Thanks. |
Beta Was this translation helpful? Give feedback.
-
|
I see you have add many more code/optim in this branch 😉 👍 static __m512 from_fp8_e4m3_avx512(__m128i fp8_vec) {
// extract exponent:
__m128i expo_8 = _mm_and_si128(fp8_vec, _mm_set1_epi8(0x78));
__m128i mant_8 = _mm_and_si128(fp8_vec, _mm_set1_epi8(0x07));
__m128i sign_8 = _mm_and_si128(fp8_vec, _mm_set1_epi8(0x80));
// mask
//> need AVX512BW + AVX512VL ?
//__mmask16 is_denorm = _mm_cmpeq_epi8_mask(expo_8, _mm_setzero_si128());
__m512i expo_32 = _mm512_cvtepu8_epi32(expo_8);
__m512i mant_32 = _mm512_cvtepu8_epi32(mant_8);
__m512i sign_32 = _mm512_cvtepu8_epi32(sign_8);
//> pure AVX512F:
__mmask16 is_denorm = _mm512_cmpeq_epi32_mask(expo_32, _mm512_setzero_epi32());
// shift
expo_32 = _mm512_slli_epi32(_mm512_add_epi32(expo_32,_mm512_set1_epi32(120<<3)), 20);
mant_32 = _mm512_slli_epi32(mant_32, 20);
sign_32 = _mm512_slli_epi32(sign_32, 24);
// correction denorm:
expo_32 = _mm512_mask_blend_epi32(is_denorm, expo_32, _mm512_set1_epi32((-6 + 127) << 23));
__m512 result = _mm512_castsi512_ps(_mm512_or_si512(expo_32,mant_32));
result = _mm512_mask_add_ps(result, is_denorm, result, _mm512_set1_ps(-1.0/64));
return _mm512_castsi512_ps(_mm512_or_si512(sign_32,_mm512_castps_si512(result)));
}This code work and is pure AVX512 I think (can really test with CPU that have no AVX512BW && AVX512VL) I expect it to be a little slower on ZEN4 or at least need a little more power but can replace you fallback AVX512 code I think. |
Beta Was this translation helpful? Give feedback.
-
|
for fp8 => bf16: // for "clean" cast.
extern __inline __m512bh
__attribute__ ((__gnu_inline__, __always_inline__, __artificial__))
_mm512_castsi512_bh (__m512i __A)
{
return (__m512bh) (__A);
}
extern __inline __m256bh
__attribute__ ((__gnu_inline__, __always_inline__, __artificial__))
_mm256_castsi256_bh (__m256i __A)
{
return (__m256bh) (__A);
}
extern __inline __m512i
__attribute__ ((__gnu_inline__, __always_inline__, __artificial__))
_mm512_castbh_si512 (__m512bh __A)
{
return (__m512i) (__A);
}
static inline __m512bh
llamafile_fp8_e4m3_to_bf16_avx512(__m256i fp8_vec){
// extract components:
__m256i expo_8 = _mm256_and_si256(fp8_vec, _mm256_set1_epi8(0x78));
__m256i mant_8 = _mm256_and_si256(fp8_vec, _mm256_set1_epi8(0x07));
__m256i sign_8 = _mm256_and_si256(fp8_vec, _mm256_set1_epi8(0x80));
// denorm mask
//> need AVX512BW + AVX512VL ?
__mmask32 is_denorm = _mm256_cmpeq_epi8_mask(expo_8, _mm256_setzero_si256());
__m512i expo_16 = _mm512_cvtepu8_epi16(expo_8);
__m512i mant_16 = _mm512_cvtepu8_epi16(mant_8);
__m512i sign_16 = _mm512_cvtepu8_epi16(sign_8);
//> pure AVX512F:
//__mmask32 is_denorm = _mm512_cmpeq_epi16_mask(expo_16, _mm512_setzero_si512());
__mmask16 is_denorm_low = is_denorm;
__mmask16 is_denorm_high = is_denorm>>16;
// shift
expo_16 = _mm512_slli_epi16(_mm512_add_epi32(expo_16,_mm512_set1_epi16(120<<3)), 4);
mant_16 = _mm512_slli_epi16(mant_16, 4);
sign_16 = _mm512_slli_epi16(sign_16, 8);
// correction denorm exp:
expo_16 = _mm512_mask_blend_epi16(is_denorm, expo_16, _mm512_set1_epi16((-6 + 127) << 7));
__m512i em = _mm512_or_si512(expo_16,mant_16);
// correction denorm mantissa using fp32 Aritmetics:
__m256bh low_bh = _mm256_castsi256_bh(_mm512_castsi512_si256(em));
__m256bh high_bh = _mm256_castsi256_bh(_mm512_extracti32x8_epi32 (em, 1));
__m512 low = _mm512_cvtpbh_ps( low_bh);
__m512 high = _mm512_cvtpbh_ps(high_bh);
low = _mm512_mask_add_ps( low, is_denorm_low , low, _mm512_set1_ps(-1.0/64));
high = _mm512_mask_add_ps(high, is_denorm_high, high, _mm512_set1_ps(-1.0/64));
__m512bh result = _mm512_cvtne2ps_pbh(high,low);
return _mm512_castsi512_bh(_mm512_or_si512(sign_16,_mm512_castbh_si512(result)));
}The result is the same as yours, so code look correct. did not test speed for now. |
Beta Was this translation helpful? Give feedback.
-
|
This fp8->bf16 conversion goes 101 tok/sec on my workstation, whereas mine goes 120 tok/sec. That intermediate fp8->fp32->bf16 cast you're doing looks costly. That's what I was trying to avoid with the selector trick. One of the bottlenecks with doing fp8->bf16/f16 is we need to return |
Beta Was this translation helpful? Give feedback.
-
Yes that what I was thinking when I finish it, Just don't have time to get it benchmark yesterday. i only they have add the possibility to make a simple ADD on this bf16. At least it can serve as an example with AVX512_FP16 and NEON_FP16. And if I can get a "correct" quantisation with a scale factor as base, I may test if it is so important to have this 16 small value or if we can replace it with 0.
But don't waste your time with it. I'll do a few tries and if I find good things I'll make a merge request on this FP8 branch. |
Beta Was this translation helpful? Give feedback.
-
|
we may take some time to read that: That how FP8 quantisation are performe on vllm (I think...) Look they have per tensor scale factor.
|
Beta Was this translation helpful? Give feedback.
-
|
Yes hence my surprise that it even works at all. I wonder if it'll work with other models. Oh wow apparently our FP8 works with Meta LLaMA 3.1 8B too. It doesn't get the best perplexity scores but they're not terrible. FP8 is getting 8.5 but Q6_K does much better at 6.95 since Kawrakow is the best at this. For Nemo the difference is 11.7748 vs. 5.6997, not to mention Q6_K gets 300 tok/sec at prompt processing on my machine. I think working on this has been really fun, but that's a tremendous gulf. |
Beta Was this translation helpful? Give feedback.
-
I know... thank for all your work, I think I have most I need to make more experiment. Thanks for all this work, It is for me a good base to study more on llama.cpp/llamafile. |
Beta Was this translation helpful? Give feedback.
-
|
Yes! Come back and collaborate any time. You have a sharp intellect. I enjoyed working with you. |
Beta Was this translation helpful? Give feedback.
-
|
Thank you... |
Beta Was this translation helpful? Give feedback.
-
|
OK after many work finaly can test with scale factor on FP8 tensor ... and it looks much better with nemo. |
Beta Was this translation helpful? Give feedback.
-
|
19/10/2024 obsolete!!! it can be "much" better! Ok some more work. I create my own backend for more control on compute... (more later o, that)
Note: the 3 V2 FP8 decoding have the same perplexity so we can go for the faster (V2.2 ...) |
Beta Was this translation helpful? Give feedback.
-
6.4760 seems a bit high? I get |
Beta Was this translation helpful? Give feedback.
-
|
on which dataset? (and yes sorry I get the wrong one... but it was for all my bench... => I may need to re-do all 😎 and that with llamafile-0.8.13 too ... |
Beta Was this translation helpful? Give feedback.
-
|
Oh, I see. I always use I think that in order to compare
I.e., If I use SOTA quantization from https://github.com/ikawrakow/ik_llama.cpp/tree/main, I have
I.e, |
Beta Was this translation helpful? Give feedback.
-
|
OK some more "reference" Perplexity:
|
Beta Was this translation helpful? Give feedback.
-
|
for now it is "juste" some experiment with fp8 ... there is more possibility to test:
and there is more "quantisation" possible:
For now I have to "finish" my new kernel that may have lower memory pressure for best scaling with high core count ... (if that work...) |
Beta Was this translation helpful? Give feedback.
-
|
Some benchmark.
The llamafile colonne is the current 0.8.13 release with BF16/Q8/Q6/Q5 compute for reference
|
Beta Was this translation helpful? Give feedback.
-
|
@Djip007 Am I understanding correctly that you got FP8 E4M3 to:
If so, that's outstanding; I am interested in including your work in llamafile. What especially interests me is that you figured out how to make it work with flush to zero. Were those perplexity scores measured while flushing subnormals to zero? What about the benchmarks? For your kernel shape, what you want to do is pick whatever shape uses the most vector registers without spilling to the stack. For BF16 on AVX512 and NEON which have 32 registers that shape was either 5x5 or 8x3. Also can you share your code? |
Beta Was this translation helpful? Give feedback.
-
|
I am really stupid we can have much faster E4M3 decoding !!!! now if we use for BF16 decoding: we have that: so if we do simple decoding we only have to get a 2^120 factor to get the good result ... in our case we can do that with alfa at the end... on the FP32 / mm512 ... hop I am correct!!! (in fact we may "loose" some value for small value on B... but need more time to figure it... may be not that much with FP32 acc on BF16...) Edit (not completely sure):
|
Beta Was this translation helpful? Give feedback.
-
|
Look promising: I use: static inline __m512bh
llamafile_fp8_e4m3_to_bf16_avx512(__m256i fp8_vec) {
__m512i fp8_v16 = _mm512_cvtepu8_epi16(fp8_vec);
// extract sign and expo+mantice
__m512i em_16 = _mm512_and_si512(fp8_v16, _mm512_set1_epi16(0x7F));
__m512i sign_16 = _mm512_and_si512(fp8_v16, _mm512_set1_epi16(0x80));
// shift them
em_16 = _mm512_slli_epi16(em_16, 4);
sign_16 = _mm512_slli_epi16(sign_16, 8);
return _mm512_castsi512_bh(_mm512_or_si512(sign_16,em_16));
}and the matmul is compute like that: the bench is good: with the 4x6 kernel:
and perplexity look good: |
Beta Was this translation helpful? Give feedback.
-
OK more info on that...
in all case for tg I use 8x3 kernel (and for pp1,pp2,pp3) ... for higher ppN I use 5x5 or 4x6. => I may have call it 5x5_32x1 (MxN_KxM) ...
=> I may have call it 16x16_2x16 (MxN_KxM) ... I can give more element on that, but may be on a other "Discussion"... it may be interesting for the llama-matmul branch. I have a look on it and I think have some "comment" for that ;) (ps: I update the benchmark result with more correct name and only best E4M3 decoding speed...) |
Beta Was this translation helpful? Give feedback.
-
|
I tried replacing |
Beta Was this translation helpful? Give feedback.
-
|
yes with the last one the converted value is "wrong" we need to apply a CORRECTION=2^120 at the end: for (j = 0; j < RN; ++j)
for (i = 0; i < RM; ++i)
store(INDEX(C, ldc, jj + j, ii + i), Cf[j][i]);
// have to be replace with:
for (j = 0; j < RN; ++j)
for (i = 0; i < RM; ++i)
store(INDEX(C, ldc, jj + j, ii + i), Cf[j][i]*CORRECTION<TA>());
# with something like that.
template<typename T> float constexpr CORRECTION() { return 1;}
template<> float constexpr CORRECTION<ggml_fp8_bf16_t> { return 2^120};
template<> float constexpr CORRECTION<ggml_fp8_t> { return 2^120};but I don not test it , may be more complicated / case than that. For good perplexity it have to be more like: for (j = 0; j < RN; ++j)
for (i = 0; i < RM; ++i)
store(INDEX(C, ldc, jj + j, ii + i), Cf[j][i]*scale[ii+i]*CORRECTION<TA>());I try to publish my code... |
Beta Was this translation helpful? Give feedback.
-
OK I publish it, it is not finish and more like a POC..
The load for fp8 can be found here: load_fp8
But it is really experimental, I want to do much more work... |
Beta Was this translation helpful? Give feedback.
-
|
Some more news, I have added more FP8 format on my backend. FP8: (19/10/2024: obsolete !! )
for reference:
G => scale global A[k,m] => scale (float) |
Beta Was this translation helpful? Give feedback.
-
|
I update my branch if someone want to bench other model... (hop it is usable...) |
Beta Was this translation helpful? Give feedback.
-
|
Some more perplexity (with some surprise) for Mistral-7B-Instruct-v0.3 Yes I check it twice ... le FP8_E3M4_C has better perplexity than original BF16 |
Beta Was this translation helpful? Give feedback.
-
|
Meta-Llama-3.1-8B-Instruct.BF16: same for this one, the best perplexity is with FP8_E3M4_C |
Beta Was this translation helpful? Give feedback.
-
|
I was so focused on the speed that I forgot to look at the quality of the quantization. With a single scale for a bloc of 256 weight: FP8_E3M4_K2 : PPL = 6.346961 ± 0.039010 In this case I even have add "output.weight" fp8 quantisation. Now I have to re-compute all the perplexity 😎 Q8_0 : PPL = 6.3445 +/- 0.03898 |
Beta Was this translation helpful? Give feedback.
-
template<int _E>
struct FP8 {
uint8_t bits;
using type = FP8<_E>;
static constexpr int E=_E;
static constexpr int M=7-_E;
static constexpr int E_BIAS=EXP2<_E-1>()-1;
static constexpr float MAX() { return (2-EXP2<-M+1>())*EXP2<EXP_I2<_E-1>()>(); }
static constexpr float MIN() { return EXP2<-M>()*EXP2<2-EXP_I2<_E-1>()>(); }
void operator=(fp32_t value) {
union {
float f;
uint32_t bits;
} in = {value};
bits = (in.bits >> 24) & 0x80;
in.bits &= 0x7fffffff;
GGML_ASSERT(in.bits < 0x7f800000);
if (in.f >= MAX()) {
bits |= 0x7E;
} else if (in.f<MIN()) {
// OK
} else {
in.f *= EXP2<E_BIAS-127>();
in.bits += 1<<(22-M); // for rounding!
bits |= (in.bits >> (23-M)) & 0x7F;
}
}
}; |
Beta Was this translation helpful? Give feedback.
-
|
OK this is better with "correct" rounding: #> zen4 / BF16 / wiki.test.raw / rounding / Mistral-Nemo-Instruct.
Now most are from Q6_K and Q8_0. (some more result in progress) But as you see I have best result with E3M4 format... that may not be a good point for AMD/Nvidia that have implement E4M3 in hardware. (but to be fare only the weight are quantize, the compute is done after convert on BF16.) liste of quantized tensor for the FP8 are : static constexpr std::list<std::string> LIST_WEIGHT_CONVERT() { return {
"ffn_down.weight",
"ffn_gate.weight",
"ffn_up.weight",
"ttn_k.weight",
"ttn_q.weight",
"ttn_v.weight",
"ttn_output.weight",
"output.weight",
};}(before I do not quantize the "output.weight") |
Beta Was this translation helpful? Give feedback.
-
|
And for the story, I had some more complet result for "trunc" implementation and for example I had: FP8_E3M4_K2 (trunc / without "output.weight"): if we compare with: PPL / Mean ln(PPL(Q)/PPL(base)) is close, but KLD / top P are really different. For Q6_K I get: |
Beta Was this translation helpful? Give feedback.
-
|
some more... Meta-Llama-3.1-8B-Instruct(zen4 / BF16 / wiki.test.raw / rounding / output-weight)
when we see what we get when model are not train with FP8 quantize award ? Mistral-7B-Instruct-v0.3( zen4 / BF16 / wiki.test.raw / rounding / output-weight )
this one is better! |
Beta Was this translation helpful? Give feedback.
-
|
OK for now I create a branch with FP8 base support. It is not for speed, I only add base part (quantize / convert / dot) Djip007@cbd3abd I add 4 FP8 format:
Now I need to make it fast..
|
Beta Was this translation helpful? Give feedback.
-
As discution start here: #543 (reply in thread):
There is 2 (or 3 E3M4 is give in some paper...)
I really like to help. Have a look what we can have.
I think E5M2 format can be convert to FP16 the same way we do with BF16 <-> FP32. I'll have a look and get back 😎
Beta Was this translation helpful? Give feedback.
All reactions