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]

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.

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.

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.

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.

Sure, that sounds even better.