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

cuda : fix out-of-bounds-access in mul_mat_vec_q

Browse source 
ggml-ci
This commit is contained in:
Georgi Gerganov 2024-01-16 23:06:18 +02:00
parent 36feaeb401
commit b7ddc8bf12
No known key found for this signature in database
GPG key ID: 449E073F9DC10735
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

@ -5131,10 +5131,10 @@ static __global__ void mul_mat_vec_q(const void * __restrict__ vx, const void *
const block_q_t * x = (const block_q_t *) vx; const block_q_t * x = (const block_q_t *) vx;
const block_q8_1 * y = (const block_q8_1 *) vy; const block_q8_1 * y = (const block_q8_1 *) vy;
for (int i = 0; i < blocks_per_row; i += blocks_per_warp) { for (int i = threadIdx.x / (qi/vdr); i < blocks_per_row; i += blocks_per_warp) {
const int ibx = row*blocks_per_row + i + threadIdx.x / (qi/vdr); // x block index const int ibx = row*blocks_per_row + i; // x block index
const int iby = (i + threadIdx.x / (qi/vdr)) * (qk/QK8_1); // y block index that aligns with ibx const int iby = i * (qk/QK8_1); // y block index that aligns with ibx
const int iqs = vdr * (threadIdx.x % (qi/vdr)); // x block quant index when casting the quants to int const int iqs = vdr * (threadIdx.x % (qi/vdr)); // x block quant index when casting the quants to int