vbatts/llama.cpp
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

Fix merge

Browse source 
---------

Co-authored-by: ardfork <134447697+ardfork@users.noreply.github.com>
Co-authored-by: Kerfuffle <44031344+KerfuffleV2@users.noreply.github.com>
This commit is contained in:
Henri Vasserman 2023-08-10 12:09:28 +03:00
parent 8f8ab6c4c0
commit 29a59b5f07
No known key found for this signature in database
GPG key ID: 2995FC0F58B1A986
1 changed files with 3 additions and 3 deletions
Download patch file Download diff file Expand all files Collapse all files

6
ggml-cuda.cu
View file

@ -1643,8 +1643,8 @@ template <int vdr> static __device__ __forceinline__ float vec_dot_q8_1_q8_1_imp
#else #else
const float2 dm8f = __half22float2(dm8); const float2 dm8f = __half22float2(dm8);
const float2 ds8f = __half22float2(ds8); const float2 ds8f = __half22float2(ds8);
const float d8d8 = dm8f.x * ds8f.x; const float d8d8 = __low2float(dm8) * __low2float(ds8);
const float m8s8 = dm8f.y * ds8f.y; const float m8s8 = __high2float(dm8) * __high2float(ds8);
#endif // GGML_CUDA_F16 #endif // GGML_CUDA_F16
// scale second part of sum by QI8_1/ vdr to compensate for multiple threads adding it // scale second part of sum by QI8_1/ vdr to compensate for multiple threads adding it
@ -3283,7 +3283,7 @@ static __global__ void mul_mat_q(
*dsi_dst = *dsi_src; *dsi_dst = *dsi_src;
} else { } else {
float * dfi_dst = (float *) dsi_dst; float * dfi_dst = (float *) dsi_dst;
*dfi_dst = (*dsi_src).x; *dfi_dst = __low2half(*dsi_src);
} }
} }