Mirrors/llama.cpp
1
0
Fork 0
mirror of https://github.com/ggerganov/llama.cpp.git synced 2024-12-27 06:39:25 +01:00
Code Issues Packages Projects Releases Wiki Activity

iq3_s_mult_shuffle: use new multiplier and cleanup

Browse Source
This commit is contained in:
Iwan Kawrakow 2024-03-05 08:36:57 +02:00
parent b1d753be34
commit 6d15da1ec0
3 changed files with 3 additions and 51 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

49
ggml-cuda.cu
View File

@ -2371,17 +2371,8 @@ static __global__ void dequantize_block_iq3_xxs(const void * __restrict__ vx, ds
}
#ifdef IQ3S_SLOW_MULT
// Better (lower PPL), but requires more bit twidling, so slower
#define IQ3S_MULTIPLIER 190842953LL
#else
#define IQ3S_MULTIPLIER 72968561ULL
//#define IQ3S_MULTIPLIER 540201
//#define IQ3S_MULTIPLIER 1378231
//#define IQ3S_MULTIPLIER 898886
//#define IQ3S_MULTIPLIER 842866
#define IQ3S_MULTIPLIER 518559
static const __device__ uint8_t iq3s_values[16] = {1, 1, 1, 3, 3, 3, 5, 5, 5, 7, 7, 9, 9, 11, 13, 15};
#endif
template<typename dst_t>
static __global__ void dequantize_block_iq3_s(const void * __restrict__ vx, dst_t * __restrict__ yy) {
@ -2400,40 +2391,11 @@ static __global__ void dequantize_block_iq3_s(const void * __restrict__ vx, dst_
const int is = (32*ib + 8*il)/IQ3S_BLOCK_SIZE;
const float d = (float)x[i].d * (1 + 2*((x[i].scales[is/2] >> 4*(is%2)) & 0xf));
const uint8_t signs = x[i].signs[4*ib + il];
#ifdef IQ3S_SLOW_MULT
aux32[0] = ((qs[2*il+0] | ((x[i].qh[ib] << (8-2*il)) & 256)) * IQ3S_MULTIPLIER) & 0x0f0f0f0f;
aux32[1] = ((qs[2*il+1] | ((x[i].qh[ib] << (7-2*il)) & 256)) * IQ3S_MULTIPLIER) & 0x0f0f0f0f;
#else
//aux32[0] = (((qs[2*il+0] | ((x[i].qh[ib] << (8-2*il)) & 256)) * IQ3S_MULTIPLIER) & 0x0f0f0f0f) | 0x01010101;
//aux32[1] = (((qs[2*il+1] | ((x[i].qh[ib] << (7-2*il)) & 256)) * IQ3S_MULTIPLIER) & 0x0f0f0f0f) | 0x01010101;
aux32[0] = (((qs[2*il+0] | ((x[i].qh[ib] << (8-2*il)) & 256)) * IQ3S_MULTIPLIER) & 0x0f0f0f0f);
aux32[1] = (((qs[2*il+1] | ((x[i].qh[ib] << (7-2*il)) & 256)) * IQ3S_MULTIPLIER) & 0x0f0f0f0f);
#endif
//#if __CUDA_ARCH__ >= MIN_CC_DP4A // lowest compute capability for integer intrinsics
//#ifdef IQ3S_SLOW_MULT
// aux32[0] = (((__vmaxs4(__vsub4(aux32[0], 0x01010101), 0x00000000) >> 1) & 0x07070707) << 1) | 0x01010101;
// aux32[1] = (((__vmaxs4(__vsub4(aux32[1], 0x01010101), 0x00000000) >> 1) & 0x07070707) << 1) | 0x01010101;
//#endif
// uint32_t signs0 = __vcmpeq4(((signs & 0xf) * 0x01010101) & 0x08040201, 0x08040201);
// uint32_t signs1 = __vcmpeq4(((signs >> 4) * 0x01010101) & 0x08040201, 0x08040201);
// aux32[0] = __vsub4(aux32[0] ^ signs0, signs0);
// aux32[1] = __vsub4(aux32[1] ^ signs1, signs1);
// for (int j = 0; j < 8; ++j) {
// y[j] = d * grid[j];
// }
//#else
#ifdef IQ3S_SLOW_MULT
for (int j = 0; j < 8; ++j) {
y[j] = d * (2*((grid[j]-1)/2) + 1) * (signs & kmask_iq2xs[j] ? -1.f : 1.f);
}
#else
//static const uint8_t k_values[16] = {1, 1, 1, 3, 3, 3, 5, 5, 7, 7, 9, 9, 11, 11, 13, 15};
for (int j = 0; j < 8; ++j) {
//y[j] = d * grid[j] * (signs & kmask_iq2xs[j] ? -1.f : 1.f);
y[j] = d * iq3s_values[grid[j]] * (signs & kmask_iq2xs[j] ? -1.f : 1.f);
}
#endif
//#endif
#else
assert(false);
#endif
@ -5251,18 +5213,9 @@ static __device__ __forceinline__ float vec_dot_iq3_s_q8_1(
int sumi[2] = {0, 0};
#endif
for (int l = 0; l < 4; ++l) {
#ifdef IQ3S_SLOW_MULT
aux32[0] = ((qs[2*l+0] | ((bq2->qh[ib32] << (8 - 2*l)) & 256)) * IQ3S_MULTIPLIER) & 0x0f0f0f0f;
aux32[1] = ((qs[2*l+1] | ((bq2->qh[ib32] << (7 - 2*l)) & 256)) * IQ3S_MULTIPLIER) & 0x0f0f0f0f;
aux32[0] = (((__vmaxs4(__vsub4(aux32[0], 0x01010101), 0) >> 1) & 0x07070707) << 1) | 0x01010101;
aux32[1] = (((__vmaxs4(__vsub4(aux32[1], 0x01010101), 0) >> 1) & 0x07070707) << 1) | 0x01010101;
#else
//aux32[0] = (((qs[2*l+0] | ((bq2->qh[ib32] << (8 - 2*l)) & 256)) * IQ3S_MULTIPLIER) & 0x0f0f0f0f) | 0x01010101;
//aux32[1] = (((qs[2*l+1] | ((bq2->qh[ib32] << (7 - 2*l)) & 256)) * IQ3S_MULTIPLIER) & 0x0f0f0f0f) | 0x01010101;
aux32[0] = (((qs[2*l+0] | ((bq2->qh[ib32] << (8 - 2*l)) & 256)) * IQ3S_MULTIPLIER) & 0x0f0f0f0f);
aux32[1] = (((qs[2*l+1] | ((bq2->qh[ib32] << (7 - 2*l)) & 256)) * IQ3S_MULTIPLIER) & 0x0f0f0f0f);
for (int j = 0; j < 8; ++j) aux8[j] = iq3s_values[aux8[j]];
#endif
uint32_t signs0 = __vcmpeq4(((bq2->signs[4*ib32+l] & 0xf) * 0x01010101) & 0x08040201, 0x08040201);
uint32_t signs1 = __vcmpeq4(((bq2->signs[4*ib32+l] >> 4) * 0x01010101) & 0x08040201, 0x08040201);
const int grid_l = __vsub4(aux32[0] ^ signs0, signs0);

2
ggml-metal.metal
View File

@ -2546,7 +2546,7 @@ typedef struct {
uint8_t signs[QK_K/8];
uint8_t scales[IQ3S_N_SCALE];
} block_iq3_s;
#define IQ3S_MULTIPLIER 72968561ULL
#define IQ3S_MULTIPLIER 518559
constexpr constant static uint8_t iq3s_values[16] = {1, 1, 1, 3, 3, 3, 5, 5, 5, 7, 7, 9, 9, 11, 13, 15};
typedef struct {

3
ggml-quants.c
View File

@ -4054,8 +4054,7 @@ void dequantize_row_iq3_xxs(const block_iq3_xxs * restrict x, float * restrict y
// ====================== 3.3125 bpw (de)-quantization
#define IQ3S_MULTIPLIER 72968561ULL
#define IQ3S_MULTIPLIER 518559
#define IQ3S_BITS 3
static const uint8_t iq3s_values[16] = {1, 1, 1, 3, 3, 3, 5, 5, 5, 7, 7, 9, 9, 11, 13, 15};