diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 90ca17beb..9432c482b 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -273,3 +273,19 @@ jobs: /opt/rocm/bin/hipcc --offload-arch=gfx1151 -O2 \ -o "$RUNNER_TEMP/hip_smoke" .github/ci/hip_smoke.cpp "$RUNNER_TEMP/hip_smoke" + + - name: Build + test ROCmFP formats + run: | + cmake -S server -B "$RUNNER_TEMP/rocmfp-build" \ + -DDFLASH27B_GPU_BACKEND=hip \ + -DDFLASH27B_HIP_ARCHITECTURES=gfx1151 \ + -DDFLASH27B_SERVER=OFF \ + -DDFLASH27B_TESTS=ON \ + -DCMAKE_BUILD_TYPE=Release \ + -DCMAKE_HIP_FLAGS=-DDFLASH_WAVE_SIZE=32 + cmake --build "$RUNNER_TEMP/rocmfp-build" \ + --target test_rocmfp4 test_rocmfpx test_rocmfp4_hip_tail \ + --parallel 8 + ctest --test-dir "$RUNNER_TEMP/rocmfp-build" \ + --output-on-failure \ + -R 'rocmfp4_reference|rocmfpx_reference|rocmfp4_hip_tail' diff --git a/server/CMakeLists.txt b/server/CMakeLists.txt index 9722244a4..853cf571e 100644 --- a/server/CMakeLists.txt +++ b/server/CMakeLists.txt @@ -655,6 +655,36 @@ if(DFLASH27B_TESTS) ${DFLASH27B_GGML_BACKEND_TARGET} ggml ggml-base) add_test(NAME cuda_pool_shutdown COMMAND test_cuda_pool_shutdown) endif() + + add_executable(test_rocmfp4 deps/llama.cpp/ggml/rocmfp4/test_rocmfp4.c) + target_link_libraries(test_rocmfp4 PRIVATE ggml-base) + if(UNIX) + target_link_libraries(test_rocmfp4 PRIVATE m) + endif() + add_test(NAME rocmfp4_reference COMMAND test_rocmfp4) + + add_executable(test_rocmfpx deps/llama.cpp/ggml/rocmfpx/test_rocmfpx.c) + target_link_libraries(test_rocmfpx PRIVATE ggml-base) + if(UNIX) + target_link_libraries(test_rocmfpx PRIVATE m) + endif() + add_test(NAME rocmfpx_reference COMMAND test_rocmfpx) + + if(DFLASH27B_GPU_BACKEND STREQUAL "hip") + add_executable(test_rocmfp4_hip_tail test/test_rocmfp4_hip_tail.cpp) + set_source_files_properties(test/test_rocmfp4_hip_tail.cpp PROPERTIES LANGUAGE HIP) + set_target_properties(test_rocmfp4_hip_tail PROPERTIES HIP_ARCHITECTURES "${_dflash_archs}") + target_include_directories(test_rocmfp4_hip_tail PRIVATE + ${CMAKE_CURRENT_SOURCE_DIR}/deps/llama.cpp/ggml/include + ${CMAKE_CURRENT_SOURCE_DIR}/deps/llama.cpp/ggml/rocmfp4) + target_link_libraries(test_rocmfp4_hip_tail PRIVATE + ggml + ggml-base + ${DFLASH27B_GGML_BACKEND_TARGET} + hip::host) + add_test(NAME rocmfp4_hip_tail COMMAND test_rocmfp4_hip_tail) + endif() + if(DFLASH27B_GPU_BACKEND STREQUAL "cuda" AND _dflash_cuda_min_sm GREATER_EQUAL 80 AND EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/test/test_flashprefill_kernels.cpp") add_executable(test_flashprefill_kernels test/test_flashprefill_kernels.cpp) set_target_properties(test_flashprefill_kernels PROPERTIES CUDA_ARCHITECTURES "${_dflash_archs}") diff --git a/server/deps/llama.cpp/ggml/include/ggml.h b/server/deps/llama.cpp/ggml/include/ggml.h index d76166f16..1d8f22bb8 100644 --- a/server/deps/llama.cpp/ggml/include/ggml.h +++ b/server/deps/llama.cpp/ggml/include/ggml.h @@ -430,7 +430,13 @@ extern "C" { GGML_TYPE_NVFP4 = 40, // NVFP4 (4 blocks, E4M3 scale) GGML_TYPE_Q1_0 = 41, GGML_TYPE_TQ3_0 = 42, // TurboQuant 3.5 bpv (3-bit Lloyd-Max + FWHT rotation) - GGML_TYPE_COUNT = 43, + GGML_TYPE_Q4_0_ROCMFP4 = 100, + GGML_TYPE_Q4_0_ROCMFP4_FAST = 101, + GGML_TYPE_Q6_0_ROCMFPX = 102, + GGML_TYPE_Q8_0_ROCMFPX = 103, + GGML_TYPE_Q3_0_ROCMFPX = 104, + GGML_TYPE_Q2_0_ROCMFP2 = 107, + GGML_TYPE_COUNT = 108, }; // precision @@ -468,6 +474,18 @@ extern "C" { GGML_FTYPE_MOSTLY_MXFP4 = 25, // except 1d tensors GGML_FTYPE_MOSTLY_NVFP4 = 26, // except 1d tensors GGML_FTYPE_MOSTLY_Q1_0 = 27, // except 1d tensors + GGML_FTYPE_MOSTLY_Q4_0_ROCMFP4 = 100, + GGML_FTYPE_MOSTLY_Q4_0_ROCMFP4_LEAN = 101, + GGML_FTYPE_MOSTLY_Q4_0_ROCMFP4_COHERENT = 102, + GGML_FTYPE_MOSTLY_Q4_0_ROCMFP4_FAST = 103, + GGML_FTYPE_MOSTLY_Q4_0_ROCMFP4_FAST_COHERENT = 104, + GGML_FTYPE_MOSTLY_Q4_0_ROCMFP4_STRIX = 105, + GGML_FTYPE_MOSTLY_Q4_0_ROCMFP4_STRIX_LEAN = 106, + GGML_FTYPE_MOSTLY_Q6_0_ROCMFPX = 110, + GGML_FTYPE_MOSTLY_Q8_0_ROCMFPX = 111, + GGML_FTYPE_MOSTLY_Q3_0_ROCMFPX = 112, + GGML_FTYPE_MOSTLY_Q2_0_ROCMFP2 = 118, + GGML_FTYPE_MOSTLY_Q2_0_ROCMFP2_STRIX = 119, }; // available tensor operations: @@ -583,6 +601,8 @@ extern "C" { GGML_OP_MOE_FUSED, // Fused MoE FFN: gate+up+swiglu+down+weighted_sum+shared_expert + GGML_OP_DS4_HC, // Fused DeepSeek4 hyper-connection pre/post/out mixing + GGML_OP_COUNT, }; @@ -618,6 +638,7 @@ extern "C" { GGML_GLU_OP_GEGLU, GGML_GLU_OP_SWIGLU, GGML_GLU_OP_SWIGLU_OAI, + GGML_GLU_OP_SWIGLU_DS4, GGML_GLU_OP_GEGLU_ERF, GGML_GLU_OP_GEGLU_QUICK, @@ -1339,6 +1360,12 @@ extern "C" { struct ggml_tensor * a, struct ggml_tensor * b); + GGML_API struct ggml_tensor * ggml_swiglu_ds4_split( + struct ggml_context * ctx, + struct ggml_tensor * gate, + struct ggml_tensor * up, + float clamp); + GGML_API struct ggml_tensor * ggml_geglu_erf_split( struct ggml_context * ctx, struct ggml_tensor * a, @@ -2399,6 +2426,37 @@ extern "C" { struct ggml_tensor * experts, struct ggml_tensor * expert_weights); + // Fused DeepSeek4 hyper-connection helpers (decode, n_tokens == 1). + // ggml_ds4_hc_pre: mix[2*n_hc+n_hc^2] + base + hc_state[n_embd*n_hc] -> + // dst[n_embd + 2*n_hc + n_hc^2] = { working, split(pre,post,comb) } + GGML_API struct ggml_tensor * ggml_ds4_hc_pre( + struct ggml_context * ctx, + struct ggml_tensor * mix, + struct ggml_tensor * base, + struct ggml_tensor * hc_state, + int n_hc, + int sinkhorn_iters, + float pre_scale, + float post_scale, + float comb_scale); + + // ggml_ds4_hc_post: residual hc_state + block_out + split -> new hc_state + GGML_API struct ggml_tensor * ggml_ds4_hc_post( + struct ggml_context * ctx, + struct ggml_tensor * residual_hc, + struct ggml_tensor * block_out, + struct ggml_tensor * split, + int n_hc); + + // ggml_ds4_hc_out: output-stage merge of hc streams into one embedding + GGML_API struct ggml_tensor * ggml_ds4_hc_out( + struct ggml_context * ctx, + struct ggml_tensor * mix, + struct ggml_tensor * base, + struct ggml_tensor * hc_state, + int n_hc, + float pre_scale); + // TODO: needs to be adapted to ggml_flash_attn_ext GGML_API struct ggml_tensor * ggml_flash_attn_back( struct ggml_context * ctx, diff --git a/server/deps/llama.cpp/ggml/rocmfp4/README.md b/server/deps/llama.cpp/ggml/rocmfp4/README.md new file mode 100644 index 000000000..4ae88297a --- /dev/null +++ b/server/deps/llama.cpp/ggml/rocmfp4/README.md @@ -0,0 +1,40 @@ +# ROCmFP4 + +This directory implements two GGUF weight formats used by the ROCm/HIP +backend: + +| GGML type | Block layout | Block size | Bits/weight | +|---|---|---:|---:| +| `Q4_0_ROCMFP4` | 32 packed 4-bit codes and one UE4M3 scale per 16 weights | 18 bytes | 4.50 | +| `Q4_0_ROCMFP4_FAST` | 32 packed 4-bit codes and one UE4M3 scale per 32 weights | 17 bytes | 4.25 | + +Both layouts use the signed Codebook10 levels +`0, +/-1, +/-2, +/-3, +/-4, +/-6, +/-8, +/-10`. Scale bytes are finite +unsigned UE4M3 values; validation rejects `0x7f` and values with the sign bit +set. + +## Implementation + +- `rocmfp4.c` contains deterministic CPU reference quantization, + dequantization, validation, and vector-dot functions. +- `rocmfp4_hip_codebook.cuh` and `rocmfp4_hip_scale.cuh` contain device + helpers shared by the ggml HIP kernels. +- Backend dispatch is integrated in `ggml-cuda`, which is also the source tree + used by ggml's HIP build. + +The quantizer selects UE4M3 scales by minimizing reconstruction error. When an +importance matrix is supplied, finite positive weights are used in the error +metric. Non-finite source values quantize to zero and non-finite or non-positive +importance weights are ignored. + +Vulkan kernels are not implemented in this vendored build. + +## Tests + +With `DFLASH27B_TESTS=ON`: + +- `test_rocmfp4` checks CPU reference round trips, validation, and imatrix + handling. +- `test_rocmfp4_hip_tail` compares HIP conversion with the CPU reference + byte-for-byte and checks output bounds for row sizes that are and are not + multiples of 256. diff --git a/server/deps/llama.cpp/ggml/rocmfp4/rocmfp4.c b/server/deps/llama.cpp/ggml/rocmfp4/rocmfp4.c new file mode 100644 index 000000000..b4b2cc285 --- /dev/null +++ b/server/deps/llama.cpp/ggml/rocmfp4/rocmfp4.c @@ -0,0 +1,839 @@ +#define GGML_COMMON_DECL_C +#include "../src/ggml-common.h" + +#include "rocmfp4.h" + +#include +#include +#include + +// ggml-base is compiled architecture-neutral (no -mavx2), so SIMD for the hot +// CPU dot product is enabled per-function via a target attribute plus a runtime +// CPU check. This keeps the AVX2 path in one translation unit without exporting +// internals to ggml-cpu. Non-GNU/non-x86 builds fall back to the scalar loop. +#if defined(__GNUC__) && (defined(__x86_64__) || defined(__i386__)) && !defined(ROCMFP4_NO_AVX2) +#include +#define ROCMFP4_X86_AVX2_DISPATCH 1 +#endif + +// ROCmFP4 stores a signed integer FP4-like codebook at half-scale. It is +// E2M1-derived, but the largest magnitude is retuned from 12 to 10 after +// sampling Qwen3 dense tensors; this reduces outlier pull without changing the +// packed 4-bit layout or integer dot-product path. +static const int8_t rocmfp4_codebook[16] = { + 0, 1, 2, 3, 4, 6, 8, 10, + 0, -1, -2, -3, -4, -6, -8,-10, +}; + +static inline int8_t rocmfp4_decode(uint8_t q) { + q &= 0x0f; + const int mag3 = q & 0x07; + const int mag = mag3 <= 4 ? mag3 : 2*mag3 - 4; + return (q & 0x08) ? -mag : mag; +} + +static inline int8_t rocmfp4_decode_table(uint8_t q) { + return rocmfp4_codebook[q & 0x0f]; +} + +// Finite unsigned E4M3 scale bytes decoded to the half-scale values used by +// ROCmFP4. Keeping this as a table avoids rebuilding identical FP32 values for +// every candidate during exhaustive scale search. +#define ROCMFP4_SCALE_SUB(M) ((M) * 0x1p-10f) +#define ROCMFP4_SCALE_E1(M) ((8 + (M)) * 0x1p-10f) +#define ROCMFP4_SCALE_E2(M) ((8 + (M)) * 0x1p-9f) +#define ROCMFP4_SCALE_E3(M) ((8 + (M)) * 0x1p-8f) +#define ROCMFP4_SCALE_E4(M) ((8 + (M)) * 0x1p-7f) +#define ROCMFP4_SCALE_E5(M) ((8 + (M)) * 0x1p-6f) +#define ROCMFP4_SCALE_E6(M) ((8 + (M)) * 0x1p-5f) +#define ROCMFP4_SCALE_E7(M) ((8 + (M)) * 0x1p-4f) +#define ROCMFP4_SCALE_E8(M) ((8 + (M)) * 0x1p-3f) +#define ROCMFP4_SCALE_E9(M) ((8 + (M)) * 0x1p-2f) +#define ROCMFP4_SCALE_E10(M) ((8 + (M)) * 0x1p-1f) +#define ROCMFP4_SCALE_E11(M) ((8 + (M)) * 0x1p0f) +#define ROCMFP4_SCALE_E12(M) ((8 + (M)) * 0x1p1f) +#define ROCMFP4_SCALE_E13(M) ((8 + (M)) * 0x1p2f) +#define ROCMFP4_SCALE_E14(M) ((8 + (M)) * 0x1p3f) +#define ROCMFP4_SCALE_E15(M) ((8 + (M)) * 0x1p4f) + +static const float rocmfp4_scale_ue4m3_half[127] = { + ROCMFP4_SCALE_SUB(0), ROCMFP4_SCALE_SUB(1), ROCMFP4_SCALE_SUB(2), ROCMFP4_SCALE_SUB(3), + ROCMFP4_SCALE_SUB(4), ROCMFP4_SCALE_SUB(5), ROCMFP4_SCALE_SUB(6), ROCMFP4_SCALE_SUB(7), + ROCMFP4_SCALE_E1(0), ROCMFP4_SCALE_E1(1), ROCMFP4_SCALE_E1(2), ROCMFP4_SCALE_E1(3), + ROCMFP4_SCALE_E1(4), ROCMFP4_SCALE_E1(5), ROCMFP4_SCALE_E1(6), ROCMFP4_SCALE_E1(7), + ROCMFP4_SCALE_E2(0), ROCMFP4_SCALE_E2(1), ROCMFP4_SCALE_E2(2), ROCMFP4_SCALE_E2(3), + ROCMFP4_SCALE_E2(4), ROCMFP4_SCALE_E2(5), ROCMFP4_SCALE_E2(6), ROCMFP4_SCALE_E2(7), + ROCMFP4_SCALE_E3(0), ROCMFP4_SCALE_E3(1), ROCMFP4_SCALE_E3(2), ROCMFP4_SCALE_E3(3), + ROCMFP4_SCALE_E3(4), ROCMFP4_SCALE_E3(5), ROCMFP4_SCALE_E3(6), ROCMFP4_SCALE_E3(7), + ROCMFP4_SCALE_E4(0), ROCMFP4_SCALE_E4(1), ROCMFP4_SCALE_E4(2), ROCMFP4_SCALE_E4(3), + ROCMFP4_SCALE_E4(4), ROCMFP4_SCALE_E4(5), ROCMFP4_SCALE_E4(6), ROCMFP4_SCALE_E4(7), + ROCMFP4_SCALE_E5(0), ROCMFP4_SCALE_E5(1), ROCMFP4_SCALE_E5(2), ROCMFP4_SCALE_E5(3), + ROCMFP4_SCALE_E5(4), ROCMFP4_SCALE_E5(5), ROCMFP4_SCALE_E5(6), ROCMFP4_SCALE_E5(7), + ROCMFP4_SCALE_E6(0), ROCMFP4_SCALE_E6(1), ROCMFP4_SCALE_E6(2), ROCMFP4_SCALE_E6(3), + ROCMFP4_SCALE_E6(4), ROCMFP4_SCALE_E6(5), ROCMFP4_SCALE_E6(6), ROCMFP4_SCALE_E6(7), + ROCMFP4_SCALE_E7(0), ROCMFP4_SCALE_E7(1), ROCMFP4_SCALE_E7(2), ROCMFP4_SCALE_E7(3), + ROCMFP4_SCALE_E7(4), ROCMFP4_SCALE_E7(5), ROCMFP4_SCALE_E7(6), ROCMFP4_SCALE_E7(7), + ROCMFP4_SCALE_E8(0), ROCMFP4_SCALE_E8(1), ROCMFP4_SCALE_E8(2), ROCMFP4_SCALE_E8(3), + ROCMFP4_SCALE_E8(4), ROCMFP4_SCALE_E8(5), ROCMFP4_SCALE_E8(6), ROCMFP4_SCALE_E8(7), + ROCMFP4_SCALE_E9(0), ROCMFP4_SCALE_E9(1), ROCMFP4_SCALE_E9(2), ROCMFP4_SCALE_E9(3), + ROCMFP4_SCALE_E9(4), ROCMFP4_SCALE_E9(5), ROCMFP4_SCALE_E9(6), ROCMFP4_SCALE_E9(7), + ROCMFP4_SCALE_E10(0), ROCMFP4_SCALE_E10(1), ROCMFP4_SCALE_E10(2), ROCMFP4_SCALE_E10(3), + ROCMFP4_SCALE_E10(4), ROCMFP4_SCALE_E10(5), ROCMFP4_SCALE_E10(6), ROCMFP4_SCALE_E10(7), + ROCMFP4_SCALE_E11(0), ROCMFP4_SCALE_E11(1), ROCMFP4_SCALE_E11(2), ROCMFP4_SCALE_E11(3), + ROCMFP4_SCALE_E11(4), ROCMFP4_SCALE_E11(5), ROCMFP4_SCALE_E11(6), ROCMFP4_SCALE_E11(7), + ROCMFP4_SCALE_E12(0), ROCMFP4_SCALE_E12(1), ROCMFP4_SCALE_E12(2), ROCMFP4_SCALE_E12(3), + ROCMFP4_SCALE_E12(4), ROCMFP4_SCALE_E12(5), ROCMFP4_SCALE_E12(6), ROCMFP4_SCALE_E12(7), + ROCMFP4_SCALE_E13(0), ROCMFP4_SCALE_E13(1), ROCMFP4_SCALE_E13(2), ROCMFP4_SCALE_E13(3), + ROCMFP4_SCALE_E13(4), ROCMFP4_SCALE_E13(5), ROCMFP4_SCALE_E13(6), ROCMFP4_SCALE_E13(7), + ROCMFP4_SCALE_E14(0), ROCMFP4_SCALE_E14(1), ROCMFP4_SCALE_E14(2), ROCMFP4_SCALE_E14(3), + ROCMFP4_SCALE_E14(4), ROCMFP4_SCALE_E14(5), ROCMFP4_SCALE_E14(6), ROCMFP4_SCALE_E14(7), + ROCMFP4_SCALE_E15(0), ROCMFP4_SCALE_E15(1), ROCMFP4_SCALE_E15(2), ROCMFP4_SCALE_E15(3), + ROCMFP4_SCALE_E15(4), ROCMFP4_SCALE_E15(5), ROCMFP4_SCALE_E15(6), +}; + +#undef ROCMFP4_SCALE_SUB +#undef ROCMFP4_SCALE_E1 +#undef ROCMFP4_SCALE_E2 +#undef ROCMFP4_SCALE_E3 +#undef ROCMFP4_SCALE_E4 +#undef ROCMFP4_SCALE_E5 +#undef ROCMFP4_SCALE_E6 +#undef ROCMFP4_SCALE_E7 +#undef ROCMFP4_SCALE_E8 +#undef ROCMFP4_SCALE_E9 +#undef ROCMFP4_SCALE_E10 +#undef ROCMFP4_SCALE_E11 +#undef ROCMFP4_SCALE_E12 +#undef ROCMFP4_SCALE_E13 +#undef ROCMFP4_SCALE_E14 +#undef ROCMFP4_SCALE_E15 + +static inline float rocmfp4_ue4m3_to_fp32_half(uint8_t e) { + return e <= 0x7e ? rocmfp4_scale_ue4m3_half[e] : 0.0f; +} + +static inline uint8_t rocmfp4_best_index_scaled_finite(float x, float inv_scale_half) { + // Exact nearest-neighbor thresholds for Codebook10: + // 0, +/-1, +/-2, +/-3, +/-4, +/-6, +/-8, +/-10 + // Ties intentionally choose the lower-magnitude code, matching the former + // linear scan because the positive codes and zero appear first. + const float a = fabsf(x * inv_scale_half); + if (a <= 0.5f) { + return 0; + } + + const bool neg = x < 0.0f; + if (a <= 1.5f) { + return neg ? 9 : 1; + } + if (a <= 2.5f) { + return neg ? 10 : 2; + } + if (a <= 3.5f) { + return neg ? 11 : 3; + } + if (a <= 5.0f) { + return neg ? 12 : 4; + } + if (a <= 7.0f) { + return neg ? 13 : 5; + } + if (a <= 9.0f) { + return neg ? 14 : 6; + } + + return neg ? 15 : 7; +} + +static inline uint8_t rocmfp4_best_index_scaled(float x, float inv_scale_half) { + if (!isfinite(x)) { + return 0; + } + + return rocmfp4_best_index_scaled_finite(x, inv_scale_half); +} + +// Fused best-index + decode used only inside the exhaustive scale search. The +// scale search re-scans every block element for every candidate scale byte, so +// avoiding the code -> decode round-trip on the hottest quantize path matters. +// Returns the same signed Codebook10 magnitude that +// rocmfp4_decode(rocmfp4_best_index_scaled_finite(x, inv_scale_half)) produces, +// so quantized output is bit-identical to the previous path. +static inline float rocmfp4_decoded_mag_scaled_finite(float x, float inv_scale_half) { + const float a = fabsf(x * inv_scale_half); + + float mag; + if (a <= 0.5f) { + mag = 0.0f; + } else if (a <= 1.5f) { + mag = 1.0f; + } else if (a <= 2.5f) { + mag = 2.0f; + } else if (a <= 3.5f) { + mag = 3.0f; + } else if (a <= 5.0f) { + mag = 4.0f; + } else if (a <= 7.0f) { + mag = 6.0f; + } else if (a <= 9.0f) { + mag = 8.0f; + } else { + mag = 10.0f; + } + + return x < 0.0f ? -mag : mag; +} + +static inline float rocmfp4_decoded_mag_scaled(float x, float inv_scale_half) { + if (!isfinite(x)) { + return 0.0f; + } + + return rocmfp4_decoded_mag_scaled_finite(x, inv_scale_half); +} + +static inline bool rocmfp4_scale_is_valid(uint8_t e) { + // ROCmFP4 scale bytes are unsigned finite E4M3 values. 0x7f is NaN in the + // unsigned encoding and values with the sign bit set are not valid scales. + return e <= 0x7e; +} + +static float rocmfp4_row_sigma2(const float * x, int64_t k) { + float sum_x2 = 0.0f; + bool overflow = false; + + for (int64_t i = 0; i < k; ++i) { + if (!isfinite(x[i])) { + continue; + } + + const float x2 = x[i]*x[i]; + if (!isfinite(x2) || sum_x2 > FLT_MAX - x2) { + overflow = true; + break; + } + sum_x2 += x2; + } + + if (!overflow) { + return sum_x2 / (float) k; + } + + double sum_x2_wide = 0.0; + for (int64_t i = 0; i < k; ++i) { + if (isfinite(x[i])) { + const double value = x[i]; + sum_x2_wide += value*value; + } + } + return (float) fmin(sum_x2_wide / (double) k, (double) FLT_MAX); +} + +static float rocmfp4_block_mse_for_scale_unweighted( + const float * x, int n, int e, float best_err) { + const float scale_half = rocmfp4_ue4m3_to_fp32_half((uint8_t) e); + const float inv_scale_half = 1.0f / scale_half; + float err = 0.0f; + + for (int i = 0; i < n; ++i) { + if (!isfinite(x[i])) { + continue; + } + + const float y = rocmfp4_decoded_mag_scaled(x[i], inv_scale_half) * scale_half; + const float d = x[i] - y; + + err += d*d; + if (err > best_err) { + return err; + } + } + + return err; +} + +static float rocmfp4_block_mse_for_scale_unweighted_finite( + const float * x, int n, int e, float best_err) { + const float scale_half = rocmfp4_ue4m3_to_fp32_half((uint8_t) e); + const float inv_scale_half = 1.0f / scale_half; + float err = 0.0f; + + for (int i = 0; i < n; ++i) { + const float y = rocmfp4_decoded_mag_scaled_finite(x[i], inv_scale_half) * scale_half; + const float d = x[i] - y; + + err += d*d; + if (err > best_err) { + return err; + } + } + + return err; +} + +static float rocmfp4_block_mse_for_scale_weighted( + const float * x, int n, const float * mse_weights, int e, float best_err) { + const float scale_half = rocmfp4_ue4m3_to_fp32_half((uint8_t) e); + const float inv_scale_half = 1.0f / scale_half; + float err = 0.0f; + + for (int i = 0; i < n; ++i) { + if (!isfinite(x[i])) { + continue; + } + + const float y = rocmfp4_decoded_mag_scaled(x[i], inv_scale_half) * scale_half; + const float d = x[i] - y; + + err += mse_weights[i]*d*d; + if (err > best_err) { + return err; + } + } + + return err; +} + +static float rocmfp4_block_mse_for_scale_weighted_finite( + const float * x, int n, const float * mse_weights, int e, float best_err) { + const float scale_half = rocmfp4_ue4m3_to_fp32_half((uint8_t) e); + const float inv_scale_half = 1.0f / scale_half; + float err = 0.0f; + + for (int i = 0; i < n; ++i) { + const float y = rocmfp4_decoded_mag_scaled_finite(x[i], inv_scale_half) * scale_half; + const float d = x[i] - y; + + err += mse_weights[i]*d*d; + if (err > best_err) { + return err; + } + } + + return err; +} + +static void rocmfp4_prepare_mse_weights( + float * dst, const float * x, int n, const float * quant_weights, float sigma2, + float * max_abs, float * max_abs_weight, bool * all_finite) { + *max_abs = 0.0f; + *max_abs_weight = 0.0f; + *all_finite = true; + + for (int i = 0; i < n; ++i) { + const float qw = quant_weights[i]; + const float ax = fabsf(x[i]); + float weight = 0.0f; + if (isfinite(qw) && qw > 0.0f && isfinite(x[i])) { + const float energy2 = sigma2 + x[i]*x[i]; + const float candidate = isfinite(energy2) ? qw*sqrtf(energy2) : FLT_MAX; + weight = isfinite(candidate) ? candidate : FLT_MAX; + } + *all_finite = *all_finite && isfinite(x[i]); + + if (isfinite(x[i])) { + if (ax > *max_abs) { + *max_abs = ax; + *max_abs_weight = weight; + } else if (ax == *max_abs && weight > *max_abs_weight) { + *max_abs_weight = weight; + } + } + + // Match llama.cpp's imatrix weighting style for Q4_0: calibration + // importance is scaled by row energy so large activations remain protected. + dst[i] = weight; + } +} + +static int rocmfp4_nearest_scale_ue4m3(float target_scale_half) { + if (!(target_scale_half > 0.0f) || !isfinite(target_scale_half)) { + return 1; + } + + int lo = 1; + int hi = 126; + while (lo < hi) { + const int mid = lo + (hi - lo) / 2; + if (rocmfp4_ue4m3_to_fp32_half((uint8_t) mid) < target_scale_half) { + lo = mid + 1; + } else { + hi = mid; + } + } + + if (lo == 1) { + return 1; + } + + const float hi_scale = rocmfp4_ue4m3_to_fp32_half((uint8_t) lo); + const float lo_scale = rocmfp4_ue4m3_to_fp32_half((uint8_t) (lo - 1)); + + // Match the former ascending nearest scan: exact midpoint ties keep the + // lower scale byte. + return (target_scale_half - lo_scale <= hi_scale - target_scale_half) ? lo - 1 : lo; +} + +static uint8_t rocmfp4_choose_scale_ue4m3_exhaustive_unweighted( + const float * x, int n, float max_abs, bool all_finite) { + const int start_e = rocmfp4_nearest_scale_ue4m3(max_abs / 10.0f); + + int best_e = 0; + float best_err = FLT_MAX; + bool lower_done = false; + + for (int delta = 0; delta <= 125; ++delta) { + const int e0 = start_e - delta; + if (!lower_done && e0 >= 1 && e0 <= 126) { + const float scale_half = rocmfp4_ue4m3_to_fp32_half((uint8_t) e0); + const float clip_delta = max_abs - 10.0f*scale_half; + if (clip_delta > 0.0f && clip_delta*clip_delta > best_err) { + lower_done = true; + } else { + const float err = all_finite ? + rocmfp4_block_mse_for_scale_unweighted_finite(x, n, e0, best_err) : + rocmfp4_block_mse_for_scale_unweighted(x, n, e0, best_err); + if (err < best_err || (err == best_err && e0 < best_e)) { + best_err = err; + best_e = e0; + } + } + } + + const int e1 = start_e + delta; + if (delta != 0 && e1 >= 1 && e1 <= 126) { + const float err = all_finite ? + rocmfp4_block_mse_for_scale_unweighted_finite(x, n, e1, best_err) : + rocmfp4_block_mse_for_scale_unweighted(x, n, e1, best_err); + if (err < best_err || (err == best_err && e1 < best_e)) { + best_err = err; + best_e = e1; + } + } + + if ((lower_done || e0 <= 1) && e1 >= 126) { + break; + } + } + + return (uint8_t) best_e; +} + +static uint8_t rocmfp4_choose_scale_ue4m3_exhaustive_weighted( + const float * x, int n, const float * mse_weights, float max_abs, float max_abs_weight, bool all_finite) { + const int start_e = rocmfp4_nearest_scale_ue4m3(max_abs / 10.0f); + + int best_e = 0; + float best_err = FLT_MAX; + bool lower_done = false; + + for (int delta = 0; delta <= 125; ++delta) { + const int e0 = start_e - delta; + if (!lower_done && e0 >= 1 && e0 <= 126) { + const float scale_half = rocmfp4_ue4m3_to_fp32_half((uint8_t) e0); + const float clip_delta = max_abs - 10.0f*scale_half; + if (max_abs_weight > 0.0f && clip_delta > 0.0f && max_abs_weight*clip_delta*clip_delta > best_err) { + lower_done = true; + } else { + const float err = all_finite ? + rocmfp4_block_mse_for_scale_weighted_finite(x, n, mse_weights, e0, best_err) : + rocmfp4_block_mse_for_scale_weighted(x, n, mse_weights, e0, best_err); + if (err < best_err || (err == best_err && e0 < best_e)) { + best_err = err; + best_e = e0; + } + } + } + + const int e1 = start_e + delta; + if (delta != 0 && e1 >= 1 && e1 <= 126) { + const float err = all_finite ? + rocmfp4_block_mse_for_scale_weighted_finite(x, n, mse_weights, e1, best_err) : + rocmfp4_block_mse_for_scale_weighted(x, n, mse_weights, e1, best_err); + if (err < best_err || (err == best_err && e1 < best_e)) { + best_err = err; + best_e = e1; + } + } + + if ((lower_done || e0 <= 1) && e1 >= 126) { + break; + } + } + + return (uint8_t) best_e; +} + +static uint8_t rocmfp4_choose_scale_ue4m3(const float * x, int n, const float * quant_weights, float sigma2) { + if (quant_weights) { + assert(n <= QK_ROCMFP4); + float mse_weights_buf[QK_ROCMFP4]; + float weighted_max_abs; + float max_abs_weight; + bool all_finite; + rocmfp4_prepare_mse_weights(mse_weights_buf, x, n, quant_weights, sigma2, &weighted_max_abs, &max_abs_weight, &all_finite); + if (!(weighted_max_abs > 0.0f) || !isfinite(weighted_max_abs)) { + return 0; + } + return rocmfp4_choose_scale_ue4m3_exhaustive_weighted(x, n, mse_weights_buf, weighted_max_abs, max_abs_weight, all_finite); + } + + float max_abs = 0.0f; + bool all_finite = true; + for (int i = 0; i < n; ++i) { + all_finite = all_finite && isfinite(x[i]); + if (!isfinite(x[i])) { + continue; + } + + const float ax = fabsf(x[i]); + if (ax > max_abs) { + max_abs = ax; + } + } + + if (!(max_abs > 0.0f) || !isfinite(max_abs)) { + return 0; + } + + return rocmfp4_choose_scale_ue4m3_exhaustive_unweighted(x, n, max_abs, all_finite); +} + +static void rocmfp4_quantize_row_q4_0_weighted( + const float * GGML_RESTRICT x, block_rocmfp4 * GGML_RESTRICT y, int64_t k, const float * GGML_RESTRICT quant_weights) { + assert(k % QK_ROCMFP4 == 0); + + const float sigma2 = rocmfp4_row_sigma2(x, k); + + const int64_t nb = k / QK_ROCMFP4; + for (int64_t ib = 0; ib < nb; ++ib) { + const float * xb = x + ib*QK_ROCMFP4; + const float * qw = quant_weights ? quant_weights + ib*QK_ROCMFP4 : NULL; + const uint8_t e0 = rocmfp4_choose_scale_ue4m3(xb, QK_ROCMFP4/2, qw, sigma2); + const uint8_t e1 = rocmfp4_choose_scale_ue4m3(xb + QK_ROCMFP4/2, QK_ROCMFP4/2, qw ? qw + QK_ROCMFP4/2 : NULL, sigma2); + const float scale_half0 = rocmfp4_ue4m3_to_fp32_half(e0); + const float scale_half1 = rocmfp4_ue4m3_to_fp32_half(e1); + const float inv_scale_half0 = scale_half0 > 0.0f ? 1.0f / scale_half0 : 0.0f; + const float inv_scale_half1 = scale_half1 > 0.0f ? 1.0f / scale_half1 : 0.0f; + + y[ib].e[0] = e0; + y[ib].e[1] = e1; + + for (int j = 0; j < QK_ROCMFP4/2; ++j) { + const uint8_t q0 = rocmfp4_best_index_scaled(xb[j], inv_scale_half0); + const uint8_t q1 = rocmfp4_best_index_scaled(xb[j + QK_ROCMFP4/2], inv_scale_half1); + y[ib].qs[j] = q0 | (q1 << 4); + } + } +} + +static void rocmfp4_quantize_row_q4_0_fast_weighted( + const float * GGML_RESTRICT x, block_rocmfp4_fast * GGML_RESTRICT y, int64_t k, const float * GGML_RESTRICT quant_weights) { + assert(k % QK_ROCMFP4 == 0); + + const float sigma2 = rocmfp4_row_sigma2(x, k); + + const int64_t nb = k / QK_ROCMFP4; + for (int64_t ib = 0; ib < nb; ++ib) { + const float * xb = x + ib*QK_ROCMFP4; + const float * qw = quant_weights ? quant_weights + ib*QK_ROCMFP4 : NULL; + const uint8_t e = rocmfp4_choose_scale_ue4m3(xb, QK_ROCMFP4, qw, sigma2); + const float scale_half = rocmfp4_ue4m3_to_fp32_half(e); + const float inv_scale_half = scale_half > 0.0f ? 1.0f / scale_half : 0.0f; + + y[ib].e = e; + + for (int j = 0; j < QK_ROCMFP4/2; ++j) { + const uint8_t q0 = rocmfp4_best_index_scaled(xb[j], inv_scale_half); + const uint8_t q1 = rocmfp4_best_index_scaled(xb[j + QK_ROCMFP4/2], inv_scale_half); + y[ib].qs[j] = q0 | (q1 << 4); + } + } +} + +void rocmfp4_quantize_row_q4_0_ref(const float * GGML_RESTRICT x, block_rocmfp4 * GGML_RESTRICT y, int64_t k) { + assert(k % QK_ROCMFP4 == 0); + + const int64_t nb = k / QK_ROCMFP4; + for (int64_t ib = 0; ib < nb; ++ib) { + const float * xb = x + ib*QK_ROCMFP4; + const uint8_t e0 = rocmfp4_choose_scale_ue4m3(xb, QK_ROCMFP4/2, NULL, 0.0f); + const uint8_t e1 = rocmfp4_choose_scale_ue4m3(xb + QK_ROCMFP4/2, QK_ROCMFP4/2, NULL, 0.0f); + const float scale_half0 = rocmfp4_ue4m3_to_fp32_half(e0); + const float scale_half1 = rocmfp4_ue4m3_to_fp32_half(e1); + const float inv_scale_half0 = scale_half0 > 0.0f ? 1.0f / scale_half0 : 0.0f; + const float inv_scale_half1 = scale_half1 > 0.0f ? 1.0f / scale_half1 : 0.0f; + + y[ib].e[0] = e0; + y[ib].e[1] = e1; + + for (int j = 0; j < QK_ROCMFP4/2; ++j) { + const uint8_t q0 = rocmfp4_best_index_scaled(xb[j], inv_scale_half0); + const uint8_t q1 = rocmfp4_best_index_scaled(xb[j + QK_ROCMFP4/2], inv_scale_half1); + y[ib].qs[j] = q0 | (q1 << 4); + } + } +} + +void rocmfp4_quantize_row_q4_0_fast_ref(const float * GGML_RESTRICT x, block_rocmfp4_fast * GGML_RESTRICT y, int64_t k) { + assert(k % QK_ROCMFP4 == 0); + + const int64_t nb = k / QK_ROCMFP4; + for (int64_t ib = 0; ib < nb; ++ib) { + const float * xb = x + ib*QK_ROCMFP4; + const uint8_t e = rocmfp4_choose_scale_ue4m3(xb, QK_ROCMFP4, NULL, 0.0f); + const float scale_half = rocmfp4_ue4m3_to_fp32_half(e); + const float inv_scale_half = scale_half > 0.0f ? 1.0f / scale_half : 0.0f; + + y[ib].e = e; + + for (int j = 0; j < QK_ROCMFP4/2; ++j) { + const uint8_t q0 = rocmfp4_best_index_scaled(xb[j], inv_scale_half); + const uint8_t q1 = rocmfp4_best_index_scaled(xb[j + QK_ROCMFP4/2], inv_scale_half); + y[ib].qs[j] = q0 | (q1 << 4); + } + } +} + +void rocmfp4_dequantize_row_q4_0(const block_rocmfp4 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { + assert(k % QK_ROCMFP4 == 0); + + const int64_t nb = k / QK_ROCMFP4; + for (int64_t ib = 0; ib < nb; ++ib) { + const float d0 = rocmfp4_ue4m3_to_fp32_half(x[ib].e[0]); + const float d1 = rocmfp4_ue4m3_to_fp32_half(x[ib].e[1]); + + for (int j = 0; j < QK_ROCMFP4/2; ++j) { + y[ib*QK_ROCMFP4 + j] = (float) rocmfp4_decode(x[ib].qs[j] & 0x0f) * d0; + y[ib*QK_ROCMFP4 + j + QK_ROCMFP4/2] = (float) rocmfp4_decode(x[ib].qs[j] >> 4) * d1; + } + } +} + +void rocmfp4_dequantize_row_q4_0_fast(const block_rocmfp4_fast * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { + assert(k % QK_ROCMFP4 == 0); + + const int64_t nb = k / QK_ROCMFP4; + for (int64_t ib = 0; ib < nb; ++ib) { + const float d = rocmfp4_ue4m3_to_fp32_half(x[ib].e); + + for (int j = 0; j < QK_ROCMFP4/2; ++j) { + y[ib*QK_ROCMFP4 + j] = (float) rocmfp4_decode(x[ib].qs[j] & 0x0f) * d; + y[ib*QK_ROCMFP4 + j + QK_ROCMFP4/2] = (float) rocmfp4_decode(x[ib].qs[j] >> 4) * d; + } + } +} + +void rocmfp4_quantize_row_q4_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { + rocmfp4_quantize_row_q4_0_ref(x, (block_rocmfp4 *) y, k); +} + +void rocmfp4_quantize_row_q4_0_fast(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { + rocmfp4_quantize_row_q4_0_fast_ref(x, (block_rocmfp4_fast *) y, k); +} + +size_t rocmfp4_quantize_q4_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix) { + const size_t row_size = ggml_row_size(GGML_TYPE_Q4_0_ROCMFP4, n_per_row); + + if (!imatrix) { + rocmfp4_quantize_row_q4_0_ref(src, (block_rocmfp4 *) dst, nrows*n_per_row); + return nrows * row_size; + } + + char * qrow = (char *) dst; + for (int64_t row = 0; row < nrows; ++row) { + rocmfp4_quantize_row_q4_0_weighted(src, (block_rocmfp4 *) qrow, n_per_row, imatrix); + src += n_per_row; + qrow += row_size; + } + + return nrows * row_size; +} + +size_t rocmfp4_quantize_q4_0_fast(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix) { + const size_t row_size = ggml_row_size(GGML_TYPE_Q4_0_ROCMFP4_FAST, n_per_row); + + if (!imatrix) { + rocmfp4_quantize_row_q4_0_fast_ref(src, (block_rocmfp4_fast *) dst, nrows*n_per_row); + return nrows * row_size; + } + + char * qrow = (char *) dst; + for (int64_t row = 0; row < nrows; ++row) { + rocmfp4_quantize_row_q4_0_fast_weighted(src, (block_rocmfp4_fast *) qrow, n_per_row, imatrix); + src += n_per_row; + qrow += row_size; + } + + return nrows * row_size; +} + +bool rocmfp4_validate_row_data(const void * data, size_t nbytes) { + if (nbytes % sizeof(block_rocmfp4) != 0) { + return false; + } + + const block_rocmfp4 * blocks = (const block_rocmfp4 *) data; + const size_t nblocks = nbytes / sizeof(block_rocmfp4); + for (size_t i = 0; i < nblocks; ++i) { + if (!rocmfp4_scale_is_valid(blocks[i].e[0]) || !rocmfp4_scale_is_valid(blocks[i].e[1])) { + return false; + } + } + + return true; +} + +bool rocmfp4_validate_row_data_fast(const void * data, size_t nbytes) { + if (nbytes % sizeof(block_rocmfp4_fast) != 0) { + return false; + } + + const block_rocmfp4_fast * blocks = (const block_rocmfp4_fast *) data; + const size_t nblocks = nbytes / sizeof(block_rocmfp4_fast); + for (size_t i = 0; i < nblocks; ++i) { + if (!rocmfp4_scale_is_valid(blocks[i].e)) { + return false; + } + } + + return true; +} + +#ifdef ROCMFP4_X86_AVX2_DISPATCH +__attribute__((target("avx2"))) +static inline int rocmfp4_hsum_i32_8_avx2(__m256i v) { + __m128i s = _mm_add_epi32(_mm256_castsi256_si128(v), _mm256_extracti128_si256(v, 1)); + s = _mm_add_epi32(s, _mm_shuffle_epi32(s, _MM_SHUFFLE(1, 0, 3, 2))); + s = _mm_add_epi32(s, _mm_shuffle_epi32(s, _MM_SHUFFLE(2, 3, 0, 1))); + return _mm_cvtsi128_si32(s); +} + +// Decode one 32-weight block's low and high nibble streams through the +// Codebook10 table with a single PSHUFB, then integer-dot each against its half +// of the q8_0 block. Integer sums are order-independent, so sumi0/sumi1 match +// the scalar reference exactly and the float result is bit-identical. +__attribute__((target("avx2"))) +static inline void rocmfp4_block_isums_avx2( + const uint8_t * qs, const int8_t * q8, int * sumi0, int * sumi1) { + const __m128i tbl = _mm_loadu_si128((const __m128i *) rocmfp4_codebook); + const __m128i q = _mm_loadu_si128((const __m128i *) qs); + const __m128i lo = _mm_and_si128(q, _mm_set1_epi8(0x0F)); + const __m128i hi = _mm_and_si128(_mm_srli_epi16(q, 4), _mm_set1_epi8(0x0F)); + const __m128i dlo = _mm_shuffle_epi8(tbl, lo); + const __m128i dhi = _mm_shuffle_epi8(tbl, hi); + const __m128i ylo = _mm_loadu_si128((const __m128i *) q8); + const __m128i yhi = _mm_loadu_si128((const __m128i *) (q8 + QK_ROCMFP4/2)); + const __m256i pl = _mm256_madd_epi16(_mm256_cvtepi8_epi16(dlo), _mm256_cvtepi8_epi16(ylo)); + const __m256i ph = _mm256_madd_epi16(_mm256_cvtepi8_epi16(dhi), _mm256_cvtepi8_epi16(yhi)); + *sumi0 = rocmfp4_hsum_i32_8_avx2(pl); + *sumi1 = rocmfp4_hsum_i32_8_avx2(ph); +} + +__attribute__((target("avx2"))) +static void rocmfp4_vec_dot_q4_0_q8_0_avx2( + int nb, float * GGML_RESTRICT s, const block_rocmfp4 * GGML_RESTRICT x, const block_q8_0 * GGML_RESTRICT y) { + float sumf = 0.0f; + for (int ib = 0; ib < nb; ++ib) { + const float d0 = rocmfp4_ue4m3_to_fp32_half(x[ib].e[0]) * ggml_fp16_to_fp32(y[ib].d); + const float d1 = rocmfp4_ue4m3_to_fp32_half(x[ib].e[1]) * ggml_fp16_to_fp32(y[ib].d); + int sumi0, sumi1; + rocmfp4_block_isums_avx2(x[ib].qs, y[ib].qs, &sumi0, &sumi1); + sumf += d0 * (float) sumi0 + d1 * (float) sumi1; + } + *s = sumf; +} + +__attribute__((target("avx2"))) +static void rocmfp4_vec_dot_q4_0_fast_q8_0_avx2( + int nb, float * GGML_RESTRICT s, const block_rocmfp4_fast * GGML_RESTRICT x, const block_q8_0 * GGML_RESTRICT y) { + float sumf = 0.0f; + for (int ib = 0; ib < nb; ++ib) { + const float d = rocmfp4_ue4m3_to_fp32_half(x[ib].e) * ggml_fp16_to_fp32(y[ib].d); + int sumi0, sumi1; + rocmfp4_block_isums_avx2(x[ib].qs, y[ib].qs, &sumi0, &sumi1); + sumf += d * (float) (sumi0 + sumi1); + } + *s = sumf; +} +#endif // ROCMFP4_X86_AVX2_DISPATCH + +void rocmfp4_vec_dot_q4_0_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc) { + GGML_UNUSED(bs); + GGML_UNUSED(bx); + GGML_UNUSED(by); + assert(nrc == 1); + GGML_UNUSED(nrc); + assert(n % QK_ROCMFP4 == 0); + assert(QK_ROCMFP4 == QK8_0); + + const block_rocmfp4 * GGML_RESTRICT x = (const block_rocmfp4 *) vx; + const block_q8_0 * GGML_RESTRICT y = (const block_q8_0 *) vy; + + const int nb = n / QK_ROCMFP4; + +#ifdef ROCMFP4_X86_AVX2_DISPATCH + if (__builtin_cpu_supports("avx2")) { + rocmfp4_vec_dot_q4_0_q8_0_avx2(nb, s, x, y); + return; + } +#endif + + float sumf = 0.0f; + + for (int ib = 0; ib < nb; ++ib) { + const float d0 = rocmfp4_ue4m3_to_fp32_half(x[ib].e[0]) * ggml_fp16_to_fp32(y[ib].d); + const float d1 = rocmfp4_ue4m3_to_fp32_half(x[ib].e[1]) * ggml_fp16_to_fp32(y[ib].d); + int sumi0 = 0; + int sumi1 = 0; + + for (int j = 0; j < QK_ROCMFP4/2; ++j) { + const uint8_t q = x[ib].qs[j]; + sumi0 += rocmfp4_decode_table(q) * y[ib].qs[j]; + sumi1 += rocmfp4_decode_table(q >> 4) * y[ib].qs[j + QK_ROCMFP4/2]; + } + + sumf += d0 * (float) sumi0 + d1 * (float) sumi1; + } + + *s = sumf; +} + +void rocmfp4_vec_dot_q4_0_fast_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc) { + GGML_UNUSED(bs); + GGML_UNUSED(bx); + GGML_UNUSED(by); + assert(nrc == 1); + GGML_UNUSED(nrc); + assert(n % QK_ROCMFP4 == 0); + assert(QK_ROCMFP4 == QK8_0); + + const block_rocmfp4_fast * GGML_RESTRICT x = (const block_rocmfp4_fast *) vx; + const block_q8_0 * GGML_RESTRICT y = (const block_q8_0 *) vy; + + const int nb = n / QK_ROCMFP4; + +#ifdef ROCMFP4_X86_AVX2_DISPATCH + if (__builtin_cpu_supports("avx2")) { + rocmfp4_vec_dot_q4_0_fast_q8_0_avx2(nb, s, x, y); + return; + } +#endif + + float sumf = 0.0f; + + for (int ib = 0; ib < nb; ++ib) { + const float d = rocmfp4_ue4m3_to_fp32_half(x[ib].e) * ggml_fp16_to_fp32(y[ib].d); + int sumi = 0; + + for (int j = 0; j < QK_ROCMFP4/2; ++j) { + const uint8_t q = x[ib].qs[j]; + sumi += rocmfp4_decode_table(q) * y[ib].qs[j]; + sumi += rocmfp4_decode_table(q >> 4) * y[ib].qs[j + QK_ROCMFP4/2]; + } + + sumf += d * (float) sumi; + } + + *s = sumf; +} diff --git a/server/deps/llama.cpp/ggml/rocmfp4/rocmfp4.h b/server/deps/llama.cpp/ggml/rocmfp4/rocmfp4.h new file mode 100644 index 000000000..9756f6ad4 --- /dev/null +++ b/server/deps/llama.cpp/ggml/rocmfp4/rocmfp4.h @@ -0,0 +1,58 @@ +#pragma once + +#include +#include +#include + +#include "ggml.h" + +#ifdef __cplusplus +extern "C" { +#endif + +#define QK_ROCMFP4 32 +#define QR_ROCMFP4 2 +#define QI_ROCMFP4 (QK_ROCMFP4 / (4 * QR_ROCMFP4)) +#define QS_ROCMFP4 32 + +// AMD-tuned compact layout: 16 bytes of packed E2M1-derived 4-bit codes, then +// one unsigned E4M3 scale byte per 16-weight half block. +typedef struct { + uint8_t qs[QK_ROCMFP4/2]; + uint8_t e[2]; +} block_rocmfp4; + +// Speed-focused layout: same 32 packed ROCmFP4 nibbles, but one UE4M3 scale +// for the whole block. This is a separate GGUF type so fast 4.25 BPW artifacts +// never alias the safer dual-scale format above. +typedef struct { + uint8_t qs[QK_ROCMFP4/2]; + uint8_t e; +} block_rocmfp4_fast; + +#if defined(__cplusplus) +static_assert(sizeof(block_rocmfp4) == QK_ROCMFP4/2 + 2*sizeof(uint8_t), "wrong rocmfp4 block size/padding"); +static_assert(sizeof(block_rocmfp4_fast) == QK_ROCMFP4/2 + sizeof(uint8_t), "wrong rocmfp4 fast block size/padding"); +#else +_Static_assert(sizeof(block_rocmfp4) == QK_ROCMFP4/2 + 2*sizeof(uint8_t), "wrong rocmfp4 block size/padding"); +_Static_assert(sizeof(block_rocmfp4_fast) == QK_ROCMFP4/2 + sizeof(uint8_t), "wrong rocmfp4 fast block size/padding"); +#endif + +GGML_API void rocmfp4_quantize_row_q4_0_ref(const float * GGML_RESTRICT x, block_rocmfp4 * GGML_RESTRICT y, int64_t k); +GGML_API void rocmfp4_dequantize_row_q4_0(const block_rocmfp4 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API void rocmfp4_quantize_row_q4_0_fast_ref(const float * GGML_RESTRICT x, block_rocmfp4_fast * GGML_RESTRICT y, int64_t k); +GGML_API void rocmfp4_dequantize_row_q4_0_fast(const block_rocmfp4_fast * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); + +GGML_API void rocmfp4_quantize_row_q4_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +GGML_API size_t rocmfp4_quantize_q4_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); +GGML_API void rocmfp4_quantize_row_q4_0_fast(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +GGML_API size_t rocmfp4_quantize_q4_0_fast(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); +GGML_API bool rocmfp4_validate_row_data(const void * data, size_t nbytes); +GGML_API bool rocmfp4_validate_row_data_fast(const void * data, size_t nbytes); + +GGML_API void rocmfp4_vec_dot_q4_0_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); +GGML_API void rocmfp4_vec_dot_q4_0_fast_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); + +#ifdef __cplusplus +} +#endif diff --git a/server/deps/llama.cpp/ggml/rocmfp4/rocmfp4_hip_codebook.cuh b/server/deps/llama.cpp/ggml/rocmfp4/rocmfp4_hip_codebook.cuh new file mode 100644 index 000000000..b5dea2e11 --- /dev/null +++ b/server/deps/llama.cpp/ggml/rocmfp4/rocmfp4_hip_codebook.cuh @@ -0,0 +1,101 @@ +#pragma once + +#include "rocmfp4_hip_scale.cuh" + +#include +#include + +#ifndef GGML_ROCMFP4_UNALIGNED_QS_DWORD_LOAD +#define GGML_ROCMFP4_UNALIGNED_QS_DWORD_LOAD 1 +#endif + +#if !defined(GGML_USE_HIP) +// Self-contained 16-entry table expander for the non-HIP fallback. This header +// is pulled into translation units that do not include vecdotq.cuh (where the +// generic get_int_from_table_16 lives) - e.g. fattn-chunked.cu - so relying on +// that symbol breaks the CUDA build on include order. This is the generic +// branch of get_int_from_table_16 verbatim, so results are bit-identical; only +// the HIP path (Strix) uses __builtin_amdgcn_perm and never reaches here. +static __device__ __forceinline__ int2 rocmfp4_table16_fallback(const int & q4, const int8_t * table) { + const int q0_32 = (q4 >> 0) & 0x0F0F0F0F; + const int8_t * q0_8 = (const int8_t *) &q0_32; + const char4 val0_8 = make_char4( + table[q0_8[0]], table[q0_8[1]], table[q0_8[2]], table[q0_8[3]]); + + const int q1_32 = (q4 >> 4) & 0x0F0F0F0F; + const int8_t * q1_8 = (const int8_t *) &q1_32; + const char4 val1_8 = make_char4( + table[q1_8[0]], table[q1_8[1]], table[q1_8[2]], table[q1_8[3]]); + + return make_int2(*((const int *) &val0_8), *((const int *) &val1_8)); +} +#endif + +static __device__ __forceinline__ int rocmfp4_get_qs_i32(const void * x, const int & i32) { +#if defined(GGML_USE_HIP) && GGML_ROCMFP4_UNALIGNED_QS_DWORD_LOAD + return *((const int *) ((const uint8_t *) x + 4*i32)); +#else + const uint8_t * x8 = (const uint8_t *) x; + + int x32 = x8[4*i32 + 0] << 0; + x32 |= x8[4*i32 + 1] << 8; + x32 |= x8[4*i32 + 2] << 16; + x32 |= x8[4*i32 + 3] << 24; + + return x32; +#endif +} + +// AMD-specific fast path for expanding eight packed ROCmFP4 nibbles into two +// int32 DP4A operands. This encodes the Codebook10 table directly as four +// 32-bit constants: +// [0, 1, 2, 3], [4, 6, 8, 10], [0, -1, -2, -3], [-4, -6, -8, -10] +// Avoiding the table pointer keeps the ROCm/HIP MMVQ/MMQ hot path fully local +// to this format. Non-HIP builds still use llama.cpp's generic table expander. +static __device__ __forceinline__ int2 rocmfp4_get_int_from_codebook_16(const int & q4, const int8_t * fallback_table) { +#if defined(GGML_USE_HIP) + constexpr uint32_t values0 = 0x03020100u; + constexpr uint32_t values1 = 0x0a080604u; + constexpr uint32_t values2 = 0xfdfeff00u; + constexpr uint32_t values3 = 0xf6f8fafcu; + + const uint32_t q_even = q4; + const uint32_t q_odd = q4 >> 4; + + const uint32_t v_even_low = __builtin_amdgcn_perm(values1, values0, q_even & 0x07070707u); + const uint32_t v_odd_low = __builtin_amdgcn_perm(values1, values0, q_odd & 0x07070707u); + const uint32_t v_even_high = __builtin_amdgcn_perm(values3, values2, q_even & 0x07070707u); + const uint32_t v_odd_high = __builtin_amdgcn_perm(values3, values2, q_odd & 0x07070707u); + + const uint32_t mask_even = 0x03020100u | ((q_even & 0x08080808u) >> 1); + const uint32_t mask_odd = 0x03020100u | ((q_odd & 0x08080808u) >> 1); + + return make_int2( + __builtin_amdgcn_perm(v_even_high, v_even_low, mask_even), + __builtin_amdgcn_perm(v_odd_high, v_odd_low, mask_odd)); +#else + return rocmfp4_table16_fallback(q4, fallback_table); +#endif +} + +// Variant for call sites that already selected either the low or high nibble +// stream and only need one DP4A operand. This avoids the extra odd/even table +// expansion work in ROCmFP4 FlashAttention K/V decode. +static __device__ __forceinline__ int rocmfp4_get_low_int_from_codebook_16(const int & q4, const int8_t * fallback_table) { +#if defined(GGML_USE_HIP) + constexpr uint32_t values0 = 0x03020100u; + constexpr uint32_t values1 = 0x0a080604u; + constexpr uint32_t values2 = 0xfdfeff00u; + constexpr uint32_t values3 = 0xf6f8fafcu; + + const uint32_t q = q4; + + const uint32_t v_low = __builtin_amdgcn_perm(values1, values0, q & 0x07070707u); + const uint32_t v_high = __builtin_amdgcn_perm(values3, values2, q & 0x07070707u); + const uint32_t mask = 0x03020100u | ((q & 0x08080808u) >> 1); + + return __builtin_amdgcn_perm(v_high, v_low, mask); +#else + return rocmfp4_table16_fallback(q4, fallback_table).x; +#endif +} diff --git a/server/deps/llama.cpp/ggml/rocmfp4/rocmfp4_hip_scale.cuh b/server/deps/llama.cpp/ggml/rocmfp4/rocmfp4_hip_scale.cuh new file mode 100644 index 000000000..11acc9972 --- /dev/null +++ b/server/deps/llama.cpp/ggml/rocmfp4/rocmfp4_hip_scale.cuh @@ -0,0 +1,78 @@ +#pragma once + +#include +#include +#include + +static __device__ __forceinline__ float rocmfp4_u32_as_f32(uint32_t bits) { +#if defined(GGML_USE_HIP) + return __uint_as_float(bits); +#else + float result; + memcpy(&result, &bits, sizeof(float)); + return result; +#endif +} + +// ROCmFP4 validates scale bytes before backend execution, so HIP/ROCm hot +// paths can decode finite unsigned E4M3 half-scales directly without the +// generic FP8 NaN handling used by other formats. +static __device__ __forceinline__ float rocmfp4_ue4m3_to_fp32_half_finite(uint8_t x) { + const int exp = (x >> 3) & 0xF; + const int man = x & 0x7; + + if (exp == 0) { + return (float) man * (1.0f / 1024.0f); + } + + const uint32_t bits = ((uint32_t) exp + 119u) << 23 | ((uint32_t) man << 20); + return rocmfp4_u32_as_f32(bits); +} + +static __device__ __forceinline__ float rocmfpx_ue4m3_to_fp32_finite(uint8_t x) { + if (x > 0x7e) { + return 0.0f; + } + + const int exp = (x >> 3) & 0xF; + const int man = x & 0x7; + + if (exp == 0) { + return (float) man * (1.0f / 1024.0f); + } + + const uint32_t bits = ((uint32_t) exp + 119u) << 23 | ((uint32_t) man << 20); + return rocmfp4_u32_as_f32(bits); +} + +static __device__ __forceinline__ uint8_t rocmfpx_nearest_scale_ue4m3_cuda(float target_scale) { + if (!(target_scale > 0.0f) || !isfinite(target_scale)) { + return 0; + } + + uint8_t lo = 1; + uint8_t hi = 0x7e; + while (lo < hi) { + const uint8_t mid = lo + (hi - lo) / 2; + if (rocmfpx_ue4m3_to_fp32_finite(mid) < target_scale) { + lo = mid + 1; + } else { + hi = mid; + } + } + + if (lo == 1) { + return 1; + } + + const float hi_scale = rocmfpx_ue4m3_to_fp32_finite(lo); + const float lo_scale = rocmfpx_ue4m3_to_fp32_finite((uint8_t) (lo - 1)); + return (target_scale - lo_scale <= hi_scale - target_scale) ? (uint8_t) (lo - 1) : lo; +} + +static __device__ __forceinline__ int8_t rocmfp4_decode_i8(uint8_t q) { + q &= 0x0f; + const int mag3 = q & 0x07; + const int mag = mag3 <= 4 ? mag3 : 2*mag3 - 4; + return (q & 0x08) ? -mag : mag; +} diff --git a/server/deps/llama.cpp/ggml/rocmfp4/test_rocmfp4.c b/server/deps/llama.cpp/ggml/rocmfp4/test_rocmfp4.c new file mode 100644 index 000000000..5ac092a6c --- /dev/null +++ b/server/deps/llama.cpp/ggml/rocmfp4/test_rocmfp4.c @@ -0,0 +1,119 @@ +#include "rocmfp4.h" + +#include +#include +#include +#include +#include + +enum { TEST_N = 64 }; + +static void fill_input(float * src) { + for (int i = 0; i < TEST_N; ++i) { + src[i] = 0.75f*sinf(0.37f*(float) i) + + 0.25f*cosf(0.13f*(float) i) + + 0.035f*(float) (i % 11 - 5); + } + + src[7] = 3.25f; + src[19] = -2.75f; + src[43] = 1.875f; +} + +static float mse(const float * expected, const float * actual) { + float error = 0.0f; + for (int i = 0; i < TEST_N; ++i) { + const float delta = expected[i] - actual[i]; + error += delta*delta; + } + return error / (float) TEST_N; +} + +static void check_reference_roundtrip(void) { + float src[TEST_N]; + float dual[TEST_N]; + float fast[TEST_N]; + block_rocmfp4 q_dual[TEST_N / QK_ROCMFP4]; + block_rocmfp4_fast q_fast[TEST_N / QK_ROCMFP4]; + + fill_input(src); + rocmfp4_quantize_row_q4_0_ref(src, q_dual, TEST_N); + rocmfp4_quantize_row_q4_0_fast_ref(src, q_fast, TEST_N); + + assert(rocmfp4_validate_row_data(q_dual, sizeof(q_dual))); + assert(rocmfp4_validate_row_data_fast(q_fast, sizeof(q_fast))); + + rocmfp4_dequantize_row_q4_0(q_dual, dual, TEST_N); + rocmfp4_dequantize_row_q4_0_fast(q_fast, fast, TEST_N); + + const float dual_mse = mse(src, dual); + const float fast_mse = mse(src, fast); + printf("ROCmFP4: dual_mse=%g fast_mse=%g\n", dual_mse, fast_mse); + assert(isfinite(dual_mse) && dual_mse < 0.1f); + assert(isfinite(fast_mse) && fast_mse < 0.1f); +} + +static void check_nonfinite_imatrix_input(void) { + float src[TEST_N]; + float weights[TEST_N]; + float dual[TEST_N]; + float fast[TEST_N]; + block_rocmfp4 q_dual[TEST_N / QK_ROCMFP4]; + block_rocmfp4_fast q_fast[TEST_N / QK_ROCMFP4]; + + fill_input(src); + for (int i = 0; i < TEST_N; ++i) { + weights[i] = 1.0f + (float) (i % 5); + } + + src[3] = NAN; + src[9] = INFINITY; + src[37] = -INFINITY; + src[55] = FLT_MAX; + weights[5] = NAN; + weights[11] = INFINITY; + weights[13] = -1.0f; + + assert(rocmfp4_quantize_q4_0(src, q_dual, 1, TEST_N, weights) == sizeof(q_dual)); + assert(rocmfp4_quantize_q4_0_fast(src, q_fast, 1, TEST_N, weights) == sizeof(q_fast)); + assert(rocmfp4_validate_row_data(q_dual, sizeof(q_dual))); + assert(rocmfp4_validate_row_data_fast(q_fast, sizeof(q_fast))); + + rocmfp4_dequantize_row_q4_0(q_dual, dual, TEST_N); + rocmfp4_dequantize_row_q4_0_fast(q_fast, fast, TEST_N); + for (int i = 0; i < TEST_N; ++i) { + assert(isfinite(dual[i])); + assert(isfinite(fast[i])); + } + + assert(dual[3] == 0.0f && dual[9] == 0.0f && dual[37] == 0.0f); + assert(fast[3] == 0.0f && fast[9] == 0.0f && fast[37] == 0.0f); + assert(dual[7] != 0.0f && fast[7] != 0.0f); +} + +static void check_validation_rejects_invalid_scales(void) { + block_rocmfp4 dual = { { 0 }, { 1, 1 } }; + block_rocmfp4_fast fast = { { 0 }, 1 }; + + assert(rocmfp4_validate_row_data(&dual, sizeof(dual))); + assert(rocmfp4_validate_row_data_fast(&fast, sizeof(fast))); + + dual.e[1] = 0x7f; + fast.e = 0x80; + assert(!rocmfp4_validate_row_data(&dual, sizeof(dual))); + assert(!rocmfp4_validate_row_data_fast(&fast, sizeof(fast))); + assert(!rocmfp4_validate_row_data(&dual, sizeof(dual) - 1)); + assert(!rocmfp4_validate_row_data_fast(&fast, sizeof(fast) - 1)); +} + +static void check_reserved_type_name(void) { + assert(strcmp(ggml_type_name((enum ggml_type) 50), "unknown") == 0); +} + +int main(void) { + check_reference_roundtrip(); + check_nonfinite_imatrix_input(); + check_validation_rejects_invalid_scales(); + check_reserved_type_name(); + return 0; +} diff --git a/server/deps/llama.cpp/ggml/rocmfpx/README.md b/server/deps/llama.cpp/ggml/rocmfpx/README.md new file mode 100644 index 000000000..ad90a0854 --- /dev/null +++ b/server/deps/llama.cpp/ggml/rocmfpx/README.md @@ -0,0 +1,28 @@ +# ROCmFPx formats + +This directory contains the lower- and higher-bit siblings of ROCmFP4. They +remain separate from `rocmfp4/` so the promoted 4-bit layouts can evolve +without changing these experimental formats. + +All layouts store 32 weights per block and use finite unsigned UE4M3 scale +bytes. + +| GGML type | Quantized payload | Scales | Block size | Bits/weight | +|---|---:|---:|---:|---:| +| `Q2_0_ROCMFP2` | 8 bytes | 2 | 10 bytes | 2.50 | +| `Q3_0_ROCMFPX` | 12 bytes | 2 | 14 bytes | 3.50 | +| `Q6_0_ROCMFPX` | 24 bytes | 2 | 26 bytes | 6.50 | +| `Q8_0_ROCMFPX` | 32 bytes | 1 | 33 bytes | 8.25 | + +ROCmFP2 uses the integer levels `-1, 0, 1, 2`; ROCmFP3 uses +`0, +/-1, +/-2, +/-4`; ROCmFP6 uses signed-magnitude levels up to 31; and +ROCmFP8 uses signed integer levels clamped to `[-127, 127]`. + +`rocmfpx.c` provides deterministic CPU reference quantization, +dequantization, validation, and vector-dot functions. Backend dispatch is +integrated in ggml's HIP source tree. Vulkan kernels are not implemented in +this vendored build. + +With `DFLASH27B_TESTS=ON`, `test_rocmfpx` checks reference round trips, +expected error ordering, weighted quantization, validation, and saturation of +large finite FP6/FP8 inputs. diff --git a/server/deps/llama.cpp/ggml/rocmfpx/rocmfpx.c b/server/deps/llama.cpp/ggml/rocmfpx/rocmfpx.c new file mode 100644 index 000000000..da798d23a --- /dev/null +++ b/server/deps/llama.cpp/ggml/rocmfpx/rocmfpx.c @@ -0,0 +1,1152 @@ +#include "rocmfpx.h" + +#include +#include +#include +#include + +// Finite unsigned E4M3 scale bytes decoded to FP32. Precomputed from the same +// exp/mant formula rocmfpx_ue4m3_to_fp32() used to evaluate with ldexpf(): +// exp == 0 -> mant * 2^-10 ; otherwise (8 + mant) * 2^(exp - 11). +// The scale search re-decodes candidate bytes for every block, and dequant +// decodes a scale for every element, so keeping this as a table (identical to +// the former per-call ldexpf result) removes the transcendental from both hot +// paths without changing any produced value. +#define ROCMFPX_SCALE_SUB(M) ((M) * 0x1p-10f) +#define ROCMFPX_SCALE_E(B, M) ((8 + (M)) * (B)) + +static const float rocmfpx_scale_ue4m3[127] = { + ROCMFPX_SCALE_SUB(0), ROCMFPX_SCALE_SUB(1), ROCMFPX_SCALE_SUB(2), ROCMFPX_SCALE_SUB(3), + ROCMFPX_SCALE_SUB(4), ROCMFPX_SCALE_SUB(5), ROCMFPX_SCALE_SUB(6), ROCMFPX_SCALE_SUB(7), + ROCMFPX_SCALE_E(0x1p-10f,0), ROCMFPX_SCALE_E(0x1p-10f,1), ROCMFPX_SCALE_E(0x1p-10f,2), ROCMFPX_SCALE_E(0x1p-10f,3), + ROCMFPX_SCALE_E(0x1p-10f,4), ROCMFPX_SCALE_E(0x1p-10f,5), ROCMFPX_SCALE_E(0x1p-10f,6), ROCMFPX_SCALE_E(0x1p-10f,7), + ROCMFPX_SCALE_E(0x1p-9f,0), ROCMFPX_SCALE_E(0x1p-9f,1), ROCMFPX_SCALE_E(0x1p-9f,2), ROCMFPX_SCALE_E(0x1p-9f,3), + ROCMFPX_SCALE_E(0x1p-9f,4), ROCMFPX_SCALE_E(0x1p-9f,5), ROCMFPX_SCALE_E(0x1p-9f,6), ROCMFPX_SCALE_E(0x1p-9f,7), + ROCMFPX_SCALE_E(0x1p-8f,0), ROCMFPX_SCALE_E(0x1p-8f,1), ROCMFPX_SCALE_E(0x1p-8f,2), ROCMFPX_SCALE_E(0x1p-8f,3), + ROCMFPX_SCALE_E(0x1p-8f,4), ROCMFPX_SCALE_E(0x1p-8f,5), ROCMFPX_SCALE_E(0x1p-8f,6), ROCMFPX_SCALE_E(0x1p-8f,7), + ROCMFPX_SCALE_E(0x1p-7f,0), ROCMFPX_SCALE_E(0x1p-7f,1), ROCMFPX_SCALE_E(0x1p-7f,2), ROCMFPX_SCALE_E(0x1p-7f,3), + ROCMFPX_SCALE_E(0x1p-7f,4), ROCMFPX_SCALE_E(0x1p-7f,5), ROCMFPX_SCALE_E(0x1p-7f,6), ROCMFPX_SCALE_E(0x1p-7f,7), + ROCMFPX_SCALE_E(0x1p-6f,0), ROCMFPX_SCALE_E(0x1p-6f,1), ROCMFPX_SCALE_E(0x1p-6f,2), ROCMFPX_SCALE_E(0x1p-6f,3), + ROCMFPX_SCALE_E(0x1p-6f,4), ROCMFPX_SCALE_E(0x1p-6f,5), ROCMFPX_SCALE_E(0x1p-6f,6), ROCMFPX_SCALE_E(0x1p-6f,7), + ROCMFPX_SCALE_E(0x1p-5f,0), ROCMFPX_SCALE_E(0x1p-5f,1), ROCMFPX_SCALE_E(0x1p-5f,2), ROCMFPX_SCALE_E(0x1p-5f,3), + ROCMFPX_SCALE_E(0x1p-5f,4), ROCMFPX_SCALE_E(0x1p-5f,5), ROCMFPX_SCALE_E(0x1p-5f,6), ROCMFPX_SCALE_E(0x1p-5f,7), + ROCMFPX_SCALE_E(0x1p-4f,0), ROCMFPX_SCALE_E(0x1p-4f,1), ROCMFPX_SCALE_E(0x1p-4f,2), ROCMFPX_SCALE_E(0x1p-4f,3), + ROCMFPX_SCALE_E(0x1p-4f,4), ROCMFPX_SCALE_E(0x1p-4f,5), ROCMFPX_SCALE_E(0x1p-4f,6), ROCMFPX_SCALE_E(0x1p-4f,7), + ROCMFPX_SCALE_E(0x1p-3f,0), ROCMFPX_SCALE_E(0x1p-3f,1), ROCMFPX_SCALE_E(0x1p-3f,2), ROCMFPX_SCALE_E(0x1p-3f,3), + ROCMFPX_SCALE_E(0x1p-3f,4), ROCMFPX_SCALE_E(0x1p-3f,5), ROCMFPX_SCALE_E(0x1p-3f,6), ROCMFPX_SCALE_E(0x1p-3f,7), + ROCMFPX_SCALE_E(0x1p-2f,0), ROCMFPX_SCALE_E(0x1p-2f,1), ROCMFPX_SCALE_E(0x1p-2f,2), ROCMFPX_SCALE_E(0x1p-2f,3), + ROCMFPX_SCALE_E(0x1p-2f,4), ROCMFPX_SCALE_E(0x1p-2f,5), ROCMFPX_SCALE_E(0x1p-2f,6), ROCMFPX_SCALE_E(0x1p-2f,7), + ROCMFPX_SCALE_E(0x1p-1f,0), ROCMFPX_SCALE_E(0x1p-1f,1), ROCMFPX_SCALE_E(0x1p-1f,2), ROCMFPX_SCALE_E(0x1p-1f,3), + ROCMFPX_SCALE_E(0x1p-1f,4), ROCMFPX_SCALE_E(0x1p-1f,5), ROCMFPX_SCALE_E(0x1p-1f,6), ROCMFPX_SCALE_E(0x1p-1f,7), + ROCMFPX_SCALE_E(0x1p0f,0), ROCMFPX_SCALE_E(0x1p0f,1), ROCMFPX_SCALE_E(0x1p0f,2), ROCMFPX_SCALE_E(0x1p0f,3), + ROCMFPX_SCALE_E(0x1p0f,4), ROCMFPX_SCALE_E(0x1p0f,5), ROCMFPX_SCALE_E(0x1p0f,6), ROCMFPX_SCALE_E(0x1p0f,7), + ROCMFPX_SCALE_E(0x1p1f,0), ROCMFPX_SCALE_E(0x1p1f,1), ROCMFPX_SCALE_E(0x1p1f,2), ROCMFPX_SCALE_E(0x1p1f,3), + ROCMFPX_SCALE_E(0x1p1f,4), ROCMFPX_SCALE_E(0x1p1f,5), ROCMFPX_SCALE_E(0x1p1f,6), ROCMFPX_SCALE_E(0x1p1f,7), + ROCMFPX_SCALE_E(0x1p2f,0), ROCMFPX_SCALE_E(0x1p2f,1), ROCMFPX_SCALE_E(0x1p2f,2), ROCMFPX_SCALE_E(0x1p2f,3), + ROCMFPX_SCALE_E(0x1p2f,4), ROCMFPX_SCALE_E(0x1p2f,5), ROCMFPX_SCALE_E(0x1p2f,6), ROCMFPX_SCALE_E(0x1p2f,7), + ROCMFPX_SCALE_E(0x1p3f,0), ROCMFPX_SCALE_E(0x1p3f,1), ROCMFPX_SCALE_E(0x1p3f,2), ROCMFPX_SCALE_E(0x1p3f,3), + ROCMFPX_SCALE_E(0x1p3f,4), ROCMFPX_SCALE_E(0x1p3f,5), ROCMFPX_SCALE_E(0x1p3f,6), ROCMFPX_SCALE_E(0x1p3f,7), + ROCMFPX_SCALE_E(0x1p4f,0), ROCMFPX_SCALE_E(0x1p4f,1), ROCMFPX_SCALE_E(0x1p4f,2), ROCMFPX_SCALE_E(0x1p4f,3), + ROCMFPX_SCALE_E(0x1p4f,4), ROCMFPX_SCALE_E(0x1p4f,5), ROCMFPX_SCALE_E(0x1p4f,6), +}; + +#undef ROCMFPX_SCALE_SUB +#undef ROCMFPX_SCALE_E + +float rocmfpx_ue4m3_to_fp32(uint8_t e) { + return rocmfpx_scale_is_valid(e) ? rocmfpx_scale_ue4m3[e] : 0.0f; +} + +bool rocmfpx_scale_is_valid(uint8_t e) { + return e <= 0x7e; +} + +size_t rocmfpx_row_size_fp2(int64_t k) { + assert(k % QK_ROCMFP2 == 0); + return (size_t) (k / QK_ROCMFP2) * sizeof(block_rocmfp2); +} + +size_t rocmfpx_row_size_fp3(int64_t k) { + assert(k % QK_ROCMFP3 == 0); + return (size_t) (k / QK_ROCMFP3) * sizeof(block_rocmfp3); +} + +size_t rocmfpx_row_size_fp6(int64_t k) { + assert(k % QK_ROCMFP6 == 0); + return (size_t) (k / QK_ROCMFP6) * sizeof(block_rocmfp6); +} + +size_t rocmfpx_row_size_fp8(int64_t k) { + assert(k % QK_ROCMFP8 == 0); + return (size_t) (k / QK_ROCMFP8) * sizeof(block_rocmfp8); +} + +static uint8_t rocmfpx_nearest_scale_ue4m3(float target) { + if (!(target > 0.0f) || !isfinite(target)) { + return 0; + } + + uint8_t best_e = 1; + float best_err = fabsf(rocmfpx_ue4m3_to_fp32(best_e) - target); + + for (int e = 2; e <= 0x7e; ++e) { + const float err = fabsf(rocmfpx_ue4m3_to_fp32((uint8_t) e) - target); + if (err < best_err) { + best_err = err; + best_e = (uint8_t) e; + } + } + + return best_e; +} + +static float rocmfpx_max_abs(const float * x, int n) { + float max_abs = 0.0f; + + for (int i = 0; i < n; ++i) { + if (!isfinite(x[i])) { + continue; + } + + const float ax = fabsf(x[i]); + if (ax > max_abs) { + max_abs = ax; + } + } + + return max_abs; +} + +static float rocmfpx_row_sigma2(const float * x, int64_t k) { + float sum_x2 = 0.0f; + bool overflow = false; + + for (int64_t i = 0; i < k; ++i) { + if (!isfinite(x[i])) { + continue; + } + + const float x2 = x[i]*x[i]; + if (!isfinite(x2) || sum_x2 > FLT_MAX - x2) { + overflow = true; + break; + } + sum_x2 += x2; + } + + if (!overflow) { + return sum_x2 / (float) k; + } + + double sum_x2_wide = 0.0; + for (int64_t i = 0; i < k; ++i) { + if (isfinite(x[i])) { + const double value = x[i]; + sum_x2_wide += value*value; + } + } + return (float) fmin(sum_x2_wide / (double) k, (double) FLT_MAX); +} + +static void rocmfpx_prepare_mse_weights( + float * dst, const float * x, int n, const float * quant_weights, float sigma2, + float * max_abs, float * max_abs_weight) { + *max_abs = 0.0f; + *max_abs_weight = 0.0f; + + for (int i = 0; i < n; ++i) { + const float ax = fabsf(x[i]); + const float qw = quant_weights[i]; + float weight = 0.0f; + if (isfinite(qw) && qw > 0.0f && isfinite(x[i])) { + const float energy2 = sigma2 + x[i]*x[i]; + const float candidate = isfinite(energy2) ? qw*sqrtf(energy2) : FLT_MAX; + weight = isfinite(candidate) ? candidate : FLT_MAX; + } + + if (isfinite(x[i])) { + if (ax > *max_abs) { + *max_abs = ax; + *max_abs_weight = weight; + } else if (ax == *max_abs && weight > *max_abs_weight) { + *max_abs_weight = weight; + } + } + + // Match llama.cpp imatrix weighting style: calibration importance is + // scaled by row energy so large activations stay protected. + dst[i] = weight; + } +} + +static void rocmfpx_set_bits(uint8_t * dst, int bit_pos, int nbits, uint32_t code) { + for (int bit = 0; bit < nbits; ++bit) { + const int absolute_bit = bit_pos + bit; + const int byte_index = absolute_bit >> 3; + const int bit_index = absolute_bit & 7; + + if ((code >> bit) & 1u) { + dst[byte_index] |= (uint8_t) (1u << bit_index); + } + } +} + +static uint32_t rocmfpx_get_bits(const uint8_t * src, int bit_pos, int nbits) { + uint32_t code = 0; + + for (int bit = 0; bit < nbits; ++bit) { + const int absolute_bit = bit_pos + bit; + const int byte_index = absolute_bit >> 3; + const int bit_index = absolute_bit & 7; + + code |= (uint32_t) ((src[byte_index] >> bit_index) & 1u) << bit; + } + + return code; +} + +// Starting 2-bit ROCmFP2 codebook. Keep this single definition easy to tune. +// TODO: affine scale+min would likely improve quality, but it would break the +// ROCmFPx family's unsigned-UE4M3 scale contract; revisit for a v2 layout. +static const float kvalues_rocmfp2[4] = ROCMFP2_KVALUES_INIT; + +static uint8_t rocmfpx_quantize_fp2_code(float x, float inv_scale) { + if (!isfinite(x) || inv_scale <= 0.0f) { + return 1; + } + + const float q = x * inv_scale; + uint8_t best_code = 0; + float best_err = fabsf(q - kvalues_rocmfp2[0]); + + for (uint8_t code = 1; code < 4; ++code) { + const float err = fabsf(q - kvalues_rocmfp2[code]); + if (err < best_err) { + best_err = err; + best_code = code; + } + } + + return best_code; +} + +static inline float rocmfpx_fp2_decoded_mag(float x, float inv_scale) { + return kvalues_rocmfp2[rocmfpx_quantize_fp2_code(x, inv_scale)]; +} + +static float rocmfpx_fp2_block_mse_for_scale(const float * x, int n, uint8_t e, float best_err) { + const float scale = rocmfpx_ue4m3_to_fp32(e); + const float inv_scale = scale > 0.0f ? 1.0f / scale : 0.0f; + float err = 0.0f; + + for (int i = 0; i < n; ++i) { + if (!isfinite(x[i])) { + continue; + } + + const float y = rocmfpx_fp2_decoded_mag(x[i], inv_scale) * scale; + const float d = x[i] - y; + + err += d*d; + if (err > best_err) { + return err; + } + } + + return err; +} + +static float rocmfpx_fp2_block_weighted_mse_for_scale(const float * x, int n, const float * mse_weights, uint8_t e, float best_err) { + const float scale = rocmfpx_ue4m3_to_fp32(e); + const float inv_scale = scale > 0.0f ? 1.0f / scale : 0.0f; + float err = 0.0f; + + for (int i = 0; i < n; ++i) { + if (!isfinite(x[i])) { + continue; + } + + const float y = rocmfpx_fp2_decoded_mag(x[i], inv_scale) * scale; + const float d = x[i] - y; + + err += mse_weights[i]*d*d; + if (err > best_err) { + return err; + } + } + + return err; +} + +static uint8_t rocmfpx_choose_scale_fp2_mse_impl(const float * x, int n, const float * mse_weights, float max_abs, float max_abs_weight) { + const uint8_t start_e = rocmfpx_nearest_scale_ue4m3(max_abs / 2.0f); + uint8_t best_e = start_e; + float best_err = INFINITY; + bool lower_done = false; + + for (int delta = 0; delta <= 125; ++delta) { + const int e0 = (int) start_e - delta; + if (!lower_done && e0 >= 1 && e0 <= 126) { + const float scale = rocmfpx_ue4m3_to_fp32((uint8_t) e0); + const float clip_delta = max_abs - 2.0f*scale; + const float clip_err = mse_weights ? max_abs_weight*clip_delta*clip_delta : clip_delta*clip_delta; + if (clip_delta > 0.0f && clip_err > best_err) { + lower_done = true; + } else { + const float err = mse_weights ? + rocmfpx_fp2_block_weighted_mse_for_scale(x, n, mse_weights, (uint8_t) e0, best_err) : + rocmfpx_fp2_block_mse_for_scale(x, n, (uint8_t) e0, best_err); + if (err < best_err || (err == best_err && e0 < best_e)) { + best_err = err; + best_e = (uint8_t) e0; + } + } + } + + const int e1 = (int) start_e + delta; + if (delta != 0 && e1 >= 1 && e1 <= 126) { + const float err = mse_weights ? + rocmfpx_fp2_block_weighted_mse_for_scale(x, n, mse_weights, (uint8_t) e1, best_err) : + rocmfpx_fp2_block_mse_for_scale(x, n, (uint8_t) e1, best_err); + if (err < best_err || (err == best_err && e1 < best_e)) { + best_err = err; + best_e = (uint8_t) e1; + } + } + + if ((lower_done || e0 <= 1) && e1 >= 126) { + break; + } + } + + return best_e; +} + +static uint8_t rocmfpx_choose_scale_fp2_mse(const float * x, int n) { + const float max_abs = rocmfpx_max_abs(x, n); + if (!(max_abs > 0.0f) || !isfinite(max_abs)) { + return 0; + } + + return rocmfpx_choose_scale_fp2_mse_impl(x, n, NULL, max_abs, 0.0f); +} + +static uint8_t rocmfpx_choose_scale_fp2_weighted_mse(const float * x, int n, const float * quant_weights, float sigma2) { + assert(n <= QK_ROCMFP2); + float mse_weights[QK_ROCMFP2]; + float max_abs; + float max_abs_weight; + rocmfpx_prepare_mse_weights(mse_weights, x, n, quant_weights, sigma2, &max_abs, &max_abs_weight); + if (!(max_abs > 0.0f) || !isfinite(max_abs)) { + return 0; + } + + return rocmfpx_choose_scale_fp2_mse_impl(x, n, mse_weights, max_abs, max_abs_weight); +} + +static int rocmfpx_decode_fp3_code(uint8_t code) { + static const int mag[4] = { 0, 1, 2, 4 }; + const int value = mag[code & 3u]; + return (code & 4u) ? -value : value; +} + +static uint8_t rocmfpx_quantize_fp3_code(float x, float inv_scale) { + if (!isfinite(x) || inv_scale <= 0.0f) { + return 0; + } + + const float ax = fabsf(x * inv_scale); + uint8_t mag; + + if (ax <= 0.5f) { + mag = 0; + } else if (ax <= 1.5f) { + mag = 1; + } else if (ax <= 3.0f) { + mag = 2; + } else { + mag = 3; + } + + return mag == 0 ? 0 : (uint8_t) ((x < 0.0f ? 4u : 0u) | mag); +} + +// Fused threshold + decode used only inside the exhaustive scale search, which +// re-scans every element for every candidate scale byte. Returns the same +// signed decoded magnitude that +// rocmfpx_decode_fp3_code(rocmfpx_quantize_fp3_code(x, inv_scale)) produces +// (fp3 magnitudes {0,1,2,4}), so quantized output stays bit-identical. +static inline float rocmfpx_fp3_decoded_mag(float x, float inv_scale) { + const float a = fabsf(x * inv_scale); + float mag; + if (a <= 0.5f) { + return 0.0f; + } else if (a <= 1.5f) { + mag = 1.0f; + } else if (a <= 3.0f) { + mag = 2.0f; + } else { + mag = 4.0f; + } + return x < 0.0f ? -mag : mag; +} + +static float rocmfpx_fp3_block_mse_for_scale(const float * x, int n, uint8_t e, float best_err) { + const float scale = rocmfpx_ue4m3_to_fp32(e); + const float inv_scale = scale > 0.0f ? 1.0f / scale : 0.0f; + float err = 0.0f; + + for (int i = 0; i < n; ++i) { + if (!isfinite(x[i])) { + continue; + } + + const float y = rocmfpx_fp3_decoded_mag(x[i], inv_scale) * scale; + const float d = x[i] - y; + + err += d*d; + if (err > best_err) { + return err; + } + } + + return err; +} + +static float rocmfpx_fp3_block_weighted_mse_for_scale(const float * x, int n, const float * mse_weights, uint8_t e, float best_err) { + const float scale = rocmfpx_ue4m3_to_fp32(e); + const float inv_scale = scale > 0.0f ? 1.0f / scale : 0.0f; + float err = 0.0f; + + for (int i = 0; i < n; ++i) { + if (!isfinite(x[i])) { + continue; + } + + const float y = rocmfpx_fp3_decoded_mag(x[i], inv_scale) * scale; + const float d = x[i] - y; + + err += mse_weights[i]*d*d; + if (err > best_err) { + return err; + } + } + + return err; +} + +static uint8_t rocmfpx_choose_scale_fp3_mse_impl(const float * x, int n, const float * mse_weights, float max_abs, float max_abs_weight) { + const uint8_t start_e = rocmfpx_nearest_scale_ue4m3(max_abs / 4.0f); + uint8_t best_e = start_e; + float best_err = INFINITY; + bool lower_done = false; + + for (int delta = 0; delta <= 125; ++delta) { + const int e0 = (int) start_e - delta; + if (!lower_done && e0 >= 1 && e0 <= 126) { + const float scale = rocmfpx_ue4m3_to_fp32((uint8_t) e0); + const float clip_delta = max_abs - 4.0f*scale; + const float clip_err = mse_weights ? max_abs_weight*clip_delta*clip_delta : clip_delta*clip_delta; + if (clip_delta > 0.0f && clip_err > best_err) { + lower_done = true; + } else { + const float err = mse_weights ? + rocmfpx_fp3_block_weighted_mse_for_scale(x, n, mse_weights, (uint8_t) e0, best_err) : + rocmfpx_fp3_block_mse_for_scale(x, n, (uint8_t) e0, best_err); + if (err < best_err || (err == best_err && e0 < best_e)) { + best_err = err; + best_e = (uint8_t) e0; + } + } + } + + const int e1 = (int) start_e + delta; + if (delta != 0 && e1 >= 1 && e1 <= 126) { + const float err = mse_weights ? + rocmfpx_fp3_block_weighted_mse_for_scale(x, n, mse_weights, (uint8_t) e1, best_err) : + rocmfpx_fp3_block_mse_for_scale(x, n, (uint8_t) e1, best_err); + if (err < best_err || (err == best_err && e1 < best_e)) { + best_err = err; + best_e = (uint8_t) e1; + } + } + + if ((lower_done || e0 <= 1) && e1 >= 126) { + break; + } + } + + return best_e; +} + +static uint8_t rocmfpx_choose_scale_fp3_mse(const float * x, int n) { + const float max_abs = rocmfpx_max_abs(x, n); + if (!(max_abs > 0.0f) || !isfinite(max_abs)) { + return 0; + } + + return rocmfpx_choose_scale_fp3_mse_impl(x, n, NULL, max_abs, 0.0f); +} + +static uint8_t rocmfpx_choose_scale_fp3_weighted_mse(const float * x, int n, const float * quant_weights, float sigma2) { + assert(n <= QK_ROCMFP3); + float mse_weights[QK_ROCMFP3]; + float max_abs; + float max_abs_weight; + rocmfpx_prepare_mse_weights(mse_weights, x, n, quant_weights, sigma2, &max_abs, &max_abs_weight); + if (!(max_abs > 0.0f) || !isfinite(max_abs)) { + return 0; + } + + return rocmfpx_choose_scale_fp3_mse_impl(x, n, mse_weights, max_abs, max_abs_weight); +} + +static int rocmfpx_decode_fp6_code(uint8_t code) { + const int value = code & 31u; + return (code & 32u) ? -value : value; +} + +static uint8_t rocmfpx_quantize_fp6_code(float x, float inv_scale) { + if (!isfinite(x) || inv_scale <= 0.0f) { + return 0; + } + + const float scaled = fminf(fabsf(x * inv_scale), 31.0f); + const int mag = (int) lroundf(scaled); + + return mag == 0 ? 0 : (uint8_t) ((x < 0.0f ? 32u : 0u) | (uint8_t) mag); +} + +// Fused round + clamp + decode for the fp6 scale search. Returns the same signed +// decoded magnitude as rocmfpx_decode_fp6_code(rocmfpx_quantize_fp6_code(...)) +// (nearest integer in [0,31], signed), keeping quantized output bit-identical. +static inline float rocmfpx_fp6_decoded_mag(float x, float inv_scale) { + const float scaled = fminf(fabsf(x * inv_scale), 31.0f); + const int mag = (int) lroundf(scaled); + if (mag == 0) { + return 0.0f; + } + return x < 0.0f ? -(float) mag : (float) mag; +} + +static float rocmfpx_fp6_block_mse_for_scale(const float * x, int n, uint8_t e, float best_err) { + const float scale = rocmfpx_ue4m3_to_fp32(e); + const float inv_scale = scale > 0.0f ? 1.0f / scale : 0.0f; + float err = 0.0f; + + for (int i = 0; i < n; ++i) { + if (!isfinite(x[i])) { + continue; + } + const float y = rocmfpx_fp6_decoded_mag(x[i], inv_scale) * scale; + const float d = x[i] - y; + err += d*d; + if (err > best_err) { + return err; + } + } + + return err; +} + +static float rocmfpx_fp6_block_weighted_mse_for_scale(const float * x, int n, const float * mse_weights, uint8_t e, float best_err) { + const float scale = rocmfpx_ue4m3_to_fp32(e); + const float inv_scale = scale > 0.0f ? 1.0f / scale : 0.0f; + float err = 0.0f; + + for (int i = 0; i < n; ++i) { + if (!isfinite(x[i])) { + continue; + } + const float y = rocmfpx_fp6_decoded_mag(x[i], inv_scale) * scale; + const float d = x[i] - y; + err += mse_weights[i]*d*d; + if (err > best_err) { + return err; + } + } + + return err; +} + +static uint8_t rocmfpx_choose_scale_fp6_mse_impl(const float * x, int n, const float * mse_weights, float max_abs, float max_abs_weight) { + const uint8_t start_e = rocmfpx_nearest_scale_ue4m3(max_abs / 31.0f); + uint8_t best_e = start_e; + float best_err = INFINITY; + bool lower_done = false; + + for (int delta = 0; delta <= 125; ++delta) { + const int e0 = (int) start_e - delta; + if (!lower_done && e0 >= 1 && e0 <= 126) { + const float scale = rocmfpx_ue4m3_to_fp32((uint8_t) e0); + const float clip_delta = max_abs - 31.0f*scale; + const float clip_err = mse_weights ? max_abs_weight*clip_delta*clip_delta : clip_delta*clip_delta; + if (clip_delta > 0.0f && clip_err > best_err) { + lower_done = true; + } else { + const float err = mse_weights ? + rocmfpx_fp6_block_weighted_mse_for_scale(x, n, mse_weights, (uint8_t) e0, best_err) : + rocmfpx_fp6_block_mse_for_scale(x, n, (uint8_t) e0, best_err); + if (err < best_err || (err == best_err && e0 < best_e)) { + best_err = err; + best_e = (uint8_t) e0; + } + } + } + + const int e1 = (int) start_e + delta; + if (delta != 0 && e1 >= 1 && e1 <= 126) { + const float err = mse_weights ? + rocmfpx_fp6_block_weighted_mse_for_scale(x, n, mse_weights, (uint8_t) e1, best_err) : + rocmfpx_fp6_block_mse_for_scale(x, n, (uint8_t) e1, best_err); + if (err < best_err || (err == best_err && e1 < best_e)) { + best_err = err; + best_e = (uint8_t) e1; + } + } + + if ((lower_done || e0 <= 1) && e1 >= 126) { + break; + } + } + + return best_e; +} + +static uint8_t rocmfpx_choose_scale_fp6_mse(const float * x, int n) { + const float max_abs = rocmfpx_max_abs(x, n); + if (!(max_abs > 0.0f) || !isfinite(max_abs)) { + return 0; + } + + return rocmfpx_choose_scale_fp6_mse_impl(x, n, NULL, max_abs, 0.0f); +} + +static uint8_t rocmfpx_choose_scale_fp6_weighted_mse(const float * x, int n, const float * quant_weights, float sigma2) { + assert(n <= QK_ROCMFP6); + float mse_weights[QK_ROCMFP6]; + float max_abs; + float max_abs_weight; + rocmfpx_prepare_mse_weights(mse_weights, x, n, quant_weights, sigma2, &max_abs, &max_abs_weight); + if (!(max_abs > 0.0f) || !isfinite(max_abs)) { + return 0; + } + + return rocmfpx_choose_scale_fp6_mse_impl(x, n, mse_weights, max_abs, max_abs_weight); +} + +static int8_t rocmfpx_quantize_fp8_code(float x, float inv_scale) { + if (!isfinite(x) || inv_scale <= 0.0f) { + return 0; + } + + const float scaled = fmaxf(-127.0f, fminf(x * inv_scale, 127.0f)); + const int q = (int) lroundf(scaled); + return (int8_t) q; +} + +static float rocmfpx_fp8_block_weighted_mse_for_scale(const float * x, int n, const float * mse_weights, uint8_t e, float best_err) { + const float scale = rocmfpx_ue4m3_to_fp32(e); + const float inv_scale = scale > 0.0f ? 1.0f / scale : 0.0f; + float err = 0.0f; + + for (int i = 0; i < n; ++i) { + if (!isfinite(x[i])) { + continue; + } + + const int8_t code = rocmfpx_quantize_fp8_code(x[i], inv_scale); + const float y = (float) code * scale; + const float d = x[i] - y; + + err += mse_weights[i]*d*d; + if (err > best_err) { + return err; + } + } + + return err; +} + +static uint8_t rocmfpx_choose_scale_fp8_weighted_mse(const float * x, int n, const float * quant_weights, float sigma2) { + assert(n <= QK_ROCMFP8); + float mse_weights[QK_ROCMFP8]; + float max_abs; + float max_abs_weight; + rocmfpx_prepare_mse_weights(mse_weights, x, n, quant_weights, sigma2, &max_abs, &max_abs_weight); + if (!(max_abs > 0.0f) || !isfinite(max_abs)) { + return 0; + } + + const uint8_t start_e = rocmfpx_nearest_scale_ue4m3(max_abs / 127.0f); + uint8_t best_e = start_e; + float best_err = INFINITY; + bool lower_done = false; + + for (int delta = 0; delta <= 125; ++delta) { + const int e0 = (int) start_e - delta; + if (!lower_done && e0 >= 1 && e0 <= 126) { + const float scale = rocmfpx_ue4m3_to_fp32((uint8_t) e0); + const float clip_delta = max_abs - 127.0f*scale; + if (clip_delta > 0.0f && max_abs_weight*clip_delta*clip_delta > best_err) { + lower_done = true; + } else { + const float err = rocmfpx_fp8_block_weighted_mse_for_scale(x, n, mse_weights, (uint8_t) e0, best_err); + if (err < best_err || (err == best_err && e0 < best_e)) { + best_err = err; + best_e = (uint8_t) e0; + } + } + } + + const int e1 = (int) start_e + delta; + if (delta != 0 && e1 >= 1 && e1 <= 126) { + const float err = rocmfpx_fp8_block_weighted_mse_for_scale(x, n, mse_weights, (uint8_t) e1, best_err); + if (err < best_err || (err == best_err && e1 < best_e)) { + best_err = err; + best_e = (uint8_t) e1; + } + } + + if ((lower_done || e0 <= 1) && e1 >= 126) { + break; + } + } + + return best_e; +} + +void rocmfpx_quantize_row_fp2_ref(const float * GGML_RESTRICT x, block_rocmfp2 * GGML_RESTRICT y, int64_t k) { + assert(k % QK_ROCMFP2 == 0); + + const int64_t nb = k / QK_ROCMFP2; + for (int64_t ib = 0; ib < nb; ++ib) { + const float * xb = x + ib*QK_ROCMFP2; + block_rocmfp2 * yb = y + ib; + + memset(yb->qs, 0, sizeof(yb->qs)); + + for (int half = 0; half < 2; ++half) { + const float * xh = xb + half*(QK_ROCMFP2/2); + yb->e[half] = rocmfpx_choose_scale_fp2_mse(xh, QK_ROCMFP2/2); + + const float scale = rocmfpx_ue4m3_to_fp32(yb->e[half]); + const float inv_scale = scale > 0.0f ? 1.0f / scale : 0.0f; + + for (int j = 0; j < QK_ROCMFP2/2; ++j) { + const int i = half*(QK_ROCMFP2/2) + j; + const uint8_t code = rocmfpx_quantize_fp2_code(xb[i], inv_scale); + yb->qs[i >> 2] |= (uint8_t) (code << (2*(i & 3))); + } + } + } +} + +static void rocmfpx_quantize_row_fp2_weighted( + const float * GGML_RESTRICT x, block_rocmfp2 * GGML_RESTRICT y, int64_t k, const float * GGML_RESTRICT quant_weights) { + assert(k % QK_ROCMFP2 == 0); + + const float sigma2 = rocmfpx_row_sigma2(x, k); + + const int64_t nb = k / QK_ROCMFP2; + for (int64_t ib = 0; ib < nb; ++ib) { + const float * xb = x + ib*QK_ROCMFP2; + const float * qw = quant_weights ? quant_weights + ib*QK_ROCMFP2 : NULL; + block_rocmfp2 * yb = y + ib; + + memset(yb->qs, 0, sizeof(yb->qs)); + + for (int half = 0; half < 2; ++half) { + const int half_off = half*(QK_ROCMFP2/2); + const float * xh = xb + half_off; + const float * qh = qw ? qw + half_off : NULL; + yb->e[half] = qh ? + rocmfpx_choose_scale_fp2_weighted_mse(xh, QK_ROCMFP2/2, qh, sigma2) : + rocmfpx_choose_scale_fp2_mse(xh, QK_ROCMFP2/2); + + const float scale = rocmfpx_ue4m3_to_fp32(yb->e[half]); + const float inv_scale = scale > 0.0f ? 1.0f / scale : 0.0f; + + for (int j = 0; j < QK_ROCMFP2/2; ++j) { + const int i = half_off + j; + const uint8_t code = rocmfpx_quantize_fp2_code(xb[i], inv_scale); + yb->qs[i >> 2] |= (uint8_t) (code << (2*(i & 3))); + } + } + } +} + +void rocmfpx_dequantize_row_fp2(const block_rocmfp2 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { + assert(k % QK_ROCMFP2 == 0); + + const int64_t nb = k / QK_ROCMFP2; + for (int64_t ib = 0; ib < nb; ++ib) { + const block_rocmfp2 * xb = x + ib; + float * yb = y + ib*QK_ROCMFP2; + + for (int i = 0; i < QK_ROCMFP2; ++i) { + const float scale = rocmfpx_ue4m3_to_fp32(xb->e[i >= QK_ROCMFP2/2]); + const uint8_t code = (uint8_t) ((xb->qs[i >> 2] >> (2*(i & 3))) & 3u); + yb[i] = kvalues_rocmfp2[code] * scale; + } + } +} + +void rocmfpx_quantize_row_fp2(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { + rocmfpx_quantize_row_fp2_ref(x, (block_rocmfp2 *) y, k); +} + +size_t rocmfpx_quantize_fp2(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix) { + const size_t row_size = rocmfpx_row_size_fp2(n_per_row); + char * qrow = (char *) dst; + + for (int64_t row = 0; row < nrows; ++row) { + if (imatrix) { + rocmfpx_quantize_row_fp2_weighted(src + row*n_per_row, (block_rocmfp2 *) qrow, n_per_row, imatrix); + } else { + rocmfpx_quantize_row_fp2_ref(src + row*n_per_row, (block_rocmfp2 *) qrow, n_per_row); + } + qrow += row_size; + } + + return (size_t) nrows * row_size; +} + +void rocmfpx_quantize_row_fp3_ref(const float * GGML_RESTRICT x, block_rocmfp3 * GGML_RESTRICT y, int64_t k) { + assert(k % QK_ROCMFP3 == 0); + + const int64_t nb = k / QK_ROCMFP3; + for (int64_t ib = 0; ib < nb; ++ib) { + const float * xb = x + ib*QK_ROCMFP3; + block_rocmfp3 * yb = y + ib; + + memset(yb->qs, 0, sizeof(yb->qs)); + + for (int half = 0; half < 2; ++half) { + const float * xh = xb + half*(QK_ROCMFP3/2); + yb->e[half] = rocmfpx_choose_scale_fp3_mse(xh, QK_ROCMFP3/2); + + const float scale = rocmfpx_ue4m3_to_fp32(yb->e[half]); + const float inv_scale = scale > 0.0f ? 1.0f / scale : 0.0f; + + for (int j = 0; j < QK_ROCMFP3/2; ++j) { + const int i = half*(QK_ROCMFP3/2) + j; + const uint8_t code = rocmfpx_quantize_fp3_code(xb[i], inv_scale); + rocmfpx_set_bits(yb->qs, i*3, 3, code); + } + } + } +} + +static void rocmfpx_quantize_row_fp3_weighted( + const float * GGML_RESTRICT x, block_rocmfp3 * GGML_RESTRICT y, int64_t k, const float * GGML_RESTRICT quant_weights) { + assert(k % QK_ROCMFP3 == 0); + + const float sigma2 = rocmfpx_row_sigma2(x, k); + + const int64_t nb = k / QK_ROCMFP3; + for (int64_t ib = 0; ib < nb; ++ib) { + const float * xb = x + ib*QK_ROCMFP3; + const float * qw = quant_weights ? quant_weights + ib*QK_ROCMFP3 : NULL; + block_rocmfp3 * yb = y + ib; + + memset(yb->qs, 0, sizeof(yb->qs)); + + for (int half = 0; half < 2; ++half) { + const int half_off = half*(QK_ROCMFP3/2); + const float * xh = xb + half_off; + const float * qh = qw ? qw + half_off : NULL; + yb->e[half] = qh ? + rocmfpx_choose_scale_fp3_weighted_mse(xh, QK_ROCMFP3/2, qh, sigma2) : + rocmfpx_choose_scale_fp3_mse(xh, QK_ROCMFP3/2); + + const float scale = rocmfpx_ue4m3_to_fp32(yb->e[half]); + const float inv_scale = scale > 0.0f ? 1.0f / scale : 0.0f; + + for (int j = 0; j < QK_ROCMFP3/2; ++j) { + const int i = half_off + j; + const uint8_t code = rocmfpx_quantize_fp3_code(xb[i], inv_scale); + rocmfpx_set_bits(yb->qs, i*3, 3, code); + } + } + } +} + +void rocmfpx_dequantize_row_fp3(const block_rocmfp3 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { + assert(k % QK_ROCMFP3 == 0); + + const int64_t nb = k / QK_ROCMFP3; + for (int64_t ib = 0; ib < nb; ++ib) { + const block_rocmfp3 * xb = x + ib; + float * yb = y + ib*QK_ROCMFP3; + + for (int i = 0; i < QK_ROCMFP3; ++i) { + const float scale = rocmfpx_ue4m3_to_fp32(xb->e[i >= QK_ROCMFP3/2]); + const uint8_t code = (uint8_t) rocmfpx_get_bits(xb->qs, i*3, 3); + yb[i] = (float) rocmfpx_decode_fp3_code(code) * scale; + } + } +} + +void rocmfpx_quantize_row_fp3(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { + rocmfpx_quantize_row_fp3_ref(x, (block_rocmfp3 *) y, k); +} + +size_t rocmfpx_quantize_fp3(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix) { + const size_t row_size = rocmfpx_row_size_fp3(n_per_row); + char * qrow = (char *) dst; + + for (int64_t row = 0; row < nrows; ++row) { + if (imatrix) { + rocmfpx_quantize_row_fp3_weighted(src + row*n_per_row, (block_rocmfp3 *) qrow, n_per_row, imatrix); + } else { + rocmfpx_quantize_row_fp3_ref(src + row*n_per_row, (block_rocmfp3 *) qrow, n_per_row); + } + qrow += row_size; + } + + return (size_t) nrows * row_size; +} + +static void rocmfpx_quantize_row_fp6_weighted( + const float * GGML_RESTRICT x, block_rocmfp6 * GGML_RESTRICT y, int64_t k, const float * GGML_RESTRICT quant_weights) { + assert(k % QK_ROCMFP6 == 0); + + const float sigma2 = rocmfpx_row_sigma2(x, k); + + const int64_t nb = k / QK_ROCMFP6; + for (int64_t ib = 0; ib < nb; ++ib) { + const float * xb = x + ib*QK_ROCMFP6; + const float * qw = quant_weights ? quant_weights + ib*QK_ROCMFP6 : NULL; + block_rocmfp6 * yb = y + ib; + + memset(yb->qs, 0, sizeof(yb->qs)); + + for (int half = 0; half < 2; ++half) { + const int half_off = half*(QK_ROCMFP6/2); + const float * xh = xb + half_off; + const float * qh = qw ? qw + half_off : NULL; + yb->e[half] = qh ? + rocmfpx_choose_scale_fp6_weighted_mse(xh, QK_ROCMFP6/2, qh, sigma2) : + rocmfpx_choose_scale_fp6_mse(xh, QK_ROCMFP6/2); + + const float scale = rocmfpx_ue4m3_to_fp32(yb->e[half]); + const float inv_scale = scale > 0.0f ? 1.0f / scale : 0.0f; + + for (int j = 0; j < QK_ROCMFP6/2; ++j) { + const int i = half_off + j; + const uint8_t code = rocmfpx_quantize_fp6_code(xb[i], inv_scale); + rocmfpx_set_bits(yb->qs, i*6, 6, code); + } + } + } +} + +void rocmfpx_quantize_row_fp6_ref(const float * GGML_RESTRICT x, block_rocmfp6 * GGML_RESTRICT y, int64_t k) { + assert(k % QK_ROCMFP6 == 0); + + const int64_t nb = k / QK_ROCMFP6; + for (int64_t ib = 0; ib < nb; ++ib) { + const float * xb = x + ib*QK_ROCMFP6; + block_rocmfp6 * yb = y + ib; + + memset(yb->qs, 0, sizeof(yb->qs)); + + for (int half = 0; half < 2; ++half) { + const float * xh = xb + half*(QK_ROCMFP6/2); + yb->e[half] = rocmfpx_choose_scale_fp6_mse(xh, QK_ROCMFP6/2); + + const float scale = rocmfpx_ue4m3_to_fp32(yb->e[half]); + const float inv_scale = scale > 0.0f ? 1.0f / scale : 0.0f; + + for (int j = 0; j < QK_ROCMFP6/2; ++j) { + const int i = half*(QK_ROCMFP6/2) + j; + const uint8_t code = rocmfpx_quantize_fp6_code(xb[i], inv_scale); + rocmfpx_set_bits(yb->qs, i*6, 6, code); + } + } + } +} + +void rocmfpx_dequantize_row_fp6(const block_rocmfp6 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { + assert(k % QK_ROCMFP6 == 0); + + const int64_t nb = k / QK_ROCMFP6; + for (int64_t ib = 0; ib < nb; ++ib) { + const block_rocmfp6 * xb = x + ib; + float * yb = y + ib*QK_ROCMFP6; + + for (int i = 0; i < QK_ROCMFP6; ++i) { + const float scale = rocmfpx_ue4m3_to_fp32(xb->e[i >= QK_ROCMFP6/2]); + const uint8_t code = (uint8_t) rocmfpx_get_bits(xb->qs, i*6, 6); + yb[i] = (float) rocmfpx_decode_fp6_code(code) * scale; + } + } +} + +void rocmfpx_quantize_row_fp6(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { + rocmfpx_quantize_row_fp6_ref(x, (block_rocmfp6 *) y, k); +} + +size_t rocmfpx_quantize_fp6(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix) { + const size_t row_size = rocmfpx_row_size_fp6(n_per_row); + char * qrow = (char *) dst; + + for (int64_t row = 0; row < nrows; ++row) { + if (imatrix) { + rocmfpx_quantize_row_fp6_weighted(src + row*n_per_row, (block_rocmfp6 *) qrow, n_per_row, imatrix); + } else { + rocmfpx_quantize_row_fp6_ref(src + row*n_per_row, (block_rocmfp6 *) qrow, n_per_row); + } + qrow += row_size; + } + + return (size_t) nrows * row_size; +} + +static void rocmfpx_quantize_row_fp8_weighted( + const float * GGML_RESTRICT x, block_rocmfp8 * GGML_RESTRICT y, int64_t k, const float * GGML_RESTRICT quant_weights) { + assert(k % QK_ROCMFP8 == 0); + + const float sigma2 = rocmfpx_row_sigma2(x, k); + + const int64_t nb = k / QK_ROCMFP8; + for (int64_t ib = 0; ib < nb; ++ib) { + const float * xb = x + ib*QK_ROCMFP8; + const float * qw = quant_weights ? quant_weights + ib*QK_ROCMFP8 : NULL; + block_rocmfp8 * yb = y + ib; + + yb->e = qw ? rocmfpx_choose_scale_fp8_weighted_mse(xb, QK_ROCMFP8, qw, sigma2) : + rocmfpx_nearest_scale_ue4m3(rocmfpx_max_abs(xb, QK_ROCMFP8) / 127.0f); + + const float scale = rocmfpx_ue4m3_to_fp32(yb->e); + const float inv_scale = scale > 0.0f ? 1.0f / scale : 0.0f; + + for (int i = 0; i < QK_ROCMFP8; ++i) { + yb->qs[i] = rocmfpx_quantize_fp8_code(xb[i], inv_scale); + } + } +} + +void rocmfpx_quantize_row_fp8_ref(const float * GGML_RESTRICT x, block_rocmfp8 * GGML_RESTRICT y, int64_t k) { + assert(k % QK_ROCMFP8 == 0); + + const int64_t nb = k / QK_ROCMFP8; + for (int64_t ib = 0; ib < nb; ++ib) { + const float * xb = x + ib*QK_ROCMFP8; + block_rocmfp8 * yb = y + ib; + + const float max_abs = rocmfpx_max_abs(xb, QK_ROCMFP8); + yb->e = rocmfpx_nearest_scale_ue4m3(max_abs / 127.0f); + + const float scale = rocmfpx_ue4m3_to_fp32(yb->e); + const float inv_scale = scale > 0.0f ? 1.0f / scale : 0.0f; + + for (int i = 0; i < QK_ROCMFP8; ++i) { + yb->qs[i] = rocmfpx_quantize_fp8_code(xb[i], inv_scale); + } + } +} + +void rocmfpx_dequantize_row_fp8(const block_rocmfp8 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { + assert(k % QK_ROCMFP8 == 0); + + const int64_t nb = k / QK_ROCMFP8; + for (int64_t ib = 0; ib < nb; ++ib) { + const block_rocmfp8 * xb = x + ib; + float * yb = y + ib*QK_ROCMFP8; + + const float scale = rocmfpx_ue4m3_to_fp32(xb->e); + for (int i = 0; i < QK_ROCMFP8; ++i) { + yb[i] = (float) xb->qs[i] * scale; + } + } +} + +void rocmfpx_quantize_row_fp8(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { + rocmfpx_quantize_row_fp8_ref(x, (block_rocmfp8 *) y, k); +} + +size_t rocmfpx_quantize_fp8(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix) { + const size_t row_size = rocmfpx_row_size_fp8(n_per_row); + char * qrow = (char *) dst; + + for (int64_t row = 0; row < nrows; ++row) { + if (imatrix) { + rocmfpx_quantize_row_fp8_weighted(src + row*n_per_row, (block_rocmfp8 *) qrow, n_per_row, imatrix); + } else { + rocmfpx_quantize_row_fp8_ref(src + row*n_per_row, (block_rocmfp8 *) qrow, n_per_row); + } + qrow += row_size; + } + + return (size_t) nrows * row_size; +} + +bool rocmfpx_validate_row_data_fp2(const void * data, size_t nbytes) { + if (nbytes % sizeof(block_rocmfp2) != 0) { + return false; + } + + const block_rocmfp2 * blocks = (const block_rocmfp2 *) data; + const size_t nb = nbytes / sizeof(block_rocmfp2); + + for (size_t i = 0; i < nb; ++i) { + if (!rocmfpx_scale_is_valid(blocks[i].e[0]) || !rocmfpx_scale_is_valid(blocks[i].e[1])) { + return false; + } + } + + return true; +} + +bool rocmfpx_validate_row_data_fp3(const void * data, size_t nbytes) { + if (nbytes % sizeof(block_rocmfp3) != 0) { + return false; + } + + const block_rocmfp3 * blocks = (const block_rocmfp3 *) data; + const size_t nb = nbytes / sizeof(block_rocmfp3); + + for (size_t i = 0; i < nb; ++i) { + if (!rocmfpx_scale_is_valid(blocks[i].e[0]) || !rocmfpx_scale_is_valid(blocks[i].e[1])) { + return false; + } + } + + return true; +} + +bool rocmfpx_validate_row_data_fp6(const void * data, size_t nbytes) { + if (nbytes % sizeof(block_rocmfp6) != 0) { + return false; + } + + const block_rocmfp6 * blocks = (const block_rocmfp6 *) data; + const size_t nb = nbytes / sizeof(block_rocmfp6); + + for (size_t i = 0; i < nb; ++i) { + if (!rocmfpx_scale_is_valid(blocks[i].e[0]) || !rocmfpx_scale_is_valid(blocks[i].e[1])) { + return false; + } + } + + return true; +} + +bool rocmfpx_validate_row_data_fp8(const void * data, size_t nbytes) { + if (nbytes % sizeof(block_rocmfp8) != 0) { + return false; + } + + const block_rocmfp8 * blocks = (const block_rocmfp8 *) data; + const size_t nb = nbytes / sizeof(block_rocmfp8); + + for (size_t i = 0; i < nb; ++i) { + if (!rocmfpx_scale_is_valid(blocks[i].e)) { + return false; + } + } + + return true; +} diff --git a/server/deps/llama.cpp/ggml/rocmfpx/rocmfpx.h b/server/deps/llama.cpp/ggml/rocmfpx/rocmfpx.h new file mode 100644 index 000000000..f9d6cb896 --- /dev/null +++ b/server/deps/llama.cpp/ggml/rocmfpx/rocmfpx.h @@ -0,0 +1,111 @@ +#pragma once + +#include +#include +#include + +#include "ggml.h" + +#ifdef __cplusplus +extern "C" { +#endif + +#define QK_ROCMFPX 32 + +#define QK_ROCMFP2 QK_ROCMFPX +#define QK_ROCMFP3 QK_ROCMFPX +#define QK_ROCMFP6 QK_ROCMFPX +#define QK_ROCMFP8 QK_ROCMFPX + +#define QS_ROCMFP2 ((QK_ROCMFP2 * 2) / 8) +#define QS_ROCMFP3 ((QK_ROCMFP3 * 3) / 8) +#define QS_ROCMFP6 ((QK_ROCMFP6 * 6) / 8) +#define QS_ROCMFP8 QK_ROCMFP8 + +#define QR_ROCMFP2 1 +#define QI_ROCMFP2 (QS_ROCMFP2 / 4) // 8B qs = 2 int32; decoupled from QR_ROCMFP2 (QR stays 1 for convert.cu dequant) + +#define ROCMFP2_KVALUE_0_I8 (-1) +#define ROCMFP2_KVALUE_1_I8 0 +#define ROCMFP2_KVALUE_2_I8 1 +#define ROCMFP2_KVALUE_3_I8 2 +#define ROCMFP2_KVALUES_INIT { (float) ROCMFP2_KVALUE_0_I8, (float) ROCMFP2_KVALUE_1_I8, (float) ROCMFP2_KVALUE_2_I8, (float) ROCMFP2_KVALUE_3_I8 } + +#define QR_ROCMFP3 1 +#define QI_ROCMFP3 (QK_ROCMFP3 / (4 * QR_ROCMFP3)) + +#define QR_ROCMFP6 1 +#define QI_ROCMFP6 (QK_ROCMFP6 / (4 * QR_ROCMFP6)) + +#define QR_ROCMFP8 1 +#define QI_ROCMFP8 (QK_ROCMFP8 / (4 * QR_ROCMFP8)) + +// AMD-native experimental family layouts. The GGUF types are registered, but +// the layouts stay isolated from the promoted ROCmFP4 formats while evaluated. +typedef struct { + uint8_t qs[QS_ROCMFP2]; + uint8_t e[2]; +} block_rocmfp2; + +typedef struct { + uint8_t qs[QS_ROCMFP3]; + uint8_t e[2]; +} block_rocmfp3; + +typedef struct { + uint8_t qs[QS_ROCMFP6]; + uint8_t e[2]; +} block_rocmfp6; + +typedef struct { + int8_t qs[QS_ROCMFP8]; + uint8_t e; +} block_rocmfp8; + +#if defined(__cplusplus) +static_assert(sizeof(block_rocmfp2) == QS_ROCMFP2 + 2*sizeof(uint8_t), "wrong rocmfp2 block size/padding"); +static_assert(sizeof(block_rocmfp3) == QS_ROCMFP3 + 2*sizeof(uint8_t), "wrong rocmfp3 block size/padding"); +static_assert(sizeof(block_rocmfp6) == QS_ROCMFP6 + 2*sizeof(uint8_t), "wrong rocmfp6 block size/padding"); +static_assert(sizeof(block_rocmfp8) == QS_ROCMFP8 + sizeof(uint8_t), "wrong rocmfp8 block size/padding"); +#else +_Static_assert(sizeof(block_rocmfp2) == QS_ROCMFP2 + 2*sizeof(uint8_t), "wrong rocmfp2 block size/padding"); +_Static_assert(sizeof(block_rocmfp3) == QS_ROCMFP3 + 2*sizeof(uint8_t), "wrong rocmfp3 block size/padding"); +_Static_assert(sizeof(block_rocmfp6) == QS_ROCMFP6 + 2*sizeof(uint8_t), "wrong rocmfp6 block size/padding"); +_Static_assert(sizeof(block_rocmfp8) == QS_ROCMFP8 + sizeof(uint8_t), "wrong rocmfp8 block size/padding"); +#endif + +GGML_API float rocmfpx_ue4m3_to_fp32(uint8_t e); +GGML_API bool rocmfpx_scale_is_valid(uint8_t e); +GGML_API size_t rocmfpx_row_size_fp2(int64_t k); +GGML_API size_t rocmfpx_row_size_fp3(int64_t k); +GGML_API size_t rocmfpx_row_size_fp6(int64_t k); +GGML_API size_t rocmfpx_row_size_fp8(int64_t k); + +GGML_API void rocmfpx_quantize_row_fp2_ref(const float * GGML_RESTRICT x, block_rocmfp2 * GGML_RESTRICT y, int64_t k); +GGML_API void rocmfpx_dequantize_row_fp2(const block_rocmfp2 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API void rocmfpx_quantize_row_fp2(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +GGML_API size_t rocmfpx_quantize_fp2(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); + +GGML_API void rocmfpx_quantize_row_fp3_ref(const float * GGML_RESTRICT x, block_rocmfp3 * GGML_RESTRICT y, int64_t k); +GGML_API void rocmfpx_dequantize_row_fp3(const block_rocmfp3 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API void rocmfpx_quantize_row_fp3(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +GGML_API size_t rocmfpx_quantize_fp3(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); + +GGML_API void rocmfpx_quantize_row_fp6_ref(const float * GGML_RESTRICT x, block_rocmfp6 * GGML_RESTRICT y, int64_t k); +GGML_API void rocmfpx_dequantize_row_fp6(const block_rocmfp6 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API void rocmfpx_quantize_row_fp6(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +GGML_API size_t rocmfpx_quantize_fp6(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); + +GGML_API void rocmfpx_quantize_row_fp8_ref(const float * GGML_RESTRICT x, block_rocmfp8 * GGML_RESTRICT y, int64_t k); +GGML_API void rocmfpx_dequantize_row_fp8(const block_rocmfp8 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API void rocmfpx_quantize_row_fp8(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +GGML_API size_t rocmfpx_quantize_fp8(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); + +GGML_API bool rocmfpx_validate_row_data_fp2(const void * data, size_t nbytes); +GGML_API bool rocmfpx_validate_row_data_fp3(const void * data, size_t nbytes); +GGML_API bool rocmfpx_validate_row_data_fp6(const void * data, size_t nbytes); +GGML_API bool rocmfpx_validate_row_data_fp8(const void * data, size_t nbytes); + +#ifdef __cplusplus +} +#endif diff --git a/server/deps/llama.cpp/ggml/rocmfpx/test_rocmfpx.c b/server/deps/llama.cpp/ggml/rocmfpx/test_rocmfpx.c new file mode 100644 index 000000000..a3c49ffb3 --- /dev/null +++ b/server/deps/llama.cpp/ggml/rocmfpx/test_rocmfpx.c @@ -0,0 +1,196 @@ +#include "rocmfpx.h" + +#include +#include +#include +#include + +static void fill_row(float * x, int n) { + assert(n >= 44); + + for (int i = 0; i < n; ++i) { + const float wave = 0.75f*sinf((float) i * 0.37f) + 0.25f*cosf((float) i * 0.13f); + const float ramp = ((float) (i % 11) - 5.0f) * 0.035f; + x[i] = wave + ramp; + } + + x[7] = 3.25f; + x[19] = -2.75f; + x[43] = 1.875f; +} + +static float mse(const float * a, const float * b, int n) { + float err = 0.0f; + + for (int i = 0; i < n; ++i) { + const float d = a[i] - b[i]; + err += d*d; + } + + return err / (float) n; +} + +static float weighted_mse(const float * a, const float * b, const float * w, int n) { + float err = 0.0f; + float sum_w = 0.0f; + + for (int i = 0; i < n; ++i) { + const float d = a[i] - b[i]; + err += w[i]*d*d; + sum_w += w[i]; + } + + return sum_w > 0.0f ? err / sum_w : 0.0f; +} + +static void check_weighted_imatrix_fp3(void) { + enum { N = QK_ROCMFP3 }; + + float src[N]; + float imatrix[N]; + float plain[N]; + float weighted[N]; + block_rocmfp3 q_plain[N / QK_ROCMFP3]; + block_rocmfp3 q_weighted[N / QK_ROCMFP3]; + + for (int i = 0; i < N; ++i) { + src[i] = (i % 2) ? 0.21f : -0.21f; + imatrix[i] = 100.0f; + } + + src[0] = 9.0f; + imatrix[0] = 0.0f; + + rocmfpx_quantize_fp3(src, q_plain, 1, N, NULL); + rocmfpx_quantize_fp3(src, q_weighted, 1, N, imatrix); + rocmfpx_dequantize_row_fp3(q_plain, plain, N); + rocmfpx_dequantize_row_fp3(q_weighted, weighted, N); + + const float plain_err = weighted_mse(src, plain, imatrix, N); + const float weighted_err = weighted_mse(src, weighted, imatrix, N); + + printf("ROCmFP3 imatrix weighted_mse: plain=%g weighted=%g\n", plain_err, weighted_err); + assert(weighted_err < plain_err); +} + +static void check_weighted_imatrix_fp2(void) { + enum { N = QK_ROCMFP2 }; + + float src[N]; + float imatrix[N]; + float plain[N]; + float weighted[N]; + block_rocmfp2 q_plain[N / QK_ROCMFP2]; + block_rocmfp2 q_weighted[N / QK_ROCMFP2]; + + for (int i = 0; i < N; ++i) { + src[i] = (i % 2) ? 0.21f : -0.21f; + imatrix[i] = 100.0f; + } + + src[0] = 9.0f; + imatrix[0] = 0.0f; + + rocmfpx_quantize_fp2(src, q_plain, 1, N, NULL); + rocmfpx_quantize_fp2(src, q_weighted, 1, N, imatrix); + rocmfpx_dequantize_row_fp2(q_plain, plain, N); + rocmfpx_dequantize_row_fp2(q_weighted, weighted, N); + + const float plain_err = weighted_mse(src, plain, imatrix, N); + const float weighted_err = weighted_mse(src, weighted, imatrix, N); + + printf("ROCmFP2 imatrix weighted_mse: plain=%g weighted=%g\n", plain_err, weighted_err); + assert(weighted_err < plain_err); +} + +static void check_large_finite_values(void) { + float src[QK_ROCMFPX] = { 0 }; + float imatrix[QK_ROCMFPX]; + float fp6[QK_ROCMFP6]; + float fp8[QK_ROCMFP8]; + block_rocmfp6 q6; + block_rocmfp8 q8; + + src[0] = FLT_MAX; + src[1] = -FLT_MAX; + for (int i = 0; i < QK_ROCMFPX; ++i) { + imatrix[i] = 1.0f; + } + + rocmfpx_quantize_fp6(src, &q6, 1, QK_ROCMFP6, imatrix); + rocmfpx_quantize_fp8(src, &q8, 1, QK_ROCMFP8, imatrix); + rocmfpx_dequantize_row_fp6(&q6, fp6, QK_ROCMFP6); + rocmfpx_dequantize_row_fp8(&q8, fp8, QK_ROCMFP8); + + assert(rocmfpx_validate_row_data_fp6(&q6, sizeof(q6))); + assert(rocmfpx_validate_row_data_fp8(&q8, sizeof(q8))); + assert(isfinite(fp6[0]) && isfinite(fp6[1])); + assert(isfinite(fp8[0]) && isfinite(fp8[1])); + assert(fp6[0] > 0.0f && fp6[1] < 0.0f); + assert(fp8[0] > 0.0f && fp8[1] < 0.0f); +} + +int main(void) { + enum { N = 64 }; + + float src[N]; + float fp2[N]; + float fp3[N]; + float fp6[N]; + float fp8[N]; + + block_rocmfp2 q2[N / QK_ROCMFP2]; + block_rocmfp3 q3[N / QK_ROCMFP3]; + block_rocmfp6 q6[N / QK_ROCMFP6]; + block_rocmfp8 q8[N / QK_ROCMFP8]; + + fill_row(src, N); + + rocmfpx_quantize_row_fp2_ref(src, q2, N); + rocmfpx_quantize_row_fp3_ref(src, q3, N); + rocmfpx_quantize_row_fp6_ref(src, q6, N); + rocmfpx_quantize_row_fp8_ref(src, q8, N); + + assert(rocmfpx_validate_row_data_fp2(q2, sizeof(q2))); + assert(rocmfpx_validate_row_data_fp3(q3, sizeof(q3))); + assert(rocmfpx_validate_row_data_fp6(q6, sizeof(q6))); + assert(rocmfpx_validate_row_data_fp8(q8, sizeof(q8))); + + rocmfpx_dequantize_row_fp2(q2, fp2, N); + rocmfpx_dequantize_row_fp3(q3, fp3, N); + rocmfpx_dequantize_row_fp6(q6, fp6, N); + rocmfpx_dequantize_row_fp8(q8, fp8, N); + + const float mse2 = mse(src, fp2, N); + const float mse3 = mse(src, fp3, N); + const float mse6 = mse(src, fp6, N); + const float mse8 = mse(src, fp8, N); + + printf("ROCmFP2: block=%zu row=%zu bpw=%.2f mse=%g\n", + sizeof(block_rocmfp2), rocmfpx_row_size_fp2(N), + 8.0f*(float) sizeof(block_rocmfp2)/(float) QK_ROCMFP2, mse2); + printf("ROCmFP3: block=%zu row=%zu bpw=%.2f mse=%g\n", + sizeof(block_rocmfp3), rocmfpx_row_size_fp3(N), + 8.0f*(float) sizeof(block_rocmfp3)/(float) QK_ROCMFP3, mse3); + printf("ROCmFP6: block=%zu row=%zu bpw=%.2f mse=%g\n", + sizeof(block_rocmfp6), rocmfpx_row_size_fp6(N), + 8.0f*(float) sizeof(block_rocmfp6)/(float) QK_ROCMFP6, mse6); + printf("ROCmFP8: block=%zu row=%zu bpw=%.2f mse=%g\n", + sizeof(block_rocmfp8), rocmfpx_row_size_fp8(N), + 8.0f*(float) sizeof(block_rocmfp8)/(float) QK_ROCMFP8, mse8); + + assert(sizeof(block_rocmfp2) == 10); + assert(isfinite(mse2)); + assert(isfinite(mse3)); + assert(isfinite(mse6)); + assert(isfinite(mse8)); + assert(mse8 < mse6); + assert(mse6 < mse3); + assert(mse3 < mse2); + + check_weighted_imatrix_fp2(); + check_weighted_imatrix_fp3(); + check_large_finite_values(); + + return 0; +} diff --git a/server/deps/llama.cpp/ggml/src/CMakeLists.txt b/server/deps/llama.cpp/ggml/src/CMakeLists.txt index 48fbe208d..315eafe2e 100644 --- a/server/deps/llama.cpp/ggml/src/CMakeLists.txt +++ b/server/deps/llama.cpp/ggml/src/CMakeLists.txt @@ -205,6 +205,10 @@ add_library(ggml-base ggml-threading.cpp ggml-threading.h ggml-quants.c + ../rocmfp4/rocmfp4.c + ../rocmfp4/rocmfp4.h + ../rocmfpx/rocmfpx.c + ../rocmfpx/rocmfpx.h ggml-quants.h gguf.cpp) diff --git a/server/deps/llama.cpp/ggml/src/ggml-common.h b/server/deps/llama.cpp/ggml/src/ggml-common.h index d47c97696..e52a2005d 100644 --- a/server/deps/llama.cpp/ggml/src/ggml-common.h +++ b/server/deps/llama.cpp/ggml/src/ggml-common.h @@ -1132,6 +1132,12 @@ GGML_TABLE_BEGIN(int8_t, kvalues_mxfp4, 16) 0, 1, 2, 3, 4, 6, 8, 12, 0, -1, -2, -3, -4, -6, -8, -12, GGML_TABLE_END() +// ROCmFP4 uses an E2M1-derived value set with the largest level retuned from +// 12 to 10, plus dual half-block UE4M3 scales. +GGML_TABLE_BEGIN(int8_t, kvalues_rocmfp4, 16) + 0, 1, 2, 3, 4, 6, 8, 10, 0, -1, -2, -3, -4, -6, -8, -10, +GGML_TABLE_END() + #define NGRID_IQ1S 2048 #define IQ1S_DELTA 0.125f #define IQ1M_DELTA 0.125f diff --git a/server/deps/llama.cpp/ggml/src/ggml-cpu/ggml-cpu.c b/server/deps/llama.cpp/ggml/src/ggml-cpu/ggml-cpu.c index f0e14a006..df6ee5144 100644 --- a/server/deps/llama.cpp/ggml/src/ggml-cpu/ggml-cpu.c +++ b/server/deps/llama.cpp/ggml/src/ggml-cpu/ggml-cpu.c @@ -1826,6 +1826,10 @@ static void ggml_compute_forward(struct ggml_compute_params * params, struct ggm { GGML_ABORT("GGML_OP_MOE_FUSED is only implemented for CUDA"); } + case GGML_OP_DS4_HC: + { + GGML_ABORT("GGML_OP_DS4_HC is only implemented for CUDA"); + } case GGML_OP_OUT_PROD: { ggml_compute_forward_out_prod(params, tensor); @@ -2280,6 +2284,7 @@ static int ggml_get_n_tasks(struct ggml_tensor * node, int n_threads) { case GGML_GLU_OP_GEGLU: case GGML_GLU_OP_SWIGLU: case GGML_GLU_OP_SWIGLU_OAI: + case GGML_GLU_OP_SWIGLU_DS4: case GGML_GLU_OP_GEGLU_ERF: case GGML_GLU_OP_GEGLU_QUICK: { diff --git a/server/deps/llama.cpp/ggml/src/ggml-cpu/ops.cpp b/server/deps/llama.cpp/ggml/src/ggml-cpu/ops.cpp index c8f811406..6dff19338 100644 --- a/server/deps/llama.cpp/ggml/src/ggml-cpu/ops.cpp +++ b/server/deps/llama.cpp/ggml/src/ggml-cpu/ops.cpp @@ -3271,6 +3271,80 @@ static void ggml_compute_forward_swiglu( } } +// ggml_compute_forward_swiglu_ds4 + +static void ggml_compute_forward_swiglu_ds4_f32( + const ggml_compute_params * params, + ggml_tensor * dst) { + + const ggml_tensor * src0 = dst->src[0]; + const ggml_tensor * src1 = dst->src[1]; + char * src0_d = (char *) src0->data; + char * src1_d = (char *) (src1 ? src1->data : src0->data); + const size_t src0_o = src0->nb[1]; + const size_t src1_o = src1 ? src1->nb[1] : src0->nb[1]; + + GGML_ASSERT(ggml_is_contiguous_1(src0)); + GGML_ASSERT(ggml_is_contiguous_1(dst)); + GGML_ASSERT(src1); + GGML_ASSERT(ggml_is_contiguous_1(src1)); + GGML_ASSERT(src0->type == src1->type); + + const int ith = params->ith; + const int nth = params->nth; + + const int nc = src0->ne[0]; + const int nr = ggml_nrows(src0); + + GGML_ASSERT(dst->ne[0] == nc); + GGML_ASSERT(ggml_nrows(dst) == nr); + + const float clamp = ggml_get_op_params_f32(dst, 2); + + const int dr = (nr + nth - 1)/nth; + const int ir0 = dr*ith; + const int ir1 = MIN(ir0 + dr, nr); + + for (int i1 = ir0; i1 < ir1; i1++) { + float * gate_p = (float *) (src0_d + i1*src0_o); + float * up_p = (float *) (src1_d + i1*src1_o); + float * dst_p = (float *) ((char *) dst->data + i1*(dst->nb[1])); + + for (int k = 0; k < nc; k++) { + const float gate = std::min(gate_p[k], clamp); + const float up = std::clamp(up_p[k], -clamp, clamp); + dst_p[k] = up * gate / (1.f + expf(-gate)); + } + +#ifndef NDEBUG + for (int k = 0; k < nc; k++) { + const float x = dst_p[k]; + GGML_UNUSED(x); + assert(!isnan(x)); + assert(!isinf(x)); + } +#endif // NDEBUG + } +} + +static void ggml_compute_forward_swiglu_ds4( + const ggml_compute_params * params, + ggml_tensor * dst) { + + const ggml_tensor * src0 = dst->src[0]; + + switch (src0->type) { + case GGML_TYPE_F32: + { + ggml_compute_forward_swiglu_ds4_f32(params, dst); + } break; + default: + { + GGML_ABORT("fatal error"); + } + } +} + // ggml_compute_forward_swiglu_oai static void ggml_compute_forward_swiglu_oai_f32( @@ -9801,6 +9875,10 @@ void ggml_compute_forward_glu( { ggml_compute_forward_swiglu_oai(params, dst); } break; + case GGML_GLU_OP_SWIGLU_DS4: + { + ggml_compute_forward_swiglu_ds4(params, dst); + } break; case GGML_GLU_OP_GEGLU_ERF: { ggml_compute_forward_geglu_erf(params, dst); diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/common.cuh b/server/deps/llama.cpp/ggml/src/ggml-cuda/common.cuh index 9df1461a5..59e06c049 100644 --- a/server/deps/llama.cpp/ggml/src/ggml-cuda/common.cuh +++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/common.cuh @@ -19,6 +19,8 @@ #endif #endif #include "ggml-common.h" +#include "../../rocmfp4/rocmfp4.h" +#include "../../rocmfpx/rocmfpx.h" #include #include @@ -959,6 +961,48 @@ struct ggml_cuda_type_traits { static constexpr int qi = QI8_0; }; +template<> +struct ggml_cuda_type_traits { + static constexpr int qk = QK_ROCMFP4; + static constexpr int qr = QR_ROCMFP4; + static constexpr int qi = QI_ROCMFP4; +}; + +template<> +struct ggml_cuda_type_traits { + static constexpr int qk = QK_ROCMFP4; + static constexpr int qr = QR_ROCMFP4; + static constexpr int qi = QI_ROCMFP4; +}; + +template<> +struct ggml_cuda_type_traits { + static constexpr int qk = QK_ROCMFP2; + static constexpr int qr = QR_ROCMFP2; + static constexpr int qi = QI_ROCMFP2; +}; + +template<> +struct ggml_cuda_type_traits { + static constexpr int qk = QK_ROCMFP3; + static constexpr int qr = QR_ROCMFP3; + static constexpr int qi = QI_ROCMFP3; +}; + +template<> +struct ggml_cuda_type_traits { + static constexpr int qk = QK_ROCMFP6; + static constexpr int qr = QR_ROCMFP6; + static constexpr int qi = QI_ROCMFP6; +}; + +template<> +struct ggml_cuda_type_traits { + static constexpr int qk = QK_ROCMFP8; + static constexpr int qr = QR_ROCMFP8; + static constexpr int qi = QI_ROCMFP8; +}; + template<> struct ggml_cuda_type_traits { static constexpr int qk = QK_MXFP4; @@ -1466,10 +1510,14 @@ struct ggml_cuda_mm_fusion_args_host { const ggml_tensor * gate = nullptr; const ggml_tensor * gate_bias = nullptr; ggml_glu_op glu_op; + float glu_param0 = 0.0f; + float glu_param1 = 0.0f; }; struct ggml_cuda_mm_fusion_args_device { const void * x_bias = nullptr; const void * gate = nullptr; const void * gate_bias = nullptr; ggml_glu_op glu_op; + float glu_param0 = 0.0f; + float glu_param1 = 0.0f; }; diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/concat.cu b/server/deps/llama.cpp/ggml/src/ggml-cuda/concat.cu index e9ffd274b..51a3efa53 100644 --- a/server/deps/llama.cpp/ggml/src/ggml-cuda/concat.cu +++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/concat.cu @@ -1,7 +1,8 @@ #include "concat.cuh" // contiguous kernels -static __global__ void concat_f32_dim0(const float * x, const float * y, float * dst, const int ne0, const int ne00) { +template +static __global__ void concat_dim0(const T * x, const T * y, T * dst, const int ne0, const int ne00) { int nidx = threadIdx.x + blockIdx.x * blockDim.x; if (nidx >= ne0) { return; @@ -27,7 +28,8 @@ static __global__ void concat_f32_dim0(const float * x, const float * y, float * } } -static __global__ void concat_f32_dim1(const float * x, const float * y, float * dst, const int ne0, const int ne01) { +template +static __global__ void concat_dim1(const T * x, const T * y, T * dst, const int ne0, const int ne01) { int nidx = threadIdx.x + blockIdx.x * blockDim.x; if (nidx >= ne0) { return; @@ -53,7 +55,8 @@ static __global__ void concat_f32_dim1(const float * x, const float * y, float * } } -static __global__ void concat_f32_dim2(const float * x, const float * y, float * dst, const int ne0, const int ne02) { +template +static __global__ void concat_dim2(const T * x, const T * y, T * dst, const int ne0, const int ne02) { int nidx = threadIdx.x + blockIdx.x * blockDim.x; if (nidx >= ne0) { return; @@ -79,24 +82,25 @@ static __global__ void concat_f32_dim2(const float * x, const float * y, float * } } -static void concat_f32_cuda(const float * x, const float * y, float * dst, int ne00, int ne01, int ne02, int ne0, int ne1, int ne2, int dim, cudaStream_t stream) { +template +static void concat_cuda(const T * x, const T * y, T * dst, int ne00, int ne01, int ne02, int ne0, int ne1, int ne2, int dim, cudaStream_t stream) { int num_blocks = (ne0 + CUDA_CONCAT_BLOCK_SIZE - 1) / CUDA_CONCAT_BLOCK_SIZE; dim3 gridDim(num_blocks, ne1, ne2); if (dim == 0) { - concat_f32_dim0<<>>(x, y, dst, ne0, ne00); + concat_dim0<<>>(x, y, dst, ne0, ne00); return; } if (dim == 1) { - concat_f32_dim1<<>>(x, y, dst, ne0, ne01); + concat_dim1<<>>(x, y, dst, ne0, ne01); return; } - concat_f32_dim2<<>>(x, y, dst, ne0, ne02); + concat_dim2<<>>(x, y, dst, ne0, ne02); } // non-contiguous kernel (slow) -template +template static __global__ void __launch_bounds__(CUDA_CONCAT_BLOCK_SIZE) - concat_f32_non_cont( + concat_non_cont( const char * src0, const char * src1, char * dst, @@ -130,54 +134,44 @@ static __global__ void __launch_bounds__(CUDA_CONCAT_BLOCK_SIZE) const int64_t i2 = blockIdx.y; const int64_t i1 = blockIdx.x; - const float * x; + const T * x; for (int64_t i0 = threadIdx.x; i0 < ne0; i0 += blockDim.x) { if (i0 < ne00 && i1 < ne01 && i2 < ne02 && i3 < ne03) { - x = (const float *)(src0 + (i3 )*nb03 + (i2 )*nb02 + (i1 )*nb01 + (i0 )*nb00); + x = (const T *)(src0 + (i3 )*nb03 + (i2 )*nb02 + (i1 )*nb01 + (i0 )*nb00); } else { if constexpr (dim == 0) { - x = (const float *) (src1 + i3 * nb13 + i2 * nb12 + i1 * nb11 + (i0 - ne00) * nb10); + x = (const T *) (src1 + i3 * nb13 + i2 * nb12 + i1 * nb11 + (i0 - ne00) * nb10); } else if constexpr (dim == 1) { - x = (const float *) (src1 + i3 * nb13 + i2 * nb12 + (i1 - ne01) * nb11 + i0 * nb10); + x = (const T *) (src1 + i3 * nb13 + i2 * nb12 + (i1 - ne01) * nb11 + i0 * nb10); } else if constexpr (dim == 2) { - x = (const float *) (src1 + i3 * nb13 + (i2 - ne02) * nb12 + i1 * nb11 + i0 * nb10); + x = (const T *) (src1 + i3 * nb13 + (i2 - ne02) * nb12 + i1 * nb11 + i0 * nb10); } else if constexpr (dim == 3) { - x = (const float *) (src1 + (i3 - ne03) * nb13 + i2 * nb12 + i1 * nb11 + i0 * nb10); + x = (const T *) (src1 + (i3 - ne03) * nb13 + i2 * nb12 + i1 * nb11 + i0 * nb10); } } - float * y = (float *)(dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0); + T * y = (T *)(dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0); *y = *x; } } - -void ggml_cuda_op_concat(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { - const ggml_tensor * src0 = dst->src[0]; - const ggml_tensor * src1 = dst->src[1]; - +template +static void concat_cuda_typed(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst, int dim) { cudaStream_t stream = ctx.stream(); - const int32_t dim = ((int32_t *) dst->op_params)[0]; - - GGML_ASSERT(src0->type == GGML_TYPE_F32); - GGML_ASSERT(src1->type == GGML_TYPE_F32); - GGML_ASSERT(dst->type == GGML_TYPE_F32); - if (ggml_is_contiguous(src0) && ggml_is_contiguous(src1)) { - const float * src0_d = (const float *)src0->data; - const float * src1_d = (const float *)src1->data; - - float * dst_d = (float *)dst->data; + const T * src0_d = (const T *)src0->data; + const T * src1_d = (const T *)src1->data; + T * dst_d = (T *)dst->data; if (dim != 3) { for (int i3 = 0; i3 < dst->ne[3]; i3++) { - concat_f32_cuda( - src0_d + i3 * (src0->nb[3] / 4), - src1_d + i3 * (src1->nb[3] / 4), - dst_d + i3 * ( dst->nb[3] / 4), + concat_cuda( + src0_d + i3 * (src0->nb[3] / sizeof(T)), + src1_d + i3 * (src1->nb[3] / sizeof(T)), + dst_d + i3 * ( dst->nb[3] / sizeof(T)), src0->ne[0], src0->ne[1], src0->ne[2], dst->ne[0], dst->ne[1], dst->ne[2], dim, stream); } @@ -185,13 +179,13 @@ void ggml_cuda_op_concat(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { const size_t size0 = ggml_nbytes(src0); const size_t size1 = ggml_nbytes(src1); - CUDA_CHECK(cudaMemcpyAsync(dst_d, src0_d, size0, cudaMemcpyDeviceToDevice, stream)); - CUDA_CHECK(cudaMemcpyAsync(dst_d + size0/4, src1_d, size1, cudaMemcpyDeviceToDevice, stream)); + CUDA_CHECK(cudaMemcpyAsync(dst_d, src0_d, size0, cudaMemcpyDeviceToDevice, stream)); + CUDA_CHECK(cudaMemcpyAsync((char *)dst_d + size0, src1_d, size1, cudaMemcpyDeviceToDevice, stream)); } } else { dim3 grid_dim(dst->ne[1], dst->ne[2], dst->ne[3]); auto launch_kernel = [&](auto dim) { - concat_f32_non_cont<<>>( + concat_non_cont<<>>( (const char *) src0->data, (const char *) src1->data, (char *) dst->data, src0->ne[0], src0->ne[1], src0->ne[2], src0->ne[3], src0->nb[0], src0->nb[1], src0->nb[2], src0->nb[3], @@ -219,3 +213,30 @@ void ggml_cuda_op_concat(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { } } } + +void ggml_cuda_op_concat(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { + const ggml_tensor * src0 = dst->src[0]; + const ggml_tensor * src1 = dst->src[1]; + + const int32_t dim = ((int32_t *) dst->op_params)[0]; + + GGML_ASSERT(src0->type == src1->type); + GGML_ASSERT(src0->type == dst->type); + + switch (src0->type) { + case GGML_TYPE_F32: + concat_cuda_typed(ctx, src0, src1, dst, dim); + break; + case GGML_TYPE_F16: + concat_cuda_typed(ctx, src0, src1, dst, dim); + break; + case GGML_TYPE_BF16: + concat_cuda_typed(ctx, src0, src1, dst, dim); + break; + case GGML_TYPE_I8: + concat_cuda_typed(ctx, src0, src1, dst, dim); + break; + default: + GGML_ABORT("unsupported concat type %s", ggml_type_name(src0->type)); + } +} diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/convert.cu b/server/deps/llama.cpp/ggml/src/ggml-cuda/convert.cu index c675034af..734cbb42a 100644 --- a/server/deps/llama.cpp/ggml/src/ggml-cuda/convert.cu +++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/convert.cu @@ -1,6 +1,7 @@ #include "convert.cuh" #include "dequantize.cuh" #include "tq3-quant.cuh" +#include "../../rocmfp4/rocmfp4_hip_scale.cuh" #include @@ -487,6 +488,80 @@ static __global__ void dequantize_block_mxfp4(const void * __restrict__ vx, dst_ } } +template +static __global__ void dequantize_block_rocmfp4(const void * __restrict__ vx, dst_t * __restrict__ yy) { + + const int64_t i = blockIdx.x; + const block_rocmfp4 * x = (const block_rocmfp4 *) vx + i*(QK_K/QK_ROCMFP4); + + const int64_t tid = threadIdx.x; + const int64_t il = tid/8; // 0...3 + const int64_t ib = tid%8; // 0...7 + dst_t * y = yy + i*QK_K + 32*ib + 4*il; + const uint8_t * q4 = x[ib].qs + 4*il; + const float d0 = rocmfp4_ue4m3_to_fp32_half_finite(x[ib].e[0]); + const float d1 = rocmfp4_ue4m3_to_fp32_half_finite(x[ib].e[1]); + for (int j = 0; j < 4; ++j) { + y[j+ 0] = d0 * rocmfp4_decode_i8(q4[j]); + y[j+16] = d1 * rocmfp4_decode_i8(q4[j] >> 4); + } +} + +template +static __global__ void dequantize_block_rocmfp4_fast(const void * __restrict__ vx, dst_t * __restrict__ yy) { + + const int64_t i = blockIdx.x; + const block_rocmfp4_fast * x = (const block_rocmfp4_fast *) vx + i*(QK_K/QK_ROCMFP4); + + const int64_t tid = threadIdx.x; + const int64_t il = tid/8; // 0...3 + const int64_t ib = tid%8; // 0...7 + dst_t * y = yy + i*QK_K + 32*ib + 4*il; + const uint8_t * q4 = x[ib].qs + 4*il; + const float d = rocmfp4_ue4m3_to_fp32_half_finite(x[ib].e); + for (int j = 0; j < 4; ++j) { + y[j+ 0] = d * rocmfp4_decode_i8(q4[j]); + y[j+16] = d * rocmfp4_decode_i8(q4[j] >> 4); + } +} + +template +static __global__ void dequantize_block_rocmfp4_tail( + const block_rocmfp4 * __restrict__ x, + dst_t * __restrict__ y) { + const int64_t ib = blockIdx.x; + const int64_t j = threadIdx.x; + + if (j >= QK_ROCMFP4/2) { + return; + } + + const float d0 = rocmfp4_ue4m3_to_fp32_half_finite(x[ib].e[0]); + const float d1 = rocmfp4_ue4m3_to_fp32_half_finite(x[ib].e[1]); + const uint8_t q = x[ib].qs[j]; + + y[ib*QK_ROCMFP4 + j] = d0 * rocmfp4_decode_i8(q); + y[ib*QK_ROCMFP4 + j + QK_ROCMFP4/2] = d1 * rocmfp4_decode_i8(q >> 4); +} + +template +static __global__ void dequantize_block_rocmfp4_fast_tail( + const block_rocmfp4_fast * __restrict__ x, + dst_t * __restrict__ y) { + const int64_t ib = blockIdx.x; + const int64_t j = threadIdx.x; + + if (j >= QK_ROCMFP4/2) { + return; + } + + const float d = rocmfp4_ue4m3_to_fp32_half_finite(x[ib].e); + const uint8_t q = x[ib].qs[j]; + + y[ib*QK_ROCMFP4 + j] = d * rocmfp4_decode_i8(q); + y[ib*QK_ROCMFP4 + j + QK_ROCMFP4/2] = d * rocmfp4_decode_i8(q >> 4); +} + template static void dequantize_block_cuda(const void * vx, dst_t * y, const int64_t ne00, const int64_t ne01, const int64_t ne02, const int64_t ne03, @@ -645,6 +720,30 @@ static void dequantize_row_mxfp4_cuda(const void * vx, dst_t * y, const int64_t dequantize_block_mxfp4<<>>(vx, y); } +template +static void dequantize_row_rocmfp4_hip(const void * vx, dst_t * y, const int64_t k, cudaStream_t stream) { + GGML_ASSERT(k % QK_ROCMFP4 == 0); + if (k % QK_K == 0) { + dequantize_block_rocmfp4<<>>(vx, y); + return; + } + + dequantize_block_rocmfp4_tail<<>>( + (const block_rocmfp4 *) vx, y); +} + +template +static void dequantize_row_rocmfp4_fast_hip(const void * vx, dst_t * y, const int64_t k, cudaStream_t stream) { + GGML_ASSERT(k % QK_ROCMFP4 == 0); + if (k % QK_K == 0) { + dequantize_block_rocmfp4_fast<<>>(vx, y); + return; + } + + dequantize_block_rocmfp4_fast_tail<<>>( + (const block_rocmfp4_fast *) vx, y); +} + template static __global__ void dequantize_block_nvfp4( const void * __restrict__ vx, @@ -786,6 +885,18 @@ to_fp16_cuda_t ggml_get_to_fp16_cuda(ggml_type type) { return dequantize_row_mxfp4_cuda; case GGML_TYPE_NVFP4: return dequantize_row_nvfp4_cuda; + case GGML_TYPE_Q4_0_ROCMFP4: + return dequantize_row_rocmfp4_hip; + case GGML_TYPE_Q4_0_ROCMFP4_FAST: + return dequantize_row_rocmfp4_fast_hip; + case GGML_TYPE_Q2_0_ROCMFP2: + return dequantize_block_cont_cuda; + case GGML_TYPE_Q3_0_ROCMFPX: + return dequantize_block_cont_cuda; + case GGML_TYPE_Q6_0_ROCMFPX: + return dequantize_block_cont_cuda; + case GGML_TYPE_Q8_0_ROCMFPX: + return dequantize_block_cont_cuda; case GGML_TYPE_F32: return convert_unary_cont_cuda; case GGML_TYPE_BF16: @@ -841,6 +952,18 @@ to_fp32_cuda_t ggml_get_to_fp32_cuda(ggml_type type) { return dequantize_row_mxfp4_cuda; case GGML_TYPE_NVFP4: return dequantize_row_nvfp4_cuda; + case GGML_TYPE_Q4_0_ROCMFP4: + return dequantize_row_rocmfp4_hip; + case GGML_TYPE_Q4_0_ROCMFP4_FAST: + return dequantize_row_rocmfp4_fast_hip; + case GGML_TYPE_Q2_0_ROCMFP2: + return dequantize_block_cont_cuda; + case GGML_TYPE_Q3_0_ROCMFPX: + return dequantize_block_cont_cuda; + case GGML_TYPE_Q6_0_ROCMFPX: + return dequantize_block_cont_cuda; + case GGML_TYPE_Q8_0_ROCMFPX: + return dequantize_block_cont_cuda; case GGML_TYPE_F16: return convert_unary_cont_cuda; case GGML_TYPE_BF16: @@ -864,6 +987,18 @@ to_fp16_nc_cuda_t ggml_get_to_fp16_nc_cuda(ggml_type type) { return dequantize_block_cuda; case GGML_TYPE_Q8_0: return dequantize_block_cuda; + case GGML_TYPE_Q4_0_ROCMFP4: + return dequantize_block_cuda; + case GGML_TYPE_Q4_0_ROCMFP4_FAST: + return dequantize_block_cuda; + case GGML_TYPE_Q2_0_ROCMFP2: + return dequantize_block_cuda; + case GGML_TYPE_Q3_0_ROCMFPX: + return dequantize_block_cuda; + case GGML_TYPE_Q6_0_ROCMFPX: + return dequantize_block_cuda; + case GGML_TYPE_Q8_0_ROCMFPX: + return dequantize_block_cuda; case GGML_TYPE_BF16: return convert_unary_cuda; default: @@ -885,6 +1020,18 @@ to_bf16_nc_cuda_t ggml_get_to_bf16_nc_cuda(ggml_type type) { return dequantize_block_cuda; case GGML_TYPE_Q8_0: return dequantize_block_cuda; + case GGML_TYPE_Q4_0_ROCMFP4: + return dequantize_block_cuda; + case GGML_TYPE_Q4_0_ROCMFP4_FAST: + return dequantize_block_cuda; + case GGML_TYPE_Q2_0_ROCMFP2: + return dequantize_block_cuda; + case GGML_TYPE_Q3_0_ROCMFPX: + return dequantize_block_cuda; + case GGML_TYPE_Q6_0_ROCMFPX: + return dequantize_block_cuda; + case GGML_TYPE_Q8_0_ROCMFPX: + return dequantize_block_cuda; case GGML_TYPE_F16: return convert_unary_cuda; default: @@ -906,6 +1053,18 @@ to_fp32_nc_cuda_t ggml_get_to_fp32_nc_cuda(ggml_type type) { return dequantize_block_cuda; case GGML_TYPE_Q8_0: return dequantize_block_cuda; + case GGML_TYPE_Q4_0_ROCMFP4: + return dequantize_block_cuda; + case GGML_TYPE_Q4_0_ROCMFP4_FAST: + return dequantize_block_cuda; + case GGML_TYPE_Q2_0_ROCMFP2: + return dequantize_block_cuda; + case GGML_TYPE_Q3_0_ROCMFPX: + return dequantize_block_cuda; + case GGML_TYPE_Q6_0_ROCMFPX: + return dequantize_block_cuda; + case GGML_TYPE_Q8_0_ROCMFPX: + return dequantize_block_cuda; case GGML_TYPE_BF16: return convert_unary_cuda; default: diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/cpy.cu b/server/deps/llama.cpp/ggml/src/ggml-cuda/cpy.cu index 372fc98b1..3d7a698b9 100644 --- a/server/deps/llama.cpp/ggml/src/ggml-cuda/cpy.cu +++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/cpy.cu @@ -1,6 +1,7 @@ #include "cpy.cuh" #include "dequantize.cuh" #include "cpy-utils.cuh" +#include "../../rocmfp4/rocmfp4_hip_scale.cuh" #if defined(GGML_USE_MUSA) && defined(GGML_MUSA_MUDNN_COPY) #include "ggml-musa/mudnn.cuh" #endif // GGML_USE_MUSA && GGML_MUSA_MUDNN_COPY @@ -119,6 +120,72 @@ static __device__ void cpy_blck_q_f32(const char * cxi, char * cdsti) { } } +static __device__ void cpy_blck_rocmfp4_f32(const char * cxi, char * cdsti) { + const block_rocmfp4 * x = (const block_rocmfp4 *) cxi; + float * cdstf = (float *) cdsti; + + const float d0 = rocmfp4_ue4m3_to_fp32_half_finite(x->e[0]); + const float d1 = rocmfp4_ue4m3_to_fp32_half_finite(x->e[1]); + +#pragma unroll + for (int j = 0; j < QK_ROCMFP4/2; ++j) { + const uint8_t q = x->qs[j]; + cdstf[j] = d0 * (float) rocmfp4_decode_i8(q); + cdstf[j + QK_ROCMFP4/2] = d1 * (float) rocmfp4_decode_i8(q >> 4); + } +} + +static __device__ void cpy_blck_rocmfp4_fast_f32(const char * cxi, char * cdsti) { + const block_rocmfp4_fast * x = (const block_rocmfp4_fast *) cxi; + float * cdstf = (float *) cdsti; + + const float d = rocmfp4_ue4m3_to_fp32_half_finite(x->e); + +#pragma unroll + for (int j = 0; j < QK_ROCMFP4/2; ++j) { + const uint8_t q = x->qs[j]; + cdstf[j] = d * (float) rocmfp4_decode_i8(q); + cdstf[j + QK_ROCMFP4/2] = d * (float) rocmfp4_decode_i8(q >> 4); + } +} + +static __global__ void cpy_rocmfp4_f32_contiguous(const block_rocmfp4 * cx, float * cdst, const int64_t ne) { + const int64_t packed_idx = (int64_t)blockDim.x * blockIdx.x + threadIdx.x; + const int64_t packed_count = (ne / QK_ROCMFP4) * (QK_ROCMFP4/2); + + if (packed_idx >= packed_count) { + return; + } + + const int64_t ib = packed_idx >> 4; + const int j = packed_idx & 0x0f; + const int64_t base = ib * QK_ROCMFP4; + const uint8_t q = cx[ib].qs[j]; + const float d0 = rocmfp4_ue4m3_to_fp32_half_finite(cx[ib].e[0]); + const float d1 = rocmfp4_ue4m3_to_fp32_half_finite(cx[ib].e[1]); + + cdst[base + j] = d0 * (float) rocmfp4_decode_i8(q); + cdst[base + j + QK_ROCMFP4/2] = d1 * (float) rocmfp4_decode_i8(q >> 4); +} + +static __global__ void cpy_rocmfp4_fast_f32_contiguous(const block_rocmfp4_fast * cx, float * cdst, const int64_t ne) { + const int64_t packed_idx = (int64_t)blockDim.x * blockIdx.x + threadIdx.x; + const int64_t packed_count = (ne / QK_ROCMFP4) * (QK_ROCMFP4/2); + + if (packed_idx >= packed_count) { + return; + } + + const int64_t ib = packed_idx >> 4; + const int j = packed_idx & 0x0f; + const int64_t base = ib * QK_ROCMFP4; + const uint8_t q = cx[ib].qs[j]; + const float d = rocmfp4_ue4m3_to_fp32_half_finite(cx[ib].e); + + cdst[base + j] = d * (float) rocmfp4_decode_i8(q); + cdst[base + j + QK_ROCMFP4/2] = d * (float) rocmfp4_decode_i8(q >> 4); +} + template static __global__ void cpy_f32_q(const char * cx, char * cdst, const int64_t ne, const int64_t ne00, const int64_t ne01, const int64_t ne02, const int64_t nb00, const int64_t nb01, const int64_t nb02, @@ -308,6 +375,56 @@ static void ggml_cpy_q4_1_f32_cuda( ne10, ne11, ne12, nb10, nb11, nb12, nb13); } +static void ggml_cpy_rocmfp4_f32_hip( + const char * cx, char * cdst, const int64_t ne, + const int64_t ne00, const int64_t ne01, const int64_t ne02, + const int64_t nb00, const int64_t nb01, const int64_t nb02, + const int64_t nb03, const int64_t ne10, const int64_t ne11, const int64_t ne12, + const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13, + cudaStream_t stream) { + GGML_ASSERT(ne % QK_ROCMFP4 == 0); + const int64_t num_blocks = ne / QK_ROCMFP4; + GGML_ASSERT(num_blocks < UINT_MAX); + cpy_q_f32<<>>( + cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, + ne10, ne11, ne12, nb10, nb11, nb12, nb13); +} + +static void ggml_cpy_rocmfp4_f32_contiguous_hip( + const char * cx, char * cdst, const int64_t ne, cudaStream_t stream) { + GGML_ASSERT(ne % QK_ROCMFP4 == 0); + const int64_t packed_count = (ne / QK_ROCMFP4) * (QK_ROCMFP4/2); + const int64_t num_blocks = (packed_count + CUDA_CPY_BLOCK_SIZE - 1) / CUDA_CPY_BLOCK_SIZE; + GGML_ASSERT(num_blocks < UINT_MAX); + cpy_rocmfp4_f32_contiguous<<>>( + (const block_rocmfp4 *) cx, (float *) cdst, ne); +} + +static void ggml_cpy_rocmfp4_fast_f32_hip( + const char * cx, char * cdst, const int64_t ne, + const int64_t ne00, const int64_t ne01, const int64_t ne02, + const int64_t nb00, const int64_t nb01, const int64_t nb02, + const int64_t nb03, const int64_t ne10, const int64_t ne11, const int64_t ne12, + const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13, + cudaStream_t stream) { + GGML_ASSERT(ne % QK_ROCMFP4 == 0); + const int64_t num_blocks = ne / QK_ROCMFP4; + GGML_ASSERT(num_blocks < UINT_MAX); + cpy_q_f32<<>>( + cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, + ne10, ne11, ne12, nb10, nb11, nb12, nb13); +} + +static void ggml_cpy_rocmfp4_fast_f32_contiguous_hip( + const char * cx, char * cdst, const int64_t ne, cudaStream_t stream) { + GGML_ASSERT(ne % QK_ROCMFP4 == 0); + const int64_t packed_count = (ne / QK_ROCMFP4) * (QK_ROCMFP4/2); + const int64_t num_blocks = (packed_count + CUDA_CPY_BLOCK_SIZE - 1) / CUDA_CPY_BLOCK_SIZE; + GGML_ASSERT(num_blocks < UINT_MAX); + cpy_rocmfp4_fast_f32_contiguous<<>>( + (const block_rocmfp4_fast *) cx, (float *) cdst, ne); +} + static void ggml_cpy_f32_q5_0_cuda( const char * cx, char * cdst, const int64_t ne, const int64_t ne00, const int64_t ne01, const int64_t ne02, const int64_t nb00, const int64_t nb01, const int64_t nb02, @@ -556,6 +673,20 @@ void ggml_cuda_cpy(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, gg } else if (src0->type == GGML_TYPE_Q4_0 && src1->type == GGML_TYPE_F32) { ggml_cpy_q4_0_f32_cuda (src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream); + } else if (src0->type == GGML_TYPE_Q4_0_ROCMFP4 && src1->type == GGML_TYPE_F32) { + if (contiguous_srcs) { + ggml_cpy_rocmfp4_f32_contiguous_hip(src0_ddc, src1_ddc, ne, main_stream); + } else { + ggml_cpy_rocmfp4_f32_hip + (src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream); + } + } else if (src0->type == GGML_TYPE_Q4_0_ROCMFP4_FAST && src1->type == GGML_TYPE_F32) { + if (contiguous_srcs) { + ggml_cpy_rocmfp4_fast_f32_contiguous_hip(src0_ddc, src1_ddc, ne, main_stream); + } else { + ggml_cpy_rocmfp4_fast_f32_hip + (src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream); + } } else if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_Q4_1) { ggml_cpy_f32_q4_1_cuda (src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream); diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/dequantize.cuh b/server/deps/llama.cpp/ggml/src/ggml-cuda/dequantize.cuh index 1801bd048..3ed52e12e 100644 --- a/server/deps/llama.cpp/ggml/src/ggml-cuda/dequantize.cuh +++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/dequantize.cuh @@ -1,5 +1,6 @@ #include "common.cuh" #include "tq3-quant.cuh" +#include "../../rocmfp4/rocmfp4_hip_scale.cuh" static __device__ __forceinline__ void dequantize_q4_0(const void * vx, const int64_t ib, const int iqs, float2 & v){ const block_q4_0 * x = (const block_q4_0 *) vx; @@ -76,3 +77,120 @@ static __device__ __forceinline__ void dequantize_q8_0(const void * vx, const in v.x *= d; v.y *= d; } + +static __device__ __forceinline__ void dequantize_rocmfp4(const void * vx, const int64_t ib, const int iqs, float2 & v) { + const block_rocmfp4 * x = (const block_rocmfp4 *) vx; + + const int q = x[ib].qs[iqs]; + const float d0 = rocmfp4_ue4m3_to_fp32_half_finite(x[ib].e[0]); + const float d1 = rocmfp4_ue4m3_to_fp32_half_finite(x[ib].e[1]); + + v.x = d0 * rocmfp4_decode_i8(q); + v.y = d1 * rocmfp4_decode_i8(q >> 4); +} + +static __device__ __forceinline__ void dequantize_rocmfp4_fast(const void * vx, const int64_t ib, const int iqs, float2 & v) { + const block_rocmfp4_fast * x = (const block_rocmfp4_fast *) vx; + + const int q = x[ib].qs[iqs]; + const float d = rocmfp4_ue4m3_to_fp32_half_finite(x[ib].e); + + v.x = d * rocmfp4_decode_i8(q); + v.y = d * rocmfp4_decode_i8(q >> 4); +} + +template +static __device__ __forceinline__ uint32_t rocmfpx_load_qs_window_cuda(const uint8_t * src, const int byte_pos) { + uint32_t v = (uint32_t) src[byte_pos + 0]; + + if (byte_pos + 1 < qs) { + v |= (uint32_t) src[byte_pos + 1] << 8; + } + if (byte_pos + 2 < qs) { + v |= (uint32_t) src[byte_pos + 2] << 16; + } + + return v; +} + +static __device__ __forceinline__ uint32_t rocmfpx_get_fp3_code_cuda(const uint8_t * src, const int i) { + const int bit_pos = i * 3; + const int byte_pos = bit_pos >> 3; + const int shift = bit_pos & 7; + return (rocmfpx_load_qs_window_cuda(src, byte_pos) >> shift) & 7u; +} + +static __device__ __forceinline__ uint32_t rocmfpx_get_fp2_code_cuda(const uint8_t * src, const int i) { + return (src[i >> 2] >> (2*(i & 3))) & 3u; +} + +static __device__ __forceinline__ uint32_t rocmfpx_get_fp6_code_cuda(const uint8_t * src, const int i) { + const int bit_pos = i * 6; + const int byte_pos = bit_pos >> 3; + const int shift = bit_pos & 7; + return (rocmfpx_load_qs_window_cuda(src, byte_pos) >> shift) & 63u; +} + +static __device__ __forceinline__ int rocmfpx_decode_fp3_code_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; +} + +static __device__ __forceinline__ float rocmfpx_decode_fp2_code_cuda(const uint32_t code) { + switch (code & 3u) { + case 0: return (float) ROCMFP2_KVALUE_0_I8; + case 1: return (float) ROCMFP2_KVALUE_1_I8; + case 2: return (float) ROCMFP2_KVALUE_2_I8; + default: return (float) ROCMFP2_KVALUE_3_I8; + } +} + +static __device__ __forceinline__ int rocmfpx_decode_fp6_code_cuda(const uint32_t code) { + const int mag = (int) (code & 31u); + return (code & 32u) ? -mag : mag; +} + +static __device__ __forceinline__ void dequantize_rocmfpx_fp3(const void * vx, const int64_t ib, const int iqs, float2 & v) { + const block_rocmfp3 * x = (const block_rocmfp3 *) vx; + + const int i0 = iqs + 0; + const int i1 = iqs + 1; + const float d0 = rocmfpx_ue4m3_to_fp32_finite(x[ib].e[i0 >= QK_ROCMFP3/2]); + const float d1 = rocmfpx_ue4m3_to_fp32_finite(x[ib].e[i1 >= QK_ROCMFP3/2]); + + v.x = d0 * (float) rocmfpx_decode_fp3_code_cuda(rocmfpx_get_fp3_code_cuda(x[ib].qs, i0)); + v.y = d1 * (float) rocmfpx_decode_fp3_code_cuda(rocmfpx_get_fp3_code_cuda(x[ib].qs, i1)); +} + +static __device__ __forceinline__ void dequantize_rocmfpx_fp2(const void * vx, const int64_t ib, const int iqs, float2 & v) { + const block_rocmfp2 * x = (const block_rocmfp2 *) vx; + + const int i0 = iqs + 0; + const int i1 = iqs + 1; + const float d0 = rocmfpx_ue4m3_to_fp32_finite(x[ib].e[i0 >= QK_ROCMFP2/2]); + const float d1 = rocmfpx_ue4m3_to_fp32_finite(x[ib].e[i1 >= QK_ROCMFP2/2]); + + v.x = d0 * rocmfpx_decode_fp2_code_cuda(rocmfpx_get_fp2_code_cuda(x[ib].qs, i0)); + v.y = d1 * rocmfpx_decode_fp2_code_cuda(rocmfpx_get_fp2_code_cuda(x[ib].qs, i1)); +} + +static __device__ __forceinline__ void dequantize_rocmfpx_fp6(const void * vx, const int64_t ib, const int iqs, float2 & v) { + const block_rocmfp6 * x = (const block_rocmfp6 *) vx; + + const int i0 = iqs + 0; + const int i1 = iqs + 1; + const float d0 = rocmfpx_ue4m3_to_fp32_finite(x[ib].e[i0 >= QK_ROCMFP6/2]); + const float d1 = rocmfpx_ue4m3_to_fp32_finite(x[ib].e[i1 >= QK_ROCMFP6/2]); + + v.x = d0 * (float) rocmfpx_decode_fp6_code_cuda(rocmfpx_get_fp6_code_cuda(x[ib].qs, i0)); + v.y = d1 * (float) rocmfpx_decode_fp6_code_cuda(rocmfpx_get_fp6_code_cuda(x[ib].qs, i1)); +} + +static __device__ __forceinline__ void dequantize_rocmfpx_fp8(const void * vx, const int64_t ib, const int iqs, float2 & v) { + const block_rocmfp8 * x = (const block_rocmfp8 *) vx; + + const float d = rocmfpx_ue4m3_to_fp32_finite(x[ib].e); + v.x = d * (float) x[ib].qs[iqs + 0]; + v.y = d * (float) x[ib].qs[iqs + 1]; +} diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/ds4-hc.cu b/server/deps/llama.cpp/ggml/src/ggml-cuda/ds4-hc.cu new file mode 100644 index 000000000..1c6b86972 --- /dev/null +++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/ds4-hc.cu @@ -0,0 +1,367 @@ +#include "ds4-hc.cuh" + +// Fused DeepSeek4 hyper-connection ops. +// +// mode 0 (pre): src0 = mix [mix_dim] (f32, from fn @ rms_norm(hc_state)) +// src1 = base [mix_dim] (f32) +// src2 = hc_state [n_embd*n_hc] (f32, raw residual streams) +// dst = [n_embd + mix_dim]: +// dst[0..n_embd) = working vector (pre-mixed input) +// dst[n_embd..n_embd+mix) = split = {pre[n_hc], post[n_hc], comb[n_hc*n_hc]} +// Math matches cpu_hc_sinkhorn + finish_hc_pre_from_mix_into in +// deepseek4_graph.cpp (sigmoid gates + Sinkhorn-normalized combine). +// +// mode 1 (post): src0 = residual hc_state [n_embd*n_hc] +// src1 = block_out [n_embd] +// src2 = split [mix_dim] (view of a mode-0 dst tail) +// dst = new hc_state [n_embd*n_hc]: +// dst[h*n_embd+d] = post[h]*block_out[d] +// + sum_src comb[h + src*n_hc] * residual[src*n_embd+d] +// +// mode 2 (out): src0 = mix [n_hc] +// src1 = base [n_hc] +// src2 = hc_state [n_embd*n_hc] +// dst = [n_embd]: weights[h] = sigmoid(mix[h]*s0+base[h]) + 1e-6; +// dst[d] = sum_h weights[h]*hc_state[h*n_embd+d] + +#define DS4_HC_SINKHORN_EPS 1.0e-6f +#define DS4_HC_MAX_HC 8 +#define DS4_HC_MAX_MIX (2*DS4_HC_MAX_HC + DS4_HC_MAX_HC*DS4_HC_MAX_HC) + +static __device__ __forceinline__ float ds4_hc_sigmoid(float x) { + return 1.0f / (1.0f + expf(-x)); +} + +static __device__ void ds4_hc_sinkhorn_split( + const float * mix, + const float * base, + float pre_scale, + float post_scale, + float comb_scale, + int n_hc, + int iters, + float * split) { + for (int i = 0; i < n_hc; ++i) { + split[i] = ds4_hc_sigmoid(mix[i] * pre_scale + base[i]) + DS4_HC_SINKHORN_EPS; + } + for (int i = 0; i < n_hc; ++i) { + split[n_hc + i] = 2.0f * ds4_hc_sigmoid(mix[n_hc + i] * post_scale + base[n_hc + i]); + } + + float c[DS4_HC_MAX_HC * DS4_HC_MAX_HC]; + for (int dst_i = 0; dst_i < n_hc; ++dst_i) { + float row_max = -1.0e30f; + for (int src_i = 0; src_i < n_hc; ++src_i) { + const int idx = src_i + dst_i * n_hc; + const float v = mix[2 * n_hc + idx] * comb_scale + base[2 * n_hc + idx]; + c[idx] = v; + row_max = v > row_max ? v : row_max; + } + float row_sum = 0.0f; + for (int src_i = 0; src_i < n_hc; ++src_i) { + const int idx = src_i + dst_i * n_hc; + c[idx] = expf(c[idx] - row_max); + row_sum += c[idx]; + } + const float inv = 1.0f / row_sum; + for (int src_i = 0; src_i < n_hc; ++src_i) { + c[src_i + dst_i * n_hc] = c[src_i + dst_i * n_hc] * inv + DS4_HC_SINKHORN_EPS; + } + } + for (int src_i = 0; src_i < n_hc; ++src_i) { + float sum = 0.0f; + for (int dst_i = 0; dst_i < n_hc; ++dst_i) sum += c[src_i + dst_i * n_hc]; + const float inv = 1.0f / (sum + DS4_HC_SINKHORN_EPS); + for (int dst_i = 0; dst_i < n_hc; ++dst_i) c[src_i + dst_i * n_hc] *= inv; + } + for (int iter = 1; iter < iters; ++iter) { + for (int dst_i = 0; dst_i < n_hc; ++dst_i) { + float sum = 0.0f; + for (int src_i = 0; src_i < n_hc; ++src_i) sum += c[src_i + dst_i * n_hc]; + const float inv = 1.0f / (sum + DS4_HC_SINKHORN_EPS); + for (int src_i = 0; src_i < n_hc; ++src_i) c[src_i + dst_i * n_hc] *= inv; + } + for (int src_i = 0; src_i < n_hc; ++src_i) { + float sum = 0.0f; + for (int dst_i = 0; dst_i < n_hc; ++dst_i) sum += c[src_i + dst_i * n_hc]; + const float inv = 1.0f / (sum + DS4_HC_SINKHORN_EPS); + for (int dst_i = 0; dst_i < n_hc; ++dst_i) c[src_i + dst_i * n_hc] *= inv; + } + } + for (int i = 0; i < n_hc * n_hc; ++i) { + split[2 * n_hc + i] = c[i]; + } +} + +// Fully-unrolled variant: with compile-time NHC the c[] matrix lives in +// registers. The generic version's runtime-bound loops force c[] into scratch +// (private, VRAM-backed) memory, making the serial sinkhorn ~20x slower +// (97us vs 5us measured on gfx1151). +template +static __device__ void ds4_hc_sinkhorn_split_t( + const float * mix, + const float * base, + float pre_scale, + float post_scale, + float comb_scale, + int iters, + float * split) { + #pragma unroll + for (int i = 0; i < NHC; ++i) { + split[i] = ds4_hc_sigmoid(mix[i] * pre_scale + base[i]) + DS4_HC_SINKHORN_EPS; + } + #pragma unroll + for (int i = 0; i < NHC; ++i) { + split[NHC + i] = 2.0f * ds4_hc_sigmoid(mix[NHC + i] * post_scale + base[NHC + i]); + } + float c[NHC * NHC]; + #pragma unroll + for (int dst_i = 0; dst_i < NHC; ++dst_i) { + float row_max = -1.0e30f; + #pragma unroll + for (int src_i = 0; src_i < NHC; ++src_i) { + const int idx = src_i + dst_i * NHC; + const float v = mix[2 * NHC + idx] * comb_scale + base[2 * NHC + idx]; + c[idx] = v; + row_max = v > row_max ? v : row_max; + } + float row_sum = 0.0f; + #pragma unroll + for (int src_i = 0; src_i < NHC; ++src_i) { + const int idx = src_i + dst_i * NHC; + c[idx] = expf(c[idx] - row_max); + row_sum += c[idx]; + } + const float inv = 1.0f / row_sum; + #pragma unroll + for (int src_i = 0; src_i < NHC; ++src_i) { + c[src_i + dst_i * NHC] = c[src_i + dst_i * NHC] * inv + DS4_HC_SINKHORN_EPS; + } + } + #pragma unroll + for (int src_i = 0; src_i < NHC; ++src_i) { + float sum = 0.0f; + #pragma unroll + for (int dst_i = 0; dst_i < NHC; ++dst_i) sum += c[src_i + dst_i * NHC]; + const float inv = 1.0f / (sum + DS4_HC_SINKHORN_EPS); + #pragma unroll + for (int dst_i = 0; dst_i < NHC; ++dst_i) c[src_i + dst_i * NHC] *= inv; + } + for (int iter = 1; iter < iters; ++iter) { + #pragma unroll + for (int dst_i = 0; dst_i < NHC; ++dst_i) { + float sum = 0.0f; + #pragma unroll + for (int src_i = 0; src_i < NHC; ++src_i) sum += c[src_i + dst_i * NHC]; + const float inv = 1.0f / (sum + DS4_HC_SINKHORN_EPS); + #pragma unroll + for (int src_i = 0; src_i < NHC; ++src_i) c[src_i + dst_i * NHC] *= inv; + } + #pragma unroll + for (int src_i = 0; src_i < NHC; ++src_i) { + float sum = 0.0f; + #pragma unroll + for (int dst_i = 0; dst_i < NHC; ++dst_i) sum += c[src_i + dst_i * NHC]; + const float inv = 1.0f / (sum + DS4_HC_SINKHORN_EPS); + #pragma unroll + for (int dst_i = 0; dst_i < NHC; ++dst_i) c[src_i + dst_i * NHC] *= inv; + } + } + #pragma unroll + for (int i = 0; i < NHC * NHC; ++i) { + split[2 * NHC + i] = c[i]; + } +} + +template +static __global__ void ds4_hc_pre_kernel_t( + const float * __restrict__ mix, + const float * __restrict__ base, + const float * __restrict__ hc_state, + float * __restrict__ dst, + int n_embd, + int iters, + float pre_scale, + float post_scale, + float comb_scale) { + __shared__ float split[DS4_HC_MAX_MIX]; + __shared__ float s_mix[DS4_HC_MAX_MIX]; + __shared__ float s_base[DS4_HC_MAX_MIX]; + constexpr int mix_dim = 2 * NHC + NHC * NHC; + const int tid = threadIdx.x; + + if (tid < mix_dim) { + s_mix[tid] = mix[tid]; + s_base[tid] = base[tid]; + } + __syncthreads(); + + if (tid == 0) { + ds4_hc_sinkhorn_split_t(s_mix, s_base, pre_scale, post_scale, comb_scale, iters, split); + if (blockIdx.x == 0) { + #pragma unroll + for (int i = 0; i < mix_dim; ++i) { + dst[n_embd + i] = split[i]; + } + } + } + __syncthreads(); + + const int d = (int) blockIdx.x * blockDim.x + tid; + if (d < n_embd) { + float acc = 0.0f; + #pragma unroll + for (int h = 0; h < NHC; ++h) { + acc += split[h] * hc_state[(size_t) h * n_embd + d]; + } + dst[d] = acc; + } +} + +static __global__ void ds4_hc_pre_kernel( + const float * __restrict__ mix, + const float * __restrict__ base, + const float * __restrict__ hc_state, + float * __restrict__ dst, + int n_embd, + int n_hc, + int iters, + float pre_scale, + float post_scale, + float comb_scale) { + __shared__ float split[DS4_HC_MAX_MIX]; + __shared__ float s_mix[DS4_HC_MAX_MIX]; + __shared__ float s_base[DS4_HC_MAX_MIX]; + const int mix_dim = 2 * n_hc + n_hc * n_hc; + const int tid = threadIdx.x; + + // Stage mix/base cooperatively: base lives in managed (UMA) memory where + // serial scalar loads cost ~2us each; one parallel coalesced load instead. + if (tid < mix_dim) { + s_mix[tid] = mix[tid]; + s_base[tid] = base[tid]; + } + __syncthreads(); + + // Each block redoes the (tiny) sinkhorn into shared memory so the mix + // loop below can spread across the whole GPU instead of one CU. + if (tid == 0) { + ds4_hc_sinkhorn_split(s_mix, s_base, pre_scale, post_scale, comb_scale, n_hc, iters, split); + if (blockIdx.x == 0) { + for (int i = 0; i < mix_dim; ++i) { + dst[n_embd + i] = split[i]; + } + } + } + __syncthreads(); + + const int d = (int) blockIdx.x * blockDim.x + tid; + if (d < n_embd) { + float acc = 0.0f; + for (int h = 0; h < n_hc; ++h) { + acc += split[h] * hc_state[(size_t) h * n_embd + d]; + } + dst[d] = acc; + } +} + +static __global__ void ds4_hc_post_kernel( + const float * __restrict__ residual, + const float * __restrict__ block_out, + const float * __restrict__ split, + float * __restrict__ dst, + int n_embd, + int n_hc) { + const int i = blockIdx.x * blockDim.x + threadIdx.x; + const int total = n_embd * n_hc; + if (i >= total) { + return; + } + const int h = i / n_embd; + const int d = i - h * n_embd; + const float * post = split + n_hc; + const float * comb = split + 2 * n_hc; + float acc = block_out[d] * post[h]; + for (int src = 0; src < n_hc; ++src) { + acc += comb[h + src * n_hc] * residual[(size_t) src * n_embd + d]; + } + dst[i] = acc; +} + +static __global__ void ds4_hc_out_kernel( + const float * __restrict__ mix, + const float * __restrict__ base, + const float * __restrict__ hc_state, + float * __restrict__ dst, + int n_embd, + int n_hc, + float pre_scale) { + const int d = blockIdx.x * blockDim.x + threadIdx.x; + if (d >= n_embd) { + return; + } + float acc = 0.0f; + for (int h = 0; h < n_hc; ++h) { + const float wgt = ds4_hc_sigmoid(mix[h] * pre_scale + base[h]) + DS4_HC_SINKHORN_EPS; + acc += wgt * hc_state[(size_t) h * n_embd + d]; + } + dst[d] = acc; +} + +void ggml_cuda_op_ds4_hc(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { + const ggml_tensor * src0 = dst->src[0]; + const ggml_tensor * src1 = dst->src[1]; + const ggml_tensor * src2 = dst->src[2]; + + GGML_ASSERT(src0 && src0->type == GGML_TYPE_F32); + GGML_ASSERT(src1 && src1->type == GGML_TYPE_F32); + GGML_ASSERT(src2 && src2->type == GGML_TYPE_F32); + GGML_ASSERT(dst->type == GGML_TYPE_F32); + + const int mode = ggml_get_op_params_i32(dst, 0); + const int n_embd = ggml_get_op_params_i32(dst, 1); + const int n_hc = ggml_get_op_params_i32(dst, 2); + + GGML_ASSERT(n_hc > 0 && n_hc <= DS4_HC_MAX_HC); + + cudaStream_t stream = ctx.stream(); + + switch (mode) { + case 0: { + const int iters = ggml_get_op_params_i32(dst, 3); + const float pre_scale = ggml_get_op_params_f32(dst, 4); + const float post_scale = ggml_get_op_params_f32(dst, 5); + const float comb_scale = ggml_get_op_params_f32(dst, 6); + const int pre_blocks = (n_embd + 255) / 256; + if (n_hc == 4) { + ds4_hc_pre_kernel_t<4><<>>( + (const float *) src0->data, (const float *) src1->data, + (const float *) src2->data, (float *) dst->data, + n_embd, iters, pre_scale, post_scale, comb_scale); + } else { + ds4_hc_pre_kernel<<>>( + (const float *) src0->data, (const float *) src1->data, + (const float *) src2->data, (float *) dst->data, + n_embd, n_hc, iters, pre_scale, post_scale, comb_scale); + } + } break; + case 1: { + const int total = n_embd * n_hc; + const int blocks = (total + 255) / 256; + ds4_hc_post_kernel<<>>( + (const float *) src0->data, (const float *) src1->data, + (const float *) src2->data, (float *) dst->data, + n_embd, n_hc); + } break; + case 2: { + const float pre_scale = ggml_get_op_params_f32(dst, 4); + const int blocks = (n_embd + 255) / 256; + ds4_hc_out_kernel<<>>( + (const float *) src0->data, (const float *) src1->data, + (const float *) src2->data, (float *) dst->data, + n_embd, n_hc, pre_scale); + } break; + default: + GGML_ABORT("ds4_hc: unknown mode"); + } +} diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/ds4-hc.cuh b/server/deps/llama.cpp/ggml/src/ggml-cuda/ds4-hc.cuh new file mode 100644 index 000000000..1a2cdf5eb --- /dev/null +++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/ds4-hc.cuh @@ -0,0 +1,10 @@ +#pragma once + +#include "ggml-cuda/common.cuh" + +// Fused DeepSeek4 hyper-connection (HC) decode ops. See ds4-hc.cu for the +// per-mode contract (pre / post / out). Used by the opt-in +// DFLASH_DS4_FUSED_DECODE single-graph decode path; output is deterministic +// but not bit-identical to the CPU HC reference (expf ULP differences +// amplified by the sinkhorn iterations). +void ggml_cuda_op_ds4_hc(ggml_backend_cuda_context & ctx, ggml_tensor * dst); diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/getrows.cu b/server/deps/llama.cpp/ggml/src/ggml-cuda/getrows.cu index 1d7c6d17a..73a28e529 100644 --- a/server/deps/llama.cpp/ggml/src/ggml-cuda/getrows.cu +++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/getrows.cu @@ -199,6 +199,30 @@ static void ggml_cuda_get_rows_switch_src0_type( get_rows_cuda_q(src0_d, src1_d, dst_d, ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream); break; + case GGML_TYPE_Q4_0_ROCMFP4: + get_rows_cuda_q(src0_d, src1_d, dst_d, + ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream); + break; + case GGML_TYPE_Q4_0_ROCMFP4_FAST: + get_rows_cuda_q(src0_d, src1_d, dst_d, + ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream); + break; + case GGML_TYPE_Q2_0_ROCMFP2: + get_rows_cuda_q(src0_d, src1_d, dst_d, + ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream); + break; + case GGML_TYPE_Q3_0_ROCMFPX: + get_rows_cuda_q(src0_d, src1_d, dst_d, + ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream); + break; + case GGML_TYPE_Q6_0_ROCMFPX: + get_rows_cuda_q(src0_d, src1_d, dst_d, + ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream); + break; + case GGML_TYPE_Q8_0_ROCMFPX: + get_rows_cuda_q(src0_d, src1_d, dst_d, + ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream); + break; case GGML_TYPE_TQ3_0: get_rows_cuda_q(src0_d, src1_d, dst_d, ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream); diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu b/server/deps/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu index 33f01a10a..6f887eafc 100644 --- a/server/deps/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu @@ -64,6 +64,7 @@ #include "ggml-cuda/cumsum.cuh" #include "ggml-cuda/fill.cuh" #include "ggml-cuda/moe-fused.cuh" +#include "ggml-cuda/ds4-hc.cuh" #include "ggml.h" #include @@ -2368,7 +2369,12 @@ static bool ggml_cuda_should_fuse_mul_mat(const ggml_tensor * ffn_up, return false; } - static constexpr std::array valid_glu_ops = { GGML_GLU_OP_SWIGLU, GGML_GLU_OP_GEGLU, GGML_GLU_OP_SWIGLU_OAI }; + static constexpr std::array valid_glu_ops = { + GGML_GLU_OP_SWIGLU, + GGML_GLU_OP_GEGLU, + GGML_GLU_OP_SWIGLU_OAI, + GGML_GLU_OP_SWIGLU_DS4, + }; if (std::find(valid_glu_ops.begin(), valid_glu_ops.end(), ggml_get_glu_op(glu)) == valid_glu_ops.end()) { return false; @@ -2424,6 +2430,12 @@ static bool ggml_cuda_should_fuse_mul_mat_vec_f(const ggml_tensor * tensor) { return use_mul_mat_vec_f; } +static inline void ggml_cuda_set_fusion_glu_params(ggml_cuda_mm_fusion_args_host & fusion_data, const ggml_tensor * glu) { + fusion_data.glu_op = ggml_get_glu_op(glu); + fusion_data.glu_param0 = ggml_get_op_params_f32(glu, 2); + fusion_data.glu_param1 = ggml_get_op_params_f32(glu, 3); +} + static bool ggml_cuda_should_fuse_mul_mat_vec_q(const ggml_tensor * tensor) { ggml_tensor * src0 = tensor->src[0]; ggml_tensor * src1 = tensor->src[1]; @@ -2857,6 +2869,9 @@ static bool ggml_cuda_compute_forward(ggml_backend_cuda_context & ctx, struct gg case GGML_GLU_OP_SWIGLU_OAI: ggml_cuda_op_swiglu_oai(ctx, dst); break; + case GGML_GLU_OP_SWIGLU_DS4: + ggml_cuda_op_swiglu_ds4(ctx, dst); + break; case GGML_GLU_OP_GEGLU_ERF: ggml_cuda_op_geglu_erf(ctx, dst); break; @@ -2876,6 +2891,9 @@ static bool ggml_cuda_compute_forward(ggml_backend_cuda_context & ctx, struct gg case GGML_OP_MOE_FUSED: ggml_cuda_op_moe_fused(ctx, dst); break; + case GGML_OP_DS4_HC: + ggml_cuda_op_ds4_hc(ctx, dst); + break; case GGML_OP_GROUP_NORM: ggml_cuda_op_group_norm(ctx, dst); break; @@ -4086,7 +4104,7 @@ static void ggml_cuda_graph_evaluate_and_capture(ggml_backend_cuda_context * cud fusion_data.gate = gate_n->src[0]; fusion_data.x_bias = up_bias_tensor; fusion_data.gate_bias = gate_bias_tensor; - fusion_data.glu_op = ggml_get_glu_op(glu); + ggml_cuda_set_fusion_glu_params(fusion_data, glu); ggml_cuda_mul_mat_vec_f(*cuda_ctx, src0, src1, ids, glu, &fusion_data); fused_mul_mat_vec = true; @@ -4099,7 +4117,7 @@ static void ggml_cuda_graph_evaluate_and_capture(ggml_backend_cuda_context * cud fusion_data.gate = gate_n->src[0]; fusion_data.x_bias = up_bias_tensor; fusion_data.gate_bias = gate_bias_tensor; - fusion_data.glu_op = ggml_get_glu_op(glu); + ggml_cuda_set_fusion_glu_params(fusion_data, glu); ggml_cuda_mul_mat_vec_q(*cuda_ctx, src0, src1, ids, glu, &fusion_data); fused_mul_mat_vec = true; @@ -4123,7 +4141,7 @@ static void ggml_cuda_graph_evaluate_and_capture(ggml_backend_cuda_context * cud if (ggml_cuda_should_fuse_mul_mat_vec_f(up)) { ggml_cuda_mm_fusion_args_host fusion_data{}; fusion_data.gate = gate->src[0]; - fusion_data.glu_op = ggml_get_glu_op(glu); + ggml_cuda_set_fusion_glu_params(fusion_data, glu); ggml_cuda_mul_mat_vec_f(*cuda_ctx, src0, src1, ids, glu, &fusion_data); fused_mul_mat_vec = true; @@ -4134,7 +4152,7 @@ static void ggml_cuda_graph_evaluate_and_capture(ggml_backend_cuda_context * cud if (ggml_cuda_should_fuse_mul_mat_vec_q(up)) { ggml_cuda_mm_fusion_args_host fusion_data{}; fusion_data.gate = gate->src[0]; - fusion_data.glu_op = ggml_get_glu_op(glu); + ggml_cuda_set_fusion_glu_params(fusion_data, glu); ggml_cuda_mul_mat_vec_q(*cuda_ctx, src0, src1, ids, glu, &fusion_data); fused_mul_mat_vec = true; @@ -5007,6 +5025,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g case GGML_GLU_OP_GEGLU: case GGML_GLU_OP_SWIGLU: case GGML_GLU_OP_SWIGLU_OAI: + case GGML_GLU_OP_SWIGLU_DS4: case GGML_GLU_OP_GEGLU_ERF: case GGML_GLU_OP_GEGLU_QUICK: return ggml_is_contiguous_1(op->src[0]); @@ -5020,6 +5039,8 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g return op->src[0]->nb[0] == ggml_type_size(op->src[0]->type); case GGML_OP_MOE_FUSED: return true; + case GGML_OP_DS4_HC: + return true; case GGML_OP_MUL_MAT: case GGML_OP_MUL_MAT_ID: { @@ -5065,6 +5086,12 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g case GGML_TYPE_Q8_0: case GGML_TYPE_MXFP4: case GGML_TYPE_NVFP4: + case GGML_TYPE_Q4_0_ROCMFP4: + case GGML_TYPE_Q4_0_ROCMFP4_FAST: + case GGML_TYPE_Q2_0_ROCMFP2: + case GGML_TYPE_Q3_0_ROCMFPX: + case GGML_TYPE_Q6_0_ROCMFPX: + case GGML_TYPE_Q8_0_ROCMFPX: case GGML_TYPE_Q2_K: case GGML_TYPE_Q3_K: case GGML_TYPE_Q4_K: @@ -5100,6 +5127,12 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g case GGML_TYPE_Q5_0: case GGML_TYPE_Q5_1: case GGML_TYPE_Q8_0: + case GGML_TYPE_Q4_0_ROCMFP4: + case GGML_TYPE_Q4_0_ROCMFP4_FAST: + case GGML_TYPE_Q2_0_ROCMFP2: + case GGML_TYPE_Q3_0_ROCMFPX: + case GGML_TYPE_Q6_0_ROCMFPX: + case GGML_TYPE_Q8_0_ROCMFPX: case GGML_TYPE_TQ3_0: return true; default: @@ -5206,7 +5239,10 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g case GGML_OP_CONCAT: { ggml_type src0_type = op->src[0]->type; - return src0_type != GGML_TYPE_I32 && src0_type != GGML_TYPE_I16; + return src0_type == GGML_TYPE_F32 || + src0_type == GGML_TYPE_F16 || + src0_type == GGML_TYPE_BF16 || + src0_type == GGML_TYPE_I8; } break; case GGML_OP_CONV_TRANSPOSE_1D: { diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/mmvf.cu b/server/deps/llama.cpp/ggml/src/ggml-cuda/mmvf.cu index d91472024..6e4db8e8b 100644 --- a/server/deps/llama.cpp/ggml/src/ggml-cuda/mmvf.cu +++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/mmvf.cu @@ -51,6 +51,8 @@ static __global__ void mul_mat_vec_f( bool use_bias = false; bool use_gate_bias = false; ggml_glu_op glu_op = ggml_glu_op::GGML_GLU_OP_SWIGLU; + float glu_param0 = 0.0f; + float glu_param1 = 0.0f; const T * gate_x = nullptr; const float * x_bias = nullptr; const float * gate_bias = nullptr; @@ -60,6 +62,8 @@ static __global__ void mul_mat_vec_f( use_bias = fusion.x_bias != nullptr; use_gate_bias = fusion.gate_bias != nullptr; glu_op = fusion.glu_op; + glu_param0 = fusion.glu_param0; + glu_param1 = fusion.glu_param1; if (use_gate) { gate_x = static_cast(fusion.gate); @@ -357,7 +361,11 @@ static __global__ void mul_mat_vec_f( value *= ggml_cuda_op_gelu_single(gate_value); break; case GGML_GLU_OP_SWIGLU_OAI: { - value = ggml_cuda_op_swiglu_oai_single(gate_value, value); + value = ggml_cuda_op_swiglu_oai_single(gate_value, value, glu_param0, glu_param1); + break; + } + case GGML_GLU_OP_SWIGLU_DS4: { + value = ggml_cuda_op_swiglu_ds4_single(gate_value, value, glu_param0); break; } default: @@ -369,7 +377,7 @@ static __global__ void mul_mat_vec_f( dst[tid*stride_col_dst + row] = value; if constexpr (!has_fusion) { - GGML_UNUSED_VARS(use_gate, use_bias, use_gate_bias, glu_op, gate_x, x_bias, gate_bias, sumf_gate); + GGML_UNUSED_VARS(use_gate, use_bias, use_gate_bias, glu_op, glu_param0, glu_param1, gate_x, x_bias, gate_bias, sumf_gate); } } @@ -668,6 +676,8 @@ void ggml_cuda_mul_mat_vec_f(ggml_backend_cuda_context & ctx, const ggml_tensor fusion_local.gate_bias = fusion->gate_bias->data; } fusion_local.glu_op = fusion->glu_op; + fusion_local.glu_param0 = fusion->glu_param0; + fusion_local.glu_param1 = fusion->glu_param1; } const int64_t s01 = src0->nb[1] / ts_src0; 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 befd0e731..2569d0828 100644 --- a/server/deps/llama.cpp/ggml/src/ggml-cuda/mmvq.cu +++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/mmvq.cu @@ -17,6 +17,12 @@ static constexpr __device__ vec_dot_q_cuda_t get_vec_dot_q_cuda(ggml_type type) case GGML_TYPE_Q8_0: return vec_dot_q8_0_q8_1; case GGML_TYPE_MXFP4: return vec_dot_mxfp4_q8_1; case GGML_TYPE_NVFP4: return vec_dot_nvfp4_q8_1; + case GGML_TYPE_Q4_0_ROCMFP4: return vec_dot_rocmfp4_q8_1; + case GGML_TYPE_Q4_0_ROCMFP4_FAST: return vec_dot_rocmfp4_fast_q8_1; + case GGML_TYPE_Q2_0_ROCMFP2: return vec_dot_rocmfpx_fp2_q8_1; + case GGML_TYPE_Q3_0_ROCMFPX: return vec_dot_rocmfpx_fp3_q8_1; + case GGML_TYPE_Q6_0_ROCMFPX: return vec_dot_rocmfpx_fp6_q8_1; + case GGML_TYPE_Q8_0_ROCMFPX: return vec_dot_rocmfpx_fp8_q8_1; case GGML_TYPE_Q2_K: return vec_dot_q2_K_q8_1; case GGML_TYPE_Q3_K: return vec_dot_q3_K_q8_1; case GGML_TYPE_Q4_K: return vec_dot_q4_K_q8_1; @@ -44,6 +50,12 @@ static constexpr __host__ __device__ int get_vdr_mmvq(ggml_type type) { case GGML_TYPE_Q8_0: return VDR_Q8_0_Q8_1_MMVQ; case GGML_TYPE_MXFP4: return VDR_MXFP4_Q8_1_MMVQ; case GGML_TYPE_NVFP4: return VDR_NVFP4_Q8_1_MMVQ; + case GGML_TYPE_Q4_0_ROCMFP4: return VDR_ROCMFP4_Q8_1_MMVQ; + case GGML_TYPE_Q4_0_ROCMFP4_FAST: return VDR_ROCMFP4_FAST_Q8_1_MMVQ; + case GGML_TYPE_Q2_0_ROCMFP2: return VDR_ROCMFP2_Q8_1_MMVQ; + case GGML_TYPE_Q3_0_ROCMFPX: return VDR_ROCMFP3_Q8_1_MMVQ; + case GGML_TYPE_Q6_0_ROCMFPX: return VDR_ROCMFP6_Q8_1_MMVQ; + case GGML_TYPE_Q8_0_ROCMFPX: return VDR_ROCMFP8_Q8_1_MMVQ; case GGML_TYPE_Q2_K: return VDR_Q2_K_Q8_1_MMVQ; case GGML_TYPE_Q3_K: return VDR_Q3_K_Q8_1_MMVQ; case GGML_TYPE_Q4_K: return VDR_Q4_K_Q8_1_MMVQ; @@ -536,6 +548,8 @@ static __global__ void mul_mat_vec_q( const float * x_bias = nullptr; const float * gate_bias = nullptr; ggml_glu_op active_glu; + float glu_param0 = 0.0f; + float glu_param1 = 0.0f; if constexpr (has_fusion) { use_gate = fusion.gate != nullptr; @@ -545,6 +559,8 @@ static __global__ void mul_mat_vec_q( x_bias = (const float *) fusion.x_bias; gate_bias = (const float *) fusion.gate_bias; active_glu = fusion.glu_op; + glu_param0 = fusion.glu_param0; + glu_param1 = fusion.glu_param1; } @@ -691,7 +707,11 @@ static __global__ void mul_mat_vec_q( result *= ggml_cuda_op_gelu_single(gate_value); break; case GGML_GLU_OP_SWIGLU_OAI: { - result = ggml_cuda_op_swiglu_oai_single(gate_value, result); + result = ggml_cuda_op_swiglu_oai_single(gate_value, result, glu_param0, glu_param1); + break; + } + case GGML_GLU_OP_SWIGLU_DS4: { + result = ggml_cuda_op_swiglu_ds4_single(gate_value, result, glu_param0); break; } default: @@ -705,7 +725,7 @@ static __global__ void mul_mat_vec_q( } if constexpr (!has_fusion) { - GGML_UNUSED_VARS(use_gate, use_bias, use_gate_bias, active_glu, gate_bias, x_bias, tmp_gate); + GGML_UNUSED_VARS(use_gate, use_bias, use_gate_bias, active_glu, glu_param0, glu_param1, gate_bias, x_bias, tmp_gate); } } @@ -755,6 +775,8 @@ static __global__ void mul_mat_vec_q_moe( const float * x_bias = nullptr; const float * gate_bias = nullptr; ggml_glu_op active_glu; + float glu_param0 = 0.0f; + float glu_param1 = 0.0f; if constexpr (has_fusion) { use_gate = fusion.gate != nullptr; @@ -764,6 +786,8 @@ static __global__ void mul_mat_vec_q_moe( x_bias = (const float *) fusion.x_bias; gate_bias = (const float *) fusion.gate_bias; active_glu = fusion.glu_op; + glu_param0 = fusion.glu_param0; + glu_param1 = fusion.glu_param1; } // partial sum for each thread @@ -816,7 +840,10 @@ static __global__ void mul_mat_vec_q_moe( result *= ggml_cuda_op_gelu_single(gate_value); break; case GGML_GLU_OP_SWIGLU_OAI: - result = ggml_cuda_op_swiglu_oai_single(gate_value, result); + result = ggml_cuda_op_swiglu_oai_single(gate_value, result, glu_param0, glu_param1); + break; + case GGML_GLU_OP_SWIGLU_DS4: + result = ggml_cuda_op_swiglu_ds4_single(gate_value, result, glu_param0); break; default: result = result * gate_value; @@ -828,7 +855,7 @@ static __global__ void mul_mat_vec_q_moe( } if constexpr (!has_fusion) { - GGML_UNUSED_VARS(use_gate, use_bias, use_gate_bias, vgate, x_bias, gate_bias, active_glu, tmp_gate); + GGML_UNUSED_VARS(use_gate, use_bias, use_gate_bias, vgate, x_bias, gate_bias, active_glu, glu_param0, glu_param1, tmp_gate); } } @@ -1444,6 +1471,42 @@ static void mul_mat_vec_q_switch_type( nchannels_x, nchannels_y, nchannels_dst, stride_channel_x, stride_channel_y, stride_channel_dst, nsamples_x, nsamples_dst, stride_sample_x, stride_sample_y, stride_sample_dst, ids_stride, stream); break; + case GGML_TYPE_Q4_0_ROCMFP4: + mul_mat_vec_q_switch_ncols_dst + (vx, vy, ids, fusion, dst, ncols_x, nrows_x, ncols_dst, stride_row_x, stride_col_y, stride_col_dst, + nchannels_x, nchannels_y, nchannels_dst, stride_channel_x, stride_channel_y, stride_channel_dst, + nsamples_x, nsamples_dst, stride_sample_x, stride_sample_y, stride_sample_dst, ids_stride, stream); + break; + case GGML_TYPE_Q4_0_ROCMFP4_FAST: + mul_mat_vec_q_switch_ncols_dst + (vx, vy, ids, fusion, dst, ncols_x, nrows_x, ncols_dst, stride_row_x, stride_col_y, stride_col_dst, + nchannels_x, nchannels_y, nchannels_dst, stride_channel_x, stride_channel_y, stride_channel_dst, + nsamples_x, nsamples_dst, stride_sample_x, stride_sample_y, stride_sample_dst, ids_stride, stream); + break; + case GGML_TYPE_Q2_0_ROCMFP2: + mul_mat_vec_q_switch_ncols_dst + (vx, vy, ids, fusion, dst, ncols_x, nrows_x, ncols_dst, stride_row_x, stride_col_y, stride_col_dst, + nchannels_x, nchannels_y, nchannels_dst, stride_channel_x, stride_channel_y, stride_channel_dst, + nsamples_x, nsamples_dst, stride_sample_x, stride_sample_y, stride_sample_dst, ids_stride, stream); + break; + case GGML_TYPE_Q3_0_ROCMFPX: + mul_mat_vec_q_switch_ncols_dst + (vx, vy, ids, fusion, dst, ncols_x, nrows_x, ncols_dst, stride_row_x, stride_col_y, stride_col_dst, + nchannels_x, nchannels_y, nchannels_dst, stride_channel_x, stride_channel_y, stride_channel_dst, + nsamples_x, nsamples_dst, stride_sample_x, stride_sample_y, stride_sample_dst, ids_stride, stream); + break; + case GGML_TYPE_Q6_0_ROCMFPX: + mul_mat_vec_q_switch_ncols_dst + (vx, vy, ids, fusion, dst, ncols_x, nrows_x, ncols_dst, stride_row_x, stride_col_y, stride_col_dst, + nchannels_x, nchannels_y, nchannels_dst, stride_channel_x, stride_channel_y, stride_channel_dst, + nsamples_x, nsamples_dst, stride_sample_x, stride_sample_y, stride_sample_dst, ids_stride, stream); + break; + case GGML_TYPE_Q8_0_ROCMFPX: + mul_mat_vec_q_switch_ncols_dst + (vx, vy, ids, fusion, dst, ncols_x, nrows_x, ncols_dst, stride_row_x, stride_col_y, stride_col_dst, + nchannels_x, nchannels_y, nchannels_dst, stride_channel_x, stride_channel_y, stride_channel_dst, + nsamples_x, nsamples_dst, stride_sample_x, stride_sample_y, stride_sample_dst, ids_stride, stream); + break; case GGML_TYPE_Q2_K: mul_mat_vec_q_switch_ncols_dst (vx, vy, ids, fusion, dst, ncols_x, nrows_x, ncols_dst, stride_row_x, stride_col_y, stride_col_dst, @@ -1581,6 +1644,8 @@ void ggml_cuda_mul_mat_vec_q( fusion_local.gate_bias = fusion->gate_bias->data; } fusion_local.glu_op = fusion->glu_op; + fusion_local.glu_param0 = fusion->glu_param0; + fusion_local.glu_param1 = fusion->glu_param1; } // If src0 is a temporary compute buffer, clear any potential padding. diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/unary.cu b/server/deps/llama.cpp/ggml/src/ggml-cuda/unary.cu index 4ad30fa1f..ed7f3d02a 100644 --- a/server/deps/llama.cpp/ggml/src/ggml-cuda/unary.cu +++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/unary.cu @@ -338,6 +338,63 @@ void ggml_cuda_op_swiglu(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { ggml_cuda_op_unary_gated(ctx, dst); } +// swiglu_ds4 + +template +static __global__ void swiglu_ds4_kernel(const T * gate, const T * up, T * dst, const int64_t k, const int64_t n, const int64_t o0, const int64_t o1, float limit) { + const int64_t i = int64_t(blockDim.x)*blockIdx.x + threadIdx.x; + + if (i >= k) { + return; + } + + const int64_t j0 = (i / n) * o0 + (i % n); + const int64_t j1 = o0 == o1 ? j0 : (i / n) * o1 + (i % n); + + const float gate_v = gate[j0]; + const float up_v = up[j1]; + + dst[i] = ggml_cuda_op_swiglu_ds4_single(gate_v, up_v, limit); +} + +template +static void swiglu_ds4_cuda(const T * gate, const T * up, T * dst, const int64_t k, const int64_t n, const int64_t o0, const int64_t o1, const float limit, cudaStream_t stream) { + const int64_t num_blocks = (k + CUDA_GLU_BLOCK_SIZE - 1) / CUDA_GLU_BLOCK_SIZE; + swiglu_ds4_kernel<<>>(gate, up, dst, k, n, o0, o1, limit); +} + +void ggml_cuda_op_swiglu_ds4(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { + const ggml_tensor * src0 = dst->src[0]; + const ggml_tensor * src1 = dst->src[1]; + void * src0_d = src0->data; + void * src1_d = src1 ? src1->data : src0->data; + const int64_t src0_o = src0->nb[1]; + const int64_t src1_o = src1 ? src1->nb[1] : src0->nb[1]; + void * dst_d = dst->data; + const int64_t nc = src1 ? src0->ne[0] : src0->ne[0] / 2; + cudaStream_t stream = ctx.stream(); + + GGML_ASSERT(ggml_is_contiguous_1(src0)); + GGML_ASSERT(src0->nb[0] == ggml_element_size(src0)); + GGML_ASSERT(ggml_is_contiguous(dst)); + + GGML_ASSERT(src0->type == GGML_TYPE_F32); + GGML_ASSERT( dst->type == GGML_TYPE_F32); + GGML_ASSERT(src0->type == dst->type); + GGML_ASSERT(dst->ne[0] == nc); + GGML_ASSERT(ggml_nrows(dst) == ggml_nrows(src0)); + + GGML_ASSERT(src1); + GGML_ASSERT(ggml_is_contiguous_1(src1)); + GGML_ASSERT(src1->nb[0] == ggml_element_size(src1)); + GGML_ASSERT(src1->ne[0] == nc); + GGML_ASSERT(src0->type == src1->type); + + const float limit = ggml_get_op_params_f32(dst, 2); + + swiglu_ds4_cuda((float *) src0_d, (float *) src1_d, (float *)dst_d, ggml_nelements(dst), nc, src0_o / sizeof(float), src1_o / sizeof(float), limit, stream); +} + void ggml_cuda_op_geglu_erf(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { ggml_cuda_op_unary_gated(ctx, dst); } diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/unary.cuh b/server/deps/llama.cpp/ggml/src/ggml-cuda/unary.cuh index f1dd2183a..1864ce2dc 100644 --- a/server/deps/llama.cpp/ggml/src/ggml-cuda/unary.cuh +++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/unary.cuh @@ -81,6 +81,8 @@ void ggml_cuda_op_geglu(ggml_backend_cuda_context & ctx, ggml_tensor * dst); void ggml_cuda_op_swiglu(ggml_backend_cuda_context & ctx, ggml_tensor * dst); +void ggml_cuda_op_swiglu_ds4(ggml_backend_cuda_context & ctx, ggml_tensor * dst); + void ggml_cuda_op_swiglu_oai(ggml_backend_cuda_context & ctx, ggml_tensor * dst); void ggml_cuda_op_geglu_erf(ggml_backend_cuda_context & ctx, ggml_tensor * dst); @@ -110,3 +112,11 @@ __device__ __forceinline__ float ggml_cuda_op_swiglu_oai_single(float x, float g out_glu = out_glu * (1.0f + g); return out_glu; } + +__device__ __forceinline__ float ggml_cuda_op_swiglu_ds4_single(float gate, float up, float limit) { + gate = fminf(gate, limit); + up = fmaxf(fminf(up, limit), -limit); + + const float silu = gate / (1.0f + expf(-gate)); + return silu * up; +} 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 40b2b41e7..bb75e7977 100644 --- a/server/deps/llama.cpp/ggml/src/ggml-cuda/vecdotq.cuh +++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/vecdotq.cuh @@ -1,6 +1,7 @@ #pragma once #include "common.cuh" +#include "../../rocmfp4/rocmfp4_hip_codebook.cuh" #include @@ -322,6 +323,311 @@ static __device__ __forceinline__ float vec_dot_mxfp4_q8_1( return d * sumi; } +// === ROCMFP vec_dot device functions (ported from ROCmFPX) === +#ifndef GGML_ROCMFP4_Q8_1_MMQ_VDR +#define GGML_ROCMFP4_Q8_1_MMQ_VDR 8 +#endif + +#ifndef GGML_ROCMFP4_FAST_Q8_1_MMQ_VDR +#define GGML_ROCMFP4_FAST_Q8_1_MMQ_VDR GGML_ROCMFP4_Q8_1_MMQ_VDR +#endif + +#ifndef GGML_ROCMFP4_FAST_Q8_1_MMVQ_VDR +#define GGML_ROCMFP4_FAST_Q8_1_MMVQ_VDR 2 +#endif + +#if GGML_ROCMFP4_FAST_Q8_1_MMVQ_VDR != 1 && \ + GGML_ROCMFP4_FAST_Q8_1_MMVQ_VDR != 2 && \ + GGML_ROCMFP4_FAST_Q8_1_MMVQ_VDR != 4 +#error "GGML_ROCMFP4_FAST_Q8_1_MMVQ_VDR must be 1, 2, or 4" +#endif + +#define VDR_ROCMFP4_Q8_1_MMVQ 4 +#define VDR_ROCMFP4_Q8_1_MMQ GGML_ROCMFP4_Q8_1_MMQ_VDR +#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 +#ifndef VDR_ROCMFP6_Q8_1_MMVQ +#define VDR_ROCMFP6_Q8_1_MMVQ 4 +#endif +#define VDR_ROCMFP8_Q8_1_MMVQ 2 + +#define VDR_ROCMFP2_Q8_1_MMQ 4 +#define VDR_ROCMFP3_Q8_1_MMQ 4 +#ifndef VDR_ROCMFP6_Q8_1_MMQ +#define VDR_ROCMFP6_Q8_1_MMQ 4 +#endif +#define VDR_ROCMFP8_Q8_1_MMQ 8 + +static __device__ __forceinline__ uint32_t rocmfpx_get_bits_vec_cuda(const uint8_t * src, const int bit_pos, const int nbits) { + uint32_t code = 0; + +#pragma unroll + for (int bit = 0; bit < nbits; ++bit) { + const int src_bit = bit_pos + bit; + code |= ((uint32_t) ((src[src_bit >> 3] >> (src_bit & 7)) & 1u)) << bit; + } + + return code; +} + +static __device__ __forceinline__ 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; +} + +static __device__ __forceinline__ int rocmfpx_decode_fp2_code_vec_cuda(const uint32_t code) { + switch (code & 3u) { + case 0: return ROCMFP2_KVALUE_0_I8; + case 1: return ROCMFP2_KVALUE_1_I8; + case 2: return ROCMFP2_KVALUE_2_I8; + default: return ROCMFP2_KVALUE_3_I8; + } +} + +static __device__ __forceinline__ int rocmfpx_decode_fp6_code_vec_cuda(const uint32_t code) { + const int mag = (int) (code & 31u); + return (code & 32u) ? -mag : mag; +} + +static __device__ __forceinline__ int rocmfpx_pack4_fp6_bits24_vec_cuda(const uint32_t bits24) { + const char4 v = make_char4( + (int8_t) rocmfpx_decode_fp6_code_vec_cuda(bits24 & 63u), + (int8_t) rocmfpx_decode_fp6_code_vec_cuda((bits24 >> 6) & 63u), + (int8_t) rocmfpx_decode_fp6_code_vec_cuda((bits24 >> 12) & 63u), + (int8_t) rocmfpx_decode_fp6_code_vec_cuda((bits24 >> 18) & 63u)); + return *((const int *) &v); +} + +static __device__ __forceinline__ int rocmfpx_pack4_fp3_vec_cuda(const uint8_t * qs, const int base) { + 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))); + return *((const int *) &v); +} + +static __device__ __forceinline__ int rocmfpx_pack4_fp2_bits8_vec_cuda(const uint32_t bits8) { +#if defined(GGML_USE_HIP) + const uint32_t values = + ((uint32_t) (uint8_t) (int8_t) ROCMFP2_KVALUE_0_I8) | + ((uint32_t) (uint8_t) (int8_t) ROCMFP2_KVALUE_1_I8 << 8) | + ((uint32_t) (uint8_t) (int8_t) ROCMFP2_KVALUE_2_I8 << 16) | + ((uint32_t) (uint8_t) (int8_t) ROCMFP2_KVALUE_3_I8 << 24); + const uint32_t selectors = + ((bits8 >> 0) & 3u) | + (((bits8 >> 2) & 3u) << 8) | + (((bits8 >> 4) & 3u) << 16) | + (((bits8 >> 6) & 3u) << 24); + return (int) __builtin_amdgcn_perm(0, values, selectors); +#else + const char4 v = make_char4( + (int8_t) rocmfpx_decode_fp2_code_vec_cuda((bits8 >> 0) & 3u), + (int8_t) rocmfpx_decode_fp2_code_vec_cuda((bits8 >> 2) & 3u), + (int8_t) rocmfpx_decode_fp2_code_vec_cuda((bits8 >> 4) & 3u), + (int8_t) rocmfpx_decode_fp2_code_vec_cuda((bits8 >> 6) & 3u)); + return *((const int *) &v); +#endif +} + +static __device__ __forceinline__ int rocmfpx_pack4_fp2_vec_cuda(const uint8_t * qs, const int base) { + return rocmfpx_pack4_fp2_bits8_vec_cuda(qs[base >> 2]); +} + +static __device__ __forceinline__ float vec_dot_rocmfp4_q8_1( + const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) { + + const block_rocmfp4 * bq4 = (const block_rocmfp4 *) vbq + kbx; + + const int * q8 = (const int *) bq8_1->qs + iqs; + + int sumi0 = 0; + int sumi1 = 0; +#pragma unroll + for (int l = 0; l < VDR_ROCMFP4_Q8_1_MMVQ; ++l) { + const int aux_q4 = rocmfp4_get_qs_i32(bq4->qs, iqs + l); + const int2 v = rocmfp4_get_int_from_codebook_16(aux_q4, kvalues_rocmfp4); + + sumi0 = ggml_cuda_dp4a(v.x, q8[l + 0], sumi0); + sumi1 = ggml_cuda_dp4a(v.y, q8[l + 4], sumi1); + } + + const float db = __low2float(bq8_1->ds); + return db * (rocmfp4_ue4m3_to_fp32_half_finite(bq4->e[0]) * sumi0 + rocmfp4_ue4m3_to_fp32_half_finite(bq4->e[1]) * sumi1); +} + +static __device__ __forceinline__ float vec_dot_rocmfp4_fast_q8_1( + const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) { + + const block_rocmfp4_fast * bq4 = (const block_rocmfp4_fast *) vbq + kbx; + + const int * q8 = (const int *) bq8_1->qs + iqs; + + int sumi = 0; +#pragma unroll + for (int l = 0; l < VDR_ROCMFP4_FAST_Q8_1_MMVQ; ++l) { + const int aux_q4 = rocmfp4_get_qs_i32(bq4->qs, iqs + l); + const int2 v = rocmfp4_get_int_from_codebook_16(aux_q4, kvalues_rocmfp4); + + sumi = ggml_cuda_dp4a(v.x, q8[l + 0], sumi); + sumi = ggml_cuda_dp4a(v.y, q8[l + 4], sumi); + } + + return __low2float(bq8_1->ds) * rocmfp4_ue4m3_to_fp32_half_finite(bq4->e) * sumi; +} + +static __device__ __forceinline__ float vec_dot_rocmfpx_fp2_q8_1( + const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) { + // int32-consistent MMVQ layout: QI_ROCMFP2=2 (8B qs = 2 int32), VDR=1. + // iqs in {0,1} selects one 16-weight half-block: qs bytes [4*iqs .. 4*iqs+3] + scale e[iqs]. + // 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); + } + + const float db = __low2float(bq8_1->ds); + return db * rocmfpx_ue4m3_to_fp32_finite(bq2->e[iqs]) * sumi; +} + +static __device__ __forceinline__ float vec_dot_rocmfpx_fp3_q8_1( + const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) { + + const block_rocmfp3 * bq3 = (const block_rocmfp3 *) vbq + kbx; + + uint32_t qs0, qs1, qs2; + memcpy(&qs0, bq3->qs + 0, 4); + memcpy(&qs1, bq3->qs + 4, 4); + memcpy(&qs2, bq3->qs + 8, 4); + + const uint32_t qs[4] = { qs0, qs1, qs2, 0 }; + + int sumi0 = 0; + int sumi1 = 0; + + // 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 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); +} + +static __device__ __forceinline__ float vec_dot_rocmfpx_fp6_q8_1( + const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) { + + const block_rocmfp6 * bq6 = (const block_rocmfp6 *) vbq + kbx; + + uint32_t qs0, qs1, qs2, qs3, qs4, qs5; + memcpy(&qs0, bq6->qs + 0, 4); + memcpy(&qs1, bq6->qs + 4, 4); + memcpy(&qs2, bq6->qs + 8, 4); + memcpy(&qs3, bq6->qs + 12, 4); + memcpy(&qs4, bq6->qs + 16, 4); + memcpy(&qs5, bq6->qs + 20, 4); + + // Trailing 0 pad: the last FP6 window reads qs[reg_idx + 1] with reg_idx==5; + // its high bits are always masked out, so 0 keeps the result bit-identical + // while avoiding a stack over-read (matches the FP3 sibling below). + const uint32_t qs[7] = { qs0, qs1, qs2, qs3, qs4, qs5, 0 }; + + int sumi0 = 0; + int sumi1 = 0; + + const bool fp6_first_half = iqs + VDR_ROCMFP6_Q8_1_MMVQ <= QK_ROCMFP6/8; + const bool fp6_second_half = iqs >= QK_ROCMFP6/8; + +#pragma unroll + for (int i = 0; i < VDR_ROCMFP6_Q8_1_MMVQ; ++i) { + const int base = 4 * (iqs + i); + const int start_bit = 6 * base; + 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 bits24 = (reg_shift == 0) ? (val_low & 0xFFFFFFu) : + (((val_low >> reg_shift) | (val_high << (32 - reg_shift))) & 0xFFFFFFu); + + const int val_packed = rocmfpx_pack4_fp6_bits24_vec_cuda(bits24); + const int u = get_int_b4(bq8_1->qs, iqs + i); + + if (fp6_first_half) { + sumi0 = ggml_cuda_dp4a(val_packed, u, sumi0); + } else if (fp6_second_half) { + sumi1 = ggml_cuda_dp4a(val_packed, u, sumi1); + } else + if (base < QK_ROCMFP6/2) { + sumi0 = ggml_cuda_dp4a(val_packed, u, sumi0); + } else { + sumi1 = ggml_cuda_dp4a(val_packed, u, sumi1); + } + } + + const float db = __low2float(bq8_1->ds); + return db * (rocmfpx_ue4m3_to_fp32_finite(bq6->e[0]) * sumi0 + rocmfpx_ue4m3_to_fp32_finite(bq6->e[1]) * sumi1); +} + +static __device__ __forceinline__ float vec_dot_rocmfpx_fp8_q8_1( + const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) { + + const block_rocmfp8 * bq8 = (const block_rocmfp8 *) vbq + kbx; + + int v[VDR_ROCMFP8_Q8_1_MMVQ]; + int u[VDR_ROCMFP8_Q8_1_MMVQ]; + +#pragma unroll + for (int i = 0; i < VDR_ROCMFP8_Q8_1_MMVQ; ++i) { + v[i] = get_int_b1(bq8->qs, iqs + i); + u[i] = get_int_b4(bq8_1->qs, iqs + i); + } + + return vec_dot_q8_0_q8_1_impl( + v, u, rocmfpx_ue4m3_to_fp32_finite(bq8->e), __low2half(bq8_1->ds)); +} + #define VDR_NVFP4_Q8_1_MMVQ 4 #define VDR_NVFP4_Q8_1_MMQ 8 diff --git a/server/deps/llama.cpp/ggml/src/ggml.c b/server/deps/llama.cpp/ggml/src/ggml.c index fda42ef76..db026f34e 100644 --- a/server/deps/llama.cpp/ggml/src/ggml.c +++ b/server/deps/llama.cpp/ggml/src/ggml.c @@ -9,6 +9,8 @@ // FIXME: required here for quantization functions #include "ggml-quants.h" +#include "../rocmfp4/rocmfp4.h" +#include "../rocmfpx/rocmfpx.h" #ifdef GGML_USE_CPU_HBM #include @@ -719,6 +721,54 @@ static const struct ggml_type_traits type_traits[GGML_TYPE_COUNT] = { .to_float = (ggml_to_float_t) dequantize_row_q8_0, .from_float_ref = (ggml_from_float_t) quantize_row_q8_0_ref, }, + [GGML_TYPE_Q4_0_ROCMFP4] = { + .type_name = "q4_0_rocmfp4", + .blck_size = QK_ROCMFP4, + .type_size = sizeof(block_rocmfp4), + .is_quantized = true, + .to_float = (ggml_to_float_t) rocmfp4_dequantize_row_q4_0, + .from_float_ref = (ggml_from_float_t) rocmfp4_quantize_row_q4_0_ref, + }, + [GGML_TYPE_Q4_0_ROCMFP4_FAST] = { + .type_name = "q4_0_rocmfp4_fast", + .blck_size = QK_ROCMFP4, + .type_size = sizeof(block_rocmfp4_fast), + .is_quantized = true, + .to_float = (ggml_to_float_t) rocmfp4_dequantize_row_q4_0_fast, + .from_float_ref = (ggml_from_float_t) rocmfp4_quantize_row_q4_0_fast_ref, + }, + [GGML_TYPE_Q3_0_ROCMFPX] = { + .type_name = "q3_0_rocmfpx", + .blck_size = QK_ROCMFP3, + .type_size = sizeof(block_rocmfp3), + .is_quantized = true, + .to_float = (ggml_to_float_t) rocmfpx_dequantize_row_fp3, + .from_float_ref = (ggml_from_float_t) rocmfpx_quantize_row_fp3_ref, + }, + [GGML_TYPE_Q2_0_ROCMFP2] = { + .type_name = "q2_0_rocmfp2", + .blck_size = QK_ROCMFP2, + .type_size = sizeof(block_rocmfp2), + .is_quantized = true, + .to_float = (ggml_to_float_t) rocmfpx_dequantize_row_fp2, + .from_float_ref = (ggml_from_float_t) rocmfpx_quantize_row_fp2_ref, + }, + [GGML_TYPE_Q6_0_ROCMFPX] = { + .type_name = "q6_0_rocmfpx", + .blck_size = QK_ROCMFP6, + .type_size = sizeof(block_rocmfp6), + .is_quantized = true, + .to_float = (ggml_to_float_t) rocmfpx_dequantize_row_fp6, + .from_float_ref = (ggml_from_float_t) rocmfpx_quantize_row_fp6_ref, + }, + [GGML_TYPE_Q8_0_ROCMFPX] = { + .type_name = "q8_0_rocmfpx", + .blck_size = QK_ROCMFP8, + .type_size = sizeof(block_rocmfp8), + .is_quantized = true, + .to_float = (ggml_to_float_t) rocmfpx_dequantize_row_fp8, + .from_float_ref = (ggml_from_float_t) rocmfpx_quantize_row_fp8_ref, + }, [GGML_TYPE_Q8_1] = { .type_name = "q8_1", .blck_size = QK8_1, @@ -1074,9 +1124,13 @@ static const char * GGML_OP_NAME[GGML_OP_COUNT] = { "GLU", "TURBO_WHT", + + "MOE_FUSED", + + "DS4_HC", }; -static_assert(GGML_OP_COUNT == 99, "GGML_OP_COUNT != 99"); +static_assert(GGML_OP_COUNT == 100, "GGML_OP_COUNT != 100"); static const char * GGML_OP_SYMBOL[GGML_OP_COUNT] = { "none", @@ -1187,9 +1241,13 @@ static const char * GGML_OP_SYMBOL[GGML_OP_COUNT] = { "glu(x)", "turbo_wht(a)", + + "moe_fused(x)", + + "ds4_hc(x)", }; -static_assert(GGML_OP_COUNT == 99, "GGML_OP_COUNT != 99"); +static_assert(GGML_OP_COUNT == 100, "GGML_OP_COUNT != 100"); static_assert(GGML_OP_POOL_COUNT == 2, "GGML_OP_POOL_COUNT != 2"); @@ -1225,11 +1283,12 @@ static const char * GGML_GLU_OP_NAME[GGML_GLU_OP_COUNT] = { "GEGLU", "SWIGLU", "SWIGLU_OAI", + "SWIGLU_DS4", "GEGLU_ERF", "GEGLU_QUICK", }; -static_assert(GGML_GLU_OP_COUNT == 6, "GGML_GLU_OP_COUNT != 6"); +static_assert(GGML_GLU_OP_COUNT == 7, "GGML_GLU_OP_COUNT != 7"); static_assert(sizeof(struct ggml_object)%GGML_MEM_ALIGN == 0, "ggml_object size must be a multiple of GGML_MEM_ALIGN"); @@ -1325,7 +1384,8 @@ double ggml_type_sizef(enum ggml_type type) { const char * ggml_type_name(enum ggml_type type) { assert(type >= 0); assert(type < GGML_TYPE_COUNT); - return type_traits[type].type_name; + const char * name = type_traits[type].type_name; + return name ? name : "unknown"; } bool ggml_is_quantized(enum ggml_type type) { @@ -1428,6 +1488,20 @@ enum ggml_type ggml_ftype_to_ggml_type(enum ggml_ftype ftype) { case GGML_FTYPE_MOSTLY_IQ2_S: wtype = GGML_TYPE_IQ2_S; break; case GGML_FTYPE_UNKNOWN: wtype = GGML_TYPE_COUNT; break; case GGML_FTYPE_MOSTLY_Q4_1_SOME_F16: wtype = GGML_TYPE_COUNT; break; + // ROCmFPX presets are mixed-quant recipes (see ggml/rocmfp4/README.md); + // report the dominant transformer-tensor type. + case GGML_FTYPE_MOSTLY_Q4_0_ROCMFP4: + case GGML_FTYPE_MOSTLY_Q4_0_ROCMFP4_LEAN: + case GGML_FTYPE_MOSTLY_Q4_0_ROCMFP4_COHERENT: wtype = GGML_TYPE_Q4_0_ROCMFP4; break; + case GGML_FTYPE_MOSTLY_Q4_0_ROCMFP4_FAST: + case GGML_FTYPE_MOSTLY_Q4_0_ROCMFP4_FAST_COHERENT: + case GGML_FTYPE_MOSTLY_Q4_0_ROCMFP4_STRIX: + case GGML_FTYPE_MOSTLY_Q4_0_ROCMFP4_STRIX_LEAN: wtype = GGML_TYPE_Q4_0_ROCMFP4_FAST; break; + case GGML_FTYPE_MOSTLY_Q6_0_ROCMFPX: wtype = GGML_TYPE_Q6_0_ROCMFPX; break; + case GGML_FTYPE_MOSTLY_Q8_0_ROCMFPX: wtype = GGML_TYPE_Q8_0_ROCMFPX; break; + case GGML_FTYPE_MOSTLY_Q3_0_ROCMFPX: wtype = GGML_TYPE_Q3_0_ROCMFPX; break; + case GGML_FTYPE_MOSTLY_Q2_0_ROCMFP2: + case GGML_FTYPE_MOSTLY_Q2_0_ROCMFP2_STRIX: wtype = GGML_TYPE_Q2_0_ROCMFP2; break; } GGML_ASSERT(wtype != GGML_TYPE_COUNT); @@ -3025,6 +3099,16 @@ struct ggml_tensor * ggml_swiglu_split( return ggml_glu_impl(ctx, a, b, GGML_GLU_OP_SWIGLU, false); } +struct ggml_tensor * ggml_swiglu_ds4_split( + struct ggml_context * ctx, + struct ggml_tensor * gate, + struct ggml_tensor * up, + float clamp) { + struct ggml_tensor * result = ggml_glu_impl(ctx, gate, up, GGML_GLU_OP_SWIGLU_DS4, false); + ggml_set_op_params_f32(result, 2, clamp); + return result; +} + // ggml_geglu_erf struct ggml_tensor * ggml_geglu_erf( @@ -7991,3 +8075,102 @@ struct ggml_tensor * ggml_laguna_moe_combine( return result; } + +struct ggml_tensor * ggml_ds4_hc_pre( + struct ggml_context * ctx, + struct ggml_tensor * mix, + struct ggml_tensor * base, + struct ggml_tensor * hc_state, + int n_hc, + int sinkhorn_iters, + float pre_scale, + float post_scale, + float comb_scale) { + GGML_ASSERT(mix->type == GGML_TYPE_F32); + GGML_ASSERT(base->type == GGML_TYPE_F32); + GGML_ASSERT(hc_state->type == GGML_TYPE_F32); + GGML_ASSERT(ggml_is_contiguous(mix)); + GGML_ASSERT(ggml_is_contiguous(base)); + GGML_ASSERT(ggml_is_contiguous(hc_state)); + GGML_ASSERT(n_hc > 0 && n_hc <= 8); + const int64_t mix_dim = 2*(int64_t)n_hc + (int64_t)n_hc*n_hc; + GGML_ASSERT(ggml_nelements(mix) == mix_dim); + GGML_ASSERT(ggml_nelements(base) >= mix_dim); + GGML_ASSERT(ggml_nelements(hc_state) % n_hc == 0); + const int64_t n_embd = ggml_nelements(hc_state) / n_hc; + + struct ggml_tensor * result = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, n_embd + mix_dim); + result->op = GGML_OP_DS4_HC; + result->src[0] = mix; + result->src[1] = base; + result->src[2] = hc_state; + ggml_set_op_params_i32(result, 0, 0); + ggml_set_op_params_i32(result, 1, (int32_t) n_embd); + ggml_set_op_params_i32(result, 2, (int32_t) n_hc); + ggml_set_op_params_i32(result, 3, (int32_t) sinkhorn_iters); + ggml_set_op_params_f32(result, 4, pre_scale); + ggml_set_op_params_f32(result, 5, post_scale); + ggml_set_op_params_f32(result, 6, comb_scale); + return result; +} + +struct ggml_tensor * ggml_ds4_hc_post( + struct ggml_context * ctx, + struct ggml_tensor * residual_hc, + struct ggml_tensor * block_out, + struct ggml_tensor * split, + int n_hc) { + GGML_ASSERT(residual_hc->type == GGML_TYPE_F32); + GGML_ASSERT(block_out->type == GGML_TYPE_F32); + GGML_ASSERT(split->type == GGML_TYPE_F32); + GGML_ASSERT(ggml_is_contiguous(residual_hc)); + GGML_ASSERT(ggml_is_contiguous(block_out)); + GGML_ASSERT(ggml_is_contiguous(split)); + GGML_ASSERT(n_hc > 0 && n_hc <= 8); + const int64_t mix_dim = 2*(int64_t)n_hc + (int64_t)n_hc*n_hc; + GGML_ASSERT(ggml_nelements(split) == mix_dim); + GGML_ASSERT(ggml_nelements(residual_hc) % n_hc == 0); + const int64_t n_embd = ggml_nelements(residual_hc) / n_hc; + GGML_ASSERT(ggml_nelements(block_out) == n_embd); + + struct ggml_tensor * result = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, (int64_t) n_embd * n_hc); + result->op = GGML_OP_DS4_HC; + result->src[0] = residual_hc; + result->src[1] = block_out; + result->src[2] = split; + ggml_set_op_params_i32(result, 0, 1); + ggml_set_op_params_i32(result, 1, (int32_t) n_embd); + ggml_set_op_params_i32(result, 2, (int32_t) n_hc); + return result; +} + +struct ggml_tensor * ggml_ds4_hc_out( + struct ggml_context * ctx, + struct ggml_tensor * mix, + struct ggml_tensor * base, + struct ggml_tensor * hc_state, + int n_hc, + float pre_scale) { + GGML_ASSERT(mix->type == GGML_TYPE_F32); + GGML_ASSERT(base->type == GGML_TYPE_F32); + GGML_ASSERT(hc_state->type == GGML_TYPE_F32); + GGML_ASSERT(ggml_is_contiguous(mix)); + GGML_ASSERT(ggml_is_contiguous(base)); + GGML_ASSERT(ggml_is_contiguous(hc_state)); + GGML_ASSERT(n_hc > 0 && n_hc <= 8); + GGML_ASSERT(ggml_nelements(mix) >= n_hc); + GGML_ASSERT(ggml_nelements(base) >= n_hc); + GGML_ASSERT(ggml_nelements(hc_state) % n_hc == 0); + const int64_t n_embd = ggml_nelements(hc_state) / n_hc; + + struct ggml_tensor * result = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, n_embd); + result->op = GGML_OP_DS4_HC; + result->src[0] = mix; + result->src[1] = base; + result->src[2] = hc_state; + ggml_set_op_params_i32(result, 0, 2); + ggml_set_op_params_i32(result, 1, (int32_t) n_embd); + ggml_set_op_params_i32(result, 2, (int32_t) n_hc); + ggml_set_op_params_f32(result, 4, pre_scale); + return result; +} diff --git a/server/test/test_rocmfp4_hip_tail.cpp b/server/test/test_rocmfp4_hip_tail.cpp new file mode 100644 index 000000000..f3b80c06a --- /dev/null +++ b/server/test/test_rocmfp4_hip_tail.cpp @@ -0,0 +1,134 @@ +#include + +#include "ggml.h" +#include "rocmfp4.h" + +#include +#include +#include +#include +#include + +using to_fp32_cuda_t = void (*)(const void *, float *, int64_t, hipStream_t); +to_fp32_cuda_t ggml_get_to_fp32_cuda(ggml_type type); + +static void hip_check(hipError_t status, const char * operation) { + if (status != hipSuccess) { + std::fprintf(stderr, "%s: %s\n", operation, hipGetErrorString(status)); + std::exit(1); + } +} + +static std::vector make_input(int64_t size) { + std::vector input((size_t) size); + for (int64_t i = 0; i < size; ++i) { + input[(size_t) i] = ((float) ((i*37) % 101) - 50.0f)*0.03125f; + } + return input; +} + +template +static void check_layout( + ggml_type type, + int64_t size, + void (*quantize)(const float *, Block *, int64_t), + void (*dequantize)(const Block *, float *, int64_t), + const char * label) { + const std::vector input = make_input(size); + std::vector quantized((size_t) (size / QK_ROCMFP4)); + std::vector expected((size_t) size); + quantize(input.data(), quantized.data(), size); + dequantize(quantized.data(), expected.data(), size); + + constexpr int64_t guard_count = 256; + constexpr float guard_value = 12345.25f; + std::vector actual((size_t) (size + guard_count), guard_value); + + void * device_quantized = nullptr; + float * device_actual = nullptr; + hip_check(hipMalloc(&device_quantized, quantized.size()*sizeof(Block)), "hipMalloc input"); + hip_check(hipMalloc(&device_actual, actual.size()*sizeof(float)), "hipMalloc output"); + hip_check( + hipMemcpy( + device_quantized, + quantized.data(), + quantized.size()*sizeof(Block), + hipMemcpyHostToDevice), + "hipMemcpy input"); + hip_check( + hipMemcpy( + device_actual, + actual.data(), + actual.size()*sizeof(float), + hipMemcpyHostToDevice), + "hipMemcpy output"); + + const to_fp32_cuda_t convert = ggml_get_to_fp32_cuda(type); + if (convert == nullptr) { + std::fprintf(stderr, "%s size=%lld: missing converter\n", label, (long long) size); + std::exit(1); + } + + convert(device_quantized, device_actual, size, nullptr); + hip_check(hipGetLastError(), "ROCmFP4 converter launch"); + hip_check(hipDeviceSynchronize(), "ROCmFP4 converter synchronize"); + hip_check( + hipMemcpy( + actual.data(), + device_actual, + actual.size()*sizeof(float), + hipMemcpyDeviceToHost), + "hipMemcpy result"); + + if (std::memcmp(actual.data(), expected.data(), (size_t) size*sizeof(float)) != 0) { + for (int64_t i = 0; i < size; ++i) { + if (std::memcmp(&actual[(size_t) i], &expected[(size_t) i], sizeof(float)) != 0) { + std::fprintf( + stderr, + "%s size=%lld mismatch at %lld: gpu=%a cpu=%a\n", + label, + (long long) size, + (long long) i, + actual[(size_t) i], + expected[(size_t) i]); + break; + } + } + std::exit(1); + } + + for (int64_t i = size; i < size + guard_count; ++i) { + if (actual[(size_t) i] != guard_value) { + std::fprintf( + stderr, + "%s size=%lld overwrote guard at +%lld\n", + label, + (long long) size, + (long long) (i - size)); + std::exit(1); + } + } + + hip_check(hipFree(device_actual), "hipFree output"); + hip_check(hipFree(device_quantized), "hipFree input"); + std::printf("%s size=%lld: byte-identical, guard intact\n", label, (long long) size); +} + +int main() { + const int64_t sizes[] = { 32, 64, 96, 224, 256, 288 }; + for (const int64_t size : sizes) { + check_layout( + GGML_TYPE_Q4_0_ROCMFP4, + size, + rocmfp4_quantize_row_q4_0_ref, + rocmfp4_dequantize_row_q4_0, + "dual"); + check_layout( + GGML_TYPE_Q4_0_ROCMFP4_FAST, + size, + rocmfp4_quantize_row_q4_0_fast_ref, + rocmfp4_dequantize_row_q4_0_fast, + "fast"); + } + return 0; +} diff --git a/server/tests/test_deepseek4_unit.cpp b/server/tests/test_deepseek4_unit.cpp index bb32aeed2..f36b07d33 100644 --- a/server/tests/test_deepseek4_unit.cpp +++ b/server/tests/test_deepseek4_unit.cpp @@ -294,6 +294,59 @@ static void test_compressor_pooling_correctness(ggml_backend_t backend) { std::fprintf(stderr, g_failures ? " done\n" : " ok\n"); } +static void test_swiglu_ds4_cpu_correctness(ggml_backend_t backend) { + std::fprintf(stderr, " test_swiglu_ds4_cpu_correctness ..."); + + constexpr int dim = 17; + constexpr int rows = 3; + constexpr float clamp = 1.5f; + std::vector gate((size_t) dim * rows); + std::vector up((size_t) dim * rows); + std::vector expected((size_t) dim * rows); + for (size_t i = 0; i < gate.size(); ++i) { + gate[i] = 0.25f * (float) ((int) (i % 15) - 7); + up[i] = 0.375f * (float) ((int) (i % 11) - 5); + const float gate_clamped = std::min(gate[i], clamp); + const float up_clamped = std::clamp(up[i], -clamp, clamp); + expected[i] = up_clamped * gate_clamped / (1.0f + std::exp(-gate_clamped)); + } + + ggml_context * ctx = make_test_context(); + TEST_ASSERT_MSG(ctx != nullptr, "ggml_init failed"); + if (!ctx) { + std::fprintf(stderr, " FAIL\n"); + return; + } + + ggml_tensor * gate_t = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, dim, rows); + ggml_tensor * up_t = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, dim, rows); + ggml_set_input(gate_t); + ggml_set_input(up_t); + ggml_tensor * out_t = ggml_swiglu_ds4_split(ctx, gate_t, up_t, clamp); + ggml_set_output(out_t); + + ggml_cgraph * gf = ggml_new_graph_custom(ctx, 16, false); + ggml_build_forward_expand(gf, out_t); + ggml_gallocr_t alloc = ggml_gallocr_new(ggml_backend_cpu_buffer_type()); + TEST_ASSERT(ggml_gallocr_alloc_graph(alloc, gf)); + ggml_backend_tensor_set(gate_t, gate.data(), 0, gate.size() * sizeof(float)); + ggml_backend_tensor_set(up_t, up.data(), 0, up.size() * sizeof(float)); + TEST_ASSERT(ggml_backend_graph_compute(backend, gf) == GGML_STATUS_SUCCESS); + + std::vector actual(expected.size()); + ggml_backend_tensor_get(out_t, actual.data(), 0, actual.size() * sizeof(float)); + ggml_gallocr_free(alloc); + ggml_free(ctx); + + for (size_t i = 0; i < actual.size(); ++i) { + TEST_ASSERT_MSG( + nearly_equal(actual[i], expected[i], 1.0e-6f, 1.0e-6f), + "SWIGLU_DS4 output mismatch"); + } + + std::fprintf(stderr, g_failures ? " done\n" : " ok\n"); +} + static void test_moe_routing_correctness(ggml_backend_t backend) { std::fprintf(stderr, " test_moe_routing_correctness ..."); @@ -1551,6 +1604,7 @@ int main() { } test_compressor_pooling_correctness(backend); + test_swiglu_ds4_cpu_correctness(backend); test_moe_routing_correctness(backend); test_rmsnorm_correctness(backend); test_grouped_output_projection_shape();