rocmfp4 refactor

Author: Annie Ren

ggml/src/CMakeLists.txt

@@ -206,8 +206,6 @@ add_library(ggml-base
             ggml-threading.h
             ggml-quants.c
             ggml-quants.h
-            ../rocmfp4/rocmfp4.c
-            ../rocmfp4/rocmfp4.h
             gguf.cpp)
 
 set_target_properties(ggml-base PROPERTIES

ggml/src/ggml-common.h

@@ -112,6 +112,9 @@ typedef sycl::half2 ggml_half2;
 #define QI_NVFP4 (QK_NVFP4 / (4 * QR_NVFP4))
 #define QR_NVFP4 2
 
+#define QI_ROCMFP4 (QK_ROCMFP4 / (4 * QR_ROCMFP4))
+#define QR_ROCMFP4 2
+
 #define QI5_0 (QK5_0 / (4 * QR5_0))
 #define QR5_0 2
 
@@ -226,6 +229,23 @@ typedef struct {
 } block_nvfp4;
 static_assert(sizeof(block_nvfp4) == sizeof(uint8_t)*(QK_NVFP4/QK_NVFP4_SUB) + QK_NVFP4/2, "wrong nvfp4 block size/padding");
 
+#define QK_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;
+static_assert(sizeof(block_rocmfp4) == QK_ROCMFP4/2 + 2*sizeof(uint8_t), "wrong rocmfp4 block size/padding");
+
+// Speed-focused layout: same 32 packed ROCmFP4 nibbles, but one UE4M3 scale
+// for the whole block.
+typedef struct {
+    uint8_t qs[QK_ROCMFP4/2];
+    uint8_t e;
+} block_rocmfp4_fast;
+static_assert(sizeof(block_rocmfp4_fast) == QK_ROCMFP4/2 + sizeof(uint8_t), "wrong rocmfp4 fast block size/padding");
+
 #define QK5_0 32
 typedef struct {
     ggml_half d;           // delta
@@ -1136,6 +1156,47 @@ 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()
 
+// ROCmFP4 UE4M3 "half-scale" values for the finite scale bytes 0x00..0x7e (127
+// entries): the subnormal run (byte>>3 == 0, value M*2^-10) followed by 15
+// normal exponent groups (value (8+M)*2^(e-10), e=1..15). This is the single
+// source of truth shared by both materializations of the table: the CPU
+// quantizer scale-search table (ggml-quants.c) and the opt-in GPU
+// constant-memory LUT (ggml-cuda/common.cuh). Each backend stamps its own
+// storage-qualified array from this list so the two can never drift.
+#define GGML_ROCMFP4_SCALE_UE4M3_HALF_LIST \
+    (0) * 0x1p-10f, (1) * 0x1p-10f, (2) * 0x1p-10f, (3) * 0x1p-10f, \
+    (4) * 0x1p-10f, (5) * 0x1p-10f, (6) * 0x1p-10f, (7) * 0x1p-10f, \
+    (8 + 0) * 0x1p-10f, (8 + 1) * 0x1p-10f, (8 + 2) * 0x1p-10f, (8 + 3) * 0x1p-10f, \
+    (8 + 4) * 0x1p-10f, (8 + 5) * 0x1p-10f, (8 + 6) * 0x1p-10f, (8 + 7) * 0x1p-10f, \
+    (8 + 0) * 0x1p-9f, (8 + 1) * 0x1p-9f, (8 + 2) * 0x1p-9f, (8 + 3) * 0x1p-9f, \
+    (8 + 4) * 0x1p-9f, (8 + 5) * 0x1p-9f, (8 + 6) * 0x1p-9f, (8 + 7) * 0x1p-9f, \
+    (8 + 0) * 0x1p-8f, (8 + 1) * 0x1p-8f, (8 + 2) * 0x1p-8f, (8 + 3) * 0x1p-8f, \
+    (8 + 4) * 0x1p-8f, (8 + 5) * 0x1p-8f, (8 + 6) * 0x1p-8f, (8 + 7) * 0x1p-8f, \
+    (8 + 0) * 0x1p-7f, (8 + 1) * 0x1p-7f, (8 + 2) * 0x1p-7f, (8 + 3) * 0x1p-7f, \
+    (8 + 4) * 0x1p-7f, (8 + 5) * 0x1p-7f, (8 + 6) * 0x1p-7f, (8 + 7) * 0x1p-7f, \
+    (8 + 0) * 0x1p-6f, (8 + 1) * 0x1p-6f, (8 + 2) * 0x1p-6f, (8 + 3) * 0x1p-6f, \
+    (8 + 4) * 0x1p-6f, (8 + 5) * 0x1p-6f, (8 + 6) * 0x1p-6f, (8 + 7) * 0x1p-6f, \
+    (8 + 0) * 0x1p-5f, (8 + 1) * 0x1p-5f, (8 + 2) * 0x1p-5f, (8 + 3) * 0x1p-5f, \
+    (8 + 4) * 0x1p-5f, (8 + 5) * 0x1p-5f, (8 + 6) * 0x1p-5f, (8 + 7) * 0x1p-5f, \
+    (8 + 0) * 0x1p-4f, (8 + 1) * 0x1p-4f, (8 + 2) * 0x1p-4f, (8 + 3) * 0x1p-4f, \
+    (8 + 4) * 0x1p-4f, (8 + 5) * 0x1p-4f, (8 + 6) * 0x1p-4f, (8 + 7) * 0x1p-4f, \
+    (8 + 0) * 0x1p-3f, (8 + 1) * 0x1p-3f, (8 + 2) * 0x1p-3f, (8 + 3) * 0x1p-3f, \
+    (8 + 4) * 0x1p-3f, (8 + 5) * 0x1p-3f, (8 + 6) * 0x1p-3f, (8 + 7) * 0x1p-3f, \
+    (8 + 0) * 0x1p-2f, (8 + 1) * 0x1p-2f, (8 + 2) * 0x1p-2f, (8 + 3) * 0x1p-2f, \
+    (8 + 4) * 0x1p-2f, (8 + 5) * 0x1p-2f, (8 + 6) * 0x1p-2f, (8 + 7) * 0x1p-2f, \
+    (8 + 0) * 0x1p-1f, (8 + 1) * 0x1p-1f, (8 + 2) * 0x1p-1f, (8 + 3) * 0x1p-1f, \
+    (8 + 4) * 0x1p-1f, (8 + 5) * 0x1p-1f, (8 + 6) * 0x1p-1f, (8 + 7) * 0x1p-1f, \
+    (8 + 0) * 0x1p+0f, (8 + 1) * 0x1p+0f, (8 + 2) * 0x1p+0f, (8 + 3) * 0x1p+0f, \
+    (8 + 4) * 0x1p+0f, (8 + 5) * 0x1p+0f, (8 + 6) * 0x1p+0f, (8 + 7) * 0x1p+0f, \
+    (8 + 0) * 0x1p+1f, (8 + 1) * 0x1p+1f, (8 + 2) * 0x1p+1f, (8 + 3) * 0x1p+1f, \
+    (8 + 4) * 0x1p+1f, (8 + 5) * 0x1p+1f, (8 + 6) * 0x1p+1f, (8 + 7) * 0x1p+1f, \
+    (8 + 0) * 0x1p+2f, (8 + 1) * 0x1p+2f, (8 + 2) * 0x1p+2f, (8 + 3) * 0x1p+2f, \
+    (8 + 4) * 0x1p+2f, (8 + 5) * 0x1p+2f, (8 + 6) * 0x1p+2f, (8 + 7) * 0x1p+2f, \
+    (8 + 0) * 0x1p+3f, (8 + 1) * 0x1p+3f, (8 + 2) * 0x1p+3f, (8 + 3) * 0x1p+3f, \
+    (8 + 4) * 0x1p+3f, (8 + 5) * 0x1p+3f, (8 + 6) * 0x1p+3f, (8 + 7) * 0x1p+3f, \
+    (8 + 0) * 0x1p+4f, (8 + 1) * 0x1p+4f, (8 + 2) * 0x1p+4f, (8 + 3) * 0x1p+4f, \
+    (8 + 4) * 0x1p+4f, (8 + 5) * 0x1p+4f, (8 + 6) * 0x1p+4f
+
 #define NGRID_IQ1S 2048
 #define IQ1S_DELTA 0.125f
 #define IQ1M_DELTA 0.125f

ggml/src/ggml-cpu/ggml-cpu.c

@@ -14,7 +14,6 @@
 #include "ops.h"
 #include "ggml.h"
 #include "common.h"
-#include "../../rocmfp4/rocmfp4.h"
 
 #if defined(_MSC_VER) || defined(__MINGW32__)
 #include <malloc.h> // using malloc.h with MSC/MINGW

ggml/src/ggml-cpu/quants.c

@@ -62,6 +62,14 @@ void quantize_row_nvfp4(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, i
     quantize_row_nvfp4_ref(x, y, k);
 }
 
+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);
+}
+
 //
 // 2-6 bit quantization in super-blocks
 //
