Fixed OpenLLaMA 3b CUDA mul_mat_vec_q

[JohannesGaessler]

Fixes #2136 . The issue was that the weight tensors had row sizes that are not multiples of 128. I fixed it by padding the quantized vector and the last row of the weight tensors to a multiple of 128. This is preferable over adding checks to the CUDA kernels since it has better performance.

[slaren]

The generation looks fine, but compute-sanitizer still reports out of bounds accesses:

========= Invalid __global__ read of size 1 bytes
=========     at 0xb10 in /home/slaren/code/llama.cpp/ggml-cuda.cu:1231:vec_dot_q4_0_q8_1(const void *, const block_q8_1 *, int)
=========     by thread (16,0,0) in block (0,3199,0)
=========     Address 0xdb257e402 is out of bounds
=========     and is 3 bytes after the nearest allocation at 0xdb2000000 of size 5760000 bytes
=========     Device Frame:/home/slaren/code/llama.cpp/ggml-cuda.cu:1388:void mul_mat_vec_q<(int)32, (int)4, block_q4_0, &vec_dot_q4_0_q8_1>(const void *, const void *, float *, int, int) [0xaf0]

[JohannesGaessler]

Thank you for pointing this out, I should have checked it. The value for QUANTIZE_BLOCK_SIZE was incorrect, it should have been 256 instead of 128. One q4 block contains 4 integers so one warp accesses 8 blocks = 256 elements at once.

[JohannesGaessler]

I added another change: the padding is now memset to 0. Though unlikely, it is possible for the unset memory to encode a NaN which could make the sum over the entire row NaN.

[slaren]

So, if I understand correctly, the code depends on the value of CUDA_QUANTIZE_BLOCK_SIZE to ensure that rows sizes are always a multiple of 256. I am a bit concerned that, what previously was just a parameter that could be tuned for the hardware, now will result in incorrect results if changed. To fix this, maybe in quantize_row_q8_1_cuda, k could be rounded up explicitly to a multiple of 256, before rounding it to the block size.

[JohannesGaessler]

I see your point. How about just adding another define that controls the size to which the vector and the last row are extended? I would prefer not to increase k in quantize_row_q8_1 since I want any changes to k to be in the same place that the corresponding memory is allocated.

[slaren]

Sure, that sounds even better.