fix(cuda): use N64 centroids in turbo4_64 decode

Assisted-by: OpenAI Codex

Author: kmbandy

ggml/src/ggml-cuda/mt_pagedattn_decode.cu

@@ -315,7 +315,7 @@ static __device__ __forceinline__ void decode_coop_stage_turbo4_64(
         #pragma unroll
         for (int l = 0; l < 2; ++l) {
             const uint8_t idx_nib = (packed >> (l * 4)) & 0xF;
-            const float val = TURBO_CENTROIDS_4BIT[idx_nib] * norm_f;
+            const float val = TURBO_CENTROIDS_4BIT_N64[idx_nib] * norm_f;
             smem_dst[smem_row_base + smem_col_base + l] = __float2half(val);
         }
     }

tests/test-paged-attn-vk.cpp

@@ -1121,7 +1121,7 @@ int main() {
         for (int ctx : { 128, 512 }) {
             char l[64]; snprintf(l, sizeof l, "paged turbo4_64 hd64 decode ctx=%d", ctx);
             const paged_case d64 { 64, 8, 2, 16, 1, ctx, 1, GGML_TYPE_TURBO4_64 };
-            all_ok = compare_paged_case(l, d64, vk, cuda, 5e-2) && all_ok;
+            all_ok = compare_paged_case(l, d64, vk, cuda, 1e-2) && all_ok;
         }
     }