@@ -362,6 +370,71 @@ void ggml_vec_dot_nvfp4_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs,
     *s = sumf;
 }
 
+// ROCmFP4: Q4_0-layout FP4 (QK_ROCMFP4 == QK8_0), two UE4M3 half-scales per block.
+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) {
+    UNUSED(bs);
+    UNUSED(bx);
+    UNUSED(by);
+    assert(nrc == 1);
+    UNUSED(nrc);
+    assert(n % QK_ROCMFP4 == 0);
+    static_assert(QK_ROCMFP4 == QK8_0, "QK_ROCMFP4 and QK8_0 must be the same");
+
+    const block_rocmfp4 * GGML_RESTRICT x = vx;
+    const block_q8_0    * GGML_RESTRICT y = vy;
+
+    const int nb = n / QK_ROCMFP4;
+    float sumf = 0;
+
+    for (int ib = 0; ib < nb; ++ib) {
+        const float d0 = ggml_ue4m3_to_fp32(x[ib].e[0]) * GGML_CPU_FP16_TO_FP32(y[ib].d);
+        const float d1 = ggml_ue4m3_to_fp32(x[ib].e[1]) * GGML_CPU_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 += kvalues_rocmfp4[q & 0x0f] * y[ib].qs[j];
+            sumi1 += kvalues_rocmfp4[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) {
+    UNUSED(bs);
+    UNUSED(bx);
+    UNUSED(by);
+    assert(nrc == 1);
+    UNUSED(nrc);
+    assert(n % QK_ROCMFP4 == 0);
+    static_assert(QK_ROCMFP4 == QK8_0, "QK_ROCMFP4 and QK8_0 must be the same");
+
+    const block_rocmfp4_fast * GGML_RESTRICT x = vx;
+    const block_q8_0         * GGML_RESTRICT y = vy;
+
+    const int nb = n / QK_ROCMFP4;
+    float sumf = 0;
+
+    for (int ib = 0; ib < nb; ++ib) {
+        const float d = ggml_ue4m3_to_fp32(x[ib].e) * GGML_CPU_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 += kvalues_rocmfp4[q & 0x0f] * y[ib].qs[j];
+            sumi += kvalues_rocmfp4[q >>   4] * y[ib].qs[j + QK_ROCMFP4/2];
+        }
+
+        sumf += d * (float) sumi;
+    }
+
+    *s = sumf;
+}
+
 void ggml_vec_dot_q5_0_q8_0_generic(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) {
     const int qk = QK8_0;
     const int nb = n / qk;

ggml/src/ggml-cpu/quants.h

@@ -24,6 +24,10 @@ void quantize_row_q8_1(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, in
 void quantize_row_mxfp4(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
 void quantize_row_nvfp4(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
 
+// CPU from_float wrappers; defined in ggml-cpu/quants.c over the _ref quantizers in ggml-quants.c.
+GGML_API void rocmfp4_quantize_row_q4_0     (const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
+GGML_API void rocmfp4_quantize_row_q4_0_fast(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
+
 void quantize_row_q2_K(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
 void quantize_row_q3_K(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
 void quantize_row_q4_K(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
@@ -49,6 +53,9 @@ void ggml_vec_dot_q8_0_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const voi
 void ggml_vec_dot_mxfp4_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);
 void ggml_vec_dot_nvfp4_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_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);
+
 void ggml_vec_dot_q2_K_q8_K(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);
 void ggml_vec_dot_q3_K_q8_K(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);
 void ggml_vec_dot_q4_K_q8_K(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/src/ggml-cuda/common.cuh

@@ -21,7 +21,6 @@
 #endif
 #endif
 #include "ggml-common.h"
-#include "../../rocmfp4/rocmfp4.h"
 
 #include <array>
 #include <algorithm>
@@ -867,6 +866,36 @@ static __device__ __forceinline__ float ggml_cuda_ue4m3_to_fp32(uint8_t x) {
 #endif // defined(GGML_USE_HIP) && defined(CDNA3) && defined(FP8_AVAILABLE) && HIP_VERSION >= 60200000
 }
 
+// ============================================================================
+// ROCmFP4 (AMD gfx1151) device scale + FP4 codebook decode helpers.
+// GGML_ROCMFP4_USE_SCALE_LUT is an opt-in AMD-specific escape hatch (constant-memory table) kept for profiling.
+//  ============================================================================
+
+#ifndef GGML_ROCMFP4_USE_SCALE_LUT
+#define GGML_ROCMFP4_USE_SCALE_LUT 0
+#endif
+
+#if defined(GGML_USE_HIP) && GGML_ROCMFP4_USE_SCALE_LUT
+// Values come from the shared GGML_ROCMFP4_SCALE_UE4M3_HALF_LIST in ggml-common.h
+// (single source of truth, also used by the CPU quantizer table in ggml-quants.c).
+static __device__ __constant__ const float rocmfp4_scale_ue4m3_half_lut[127] = { GGML_ROCMFP4_SCALE_UE4M3_HALF_LIST };
+#endif
+
+static __device__ __forceinline__ float rocmfp4_ue4m3_to_fp32_half_finite(uint8_t x) {
+#if defined(GGML_USE_HIP) && GGML_ROCMFP4_USE_SCALE_LUT
+    return x <= 0x7e ? rocmfp4_scale_ue4m3_half_lut[x] : 0.0f;   // opt-in fast table
+#else
+    return ggml_cuda_ue4m3_to_fp32(x);                          // default: shared decoder
+#endif
+}
+
+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;
+}
+
 static __device__ __forceinline__ uint8_t ggml_cuda_fp32_to_ue4m3(float x) {
 #if defined(BLACKWELL_MMA_AVAILABLE) // This is used for NVFP4 subblock scale quantizations only
     if (!(x > 0.0f)) {

ggml/src/ggml-cuda/convert.cu

@@ -1,6 +1,5 @@
 #include "convert.cuh"
 #include "dequantize.cuh"
-#include "../../rocmfp4/rocmfp4_hip_scale.cuh"
 
 #include <cstdint>
 

ggml/src/ggml-cuda/dequantize.cuh

@@ -1,6 +1,5 @@
 #include "common.cuh"
 #include "convert.cuh"
-#include "../../rocmfp4/rocmfp4_hip_scale.cuh"
 
 static __device__ __forceinline__ void dequantize_q1_0(const void * vx, const int64_t ib, const int iqs, float2 & v){
     const block_q1_0 * x = (const block_q1_0 *) vx;