ikawrakow/ik_llama.cpp
1
0
Fork 0
mirror of https://github.com/ikawrakow/ik_llama.cpp.git synced 2026-02-23 22:54:10 +00:00
Code Issues Packages Projects Releases Wiki Activity

Bitnet(1.75 bpw): slightly faster CUDA dot product

Browse Source 
We get 205 t/s, so ~13% slower than 2 bit.
This commit is contained in:
Iwan Kawrakow
2024-06-18 18:42:26 +03:00
parent fece7e1db7
commit 181fd9c56e
2 changed files with 5 additions and 11 deletions
Download Patch File Download Diff File Expand all files Collapse all files

3
ggml-cuda/convert.cu
View File

@@ -437,10 +437,11 @@ static __global__ void dequantize_block_iq1_bn(const void * __restrict__ vx, dst
uint8_t u = x[i].extra & 0xff;
s.i = ((((u >> 4) | 0xf0) - 132) << 23) | ((u & 0x0f) << 19);
const float dl = x[i].extra & (1 << (4*ib + il + 8)) ? -s.f : s.f;
const float ml = -dl;
uint16_t idx = x[i].ql[4*ib + il] | ((x[i].qh[2*ib + il/2] << (8 - 4*(il%2))) & 0x0f00);
const uint16_t gp = iq1bn_grid_u16[idx];
for (int j = 0; j < 8; ++j) {
y[j] = dl * (((gp >> 2*j) & 3) - 1);
y[j] = dl * ((gp >> 2*j) & 3) + ml;
}
}

13
ggml-cuda/vecdotq.cuh
View File

@@ -1086,23 +1086,16 @@ static __device__ __forceinline__ float vec_dot_iq1_bn_q8_1(
int sumi = 0;
#if __CUDA_ARCH__ >= MIN_CC_DP4A // lowest compute capability for integer intrinsics
const int * q8 = (const int *)bq8_1[iqs].qs;
//const int minus = 0xffffffff;
for (int l = 0; l < 4; ++l) {
int sign = extra & (1 << l) ? -1 : 1;
uint16_t val = iq1bn_grid_xxx[bq1->ql[4*iqs + l] | ((bq1->qh[2*iqs + l/2] << (8 - 4*(l%2))) & 0x0f00)];
uint8_t vp = val & 0xff, vm = val >> 8;
int32_t vp1 = __vcmpeq4(((vp & 0xf) * 0x01010101) & 0x08040201, 0x08040201);
int32_t vp2 = __vcmpeq4(((vp >> 4) * 0x01010101) & 0x08040201, 0x08040201);
int32_t vm1 = __vcmpeq4(((vm & 0xf) * 0x01010101) & 0x08040201, 0x08040201);
int32_t vm2 = __vcmpeq4(((vm >> 4) * 0x01010101) & 0x08040201, 0x08040201);
sumi += (__dp4a(q8[2*l+0], vm1, __dp4a(q8[2*l+1], vm2, 0)) - __dp4a(q8[2*l+0], vp1, __dp4a(q8[2*l+1], vp2, 0)))*sign;
//int32_t vp1 = __vcmpeq4(((vp & 0xf) * 0x01010101) & 0x08040201, 0x08040201) & q8[2*l+0];
//int32_t vp2 = __vcmpeq4(((vp >> 4) * 0x01010101) & 0x08040201, 0x08040201) & q8[2*l+1];
//int32_t vm1 = __vcmpeq4(((vm & 0xf) * 0x01010101) & 0x08040201, 0x08040201) & q8[2*l+0];
//int32_t vm2 = __vcmpeq4(((vm >> 4) * 0x01010101) & 0x08040201, 0x08040201) & q8[2*l+1];
//int32_t v1 = __vsubss4(vp1, vm1);
//int32_t v2 = __vsubss4(vp2, vm2);
//sumi += __dp4a(v1, 0x01010101, __dp4a(v2, 0x01010101, 0))*sign;
int32_t pm = __dp4a(q8[2*l+0], vm1, __dp4a(q8[2*l+1], vm2, 0));
int32_t pp = __dp4a(q8[2*l+0], vp1, __dp4a(q8[2*l+1], vp2, 0));
sumi += extra & (1 << l) ? pp - pm : pm - pp;
}
#else
const int8_t * q8 = bq8_1[iqs].qs;