KBaba7/llama.cpp

1#include "common.cuh"2 3static __device__ __forceinline__ void dequantize_q4_0(const void * vx, const int64_t ib, const int iqs, dfloat2 & v){4    const block_q4_0 * x = (const block_q4_0 *) vx;5 6    const dfloat d = x[ib].d;7 8    const int vui = x[ib].qs[iqs];9 10    v.x = vui & 0xF;11    v.y = vui >> 4;12 13#ifdef GGML_CUDA_F1614    v = __hsub2(v, {8.0f, 8.0f});15    v = __hmul2(v, {d, d});16#else17    v.x = (v.x - 8.0f) * d;18    v.y = (v.y - 8.0f) * d;19#endif // GGML_CUDA_F1620}21 22static __device__ __forceinline__ void dequantize_q4_1(const void * vx, const int64_t ib, const int iqs, dfloat2 & v){23    const block_q4_1 * x = (const block_q4_1 *) vx;24 25    const dfloat d = __low2half(x[ib].dm);26    const dfloat m = __high2half(x[ib].dm);27 28    const int vui = x[ib].qs[iqs];29 30    v.x = vui & 0xF;31    v.y = vui >> 4;32 33#ifdef GGML_CUDA_F1634    v = __hmul2(v, {d, d});35    v = __hadd2(v, {m, m});36#else37    v.x = (v.x * d) + m;38    v.y = (v.y * d) + m;39#endif // GGML_CUDA_F1640}41 42static __device__ __forceinline__ void dequantize_q5_0(const void * vx, const int64_t ib, const int iqs, dfloat2 & v){43    const block_q5_0 * x = (const block_q5_0 *) vx;44 45    const dfloat d = x[ib].d;46 47    uint32_t qh;48    memcpy(&qh, x[ib].qh, sizeof(qh));49 50    const int xh_0 = ((qh >> (iqs +  0)) << 4) & 0x10;51    const int xh_1 = ((qh >> (iqs + 12))     ) & 0x10;52 53    v.x = ((x[ib].qs[iqs] & 0xf) | xh_0);54    v.y = ((x[ib].qs[iqs] >>  4) | xh_1);55 56#ifdef GGML_CUDA_F1657    v = __hsub2(v, {16.0f, 16.0f});58    v = __hmul2(v, {d, d});59#else60    v.x = (v.x - 16.0f) * d;61    v.y = (v.y - 16.0f) * d;62#endif // GGML_CUDA_F1663}64 65static __device__ __forceinline__ void dequantize_q5_1(const void * vx, const int64_t ib, const int iqs, dfloat2 & v){66    const block_q5_1 * x = (const block_q5_1 *) vx;67 68    const dfloat d = __low2half(x[ib].dm);69    const dfloat m = __high2half(x[ib].dm);70 71    uint32_t qh;72    memcpy(&qh, x[ib].qh, sizeof(qh));73 74    const int xh_0 = ((qh >> (iqs +  0)) << 4) & 0x10;75    const int xh_1 = ((qh >> (iqs + 12))     ) & 0x10;76 77    v.x = ((x[ib].qs[iqs] & 0xf) | xh_0);78    v.y = ((x[ib].qs[iqs] >>  4) | xh_1);79 80#ifdef GGML_CUDA_F1681    v = __hmul2(v, {d, d});82    v = __hadd2(v, {m, m});83#else84    v.x = (v.x * d) + m;85    v.y = (v.y * d) + m;86#endif // GGML_CUDA_F1687}88 89static __device__ __forceinline__ void dequantize_q8_0(const void * vx, const int64_t ib, const int iqs, dfloat2 & v){90    const block_q8_0 * x = (const block_q8_0 *) vx;91 92    const dfloat d = x[ib].d;93 94    v.x = x[ib].qs[iqs + 0];95    v.y = x[ib].qs[iqs + 1];96 97#ifdef GGML_CUDA_F1698    v = __hmul2(v, {d, d});99#else100    v.x *= d;101    v.y *= d;102#endif // GGML_CUDA_F16103}104