diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 90ca17beb..8462f2512 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -79,12 +79,12 @@ jobs: -DCMAKE_BUILD_TYPE=Release cmake --build build --target \ test_dflash test_generate test_flash_attn_sparse test_server_unit \ - test_deepseek4_unit -j$(nproc) + test_deepseek4_unit test_rocmfpx -j$(nproc) - name: Run C++ server unit tests run: | cd server/build - ctest --output-on-failure -R "server_unit|deepseek4_unit" --no-tests=error + ctest --output-on-failure -R "server_unit|deepseek4_unit|rocmfpx" --no-tests=error - name: Populate venv with cu128 torch + setuptools # First pass: install the workspace's default deps. dflash declares diff --git a/server/CMakeLists.txt b/server/CMakeLists.txt index 687938351..e2287e87e 100644 --- a/server/CMakeLists.txt +++ b/server/CMakeLists.txt @@ -718,6 +718,22 @@ if(DFLASH27B_TESTS) target_link_libraries(test_anchor_params PRIVATE dflash_common) add_test(NAME anchor_params COMMAND test_anchor_params) endif() + # ROCMFPX fp3/fp2 MMVQ decode-contract unit test (host-only, no GPU): pins the + # code->value tables, the packed-bit unpacking, and the half-block scale split + # that vec_dot_rocmfpx_fp3/fp2_q8_1 decode from, against the CPU dequantizer. + # Self-contained (compiles rocmfpx.c directly) so it runs on the hosted CI job. + if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/deps/llama.cpp/ggml/rocmfpx/test_rocmfpx.c") + add_executable(test_rocmfpx + deps/llama.cpp/ggml/rocmfpx/test_rocmfpx.c + deps/llama.cpp/ggml/rocmfpx/rocmfpx.c) + target_include_directories(test_rocmfpx PRIVATE + deps/llama.cpp/ggml/include + deps/llama.cpp/ggml/rocmfpx) + if(NOT WIN32) + target_link_libraries(test_rocmfpx PRIVATE m) + endif() + add_test(NAME rocmfpx_unit COMMAND test_rocmfpx) + endif() if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/test/test_drafter_early_exit_score_range.cpp") add_executable(test_drafter_early_exit_score_range test/test_drafter_early_exit_score_range.cpp) diff --git a/server/deps/llama.cpp/ggml/rocmfpx/test_rocmfpx.c b/server/deps/llama.cpp/ggml/rocmfpx/test_rocmfpx.c index 5bae72bd7..6729f0b56 100644 --- a/server/deps/llama.cpp/ggml/rocmfpx/test_rocmfpx.c +++ b/server/deps/llama.cpp/ggml/rocmfpx/test_rocmfpx.c @@ -40,6 +40,166 @@ static float weighted_mse(const float * a, const float * b, const float * w, int return sum_w > 0.0f ? err / sum_w : 0.0f; } +// ── Kernel-decode correctness ──────────────────────────────────────────── +// The GPU MMVQ kernels (vec_dot_rocmfpx_fp3_q8_1 / _fp2_q8_1) don't dequantize +// to float — they decode each packed code to an int8 kvalue, integer-dp4a it +// against the q8 activations per half-block, then scale by the half-block's +// ue4m3 e[]. These host tests pin the three things that decode contract depends +// on — the code->value table, the packed-bit unpacking, and the half-block scale +// assignment — against the independent library dequantizer. The device perm-LUT +// decode is in turn tied to this same code->value mapping by a compile-time +// static_assert in vecdotq.cuh, so a divergence fails the GPU build too. + +// fp3 code -> signed kvalue, mirroring rocmfpx_decode_fp3_code_vec_cuda: +// mag = (code&3)==3 ? 4 : code&3; value = (code&4) ? -mag : mag. => {0,1,2,4,0,-1,-2,-4} +static int fp3_code_value(unsigned code) { + const unsigned mag_code = code & 3u; + const int mag = (mag_code == 3u) ? 4 : (int) mag_code; + return (code & 4u) ? -mag : mag; +} + +static int fp2_code_value(unsigned code) { + static const int kv[4] = { + ROCMFP2_KVALUE_0_I8, ROCMFP2_KVALUE_1_I8, ROCMFP2_KVALUE_2_I8, ROCMFP2_KVALUE_3_I8 + }; + return kv[code & 3u]; +} + +// Unpack fp3 code i from the little-endian 3-bit stream (matches the kernel's +// dword-splice extraction: (stream >> 3i) & 7). 16-bit window handles a code +// straddling a byte boundary; the guard avoids reading past the 12-byte qs. +static unsigned fp3_get_code(const uint8_t * qs, int i) { + const int bit = 3 * i; + const int byte = bit >> 3; + unsigned window = qs[byte]; + if (byte + 1 < QS_ROCMFP3) { + window |= (unsigned) qs[byte + 1] << 8; + } + return (window >> (bit & 7)) & 7u; +} + +// fp2 code i: byte i/4, 2 bits at 2*(i%4) (matches rocmfpx_pack4_fp2_vec_cuda). +static unsigned fp2_get_code(const uint8_t * qs, int i) { + return (qs[i >> 2] >> (2 * (i & 3))) & 3u; +} + +// Table pinned exhaustively — this is the contract the device perm-LUT encodes. +static void check_decode_tables(void) { + const int fp3_expect[8] = { 0, 1, 2, 4, 0, -1, -2, -4 }; + for (unsigned c = 0; c < 8; ++c) { + assert(fp3_code_value(c) == fp3_expect[c]); + } + const int fp2_expect[4] = { -1, 0, 1, 2 }; + for (unsigned c = 0; c < 4; ++c) { + assert(fp2_code_value(c) == fp2_expect[c]); + } + printf("decode tables: fp3 {0,1,2,4,0,-1,-2,-4} fp2 {-1,0,1,2} OK\n"); +} + +// Reconstruct each weight from code->value * half-block scale using the kernel's +// unpack, and assert it equals the library dequantizer bit-for-bit. This pins +// the bit order, the code table, AND the e[i>=QK/2] half-block scale split. +static void check_decode_contract_fp3(void) { + enum { N = 2 * QK_ROCMFP3 }; + float src[N], ref[N]; + block_rocmfp3 q[N / QK_ROCMFP3]; + fill_row(src, N); + rocmfpx_quantize_fp3(src, q, 1, N, NULL); + rocmfpx_dequantize_row_fp3(q, ref, N); + for (int b = 0; b < N / QK_ROCMFP3; ++b) { + const block_rocmfp3 * blk = &q[b]; + const float e0 = rocmfpx_ue4m3_to_fp32(blk->e[0]); + const float e1 = rocmfpx_ue4m3_to_fp32(blk->e[1]); + for (int i = 0; i < QK_ROCMFP3; ++i) { + const float e = (i >= QK_ROCMFP3 / 2) ? e1 : e0; + const float recon = (float) fp3_code_value(fp3_get_code(blk->qs, i)) * e; + assert(recon == ref[b * QK_ROCMFP3 + i]); + } + } + printf("fp3 decode contract: unpack + table + half-block scale == dequantize_row (exact) OK\n"); +} + +static void check_decode_contract_fp2(void) { + enum { N = 2 * QK_ROCMFP2 }; + float src[N], ref[N]; + block_rocmfp2 q[N / QK_ROCMFP2]; + fill_row(src, N); + rocmfpx_quantize_fp2(src, q, 1, N, NULL); + rocmfpx_dequantize_row_fp2(q, ref, N); + for (int b = 0; b < N / QK_ROCMFP2; ++b) { + const block_rocmfp2 * blk = &q[b]; + const float e0 = rocmfpx_ue4m3_to_fp32(blk->e[0]); + const float e1 = rocmfpx_ue4m3_to_fp32(blk->e[1]); + for (int i = 0; i < QK_ROCMFP2; ++i) { + const float e = (i >= QK_ROCMFP2 / 2) ? e1 : e0; + const float recon = (float) fp2_code_value(fp2_get_code(blk->qs, i)) * e; + assert(recon == ref[b * QK_ROCMFP2 + i]); + } + } + printf("fp2 decode contract: unpack + table + half-block scale == dequantize_row (exact) OK\n"); +} + +// The kernel's dot is db * sum_half( e_half * sum_i code_value[i]*q8[i] ) with an +// exact integer inner sum per half-block (the VDR=4 half-block structure for fp3 +// / the two-chain fp2 form). Assert it matches the float reference dot to tol. +static void check_dot_structure_fp3(void) { + enum { N = QK_ROCMFP3 }; + float src[N], w[N]; + block_rocmfp3 q[1]; + fill_row(src, N); + rocmfpx_quantize_fp3(src, q, 1, N, NULL); + rocmfpx_dequantize_row_fp3(q, w, N); + + int8_t a8[N]; + for (int i = 0; i < N; ++i) { a8[i] = (int8_t) ((i * 37) % 251 - 125); } + const float db = 0.031f; + + double ref = 0.0; + for (int i = 0; i < N; ++i) { ref += (double) w[i] * (double) (db * (float) a8[i]); } + + const block_rocmfp3 * blk = &q[0]; + long s0 = 0, s1 = 0; + for (int i = 0; i < N; ++i) { + const int c = fp3_code_value(fp3_get_code(blk->qs, i)); + if (i < N / 2) { s0 += (long) c * a8[i]; } else { s1 += (long) c * a8[i]; } + } + const double e0 = rocmfpx_ue4m3_to_fp32(blk->e[0]); + const double e1 = rocmfpx_ue4m3_to_fp32(blk->e[1]); + const double kern = (double) db * (e0 * (double) s0 + e1 * (double) s1); + const double rel = fabs(kern - ref) / (fabs(ref) + 1e-9); + printf("fp3 dot structure: ref=%g kernel=%g rel=%g\n", ref, kern, rel); + assert(rel < 1e-5); +} + +static void check_dot_structure_fp2(void) { + enum { N = QK_ROCMFP2 }; + float src[N], w[N]; + block_rocmfp2 q[1]; + fill_row(src, N); + rocmfpx_quantize_fp2(src, q, 1, N, NULL); + rocmfpx_dequantize_row_fp2(q, w, N); + + int8_t a8[N]; + for (int i = 0; i < N; ++i) { a8[i] = (int8_t) ((i * 37) % 251 - 125); } + const float db = 0.031f; + + double ref = 0.0; + for (int i = 0; i < N; ++i) { ref += (double) w[i] * (double) (db * (float) a8[i]); } + + const block_rocmfp2 * blk = &q[0]; + long s0 = 0, s1 = 0; + for (int i = 0; i < N; ++i) { + const int c = fp2_code_value(fp2_get_code(blk->qs, i)); + if (i < N / 2) { s0 += (long) c * a8[i]; } else { s1 += (long) c * a8[i]; } + } + const double e0 = rocmfpx_ue4m3_to_fp32(blk->e[0]); + const double e1 = rocmfpx_ue4m3_to_fp32(blk->e[1]); + const double kern = (double) db * (e0 * (double) s0 + e1 * (double) s1); + const double rel = fabs(kern - ref) / (fabs(ref) + 1e-9); + printf("fp2 dot structure: ref=%g kernel=%g rel=%g\n", ref, kern, rel); + assert(rel < 1e-5); +} + static void check_weighted_imatrix_fp3(void) { enum { N = QK_ROCMFP3 }; @@ -160,5 +320,12 @@ int main(void) { check_weighted_imatrix_fp2(); check_weighted_imatrix_fp3(); + // Kernel-decode correctness (MMVQ fp3/fp2 vec-dot contract). + check_decode_tables(); + check_decode_contract_fp3(); + check_decode_contract_fp2(); + check_dot_structure_fp3(); + check_dot_structure_fp2(); + return 0; } diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/mmvq.cu b/server/deps/llama.cpp/ggml/src/ggml-cuda/mmvq.cu index 620a1b3fb..a3a4f4415 100644 --- a/server/deps/llama.cpp/ggml/src/ggml-cuda/mmvq.cu +++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/mmvq.cu @@ -415,6 +415,51 @@ static constexpr __host__ __device__ int calc_rows_per_block(int ncols_dst, int return 1; } +#if defined(GGML_USE_HIP) +// ROCMFPX (fp2/fp3) cross-lane float reduction. On RDNA, HIP's warp_reduce_sum(float) +// lowers its __shfl_xor butterfly to ds_bpermute, which builds a per-lane +// address VGPR (lane ^ offset) that sits on the dependent shuffle chain. For an +// xor mask inside a 32-lane group the SAME permutation is one ds_swizzle_b32 +// bitmask-mode op with the xor encoded as an instruction immediate +// ((offset<<10)|0x1f: xor=offset, or=0, and=0x1f) — no address VGPR. It is +// bit-identical to warp_reduce_sum's ladder: identical {16,8,4,2,1} reduction +// order, identical float adds (ds_swizzle by mask m gives src lane (lane&0x1f)^m, +// exactly __shfl_xor's partner; verified 0 lane mismatches on gfx1201 for those +// masks). warp==32 on RDNA so every offset is a valid intra-group swizzle. +template +static __device__ __forceinline__ float warp_reduce_sum_rocmfpx_dsswizzle(float x) { + static_assert(width == 32, "ROCMFPX ds_swizzle reduce assumes a 32-lane RDNA warp"); + x += __int_as_float(__builtin_amdgcn_ds_swizzle(__float_as_int(x), (16 << 10) | 0x1f)); + x += __int_as_float(__builtin_amdgcn_ds_swizzle(__float_as_int(x), ( 8 << 10) | 0x1f)); + x += __int_as_float(__builtin_amdgcn_ds_swizzle(__float_as_int(x), ( 4 << 10) | 0x1f)); + x += __int_as_float(__builtin_amdgcn_ds_swizzle(__float_as_int(x), ( 2 << 10) | 0x1f)); + x += __int_as_float(__builtin_amdgcn_ds_swizzle(__float_as_int(x), ( 1 << 10) | 0x1f)); + return x; +} +#endif + +// MMVQ warp reduction, specialised to the ds_swizzle ladder for the ROCMFPX +// fp2 (GGML_TYPE_Q2_0_ROCMFP2) and fp3 (GGML_TYPE_Q3_0_ROCMFPX) instantiations +// only — every other quant type, and all non-HIP builds, keep the shared +// warp_reduce_sum unchanged. Gating on `type` keeps the change local to these +// kernels' codegen (no backend-wide effect). The `width == 32` guard falls back +// to warp_reduce_sum for any non-wave32 HIP target (e.g. wave64 CDNA): +// ds_swizzle is a 32-lane intra-group op, so the fast path only applies on RDNA +// wave32 and everything else stays generic. The ladder is bit-identical to +// warp_reduce_sum for every type (same {16,8,4,2,1} order, same float adds). +template +static __device__ __forceinline__ float warp_reduce_sum_mmvq(float x) { +#if defined(GGML_USE_HIP) + if constexpr ((type == GGML_TYPE_Q3_0_ROCMFPX || type == GGML_TYPE_Q2_0_ROCMFP2) && width == 32) { + return warp_reduce_sum_rocmfpx_dsswizzle(x); + } else { + return warp_reduce_sum(x); + } +#else + return warp_reduce_sum(x); +#endif +} + template __launch_bounds__(calc_nwarps(type, ncols_dst, get_device_table_id())*ggml_cuda_get_physical_warp_size(), 1) static __global__ void mul_mat_vec_q( @@ -585,10 +630,10 @@ static __global__ void mul_mat_vec_q( } } } - tmp[j][i] = warp_reduce_sum(tmp[j][i]); + tmp[j][i] = warp_reduce_sum_mmvq(tmp[j][i]); if constexpr (has_fusion) { if (use_gate) { - tmp_gate[j][i] = warp_reduce_sum(tmp_gate[j][i]); + tmp_gate[j][i] = warp_reduce_sum_mmvq(tmp_gate[j][i]); } } } diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/vecdotq.cuh b/server/deps/llama.cpp/ggml/src/ggml-cuda/vecdotq.cuh index 8e8e79eb2..de540b8ec 100644 --- a/server/deps/llama.cpp/ggml/src/ggml-cuda/vecdotq.cuh +++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/vecdotq.cuh @@ -347,7 +347,7 @@ static __device__ __forceinline__ float vec_dot_mxfp4_q8_1( #define VDR_ROCMFP4_FAST_Q8_1_MMVQ GGML_ROCMFP4_FAST_Q8_1_MMVQ_VDR #define VDR_ROCMFP4_FAST_Q8_1_MMQ GGML_ROCMFP4_FAST_Q8_1_MMQ_VDR #define VDR_ROCMFP2_Q8_1_MMVQ 1 -#define VDR_ROCMFP3_Q8_1_MMVQ 2 +#define VDR_ROCMFP3_Q8_1_MMVQ 4 #ifndef VDR_ROCMFP6_Q8_1_MMVQ #define VDR_ROCMFP6_Q8_1_MMVQ 4 #endif @@ -382,7 +382,7 @@ static __device__ __forceinline__ uint32_t rocmfpx_get_bits_vec_cuda(const uint8 return code; } -static __device__ __forceinline__ int rocmfpx_decode_fp3_code_vec_cuda(const uint32_t code) { +static __host__ __device__ __forceinline__ constexpr int rocmfpx_decode_fp3_code_vec_cuda(const uint32_t code) { const uint32_t mag_code = code & 3u; const int mag = mag_code == 3u ? 4 : (int) mag_code; return (code & 4u) ? -mag : mag; @@ -417,13 +417,81 @@ static __device__ __forceinline__ int rocmfpx_pack4_fp6_bits24_vec_cuda(const ui return *((const int *) &v); } -static __device__ __forceinline__ int rocmfpx_pack4_fp3_vec_cuda(const uint8_t * qs, const int base) { +// Decode 4 fp3 codes packed in the low 12 bits of `bits12` into 4 int8 kvalues, +// returned as a packed int32. On HIP this replaces the scalar per-code decode + +// char4 assembly with a single v_perm_b32: rocmfpx_decode_fp3_code_vec_cuda maps +// the 8 possible 3-bit codes to {0,1,2,4,0,-1,-2,-4}, which fits an 8-byte LUT +// split across two registers. __builtin_amdgcn_perm selects byte n<4 from the +// low operand and n>=4 from the high operand, so each 3-bit code (0..7) indexes +// the table directly as a perm selector. Bit-identical to the make_char4 form. +static __device__ __forceinline__ int rocmfpx_pack4_fp3_bits12_vec_cuda(const uint32_t bits12) { +#if defined(GGML_USE_HIP) + constexpr uint32_t lut_lo = 0x04020100u; // codes 0,1,2,3 -> 0,1,2,4 + constexpr uint32_t lut_hi = 0xFCFEFF00u; // codes 4,5,6,7 -> 0,-1,-2,-4 + // Keep the packed LUT in lock-step with the scalar decode (and thus the CUDA + // path): a change to the code->value mapping fails the build here instead of + // silently diverging. Single source of truth = rocmfpx_decode_fp3_code_vec_cuda. + static_assert( + (int8_t)( lut_lo & 0xffu) == (int8_t) rocmfpx_decode_fp3_code_vec_cuda(0) && + (int8_t)((lut_lo >> 8) & 0xffu) == (int8_t) rocmfpx_decode_fp3_code_vec_cuda(1) && + (int8_t)((lut_lo >> 16) & 0xffu) == (int8_t) rocmfpx_decode_fp3_code_vec_cuda(2) && + (int8_t)((lut_lo >> 24) & 0xffu) == (int8_t) rocmfpx_decode_fp3_code_vec_cuda(3) && + (int8_t)( lut_hi & 0xffu) == (int8_t) rocmfpx_decode_fp3_code_vec_cuda(4) && + (int8_t)((lut_hi >> 8) & 0xffu) == (int8_t) rocmfpx_decode_fp3_code_vec_cuda(5) && + (int8_t)((lut_hi >> 16) & 0xffu) == (int8_t) rocmfpx_decode_fp3_code_vec_cuda(6) && + (int8_t)((lut_hi >> 24) & 0xffu) == (int8_t) rocmfpx_decode_fp3_code_vec_cuda(7), + "fp3 perm LUT is out of sync with rocmfpx_decode_fp3_code_vec_cuda"); + const uint32_t selectors = + ((bits12 >> 0) & 7u) | + (((bits12 >> 3) & 7u) << 8) | + (((bits12 >> 6) & 7u) << 16) | + (((bits12 >> 9) & 7u) << 24); + return (int) __builtin_amdgcn_perm(lut_hi, lut_lo, selectors); +#else const char4 v = make_char4( - (int8_t) rocmfpx_decode_fp3_code_vec_cuda(rocmfpx_get_bits_vec_cuda(qs, (base + 0)*3, 3)), - (int8_t) rocmfpx_decode_fp3_code_vec_cuda(rocmfpx_get_bits_vec_cuda(qs, (base + 1)*3, 3)), - (int8_t) rocmfpx_decode_fp3_code_vec_cuda(rocmfpx_get_bits_vec_cuda(qs, (base + 2)*3, 3)), - (int8_t) rocmfpx_decode_fp3_code_vec_cuda(rocmfpx_get_bits_vec_cuda(qs, (base + 3)*3, 3))); + (int8_t) rocmfpx_decode_fp3_code_vec_cuda((bits12 >> 0) & 7u), + (int8_t) rocmfpx_decode_fp3_code_vec_cuda((bits12 >> 3) & 7u), + (int8_t) rocmfpx_decode_fp3_code_vec_cuda((bits12 >> 6) & 7u), + (int8_t) rocmfpx_decode_fp3_code_vec_cuda((bits12 >> 9) & 7u)); return *((const int *) &v); +#endif +} + +// Decode 4 consecutive fp3 codes starting at code index `base` from the packed +// qs byte array. Used by the MMQ load-tile path (load_tiles_rocmfp3). The 4 codes +// occupy the 12 contiguous bits [base*3, base*3+12), so extract that field once +// and hand it to the perm-LUT packer instead of doing 4 branchy scalar decodes + +// a make_char4. On HIP the four rocmfpx_decode_fp3_code_vec_cuda + make_char4 +// collapse to selector assembly + one v_perm_b32 (see rocmfpx_pack4_fp3_bits12_vec_cuda). +// +// The 12-bit field is extracted with a dword-splice idiom instead of the +// bit-by-bit rocmfpx_get_bits_vec_cuda (a 12-iteration single-bit OR loop): load +// the 12 qs bytes as 3 little-endian dwords once, then splice the field out with a +// shift + a single cross-dword OR. word_in_block ∈ [0,7] → base ∈ {0,4,..,28} → +// bit_start = base*3 ∈ {0,12,..,84}, so the field never exceeds bit 96 (the last +// dword only ever contributes its low bits; the {..,0} pad covers idx+1 at idx==2). +// Bit-identical to get_bits: byte b provides stream bits [8b,8b+8) LSB-first, and a +// little-endian memcpy packs byte b at bit 8b of the dword, so the concatenated +// dwords are the same bit stream — (stream >> bit_start) & 0xFFF == get_bits(base*3,12). +// Note: this reads the full 12-byte fp3 qs field (3 dwords) unconditionally, +// regardless of `base`; callers pass the block's 12-byte qs array (QK_ROCMFP3=32 +// weights * 3 bits = 12 bytes), so the load is always in-bounds. +static __device__ __forceinline__ int rocmfpx_pack4_fp3_vec_cuda(const uint8_t * qs, const int base) { + uint32_t w0, w1, w2; + memcpy(&w0, qs + 0, 4); + memcpy(&w1, qs + 4, 4); + memcpy(&w2, qs + 8, 4); + const uint32_t words[4] = { w0, w1, w2, 0 }; + + const int bit_start = base * 3; + const int idx = bit_start >> 5; + const int shift = bit_start & 31; + const uint32_t lo = words[idx]; + const uint32_t hi = words[idx + 1]; + const uint32_t bits12 = (shift == 0) ? (lo & 0xFFFu) + : (((lo >> shift) | (hi << (32 - shift))) & 0xFFFu); + + return rocmfpx_pack4_fp3_bits12_vec_cuda(bits12); } static __device__ __forceinline__ int rocmfpx_pack4_fp2_bits8_vec_cuda(const uint32_t bits8) { @@ -545,13 +613,22 @@ static __device__ __forceinline__ float vec_dot_rocmfpx_fp2_q8_1( // qs is a 10-byte-strided, byte-aligned array -> read byte-wise (uint8_t index); NEVER cast qs to int*/int2*. const block_rocmfp2 * bq2 = (const block_rocmfp2 *) vbq + kbx; - int sumi = 0; -#pragma unroll - for (int j = 0; j < 4; ++j) { - const int val_packed = rocmfpx_pack4_fp2_bits8_vec_cuda((uint32_t) bq2->qs[4*iqs + j]); - const int u = get_int_b4(bq8_1->qs, 4*iqs + j); - sumi = ggml_cuda_dp4a(val_packed, u, sumi); - } + // Two independent dp4a chains (bytes {0,2} and {1,3}) to hide dp4a latency. + // Integer accumulation is associative and exact, so sumi_a + sumi_b is + // BIT-IDENTICAL to the serial single-accumulator chain: this touches only the + // in-lane int ordering, not the float reduction or the lane assignment (VDR=1 + // already covers one half-block per lane), so the result is byte-identical. + const int vp0 = rocmfpx_pack4_fp2_bits8_vec_cuda((uint32_t) bq2->qs[4*iqs + 0]); + const int vp1 = rocmfpx_pack4_fp2_bits8_vec_cuda((uint32_t) bq2->qs[4*iqs + 1]); + const int vp2 = rocmfpx_pack4_fp2_bits8_vec_cuda((uint32_t) bq2->qs[4*iqs + 2]); + const int vp3 = rocmfpx_pack4_fp2_bits8_vec_cuda((uint32_t) bq2->qs[4*iqs + 3]); + int sumi_a = 0; + int sumi_b = 0; + sumi_a = ggml_cuda_dp4a(vp0, get_int_b4(bq8_1->qs, 4*iqs + 0), sumi_a); + sumi_b = ggml_cuda_dp4a(vp1, get_int_b4(bq8_1->qs, 4*iqs + 1), sumi_b); + sumi_a = ggml_cuda_dp4a(vp2, get_int_b4(bq8_1->qs, 4*iqs + 2), sumi_a); + sumi_b = ggml_cuda_dp4a(vp3, get_int_b4(bq8_1->qs, 4*iqs + 3), sumi_b); + const int sumi = sumi_a + sumi_b; const float db = __low2float(bq8_1->ds); return db * rocmfpx_ue4m3_to_fp32_finite(bq2->e[iqs]) * sumi; @@ -569,50 +646,54 @@ static __device__ __forceinline__ float vec_dot_rocmfpx_fp3_q8_1( const uint32_t qs[4] = { qs0, qs1, qs2, 0 }; - int sumi0 = 0; - int sumi1 = 0; + // VDR=4 half-block-aligned MMVQ (structural rewrite of the VDR=2 form). + // The MMVQ launchers dispatch iqs = VDR*(tid % (QI_ROCMFP3/VDR)) = 4*{0,1} = + // {0,4} (QI_ROCMFP3=8, VDR=4), so each vec_dot covers exactly ONE 16-weight + // half-block = 4 adjacent 4-weight groups with a SINGLE scale: iqs==0 -> + // weights 0..15 (e[0]); iqs==4 -> weights 16..31 (e[1]). No half-block + // straddle is ever possible, so the scale is selected once with no per-element + // branch. Versus the VDR=2 kernel this doubles work-per-lane (4 dp4a vs 2), + // halves lanes-per-block (2 vs 4) and the K-loop trip count, and amortises the + // bit-window extraction over 4 groups. + // + // NOT bit-identical to the VDR=2 form: one lane now sums 4 groups instead of + // 2, so the cross-lane float reduction tree is regrouped. The int dp4a sums + // stay exact; the only difference is <=~1 ULP per block in the int->float + // product ordering, far below the argmax margin. Validated by logit-tol + + // greedy-output SHA + spec accept_rate rather than byte-identity. + static_assert(VDR_ROCMFP3_Q8_1_MMVQ == 4, "half-block window assumes VDR==4"); + assert((iqs == 0 || iqs == QK_ROCMFP3/8) && + "vec_dot_rocmfpx_fp3_q8_1: VDR=4 expects half-block-aligned iqs (0 or 4)"); + + // The 4-group window is 48 contiguous bits starting at 12*iqs. For the two + // legal iqs (0 -> bit 0, 4 -> bit 48) the field lies within two adjacent + // dwords after the intra-dword shift (shift is 0 or 16, and shift+48 <= 64), + // so a single 64-bit splice extracts all 48 bits at once. + const int win_start = 12 * iqs; + const int win_idx = win_start >> 5; // 0 or 1 + const int win_shift = win_start & 31; // 0 or 16 + const uint64_t chunk = (uint64_t) qs[win_idx] | ((uint64_t) qs[win_idx + 1] << 32); + const uint64_t bits48 = (chunk >> win_shift) & 0xFFFFFFFFFFFFull; + + // Two independent dp4a chains (groups {0,2} and {1,3}) to hide dp4a latency; + // the int sums are exact so the final int add is order-independent. + const int vp0 = rocmfpx_pack4_fp3_bits12_vec_cuda((uint32_t)( bits48 & 0xFFFu)); + const int vp1 = rocmfpx_pack4_fp3_bits12_vec_cuda((uint32_t)((bits48 >> 12) & 0xFFFu)); + const int vp2 = rocmfpx_pack4_fp3_bits12_vec_cuda((uint32_t)((bits48 >> 24) & 0xFFFu)); + const int vp3 = rocmfpx_pack4_fp3_bits12_vec_cuda((uint32_t)((bits48 >> 36) & 0xFFFu)); + int sumi_a = 0; + int sumi_b = 0; + sumi_a = ggml_cuda_dp4a(vp0, get_int_b4(bq8_1->qs, iqs + 0), sumi_a); + sumi_b = ggml_cuda_dp4a(vp1, get_int_b4(bq8_1->qs, iqs + 1), sumi_b); + sumi_a = ggml_cuda_dp4a(vp2, get_int_b4(bq8_1->qs, iqs + 2), sumi_a); + sumi_b = ggml_cuda_dp4a(vp3, get_int_b4(bq8_1->qs, iqs + 3), sumi_b); + const int sumi = sumi_a + sumi_b; - // The two half-block scales (e[0]/e[1]) split at element QK_ROCMFP3/2. base - // < QK_ROCMFP3/2 is equivalent to (iqs+i) < QK_ROCMFP3/8, so for a VDR - // window that lies entirely in one half the accumulator choice is loop - // invariant and can be hoisted out of the unrolled loop. A straddling window - // (only possible for VDRs that cross the midpoint) still uses the exact - // per-element branch, so results are bit-identical either way. - const bool fp3_first_half = iqs + VDR_ROCMFP3_Q8_1_MMVQ <= QK_ROCMFP3/8; const bool fp3_second_half = iqs >= QK_ROCMFP3/8; - -#pragma unroll - for (int i = 0; i < VDR_ROCMFP3_Q8_1_MMVQ; ++i) { - const int base = 4 * (iqs + i); - const int start_bit = 12 * (iqs + i); - const int reg_idx = start_bit >> 5; - const int reg_shift = start_bit & 31; - const uint32_t val_low = qs[reg_idx]; - const uint32_t val_high = qs[reg_idx + 1]; - const uint32_t bits12 = (reg_shift == 0) ? (val_low & 0xFFFu) : (((val_low >> reg_shift) | (val_high << (32 - reg_shift))) & 0xFFFu); - - const char4 v = make_char4( - (int8_t) rocmfpx_decode_fp3_code_vec_cuda(bits12 & 7u), - (int8_t) rocmfpx_decode_fp3_code_vec_cuda((bits12 >> 3) & 7u), - (int8_t) rocmfpx_decode_fp3_code_vec_cuda((bits12 >> 6) & 7u), - (int8_t) rocmfpx_decode_fp3_code_vec_cuda((bits12 >> 9) & 7u)); - const int val_packed = *((const int *) &v); - - const int u = get_int_b4(bq8_1->qs, iqs + i); - - if (fp3_first_half) { - sumi0 = ggml_cuda_dp4a(val_packed, u, sumi0); - } else if (fp3_second_half) { - sumi1 = ggml_cuda_dp4a(val_packed, u, sumi1); - } else if (base < QK_ROCMFP3/2) { - sumi0 = ggml_cuda_dp4a(val_packed, u, sumi0); - } else { - sumi1 = ggml_cuda_dp4a(val_packed, u, sumi1); - } - } + const float e_sel = rocmfpx_ue4m3_to_fp32_finite(bq3->e[fp3_second_half]); const float db = __low2float(bq8_1->ds); - return db * (rocmfpx_ue4m3_to_fp32_finite(bq3->e[0]) * sumi0 + rocmfpx_ue4m3_to_fp32_finite(bq3->e[1]) * sumi1); + return db * (e_sel * sumi); } static __device__ __forceinline__ float vec_dot_rocmfpx_fp6_q8_1(