root/llama.cpp
1
0
Fork 0
mirror of https://github.com/ggerganov/llama.cpp.git synced 2024-12-28 12:24:35 +00:00
Code Issues Actions 2 Packages Projects Releases Wiki Activity

Fix misaligned memory access in Q4_1 kernel

Browse Source
This commit is contained in:
Iwan Kawrakow 2023-07-28 17:27:01 +03:00
parent 72af25998c
commit dead8f4b5b
1 changed files with 3 additions and 1 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
ggml-cuda.cu
View File

@ -1433,7 +1433,9 @@ static __device__ __forceinline__ float vec_dot_q5_1_q8_1(
const block_q5_1 * bq5_1 = (const block_q5_1 *) vbq; const block_q5_1 * bq5_1 = (const block_q5_1 *) vbq;
// TODO: fix misaligned access // TODO: fix misaligned access
const int qs = *((int *) &bq5_1->qs[sizeof(int) * (iqs + 0)]); int qs;
memcpy(&qs, &bq5_1->qs[sizeof(int) * (iqs + 0)], sizeof(qs));
//const int qs = *((int *) &bq5_1->qs[sizeof(int) * (iqs + 0)]);
const int qh0 = bq5_1->qh[iqs/2 + 0] >> 4*(iqs%2); const int qh0 = bq5_1->qh[iqs/2 + 0] >> 4*(iqs%2);
const int qh1 = bq5_1->qh[iqs/2 + 2] >> 4*(iqs%2); const int qh1 = bq5_1->qh[iqs/2 + 2] >> 4*(iqs%2);
const int ui0 = *((int *) &bq8_1->qs[sizeof(int) * (iqs + 0)]); const int ui0 = *((int *) &bq8_1->qs[sizeof(int) * (iqs + 0)]);