electron-rare/prima-cpp
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

feat: first draft of q2_k quantization support

Browse Source
This commit is contained in:
DandinPower
2025-11-03 18:24:41 +08:00
parent 25e5dd4645
commit 9e1a4f2942
8 changed files with 972 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files
Makefile
+6
View File
@@ -957,6 +957,7 @@ OBJ_LLAMA = \
src/network-utils.o \
src/quantization.o \
src/sparsity.o \
src/k_quantization.o \
OBJ_COMMON = \
common/profiler.o \
@@ -1187,6 +1188,11 @@ src/sparsity.o: \
src/sparsity.h
$(CXX) $(CXXFLAGS) -c $< -o $@ -fopenmp
src/k_quantization.o: \
src/k_quantization.cpp \
src/k_quantization.h
$(CXX) $(CXXFLAGS) -c $< -o $@
src/llama-grammar.o: \
src/llama-grammar.cpp \
src/llama-grammar.h \
common/arg.cpp
+1 -1
View File
@@ -2120,7 +2120,7 @@ gpt_params_context gpt_params_parser_init(gpt_params & params, llama_example ex,
));
add_opt(llama_arg(
{"--comm-datatype"}, "TYPE",
format("Datatype for communication, currently support f32, q8_0, q4_0 or f32_sparsity (default: %s)", params.comm_datatype.c_str()),
format("Datatype for communication, currently support f32, q8_0, q4_0, q2_k or f32_sparsity (default: %s)", params.comm_datatype.c_str()),
[](gpt_params & params, const std::string & value) {
params.comm_datatype = value;
}
src/fp16/bitcasts.h
+96
View File
@@ -0,0 +1,96 @@
#pragma once
#ifndef FP16_BITCASTS_H
#define FP16_BITCASTS_H
#if defined(__cplusplus) && (__cplusplus >= 201103L)
#include <cstdint>
#elif !defined(__OPENCL_VERSION__)
#include <stdint.h>
#endif
#if defined(__INTEL_COMPILER) || defined(_MSC_VER) && (_MSC_VER >= 1932) && (defined(_M_IX86) || defined(_M_X64))
#include <immintrin.h>
#endif
#if defined(_MSC_VER) && !defined(__clang__) && (defined(_M_ARM) || defined(_M_ARM64))
#include <intrin.h>
#endif
#if defined(__clang__) && (defined(_M_IX86) || defined(_M_X64))
#include <x86intrin.h>
#endif
static inline float fp32_from_bits(uint32_t w) {
#if defined(__OPENCL_VERSION__)
return as_float(w);
#elif defined(__CUDA_ARCH__)
return __uint_as_float((unsigned int) w);
#elif defined(__INTEL_COMPILER) || defined(_MSC_VER) && (_MSC_VER >= 1932) && (defined(_M_IX86) || defined(_M_X64))
return _castu32_f32(w);
#elif defined(_MSC_VER) && !defined(__clang__) && (defined(_M_ARM) || defined(_M_ARM64))
return _CopyFloatFromInt32((__int32) w);
#else
union {
uint32_t as_bits;
float as_value;
} fp32 = { w };
return fp32.as_value;
#endif
}
static inline uint32_t fp32_to_bits(float f) {
#if defined(__OPENCL_VERSION__)
return as_uint(f);
#elif defined(__CUDA_ARCH__)
return (uint32_t) __float_as_uint(f);
#elif defined(__INTEL_COMPILER) || defined(_MSC_VER) && (_MSC_VER >= 1932) && (defined(_M_IX86) || defined(_M_X64))
return _castf32_u32(f);
#elif defined(_MSC_VER) && !defined(__clang__) && (defined(_M_ARM) || defined(_M_ARM64))
return (uint32_t) _CopyInt32FromFloat(f);
#else
union {
float as_value;
uint32_t as_bits;
} fp32 = { f };
return fp32.as_bits;
#endif
}
static inline double fp64_from_bits(uint64_t w) {
#if defined(__OPENCL_VERSION__)
return as_double(w);
#elif defined(__CUDA_ARCH__)
return __longlong_as_double((long long) w);
#elif defined(__INTEL_COMPILER) || defined(_MSC_VER) && (_MSC_VER >= 1932) && (defined(_M_IX86) || defined(_M_X64))
return _castu64_f64(w);
#elif defined(_MSC_VER) && !defined(__clang__) && (defined(_M_ARM) || defined(_M_ARM64))
return _CopyDoubleFromInt64((__int64) w);
#else
union {
uint64_t as_bits;
double as_value;
} fp64 = { w };
return fp64.as_value;
#endif
}
static inline uint64_t fp64_to_bits(double f) {
#if defined(__OPENCL_VERSION__)
return as_ulong(f);
#elif defined(__CUDA_ARCH__)
return (uint64_t) __double_as_longlong(f);
#elif defined(__INTEL_COMPILER) || defined(_MSC_VER) && (_MSC_VER >= 1932) && (defined(_M_IX86) || defined(_M_X64))
return _castf64_u64(f);
#elif defined(_MSC_VER) && !defined(__clang__) && (defined(_M_ARM) || defined(_M_ARM64))
return (uint64_t) _CopyInt64FromDouble(f);
#else
union {
double as_value;
uint64_t as_bits;
} fp64 = { f };
return fp64.as_bits;
#endif
}
#endif /* FP16_BITCASTS_H */
src/fp16/fp16.h
+515
View File
@@ -0,0 +1,515 @@
#pragma once
#ifndef FP16_FP16_H
#define FP16_FP16_H
#if defined(__cplusplus) && (__cplusplus >= 201103L)
#include <cstdint>
#include <cmath>
#elif !defined(__OPENCL_VERSION__)
#include <stdint.h>
#include <math.h>
#endif
#include <fp16/bitcasts.h>
#include <fp16/macros.h>
#if defined(_MSC_VER)
#include <intrin.h>
#endif
#if defined(__F16C__) && FP16_USE_NATIVE_CONVERSION && !FP16_USE_FLOAT16_TYPE && !FP16_USE_FP16_TYPE
#include <immintrin.h>
#endif
#if (defined(__aarch64__) || defined(_M_ARM64)) && FP16_USE_NATIVE_CONVERSION && !FP16_USE_FLOAT16_TYPE && !FP16_USE_FP16_TYPE
#include <arm_neon.h>
#endif
/*
* Convert a 16-bit floating-point number in IEEE half-precision format, in bit representation, to
* a 32-bit floating-point number in IEEE single-precision format, in bit representation.
*
* @note The implementation doesn't use any floating-point operations.
*/
static inline uint32_t fp16_ieee_to_fp32_bits(uint16_t h) {
/*
* Extend the half-precision floating-point number to 32 bits and shift to the upper part of the 32-bit word:
* +---+-----+------------+-------------------+
* | S |EEEEE|MM MMMM MMMM|0000 0000 0000 0000|
* +---+-----+------------+-------------------+
* Bits 31 26-30 16-25 0-15
*
* S - sign bit, E - bits of the biased exponent, M - bits of the mantissa, 0 - zero bits.
*/
const uint32_t w = (uint32_t) h << 16;
/*
* Extract the sign of the input number into the high bit of the 32-bit word:
*
* +---+----------------------------------+
* | S |0000000 00000000 00000000 00000000|
* +---+----------------------------------+
* Bits 31 0-31
*/
const uint32_t sign = w & UINT32_C(0x80000000);
/*
* Extract mantissa and biased exponent of the input number into the bits 0-30 of the 32-bit word:
*
* +---+-----+------------+-------------------+
* | 0 |EEEEE|MM MMMM MMMM|0000 0000 0000 0000|
* +---+-----+------------+-------------------+
* Bits 30 27-31 17-26 0-16
*/
const uint32_t nonsign = w & UINT32_C(0x7FFFFFFF);
/*
* Renorm shift is the number of bits to shift mantissa left to make the half-precision number normalized.
* If the initial number is normalized, some of its high 6 bits (sign == 0 and 5-bit exponent) equals one.
* In this case renorm_shift == 0. If the number is denormalize, renorm_shift > 0. Note that if we shift
* denormalized nonsign by renorm_shift, the unit bit of mantissa will shift into exponent, turning the
* biased exponent into 1, and making mantissa normalized (i.e. without leading 1).
*/
#ifdef _MSC_VER
unsigned long nonsign_bsr;
_BitScanReverse(&nonsign_bsr, (unsigned long) nonsign);
uint32_t renorm_shift = (uint32_t) nonsign_bsr ^ 31;
#else
uint32_t renorm_shift = __builtin_clz(nonsign);
#endif
renorm_shift = renorm_shift > 5 ? renorm_shift - 5 : 0;
/*
* Iff half-precision number has exponent of 15, the addition overflows it into bit 31,
* and the subsequent shift turns the high 9 bits into 1. Thus
* inf_nan_mask ==
* 0x7F800000 if the half-precision number had exponent of 15 (i.e. was NaN or infinity)
* 0x00000000 otherwise
*/
const int32_t inf_nan_mask = ((int32_t) (nonsign + 0x04000000) >> 8) & INT32_C(0x7F800000);
/*
* Iff nonsign is 0, it overflows into 0xFFFFFFFF, turning bit 31 into 1. Otherwise, bit 31 remains 0.
* The signed shift right by 31 broadcasts bit 31 into all bits of the zero_mask. Thus
* zero_mask ==
* 0xFFFFFFFF if the half-precision number was zero (+0.0h or -0.0h)
* 0x00000000 otherwise
*/
const int32_t zero_mask = (int32_t) (nonsign - 1) >> 31;
/*
* 1. Shift nonsign left by renorm_shift to normalize it (if the input was denormal)
* 2. Shift nonsign right by 3 so the exponent (5 bits originally) becomes an 8-bit field and 10-bit mantissa
* shifts into the 10 high bits of the 23-bit mantissa of IEEE single-precision number.
* 3. Add 0x70 to the exponent (starting at bit 23) to compensate the different in exponent bias
* (0x7F for single-precision number less 0xF for half-precision number).
* 4. Subtract renorm_shift from the exponent (starting at bit 23) to account for renormalization. As renorm_shift
* is less than 0x70, this can be combined with step 3.
* 5. Binary OR with inf_nan_mask to turn the exponent into 0xFF if the input was NaN or infinity.
* 6. Binary ANDNOT with zero_mask to turn the mantissa and exponent into zero if the input was zero.
* 7. Combine with the sign of the input number.
*/
return sign | ((((nonsign << renorm_shift >> 3) + ((0x70 - renorm_shift) << 23)) | inf_nan_mask) & ~zero_mask);
}
/*
* Convert a 16-bit floating-point number in IEEE half-precision format, in bit representation, to
* a 32-bit floating-point number in IEEE single-precision format.
*
* @note The implementation relies on IEEE-like (no assumption about rounding mode and no operations on denormals)
* floating-point operations and bitcasts between integer and floating-point variables.
*/
static inline float fp16_ieee_to_fp32_value(uint16_t h) {
#if FP16_USE_NATIVE_CONVERSION
#if FP16_USE_FLOAT16_TYPE
union {
uint16_t as_bits;
_Float16 as_value;
} fp16 = { h };
return (float) fp16.as_value;
#elif FP16_USE_FP16_TYPE
union {
uint16_t as_bits;
__fp16 as_value;
} fp16 = { h };
return (float) fp16.as_value;
#else
#if (defined(__INTEL_COMPILER) || defined(__GNUC__)) && defined(__F16C__)
return _cvtsh_ss((unsigned short) h);
#elif defined(_MSC_VER) && (defined(_M_IX86) || defined(_M_X64)) && defined(__AVX2__)
return _mm_cvtss_f32(_mm_cvtph_ps(_mm_cvtsi32_si128((int) (unsigned int) h)));
#elif defined(_M_ARM64) || defined(__aarch64__)
return vgetq_lane_f32(vcvt_f32_f16(vreinterpret_f16_u16(vdup_n_u16(h))), 0);
#else
#error "Archtecture- or compiler-specific implementation required"
#endif
#endif
#else
/*
* Extend the half-precision floating-point number to 32 bits and shift to the upper part of the 32-bit word:
* +---+-----+------------+-------------------+
* | S |EEEEE|MM MMMM MMMM|0000 0000 0000 0000|
* +---+-----+------------+-------------------+
* Bits 31 26-30 16-25 0-15
*
* S - sign bit, E - bits of the biased exponent, M - bits of the mantissa, 0 - zero bits.
*/
const uint32_t w = (uint32_t) h << 16;
/*
* Extract the sign of the input number into the high bit of the 32-bit word:
*
* +---+----------------------------------+
* | S |0000000 00000000 00000000 00000000|
* +---+----------------------------------+
* Bits 31 0-31
*/
const uint32_t sign = w & UINT32_C(0x80000000);
/*
* Extract mantissa and biased exponent of the input number into the high bits of the 32-bit word:
*
* +-----+------------+---------------------+
* |EEEEE|MM MMMM MMMM|0 0000 0000 0000 0000|
* +-----+------------+---------------------+
* Bits 27-31 17-26 0-16
*/
const uint32_t two_w = w + w;
/*
* Shift mantissa and exponent into bits 23-28 and bits 13-22 so they become mantissa and exponent
* of a single-precision floating-point number:
*
* S|Exponent | Mantissa
* +-+---+-----+------------+----------------+
* |0|000|EEEEE|MM MMMM MMMM|0 0000 0000 0000|
* +-+---+-----+------------+----------------+
* Bits | 23-31 | 0-22
*
* Next, there are some adjustments to the exponent:
* - The exponent needs to be corrected by the difference in exponent bias between single-precision and half-precision
* formats (0x7F - 0xF = 0x70)
* - Inf and NaN values in the inputs should become Inf and NaN values after conversion to the single-precision number.
* Therefore, if the biased exponent of the half-precision input was 0x1F (max possible value), the biased exponent
* of the single-precision output must be 0xFF (max possible value). We do this correction in two steps:
* - First, we adjust the exponent by (0xFF - 0x1F) = 0xE0 (see exp_offset below) rather than by 0x70 suggested
* by the difference in the exponent bias (see above).
* - Then we multiply the single-precision result of exponent adjustment by 2**(-112) to reverse the effect of
* exponent adjustment by 0xE0 less the necessary exponent adjustment by 0x70 due to difference in exponent bias.
* The floating-point multiplication hardware would ensure than Inf and NaN would retain their value on at least
* partially IEEE754-compliant implementations.
*
* Note that the above operations do not handle denormal inputs (where biased exponent == 0). However, they also do not
* operate on denormal inputs, and do not produce denormal results.
*/
const uint32_t exp_offset = UINT32_C(0xE0) << 23;
#if defined(__STDC_VERSION__) && (__STDC_VERSION__ >= 199901L) || defined(__GNUC__) && !defined(__STRICT_ANSI__)
const float exp_scale = 0x1.0p-112f;
#else
const float exp_scale = fp32_from_bits(UINT32_C(0x7800000));
#endif
const float normalized_value = fp32_from_bits((two_w >> 4) + exp_offset) * exp_scale;
/*
* Convert denormalized half-precision inputs into single-precision results (always normalized).
* Zero inputs are also handled here.
*
* In a denormalized number the biased exponent is zero, and mantissa has on-zero bits.
* First, we shift mantissa into bits 0-9 of the 32-bit word.
*
* zeros | mantissa
* +---------------------------+------------+
* |0000 0000 0000 0000 0000 00|MM MMMM MMMM|
* +---------------------------+------------+
* Bits 10-31 0-9
*
* Now, remember that denormalized half-precision numbers are represented as:
* FP16 = mantissa * 2**(-24).
* The trick is to construct a normalized single-precision number with the same mantissa and thehalf-precision input
* and with an exponent which would scale the corresponding mantissa bits to 2**(-24).
* A normalized single-precision floating-point number is represented as:
* FP32 = (1 + mantissa * 2**(-23)) * 2**(exponent - 127)
* Therefore, when the biased exponent is 126, a unit change in the mantissa of the input denormalized half-precision
* number causes a change of the constructud single-precision number by 2**(-24), i.e. the same ammount.
*
* The last step is to adjust the bias of the constructed single-precision number. When the input half-precision number
* is zero, the constructed single-precision number has the value of
* FP32 = 1 * 2**(126 - 127) = 2**(-1) = 0.5
* Therefore, we need to subtract 0.5 from the constructed single-precision number to get the numerical equivalent of
* the input half-precision number.
*/
const uint32_t magic_mask = UINT32_C(126) << 23;
const float magic_bias = 0.5f;
const float denormalized_value = fp32_from_bits((two_w >> 17) | magic_mask) - magic_bias;
/*
* - Choose either results of conversion of input as a normalized number, or as a denormalized number, depending on the
* input exponent. The variable two_w contains input exponent in bits 27-31, therefore if its smaller than 2**27, the
* input is either a denormal number, or zero.
* - Combine the result of conversion of exponent and mantissa with the sign of the input number.
*/
const uint32_t denormalized_cutoff = UINT32_C(1) << 27;
const uint32_t result = sign |
(two_w < denormalized_cutoff ? fp32_to_bits(denormalized_value) : fp32_to_bits(normalized_value));
return fp32_from_bits(result);
#endif
}
/*
* Convert a 32-bit floating-point number in IEEE single-precision format to a 16-bit floating-point number in
* IEEE half-precision format, in bit representation.
*
* @note The implementation relies on IEEE-like (no assumption about rounding mode and no operations on denormals)
* floating-point operations and bitcasts between integer and floating-point variables.
*/
static inline uint16_t fp16_ieee_from_fp32_value(float f) {
#if FP16_USE_NATIVE_CONVERSION
#if FP16_USE_FLOAT16_TYPE
union {
_Float16 as_value;
uint16_t as_bits;
} fp16 = { (_Float16) f };
return fp16.as_bits;
#elif FP16_USE_FP16_TYPE
union {
__fp16 as_value;
uint16_t as_bits;
} fp16 = { (__fp16) f };
return fp16.as_bits;
#else
#if (defined(__INTEL_COMPILER) || defined(__GNUC__)) && defined(__F16C__)
return _cvtss_sh(f, _MM_FROUND_CUR_DIRECTION);
#elif defined(_MSC_VER) && (defined(_M_IX86) || defined(_M_X64)) && defined(__AVX2__)
return (uint16_t) _mm_cvtsi128_si32(_mm_cvtps_ph(_mm_set_ss(f), _MM_FROUND_CUR_DIRECTION));
#elif defined(_M_ARM64) || defined(__aarch64__)
return vget_lane_u16(vcvt_f16_f32(vdupq_n_f32(f)), 0);
#else
#error "Archtecture- or compiler-specific implementation required"
#endif
#endif
#else
#if defined(__STDC_VERSION__) && (__STDC_VERSION__ >= 199901L) || defined(__GNUC__) && !defined(__STRICT_ANSI__)
const float scale_to_inf = 0x1.0p+112f;
const float scale_to_zero = 0x1.0p-110f;
#else
const float scale_to_inf = fp32_from_bits(UINT32_C(0x77800000));
const float scale_to_zero = fp32_from_bits(UINT32_C(0x08800000));
#endif
#if defined(_MSC_VER) && defined(_M_IX86_FP) && (_M_IX86_FP == 0) || defined(__GNUC__) && defined(__FLT_EVAL_METHOD__) && (__FLT_EVAL_METHOD__ != 0)
const volatile float saturated_f = fabsf(f) * scale_to_inf;
#else
const float saturated_f = fabsf(f) * scale_to_inf;
#endif
float base = saturated_f * scale_to_zero;
const uint32_t w = fp32_to_bits(f);
const uint32_t shl1_w = w + w;
const uint32_t sign = w & UINT32_C(0x80000000);
uint32_t bias = shl1_w & UINT32_C(0xFF000000);
if (bias < UINT32_C(0x71000000)) {
bias = UINT32_C(0x71000000);
}
base = fp32_from_bits((bias >> 1) + UINT32_C(0x07800000)) + base;
const uint32_t bits = fp32_to_bits(base);
const uint32_t exp_bits = (bits >> 13) & UINT32_C(0x00007C00);
const uint32_t mantissa_bits = bits & UINT32_C(0x00000FFF);
const uint32_t nonsign = exp_bits + mantissa_bits;
return (sign >> 16) | (shl1_w > UINT32_C(0xFF000000) ? UINT16_C(0x7E00) : nonsign);
#endif
}
/*
* Convert a 16-bit floating-point number in ARM alternative half-precision format, in bit representation, to
* a 32-bit floating-point number in IEEE single-precision format, in bit representation.
*
* @note The implementation doesn't use any floating-point operations.
*/
static inline uint32_t fp16_alt_to_fp32_bits(uint16_t h) {
/*
* Extend the half-precision floating-point number to 32 bits and shift to the upper part of the 32-bit word:
* +---+-----+------------+-------------------+
* | S |EEEEE|MM MMMM MMMM|0000 0000 0000 0000|
* +---+-----+------------+-------------------+
* Bits 31 26-30 16-25 0-15
*
* S - sign bit, E - bits of the biased exponent, M - bits of the mantissa, 0 - zero bits.
*/
const uint32_t w = (uint32_t) h << 16;
/*
* Extract the sign of the input number into the high bit of the 32-bit word:
*
* +---+----------------------------------+
* | S |0000000 00000000 00000000 00000000|
* +---+----------------------------------+
* Bits 31 0-31
*/
const uint32_t sign = w & UINT32_C(0x80000000);
/*
* Extract mantissa and biased exponent of the input number into the bits 0-30 of the 32-bit word:
*
* +---+-----+------------+-------------------+
* | 0 |EEEEE|MM MMMM MMMM|0000 0000 0000 0000|
* +---+-----+------------+-------------------+
* Bits 30 27-31 17-26 0-16
*/
const uint32_t nonsign = w & UINT32_C(0x7FFFFFFF);
/*
* Renorm shift is the number of bits to shift mantissa left to make the half-precision number normalized.
* If the initial number is normalized, some of its high 6 bits (sign == 0 and 5-bit exponent) equals one.
* In this case renorm_shift == 0. If the number is denormalize, renorm_shift > 0. Note that if we shift
* denormalized nonsign by renorm_shift, the unit bit of mantissa will shift into exponent, turning the
* biased exponent into 1, and making mantissa normalized (i.e. without leading 1).
*/
#ifdef _MSC_VER
unsigned long nonsign_bsr;
_BitScanReverse(&nonsign_bsr, (unsigned long) nonsign);
uint32_t renorm_shift = (uint32_t) nonsign_bsr ^ 31;
#else
uint32_t renorm_shift = __builtin_clz(nonsign);
#endif
renorm_shift = renorm_shift > 5 ? renorm_shift - 5 : 0;
/*
* Iff nonsign is 0, it overflows into 0xFFFFFFFF, turning bit 31 into 1. Otherwise, bit 31 remains 0.
* The signed shift right by 31 broadcasts bit 31 into all bits of the zero_mask. Thus
* zero_mask ==
* 0xFFFFFFFF if the half-precision number was zero (+0.0h or -0.0h)
* 0x00000000 otherwise
*/
const int32_t zero_mask = (int32_t) (nonsign - 1) >> 31;
/*
* 1. Shift nonsign left by renorm_shift to normalize it (if the input was denormal)
* 2. Shift nonsign right by 3 so the exponent (5 bits originally) becomes an 8-bit field and 10-bit mantissa
* shifts into the 10 high bits of the 23-bit mantissa of IEEE single-precision number.
* 3. Add 0x70 to the exponent (starting at bit 23) to compensate the different in exponent bias
* (0x7F for single-precision number less 0xF for half-precision number).
* 4. Subtract renorm_shift from the exponent (starting at bit 23) to account for renormalization. As renorm_shift
* is less than 0x70, this can be combined with step 3.
* 5. Binary ANDNOT with zero_mask to turn the mantissa and exponent into zero if the input was zero.
* 6. Combine with the sign of the input number.
*/
return sign | (((nonsign << renorm_shift >> 3) + ((0x70 - renorm_shift) << 23)) & ~zero_mask);
}
/*
* Convert a 16-bit floating-point number in ARM alternative half-precision format, in bit representation, to
* a 32-bit floating-point number in IEEE single-precision format.
*
* @note The implementation relies on IEEE-like (no assumption about rounding mode and no operations on denormals)
* floating-point operations and bitcasts between integer and floating-point variables.
*/
static inline float fp16_alt_to_fp32_value(uint16_t h) {
/*
* Extend the half-precision floating-point number to 32 bits and shift to the upper part of the 32-bit word:
* +---+-----+------------+-------------------+
* | S |EEEEE|MM MMMM MMMM|0000 0000 0000 0000|
* +---+-----+------------+-------------------+
* Bits 31 26-30 16-25 0-15
*
* S - sign bit, E - bits of the biased exponent, M - bits of the mantissa, 0 - zero bits.
*/
const uint32_t w = (uint32_t) h << 16;
/*
* Extract the sign of the input number into the high bit of the 32-bit word:
*
* +---+----------------------------------+
* | S |0000000 00000000 00000000 00000000|
* +---+----------------------------------+
* Bits 31 0-31
*/
const uint32_t sign = w & UINT32_C(0x80000000);
/*
* Extract mantissa and biased exponent of the input number into the high bits of the 32-bit word:
*
* +-----+------------+---------------------+
* |EEEEE|MM MMMM MMMM|0 0000 0000 0000 0000|
* +-----+------------+---------------------+
* Bits 27-31 17-26 0-16
*/
const uint32_t two_w = w + w;
/*
* Shift mantissa and exponent into bits 23-28 and bits 13-22 so they become mantissa and exponent
* of a single-precision floating-point number:
*
* S|Exponent | Mantissa
* +-+---+-----+------------+----------------+
* |0|000|EEEEE|MM MMMM MMMM|0 0000 0000 0000|
* +-+---+-----+------------+----------------+
* Bits | 23-31 | 0-22
*
* Next, the exponent is adjusted for the difference in exponent bias between single-precision and half-precision
* formats (0x7F - 0xF = 0x70). This operation never overflows or generates non-finite values, as the largest
* half-precision exponent is 0x1F and after the adjustment is can not exceed 0x8F < 0xFE (largest single-precision
* exponent for non-finite values).
*
* Note that this operation does not handle denormal inputs (where biased exponent == 0). However, they also do not
* operate on denormal inputs, and do not produce denormal results.
*/
const uint32_t exp_offset = UINT32_C(0x70) << 23;
const float normalized_value = fp32_from_bits((two_w >> 4) + exp_offset);
/*
* Convert denormalized half-precision inputs into single-precision results (always normalized).
* Zero inputs are also handled here.
*
* In a denormalized number the biased exponent is zero, and mantissa has on-zero bits.
* First, we shift mantissa into bits 0-9 of the 32-bit word.
*
* zeros | mantissa
* +---------------------------+------------+
* |0000 0000 0000 0000 0000 00|MM MMMM MMMM|
* +---------------------------+------------+
* Bits 10-31 0-9
*
* Now, remember that denormalized half-precision numbers are represented as:
* FP16 = mantissa * 2**(-24).
* The trick is to construct a normalized single-precision number with the same mantissa and thehalf-precision input
* and with an exponent which would scale the corresponding mantissa bits to 2**(-24).
* A normalized single-precision floating-point number is represented as:
* FP32 = (1 + mantissa * 2**(-23)) * 2**(exponent - 127)
* Therefore, when the biased exponent is 126, a unit change in the mantissa of the input denormalized half-precision
* number causes a change of the constructud single-precision number by 2**(-24), i.e. the same ammount.
*
* The last step is to adjust the bias of the constructed single-precision number. When the input half-precision number
* is zero, the constructed single-precision number has the value of
* FP32 = 1 * 2**(126 - 127) = 2**(-1) = 0.5
* Therefore, we need to subtract 0.5 from the constructed single-precision number to get the numerical equivalent of
* the input half-precision number.
*/
const uint32_t magic_mask = UINT32_C(126) << 23;
const float magic_bias = 0.5f;
const float denormalized_value = fp32_from_bits((two_w >> 17) | magic_mask) - magic_bias;
/*
* - Choose either results of conversion of input as a normalized number, or as a denormalized number, depending on the
* input exponent. The variable two_w contains input exponent in bits 27-31, therefore if its smaller than 2**27, the
* input is either a denormal number, or zero.
* - Combine the result of conversion of exponent and mantissa with the sign of the input number.
*/
const uint32_t denormalized_cutoff = UINT32_C(1) << 27;
const uint32_t result = sign |
(two_w < denormalized_cutoff ? fp32_to_bits(denormalized_value) : fp32_to_bits(normalized_value));
return fp32_from_bits(result);
}
/*
* Convert a 32-bit floating-point number in IEEE single-precision format to a 16-bit floating-point number in
* ARM alternative half-precision format, in bit representation.
*
* @note The implementation relies on IEEE-like (no assumption about rounding mode and no operations on denormals)
* floating-point operations and bitcasts between integer and floating-point variables.
*/
static inline uint16_t fp16_alt_from_fp32_value(float f) {
const uint32_t w = fp32_to_bits(f);
const uint32_t sign = w & UINT32_C(0x80000000);
const uint32_t shl1_w = w + w;
const uint32_t shl1_max_fp16_fp32 = UINT32_C(0x8FFFC000);
const uint32_t shl1_base = shl1_w > shl1_max_fp16_fp32 ? shl1_max_fp16_fp32 : shl1_w;
uint32_t shl1_bias = shl1_base & UINT32_C(0xFF000000);
const uint32_t exp_difference = 23 - 10;
const uint32_t shl1_bias_min = (127 - 1 - exp_difference) << 24;
if (shl1_bias < shl1_bias_min) {
shl1_bias = shl1_bias_min;
}
const float bias = fp32_from_bits((shl1_bias >> 1) + ((exp_difference + 2) << 23));
const float base = fp32_from_bits((shl1_base >> 1) + (2 << 23)) + bias;
const uint32_t exp_f = fp32_to_bits(base) >> 13;
return (sign >> 16) | ((exp_f & UINT32_C(0x00007C00)) + (fp32_to_bits(base) & UINT32_C(0x00000FFF)));
}
#endif /* FP16_FP16_H */
src/fp16/macros.h
+46
View File
@@ -0,0 +1,46 @@
#pragma once
#ifndef FP16_MACROS_H
#define FP16_MACROS_H
#ifndef FP16_USE_NATIVE_CONVERSION
#if (defined(__INTEL_COMPILER) || defined(__GNUC__)) && defined(__F16C__)
#define FP16_USE_NATIVE_CONVERSION 1
#elif defined(_MSC_VER) && (defined(_M_IX86) || defined(_M_X64)) && defined(__AVX2__)
#define FP16_USE_NATIVE_CONVERSION 1
#elif defined(_MSC_VER) && defined(_M_ARM64)
#define FP16_USE_NATIVE_CONVERSION 1
#elif defined(__GNUC__) && defined(__aarch64__)
#define FP16_USE_NATIVE_CONVERSION 1
#endif
#if !defined(FP16_USE_NATIVE_CONVERSION)
#define FP16_USE_NATIVE_CONVERSION 0
#endif // !defined(FP16_USE_NATIVE_CONVERSION)
#endif // !define(FP16_USE_NATIVE_CONVERSION)
#ifndef FP16_USE_FLOAT16_TYPE
#if !defined(__clang__) && !defined(__INTEL_COMPILER) && defined(__GNUC__) && (__GNUC__ >= 12)
#if defined(__F16C__)
#define FP16_USE_FLOAT16_TYPE 1
#endif
#endif
#if !defined(FP16_USE_FLOAT16_TYPE)
#define FP16_USE_FLOAT16_TYPE 0
#endif // !defined(FP16_USE_FLOAT16_TYPE)
#endif // !defined(FP16_USE_FLOAT16_TYPE)
#ifndef FP16_USE_FP16_TYPE
#if defined(__clang__)
#if defined(__F16C__) || defined(__aarch64__)
#define FP16_USE_FP16_TYPE 1
#endif
#elif defined(__GNUC__)
#if defined(__aarch64__)
#define FP16_USE_FP16_TYPE 1
#endif
#endif
#if !defined(FP16_USE_FP16_TYPE)
#define FP16_USE_FP16_TYPE 0
#endif // !defined(FP16_USE_FP16_TYPE)
#endif // !defined(FP16_USE_FP16_TYPE)
#endif /* FP16_MACROS_H */
src/k_quantization.cpp
+218
View File
@@ -0,0 +1,218 @@
#include "k_quantization.h"
#include <stdio.h>
#define MAX(a, b) ((a) > (b) ? (a) : (b))
#define MIN(a, b) ((a) < (b) ? (a) : (b))
// The implementation is refer to https://github.com/ggml-org/llama.cpp/blob/master/ggml/src/ggml-quants.c#L622
quantized_array_q2_k_t *allocate_q2_k_array(uint64_t num_elements) {
if (!num_elements) return NULL;
uint64_t num_elements_aligned = (num_elements % WEIGHT_PER_SUPER_BLOCK == 0) ? num_elements : num_elements + (WEIGHT_PER_SUPER_BLOCK - (num_elements % WEIGHT_PER_SUPER_BLOCK));
uint64_t num_super_blocks = num_elements_aligned / WEIGHT_PER_SUPER_BLOCK;
size_t total = sizeof(quantized_array_q2_k_t) + num_super_blocks * sizeof(super_block_q2_k);
quantized_array_q2_k_t *qa = (quantized_array_q2_k_t*)calloc(1, total);
if (!qa) return NULL;
qa->num_elements = num_elements;
qa->num_elements_aligned = num_elements_aligned;
qa->num_super_blocks = num_super_blocks;
qa->super_blocks = (super_block_q2_k*)(qa + 1);
return qa;
}
void free_quantized_q2_k_array(quantized_array_q2_k_t *quantized_array_q2_k) {
if (!quantized_array_q2_k) return;
free(quantized_array_q2_k);
}
int64_t get_quantized_q2_k_array_size(const quantized_array_q2_k_t *quantized_array_q2_k) {
if (!quantized_array_q2_k) return 0;
return sizeof(quantized_array_q2_k_t) + quantized_array_q2_k->num_super_blocks * sizeof(super_block_q2_k);
}
quantized_array_q2_k_t *load_quantized_q2_k_array_from_buffer(const void *buffer, int64_t buffer_size) {
quantized_array_q2_k_t *quantized_array = (quantized_array_q2_k_t*)calloc(1, buffer_size);
if (!quantized_array) return NULL;
memcpy(quantized_array, buffer, buffer_size);
quantized_array->super_blocks = (super_block_q2_k*)(quantized_array + 1);
return quantized_array;
}
static void find_optimal_scale_and_min(int curr_block_index, float *weights, float *scales, float*mins){
// naive approach
const float q2_scale = 3.f;
float min_val = INFINITY;
float max_val = -INFINITY;
for (int l = 0; l < Q2_K_BLOCK_SIZE; l++) {
if (weights[l] < min_val) min_val = weights[l];
}
for (int l = 0; l < Q2_K_BLOCK_SIZE; l++) {
weights[l] -= min_val;
}
for (int l = 0; l < Q2_K_BLOCK_SIZE; l++) {
if (weights[l] > max_val) max_val = weights[l];
}
scales[curr_block_index] = max_val / q2_scale;
mins[curr_block_index] = min_val;
}
int k_quantize(const float *float_array, uint64_t num_elements, quantized_array_q2_k_t **quantized_array_q2_k) {
const float q4_scale = 15.f;
uint8_t L[WEIGHT_PER_SUPER_BLOCK];
float weights[Q2_K_BLOCK_SIZE];
float mins[Q2_K_SUPER_BLOCK_SIZE];
float scales[Q2_K_SUPER_BLOCK_SIZE];
if (!float_array || num_elements == 0 || *quantized_array_q2_k) {
return 1;
}
*quantized_array_q2_k = allocate_q2_k_array(num_elements);
if (!*quantized_array_q2_k) {
return 1;
}
quantized_array_q2_k_t *qa = *quantized_array_q2_k;
float *float_array_aligned = (float*) calloc(1, sizeof(float) * qa->num_elements_aligned);
memcpy(float_array_aligned, float_array, (qa->num_elements) * sizeof(float));
for (uint32_t curr_super_block_index = 0; curr_super_block_index < qa->num_super_blocks; curr_super_block_index++) {
super_block_q2_k *curr_super_block = &qa->super_blocks[curr_super_block_index];
float max_scale = -INFINITY;
float max_abs_min = 0.f;
for (int j = 0; j < Q2_K_SUPER_BLOCK_SIZE; j++) {
for (int l = 0; l < Q2_K_BLOCK_SIZE; l++) {
weights[l] = float_array_aligned[j * Q2_K_BLOCK_SIZE + l];
}
find_optimal_scale_and_min(j, weights, scales, mins);
if (scales[j] > max_scale) {
max_scale = scales[j];
}
if (fabsf(mins[j]) > max_abs_min) {
max_abs_min = fabsf(mins[j]);
}
}
if (max_scale > 0) {
float iscale = q4_scale / max_scale;
for (int j = 0; j < Q2_K_SUPER_BLOCK_SIZE; j++) {
int l = (int)lrintf(iscale*scales[j]);
curr_super_block->scales[j] = l;
}
curr_super_block->super_scale = fp16_ieee_from_fp32_value(max_scale / q4_scale);
} else {
for (int j = 0; j < Q2_K_SUPER_BLOCK_SIZE; j++) curr_super_block->scales[j] = 0;
curr_super_block->super_scale = fp16_ieee_from_fp32_value(0.f);
}
if (max_abs_min > 0) {
const float iscale = 7.f / max_abs_min;
for (int j = 0; j < Q2_K_SUPER_BLOCK_SIZE; j++) {
int l = (int)lrintf(iscale * mins[j]);
l = MAX(-8, MIN(7, l));
curr_super_block->scales[j] |= ((l & 0xF) << 4);
}
curr_super_block->super_min = fp16_ieee_from_fp32_value(max_abs_min / 7.f);
} else {
curr_super_block->super_min = fp16_ieee_from_fp32_value(0.f);
}
for (int j = 0; j < Q2_K_SUPER_BLOCK_SIZE; j++) {
const float temp_scale = fp16_ieee_to_fp32_value(curr_super_block->super_scale) * (curr_super_block->scales[j] & 0xF);
const float m = fp16_ieee_to_fp32_value(curr_super_block->super_min);
const int8_t min_q = (curr_super_block->scales[j] >> 4);
const float temp_min = m * ((int8_t)(min_q << 4) >> 4);
for (int ii = 0; ii < Q2_K_BLOCK_SIZE; ii++) {
float val = (temp_scale > 0.f) ? (float_array_aligned[j * Q2_K_BLOCK_SIZE + ii] - temp_min) / temp_scale : 0.f;
int l = (int)lrintf(val);
l = MAX(0, MIN(3, l));
L[j * Q2_K_BLOCK_SIZE + ii] = l;
}
}
uint32_t packed_run = WEIGHT_PER_SUPER_BLOCK / 2; // 128
for (int j = 0; j < WEIGHT_PER_SUPER_BLOCK; j += packed_run) {
for (int l = 0; l < Q2_K_BLOCK_SIZE * 2; l++) { // l = 0..31
uint8_t b0 = L[j + l + 0];
uint8_t b1 = L[j + l + 32];
uint8_t b2 = L[j + l + 64];
uint8_t b3 = L[j + l + 96];
curr_super_block->data[j / 4 + l] = b0 | (b1 << 2) | (b2 << 4) | (b3 << 6);
}
}
float_array_aligned += WEIGHT_PER_SUPER_BLOCK;
}
return 0;
}
int k_dequantize(const quantized_array_q2_k_t *quantized_array_q2_k, float *float_array) {
if (!quantized_array_q2_k || !float_array || quantized_array_q2_k->num_super_blocks == 0) {
return 1;
}
for (uint32_t s = 0; s < quantized_array_q2_k->num_super_blocks; ++s) {
const super_block_q2_k *curr_super_block = &quantized_array_q2_k->super_blocks[s];
const float super_scale = fp16_ieee_to_fp32_value(curr_super_block->super_scale);
const float super_min = fp16_ieee_to_fp32_value(curr_super_block->super_min);
float scales[Q2_K_SUPER_BLOCK_SIZE];
float mins[Q2_K_SUPER_BLOCK_SIZE];
for(int i = 0; i < Q2_K_SUPER_BLOCK_SIZE; ++i) {
uint8_t packed_val = curr_super_block->scales[i];
scales[i] = super_scale * (packed_val & 0x0F);
int8_t min_q = (packed_val >> 4);
mins[i] = super_min * ((int8_t)(min_q << 4) >> 4);
}
const uint8_t *q = curr_super_block->data;
for (int l = 0; l < 32; ++l) {
uint8_t packed_byte = q[l];
int idx0 = l;
int idx1 = l + 32;
int idx2 = l + 64;
int idx3 = l + 96;
float_array[idx0] = mins[idx0/16] + scales[idx0/16] * ((packed_byte >> 0) & 3);
float_array[idx1] = mins[idx1/16] + scales[idx1/16] * ((packed_byte >> 2) & 3);
float_array[idx2] = mins[idx2/16] + scales[idx2/16] * ((packed_byte >> 4) & 3);
float_array[idx3] = mins[idx3/16] + scales[idx3/16] * ((packed_byte >> 6) & 3);
}
for (int l = 0; l < 32; ++l) {
uint8_t packed_byte = q[32 + l];
int idx0 = 128 + l;
int idx1 = 160 + l;
int idx2 = 192 + l;
int idx3 = 224 + l;
float_array[idx0] = mins[idx0/16] + scales[idx0/16] * ((packed_byte >> 0) & 3);
float_array[idx1] = mins[idx1/16] + scales[idx1/16] * ((packed_byte >> 2) & 3);
float_array[idx2] = mins[idx2/16] + scales[idx2/16] * ((packed_byte >> 4) & 3);
float_array[idx3] = mins[idx3/16] + scales[idx3/16] * ((packed_byte >> 6) & 3);
}
float_array += WEIGHT_PER_SUPER_BLOCK;
}
return 0;
}
src/k_quantization.h
+47
View File
@@ -0,0 +1,47 @@
#ifndef K_QUANTIZATION_H
#define K_QUANTIZATION_H
#include <stdint.h>
#include <stdlib.h>
#include <string.h>
#include <math.h>
#include "fp16/fp16.h"
// The setting is refer to https://github.com/ggml-org/llama.cpp/blob/master/ggml/src/ggml-common.h
// fp16 implementation is refer to https://github.com/Maratyszcza/FP16/tree/master/include/fp16
#define Q2_K_BLOCK_SIZE 16
#define Q2_K_SUPER_BLOCK_SIZE 16
#define WEIGHT_PER_SUPER_BLOCK (Q2_K_BLOCK_SIZE*Q2_K_SUPER_BLOCK_SIZE)
// Q2_K 2-bit quantization
// weight is represented as x = a * q + b
// 16 blocks of 16 elements each
// 2.625 bits per weight ((16 * 4 * 2) + (256 * 2) + (16 * 2)) / 256 = 2.625
typedef struct {
uint16_t super_scale; // super-block scale for quantized scales (fp16)
uint16_t super_min; // super-block min for quantized scales (fp16)
uint8_t scales[Q2_K_SUPER_BLOCK_SIZE]; // scales and mins, quantized with 4 bits (length: Q2_K_SUPER_BLOCK_SIZE)
uint8_t data[WEIGHT_PER_SUPER_BLOCK / 4]; // quants with 2 bits (length: WEIGHT_PER_SUPER_BLOCK / 4)
} super_block_q2_k;
typedef struct {
uint64_t num_elements; /* total elements in the original float array */
uint64_t num_elements_aligned; /* aligned (padding) total elements for SUPER_BLOCK ELEMENTS */
uint32_t num_super_blocks;
super_block_q2_k *super_blocks;
} quantized_array_q2_k_t;
quantized_array_q2_k_t *allocate_q2_k_array(uint64_t num_elements);
void free_quantized_q2_k_array(quantized_array_q2_k_t *quantized_array_q2_k);
int64_t get_quantized_q2_k_array_size(const quantized_array_q2_k_t *quantized_array_q2_k);
quantized_array_q2_k_t *load_quantized_q2_k_array_from_buffer(const void *buffer, int64_t buffer_size);
int k_quantize(const float *float_array, uint64_t num_elements, quantized_array_q2_k_t **quantized_array_q2_k);
int k_dequantize(const quantized_array_q2_k_t *quantized_array_q2_k, float *float_array);
#endif
src/llama.cpp
+43
View File
@@ -15,6 +15,7 @@
#include "quantization.h"
#include "sparsity.h"
#include "k_quantization.h"
#ifdef GGML_USE_RPC
# include "ggml-rpc.h"
@@ -18183,6 +18184,30 @@ static void llama_send_tensors(zmq::socket_t & socket, struct llama_ubatch * uba
LLAMA_LOG_INFO("[%d][%s][compute][start][send_tensors][quantize]\n", my_rank, start_compute_time.c_str());
LLAMA_LOG_INFO("[%d][%s][compute][end][send_tensors][quantize]\n", my_rank, end_compute_time.c_str());
}
} else if (comm_datatype_string == "q2_k") {
start_compute_time = get_iso8601_ms_timestamp();
quantized_array_q2_k_t *quantized_array = NULL;
if (k_quantize(ubatch->backend_embd, num_elements,
&quantized_array) || !quantized_array) {
LLAMA_LOG_INFO("Failed to allocate space or do quantization\n");
return;
}
end_compute_time = get_iso8601_ms_timestamp();
buf_size = get_quantized_q2_k_array_size(quantized_array);
send_msgs.emplace_back("sub_gf_out", strlen("sub_gf_out"));
send_msgs.emplace_back("k_quantized", strlen("k_quantized"));
send_msgs.emplace_back(tensors->sub_gf_out->ne,
sizeof(tensors->sub_gf_out->ne));
send_msgs.emplace_back(quantized_array, buf_size);
send_msgs.emplace_back(&buf_size, sizeof(buf_size));
free_quantized_q2_k_array(quantized_array);
if (enable_comm_compute_log) {
LLAMA_LOG_INFO("[%d][%s][compute][start][send_tensors][k_quantize]\n", my_rank, start_compute_time.c_str());
LLAMA_LOG_INFO("[%d][%s][compute][end][send_tensors][k_quantize]\n", my_rank, end_compute_time.c_str());
}
} else if (comm_datatype_string == "f32_sparsity") {
if (comm_sparse_percentage < 1 && comm_sparse_percentage > 100) {
fprintf(stderr, "Sparse percentage %d should between 1~100\n", comm_sparse_percentage);
@@ -18285,6 +18310,24 @@ static void llama_recv_tensors(zmq::socket_t & socket, struct llama_ubatch * uba
LLAMA_LOG_INFO("[%d][%s][compute][end][recv_tensors][dequantize]\n", my_rank, end_compute_time.c_str());
}
}
else if (comm_type == "k_quantized") {
quantized_array_q2_k_t *quantized_array = load_quantized_q2_k_array_from_buffer(data_msg.data(), *buf_size);
if (!quantized_array) {
LLAMA_LOG_INFO("Failed to load quantized array from buffer.\n");
return;
}
std::string start_compute_time = get_iso8601_ms_timestamp();
k_dequantize(quantized_array, batch_embd);
std::string end_compute_time = get_iso8601_ms_timestamp();
free_quantized_q2_k_array(quantized_array);
if (enable_comm_compute_log) {
LLAMA_LOG_INFO("[%d][%s][compute][start][recv_tensors][k_dequantize]\n", my_rank, start_compute_time.c_str());
LLAMA_LOG_INFO("[%d][%s][compute][end][recv_tensors][k_dequantize]\n", my_rank, end_compute_time.c_str());
}
}
else if (comm_type == "sparse") {
sparse_array_t *sparse_array = load_sparse_array_from_buffer(data_msg.data(), *buf_size);
if (!sparse_array) {