From 2fd4c2a770c801dab722ceb37151c9e6ab173ff6 Mon Sep 17 00:00:00 2001 From: Baraxanx Date: Thu, 10 Sep 2026 02:12:41 -0400 Subject: [PATCH 1/5] cuda: ROCmFPx load/dequant/dispatch paths for HIP builds Loads Q4_0_ROCMFP4, Q4_0_ROCMFP4_FAST and the Q2/Q3/Q6/Q8_0_ROCMFPX GGUF tensor types on ggml-cuda/HIP (previously CPU+Vulkan only). Adds dequant, convert, get_rows support and MMVQ dispatch; restores the *_hip_* helper headers (scale LUT + codebooks) that the format import pruned. Part 1/2 of the FP4-on-HIP work; part 2 adds the MMQ tile path. --- ggml/rocmfp4/rocmfp4_hip.cu | 93 +++++++++++++++ ggml/rocmfp4/rocmfp4_hip_codebook.cuh | 87 ++++++++++++++ ggml/rocmfp4/rocmfp4_hip_scale.cuh | 165 ++++++++++++++++++++++++++ ggml/rocmfpx/rocmfpx_hip_codebook.cuh | 53 +++++++++ ggml/src/ggml-cuda/common.cuh | 18 +++ ggml/src/ggml-cuda/convert.cu | 38 ++++++ ggml/src/ggml-cuda/dequantize.cuh | 34 ++++++ ggml/src/ggml-cuda/getrows.cu | 8 ++ ggml/src/ggml-cuda/ggml-cuda.cu | 4 + ggml/src/ggml-cuda/mmvq.cu | 16 +++ ggml/src/ggml-cuda/vecdotq.cuh | 50 ++++++++ 11 files changed, 566 insertions(+) create mode 100644 ggml/rocmfp4/rocmfp4_hip.cu create mode 100644 ggml/rocmfp4/rocmfp4_hip_codebook.cuh create mode 100644 ggml/rocmfp4/rocmfp4_hip_scale.cuh create mode 100644 ggml/rocmfpx/rocmfpx_hip_codebook.cuh diff --git a/ggml/rocmfp4/rocmfp4_hip.cu b/ggml/rocmfp4/rocmfp4_hip.cu new file mode 100644 index 000000000000..3174743cc41d --- /dev/null +++ b/ggml/rocmfp4/rocmfp4_hip.cu @@ -0,0 +1,93 @@ +// ROCmFPx quant formats, hand-ported into this fork. +// +// Origin: https://github.com/charlie12345/ROCmFPX - creator of the ROCmFP4 format. +// Ported from: https://github.com/ciru-ai/ROCmFPX - a fork of the above. +// +// Both upstream projects are MIT licensed and based on llama.cpp; upstream authors +// retain their authorship and MIT license credit. See LICENSE. + +#include "rocmfp4.h" + +#include + +#include "rocmfp4_hip_scale.cuh" + +// Standalone ROCm/HIP dequant kernel for integration tests and future fused +// paths. One lane owns one packed byte and writes the matching low/high +// half-block values, so each byte is read once. +extern "C" __global__ void rocmfp4_dequantize_q4_0_f32_kernel( + const block_rocmfp4 * __restrict__ x, + float * __restrict__ y, + int64_t k) { + const int64_t packed_idx = (int64_t) blockIdx.x*blockDim.x + threadIdx.x; + const int64_t nblocks = (k + QK_ROCMFP4 - 1) / QK_ROCMFP4; + const int64_t packed_count = nblocks * (QK_ROCMFP4/2); + + if (packed_idx >= packed_count) { + return; + } + + const int64_t ib = packed_idx / (QK_ROCMFP4/2); + const int tid = packed_idx - ib*(QK_ROCMFP4/2); + const int64_t base = ib*QK_ROCMFP4; + const uint8_t packed = x[ib].qs[tid]; + 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]); + + if (base + tid < k) { + y[base + tid] = (float) rocmfp4_decode_i8(packed & 0x0f) * d0; + } + if (base + tid + QK_ROCMFP4/2 < k) { + y[base + tid + QK_ROCMFP4/2] = (float) rocmfp4_decode_i8(packed >> 4) * d1; + } +} + +extern "C" __global__ void rocmfp4_dequantize_q4_0_fast_f32_kernel( + const block_rocmfp4_fast * __restrict__ x, + float * __restrict__ y, + int64_t k) { + const int64_t packed_idx = (int64_t) blockIdx.x*blockDim.x + threadIdx.x; + const int64_t nblocks = (k + QK_ROCMFP4 - 1) / QK_ROCMFP4; + const int64_t packed_count = nblocks * (QK_ROCMFP4/2); + + if (packed_idx >= packed_count) { + return; + } + + const int64_t ib = packed_idx / (QK_ROCMFP4/2); + const int tid = packed_idx - ib*(QK_ROCMFP4/2); + const int64_t base = ib*QK_ROCMFP4; + const uint8_t packed = x[ib].qs[tid]; + const float d = rocmfp4_ue4m3_to_fp32_half_finite(x[ib].e); + + if (base + tid < k) { + y[base + tid] = (float) rocmfp4_decode_i8(packed & 0x0f) * d; + } + if (base + tid + QK_ROCMFP4/2 < k) { + y[base + tid + QK_ROCMFP4/2] = (float) rocmfp4_decode_i8(packed >> 4) * d; + } +} + +extern "C" void rocmfp4_hip_dequantize_q4_0_to_f32( + const void * src, + float * dst, + int64_t k, + hipStream_t stream) { + const int64_t nblocks = (k + QK_ROCMFP4 - 1) / QK_ROCMFP4; + const int64_t packed_count = nblocks * (QK_ROCMFP4/2); + const dim3 block(256); + const dim3 grid((unsigned int) ((packed_count + block.x - 1) / block.x)); + rocmfp4_dequantize_q4_0_f32_kernel<<>>((const block_rocmfp4 *) src, dst, k); +} + +extern "C" void rocmfp4_hip_dequantize_q4_0_fast_to_f32( + const void * src, + float * dst, + int64_t k, + hipStream_t stream) { + const int64_t nblocks = (k + QK_ROCMFP4 - 1) / QK_ROCMFP4; + const int64_t packed_count = nblocks * (QK_ROCMFP4/2); + const dim3 block(256); + const dim3 grid((unsigned int) ((packed_count + block.x - 1) / block.x)); + rocmfp4_dequantize_q4_0_fast_f32_kernel<<>>((const block_rocmfp4_fast *) src, dst, k); +} diff --git a/ggml/rocmfp4/rocmfp4_hip_codebook.cuh b/ggml/rocmfp4/rocmfp4_hip_codebook.cuh new file mode 100644 index 000000000000..7fa1357f8e8d --- /dev/null +++ b/ggml/rocmfp4/rocmfp4_hip_codebook.cuh @@ -0,0 +1,87 @@ +// ROCmFPx quant formats, hand-ported into this fork. +// +// Origin: https://github.com/charlie12345/ROCmFPX - creator of the ROCmFP4 format. +// Ported from: https://github.com/ciru-ai/ROCmFPX - a fork of the above. +// +// Both upstream projects are MIT licensed and based on llama.cpp; upstream authors +// retain their authorship and MIT license credit. See LICENSE. + +#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 + +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 get_int_from_table_16(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 get_int_from_table_16(q4, fallback_table).x; +#endif +} diff --git a/ggml/rocmfp4/rocmfp4_hip_scale.cuh b/ggml/rocmfp4/rocmfp4_hip_scale.cuh new file mode 100644 index 000000000000..a783be2e9493 --- /dev/null +++ b/ggml/rocmfp4/rocmfp4_hip_scale.cuh @@ -0,0 +1,165 @@ +// ROCmFPx quant formats, hand-ported into this fork. +// +// Origin: https://github.com/charlie12345/ROCmFPX - creator of the ROCmFP4 format. +// Ported from: https://github.com/ciru-ai/ROCmFPX - a fork of the above. +// +// Both upstream projects are MIT licensed and based on llama.cpp; upstream authors +// retain their authorship and MIT license credit. See LICENSE. + +#pragma once + +#include +#include +#include + +#ifndef GGML_ROCMFP4_USE_SCALE_LUT +#define GGML_ROCMFP4_USE_SCALE_LUT 0 +#endif + +#if defined(GGML_USE_HIP) && GGML_ROCMFP4_USE_SCALE_LUT +#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 __device__ __constant__ const float rocmfp4_scale_ue4m3_half_lut[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 +#endif + +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) { +#if defined(GGML_USE_HIP) && GGML_ROCMFP4_USE_SCALE_LUT + return x <= 0x7e ? rocmfp4_scale_ue4m3_half_lut[x] : 0.0f; +#else + 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); +#endif +} + +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/ggml/rocmfpx/rocmfpx_hip_codebook.cuh b/ggml/rocmfpx/rocmfpx_hip_codebook.cuh new file mode 100644 index 000000000000..b9e36528c2fc --- /dev/null +++ b/ggml/rocmfpx/rocmfpx_hip_codebook.cuh @@ -0,0 +1,53 @@ +// ROCmFPx quant formats, hand-ported into this fork. +// +// Origin: https://github.com/charlie12345/ROCmFPX - creator of the ROCmFP4 format. +// Ported from: https://github.com/ciru-ai/ROCmFPX - a fork of the above. +// +// Both upstream projects are MIT licensed and based on llama.cpp; upstream authors +// retain their authorship and MIT license credit. See LICENSE. + +#pragma once + +#include + +static __device__ __constant__ const int8_t rocmfpx_fp3_codebook[8] = { + 0, 1, 2, 4, 0, -1, -2, -4, +}; + +static __device__ __constant__ const int8_t rocmfpx_fp6_codebook[64] = { + 0, 1, 2, 3, 4, 5, 6, 7, + 8, 9, 10, 11, 12, 13, 14, 15, + 16, 17, 18, 19, 20, 21, 22, 23, + 24, 25, 26, 27, 28, 29, 30, 31, + 0, -1, -2, -3, -4, -5, -6, -7, + -8, -9, -10, -11, -12, -13, -14, -15, + -16, -17, -18, -19, -20, -21, -22, -23, + -24, -25, -26, -27, -28, -29, -30, -31, +}; + +static __device__ __forceinline__ int rocmfpx_pack4_fp3_codes(const uint32_t bits12) { + const uint32_t c0 = (bits12 >> 0) & 7u; + const uint32_t c1 = (bits12 >> 3) & 7u; + const uint32_t c2 = (bits12 >> 6) & 7u; + const uint32_t c3 = (bits12 >> 9) & 7u; + const char4 v = make_char4( + rocmfpx_fp3_codebook[c0], + rocmfpx_fp3_codebook[c1], + rocmfpx_fp3_codebook[c2], + rocmfpx_fp3_codebook[c3]); + return *((const int *) &v); +} + +static __device__ __forceinline__ int rocmfpx_pack4_fp6_codes(const uint32_t bits24) { + const uint32_t c0 = (bits24 >> 0) & 63u; + const uint32_t c1 = (bits24 >> 6) & 63u; + const uint32_t c2 = (bits24 >> 12) & 63u; + const uint32_t c3 = (bits24 >> 18) & 63u; + + const char4 v = make_char4( + rocmfpx_fp6_codebook[c0], + rocmfpx_fp6_codebook[c1], + rocmfpx_fp6_codebook[c2], + rocmfpx_fp6_codebook[c3]); + return *((const int *) &v); +} diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index 3ab92800552e..115ad6d14a94 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -21,6 +21,7 @@ #endif #endif #include "ggml-common.h" +#include "../../rocmfp4/rocmfp4.h" #include #include @@ -40,6 +41,9 @@ #include "vendors/cuda.h" #endif // defined(GGML_USE_HIP) +// device-side UE4M3 scale decoder for ROCmFP4 (needs the runtime header above) +#include "../../rocmfp4/rocmfp4_hip_scale.cuh" + #define STRINGIZE_IMPL(...) #__VA_ARGS__ #define STRINGIZE(...) STRINGIZE_IMPL(__VA_ARGS__) @@ -1086,6 +1090,20 @@ struct ggml_cuda_type_traits { static constexpr int bs = sizeof(block_nvfp4); }; +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_K; diff --git a/ggml/src/ggml-cuda/convert.cu b/ggml/src/ggml-cuda/convert.cu index 47cab5f980d4..6677146c0cfc 100644 --- a/ggml/src/ggml-cuda/convert.cu +++ b/ggml/src/ggml-cuda/convert.cu @@ -243,6 +243,20 @@ static __global__ void dequantize_block_mxfp4(const void * __restrict__ vx, dst_ dequantize_mxfp4(vx, i, yy + i*QK_K, threadIdx.x); } +template +static __global__ void dequantize_block_rocmfp4(const void * __restrict__ vx, dst_t * __restrict__ yy) { + const int64_t i = blockIdx.x; + + dequantize_rocmfp4(vx, i, yy + i*QK_K, threadIdx.x); +} + +template +static __global__ void dequantize_block_rocmfp4_fast(const void * __restrict__ vx, dst_t * __restrict__ yy) { + const int64_t i = blockIdx.x; + + dequantize_rocmfp4_fast(vx, i, yy + i*QK_K, threadIdx.x); +} + 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, @@ -374,6 +388,18 @@ 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_cuda(const void * vx, dst_t * y, const int64_t k, cudaStream_t stream) { + const int nb = (k + QK_K - 1) / QK_K; + dequantize_block_rocmfp4<<>>(vx, y); +} + +template +static void dequantize_row_rocmfp4_fast_cuda(const void * vx, dst_t * y, const int64_t k, cudaStream_t stream) { + const int nb = (k + QK_K - 1) / QK_K; + dequantize_block_rocmfp4_fast<<>>(vx, y); +} + template static __global__ void dequantize_block_nvfp4( const void * __restrict__ vx, @@ -533,6 +559,10 @@ to_bf16_cuda_t ggml_get_to_bf16_cuda(ggml_type type) { return dequantize_row_iq3_s_cuda; case GGML_TYPE_MXFP4: return dequantize_row_mxfp4_cuda; + case GGML_TYPE_Q4_0_ROCMFP4: + return dequantize_row_rocmfp4_cuda; + case GGML_TYPE_Q4_0_ROCMFP4_FAST: + return dequantize_row_rocmfp4_fast_cuda; case GGML_TYPE_NVFP4: return dequantize_row_nvfp4_cuda; case GGML_TYPE_F32: @@ -593,6 +623,10 @@ to_fp16_cuda_t ggml_get_to_fp16_cuda(ggml_type type) { return dequantize_row_iq3_s_cuda; case GGML_TYPE_MXFP4: return dequantize_row_mxfp4_cuda; + case GGML_TYPE_Q4_0_ROCMFP4: + return dequantize_row_rocmfp4_cuda; + case GGML_TYPE_Q4_0_ROCMFP4_FAST: + return dequantize_row_rocmfp4_fast_cuda; case GGML_TYPE_NVFP4: return dequantize_row_nvfp4_cuda; case GGML_TYPE_F32: @@ -650,6 +684,10 @@ to_fp32_cuda_t ggml_get_to_fp32_cuda(ggml_type type) { return dequantize_row_iq3_s_cuda; case GGML_TYPE_MXFP4: return dequantize_row_mxfp4_cuda; + case GGML_TYPE_Q4_0_ROCMFP4: + return dequantize_row_rocmfp4_cuda; + case GGML_TYPE_Q4_0_ROCMFP4_FAST: + return dequantize_row_rocmfp4_fast_cuda; case GGML_TYPE_NVFP4: return dequantize_row_nvfp4_cuda; case GGML_TYPE_F16: diff --git a/ggml/src/ggml-cuda/dequantize.cuh b/ggml/src/ggml-cuda/dequantize.cuh index fd9b6d6e01bf..6489e65442f4 100644 --- a/ggml/src/ggml-cuda/dequantize.cuh +++ b/ggml/src/ggml-cuda/dequantize.cuh @@ -525,3 +525,37 @@ template static __device__ __forceinline__ void dequantize_mxfp4(const void * vx, const int64_t ib, dst_t * yy, const int tid) { dequantize_mxfp4(vx, ib, yy, tid); } + +template +static __device__ __forceinline__ void dequantize_rocmfp4(const void * vx, const int64_t ibs, dst_t * yy, const int tid) { + + const block_rocmfp4 * x = (const block_rocmfp4 *) vx + ibs*(QK_K/QK_ROCMFP4); + + const int64_t il = tid/8; // 0...3 + const int64_t ib = tid%8; // 0...7 + dst_t * y = yy + 32*ib + 4*il; + const uint8_t * q4 = x[ib].qs + 4*il; + // dual UE4M3 scales: low nibbles are weights j (e[0]), high nibbles are weights j+16 (e[1]) + 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] = ggml_cuda_cast(d0 * kvalues_rocmfp4[q4[j] & 0xf]); + y[j+16] = ggml_cuda_cast(d1 * kvalues_rocmfp4[q4[j] >> 4]); + } +} + +template +static __device__ __forceinline__ void dequantize_rocmfp4_fast(const void * vx, const int64_t ibs, dst_t * yy, const int tid) { + + const block_rocmfp4_fast * x = (const block_rocmfp4_fast *) vx + ibs*(QK_K/QK_ROCMFP4); + + const int64_t il = tid/8; // 0...3 + const int64_t ib = tid%8; // 0...7 + dst_t * y = yy + 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] = ggml_cuda_cast(d * kvalues_rocmfp4[q4[j] & 0xf]); + y[j+16] = ggml_cuda_cast(d * kvalues_rocmfp4[q4[j] >> 4]); + } +} diff --git a/ggml/src/ggml-cuda/getrows.cu b/ggml/src/ggml-cuda/getrows.cu index ffcb51f00912..d80ed22a172a 100644 --- a/ggml/src/ggml-cuda/getrows.cu +++ b/ggml/src/ggml-cuda/getrows.cu @@ -470,6 +470,14 @@ static void ggml_cuda_get_rows_switch_src0_type( get_rows_cuda_kq<32, dst_t, dequantize_mxfp4>(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_kq<32, dst_t, dequantize_rocmfp4>(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_kq<32, dst_t, dequantize_rocmfp4_fast>(src0_d, src1_d, dst_d, + ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream); + break; default: GGML_ABORT("%s: unsupported src0 type: %s\n", __func__, ggml_type_name(src0_type)); break; diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index c465879d7e6b..a93b5905cfec 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -7152,6 +7152,8 @@ 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_K: case GGML_TYPE_Q3_K: case GGML_TYPE_Q4_K: @@ -7205,6 +7207,8 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g return true; case GGML_TYPE_IQ4_NL: case GGML_TYPE_MXFP4: + case GGML_TYPE_Q4_0_ROCMFP4: + case GGML_TYPE_Q4_0_ROCMFP4_FAST: // 32-value sub-blocks, the row size does not guarantee // the QK_K super-blocks the get_rows kernel iterates on return op->src[0]->ne[0] % QK_K == 0; diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index 101da1f5acee..56a41ddc4216 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -49,6 +49,8 @@ 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_q4_0_rocmfp4_q8_1; + case GGML_TYPE_Q4_0_ROCMFP4_FAST: return vec_dot_q4_0_rocmfp4_fast_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; @@ -78,6 +80,8 @@ 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_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; @@ -2843,6 +2847,18 @@ 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_NVFP4: 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, diff --git a/ggml/src/ggml-cuda/vecdotq.cuh b/ggml/src/ggml-cuda/vecdotq.cuh index b3e49e0072af..7999f7011bc9 100644 --- a/ggml/src/ggml-cuda/vecdotq.cuh +++ b/ggml/src/ggml-cuda/vecdotq.cuh @@ -328,6 +328,56 @@ static __device__ __forceinline__ float vec_dot_mxfp4_q8_1( return d * sumi; } +#define VDR_ROCMFP4_Q8_1_MMVQ 2 + +// Dual UE4M3 half-block scales: low nibbles of qs[j] are weights j (scale e[0]), +// high nibbles are weights j + 16 (scale e[1]). Two int-dot accumulators, matching +// the q8_1 slot pairing used by vec_dot_mxfp4_q8_1 above. +static __device__ __forceinline__ float vec_dot_q4_0_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 = get_int_b1(bq4->qs, iqs + l); + const int2 v = get_int_from_table_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 dq8 = __low2float(bq8_1->ds); + const float d0 = rocmfp4_ue4m3_to_fp32_half_finite(bq4->e[0]) * dq8; + const float d1 = rocmfp4_ue4m3_to_fp32_half_finite(bq4->e[1]) * dq8; + return d0 * sumi0 + d1 * sumi1; +} + +static __device__ __forceinline__ float vec_dot_q4_0_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_Q8_1_MMVQ; ++l) { + const int aux_q4 = get_int_b1(bq4->qs, iqs + l); + const int2 v = get_int_from_table_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); + } + + const float d = rocmfp4_ue4m3_to_fp32_half_finite(bq4->e) * __low2float(bq8_1->ds); + return d * sumi; +} + #define VDR_NVFP4_Q8_1_MMVQ 4 #define VDR_NVFP4_Q8_1_MMQ 8 From b0b72220f7a17d63462b02148b428bbef1350a53 Mon Sep 17 00:00:00 2001 From: Baraxanx Date: Mon, 7 Sep 2026 21:00:59 -0400 Subject: [PATCH 2/5] cuda: MMQ tile path for Q4_0_ROCMFP4_FAST (dp4a + WMMA/MFMA layouts) Step 2 of the Strix Halo HIP plan: give the production FP4 format an MMQ path so prefill batches (>8, incl. spec-verify at wide ubatch and MoE routing) stop falling through to dequant+hipBLAS purely because the type was absent from the MMQ switches. Mirrors MXFP4 exactly - same block size (32), one UE4M3 scale per 32 values, SRAM_LAYOUT_Q8_1 tiles, q8_0_q8_1 vec_dot kernels (dp4a and MMA data layout), D4 y-side quantizer. Differences from MXFP4: kvalues_rocmfp4 codebook (Codebook10, max level 10 not 12) and the UE4M3 scale decoder without the e8m0 *0.5 factor. Tile rows mirror this master's current MXFP4 config set, incl. the 128/64/128 row. Not added to Blackwell configs on purpose: use_native_fp4 is false for this type there, so it falls through to the Ampere config like NVFP4-generic does. Dual-scale Q4_0_ROCMFP4 (2 scales/block) deferred to a follow-up slice (NVFP4-style per-16 convention). Validated: test-backend-ops vs CPU ref on RTX 4090 CUDA (default build -p rocmfp4: 288 OK / 0 FAIL; GGML_CUDA_FORCE_MMQ=ON -o MUL_MAT -p rocmfp4: 42 OK / 0 FAIL) and on gfx1151 HIP (-o MUL_MAT -b ROCm0 -p rocmfp4: 78 OK / 0 FAIL). Measured on Ryzen AI Max+ 395 against a rebuilt merge-base in the same session: pp512 +6.1%, pp2048 +5.4% (palindrome x2, zero overlap); MUL_MAT m=4096,n=512 kernel 5389 -> 3503 us (x1.54). --- ggml/src/ggml-cuda/mmq-config-ampere.cuh | 16 +++++ ggml/src/ggml-cuda/mmq-config-cdna.cuh | 7 ++ ggml/src/ggml-cuda/mmq-config-pascal-dp4a.cuh | 11 +++ .../src/ggml-cuda/mmq-config-pascal-older.cuh | 11 +++ ggml/src/ggml-cuda/mmq-config-rdna2.cuh | 11 +++ ggml/src/ggml-cuda/mmq-config-rdna3-5.cuh | 12 ++++ ggml/src/ggml-cuda/mmq-config-rdna3.cuh | 11 +++ ggml/src/ggml-cuda/mmq-config-rdna4.cuh | 12 ++++ ggml/src/ggml-cuda/mmq-load-tiles.cuh | 70 +++++++++++++++++++ ggml/src/ggml-cuda/mmq.cu | 4 ++ ggml/src/ggml-cuda/mmq.cuh | 16 +++++ .../template-instances/generate_cu_files.py | 2 +- .../mmq-instance-q4_0_rocmfp4_fast.cu | 5 ++ ggml/src/ggml-cuda/vecdotq.cuh | 1 + 14 files changed, 188 insertions(+), 1 deletion(-) create mode 100644 ggml/src/ggml-cuda/template-instances/mmq-instance-q4_0_rocmfp4_fast.cu diff --git a/ggml/src/ggml-cuda/mmq-config-ampere.cuh b/ggml/src/ggml-cuda/mmq-config-ampere.cuh index 9f9fd197382f..fa62f34d6876 100644 --- a/ggml/src/ggml-cuda/mmq-config-ampere.cuh +++ b/ggml/src/ggml-cuda/mmq-config-ampere.cuh @@ -346,21 +346,37 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf // --------------------------------------------------------------------------------------------- CASE(GGML_TYPE_MXFP4, 256, 1, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, true); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 1, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, true); CASE(GGML_TYPE_MXFP4, 256, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, true); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, true); CASE(GGML_TYPE_MXFP4, 256, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, true); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, true); CASE(GGML_TYPE_MXFP4, 256, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, true); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, true); CASE(GGML_TYPE_MXFP4, 256, 1, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, true); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 1, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, true); CASE(GGML_TYPE_MXFP4, 256, 1, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 1, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); CASE(GGML_TYPE_MXFP4, 256, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); CASE(GGML_TYPE_MXFP4, 256, 1, 128, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 1, 128, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); CASE(GGML_TYPE_MXFP4, 256, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); CASE(GGML_TYPE_MXFP4, 256, 1, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 1, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); CASE(GGML_TYPE_MXFP4, 256, 1, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 1, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); CASE(GGML_TYPE_MXFP4, 256, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); CASE(GGML_TYPE_MXFP4, 256, 1, 128, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 1, 128, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); CASE(GGML_TYPE_MXFP4, 256, 1, 128, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 1, 128, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); CASE(GGML_TYPE_MXFP4, 256, 1, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 1, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); CASE(GGML_TYPE_MXFP4, 256, 1, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 1, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); CASE(GGML_TYPE_NVFP4, 256, 1, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, true, true); CASE(GGML_TYPE_NVFP4, 256, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, true, true); diff --git a/ggml/src/ggml-cuda/mmq-config-cdna.cuh b/ggml/src/ggml-cuda/mmq-config-cdna.cuh index 4a8d89f72019..18d9ac6dc698 100644 --- a/ggml/src/ggml-cuda/mmq-config-cdna.cuh +++ b/ggml/src/ggml-cuda/mmq-config-cdna.cuh @@ -166,12 +166,19 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf // --------------------------------------------------------------------------------------------- CASE(GGML_TYPE_MXFP4, 512, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, true); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 512, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, true); CASE(GGML_TYPE_MXFP4, 512, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, true); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 512, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, true); CASE(GGML_TYPE_MXFP4, 512, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, true); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 512, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, true); CASE(GGML_TYPE_MXFP4, 512, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 512, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); CASE(GGML_TYPE_MXFP4, 512, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 512, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); CASE(GGML_TYPE_MXFP4, 512, 1, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 512, 1, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); CASE(GGML_TYPE_MXFP4, 512, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 512, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); CASE(GGML_TYPE_NVFP4, 512, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, true, true); CASE(GGML_TYPE_NVFP4, 512, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, true, true); diff --git a/ggml/src/ggml-cuda/mmq-config-pascal-dp4a.cuh b/ggml/src/ggml-cuda/mmq-config-pascal-dp4a.cuh index 83eb7c146e11..87b76b6c8fc0 100644 --- a/ggml/src/ggml-cuda/mmq-config-pascal-dp4a.cuh +++ b/ggml/src/ggml-cuda/mmq-config-pascal-dp4a.cuh @@ -246,16 +246,27 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf // --------------------------------------------------------------------------------------------- CASE(GGML_TYPE_MXFP4, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_NVFP4, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, true); CASE(GGML_TYPE_NVFP4, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, true); diff --git a/ggml/src/ggml-cuda/mmq-config-pascal-older.cuh b/ggml/src/ggml-cuda/mmq-config-pascal-older.cuh index 2a8dc9e1a93e..31078d1462b5 100644 --- a/ggml/src/ggml-cuda/mmq-config-pascal-older.cuh +++ b/ggml/src/ggml-cuda/mmq-config-pascal-older.cuh @@ -246,16 +246,27 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf // --------------------------------------------------------------------------------------------- CASE(GGML_TYPE_MXFP4, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_NVFP4, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, true); CASE(GGML_TYPE_NVFP4, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, true); diff --git a/ggml/src/ggml-cuda/mmq-config-rdna2.cuh b/ggml/src/ggml-cuda/mmq-config-rdna2.cuh index 8324d9e1a830..6cd7ccc17c72 100644 --- a/ggml/src/ggml-cuda/mmq-config-rdna2.cuh +++ b/ggml/src/ggml-cuda/mmq-config-rdna2.cuh @@ -246,16 +246,27 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf // --------------------------------------------------------------------------------------------- CASE(GGML_TYPE_MXFP4, 256, 2, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_NVFP4, 256, 2, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, true); CASE(GGML_TYPE_NVFP4, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, true); diff --git a/ggml/src/ggml-cuda/mmq-config-rdna3-5.cuh b/ggml/src/ggml-cuda/mmq-config-rdna3-5.cuh index f8a9e6350926..aaadfaa1f3bf 100644 --- a/ggml/src/ggml-cuda/mmq-config-rdna3-5.cuh +++ b/ggml/src/ggml-cuda/mmq-config-rdna3-5.cuh @@ -261,17 +261,29 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf // --------------------------------------------------------------------------------------------- CASE(GGML_TYPE_MXFP4, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 128, 2, 64, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 128, 2, 64, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_NVFP4, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, true); CASE(GGML_TYPE_NVFP4, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, true); diff --git a/ggml/src/ggml-cuda/mmq-config-rdna3.cuh b/ggml/src/ggml-cuda/mmq-config-rdna3.cuh index 3a3ef7bd9c09..8af3423d7d68 100644 --- a/ggml/src/ggml-cuda/mmq-config-rdna3.cuh +++ b/ggml/src/ggml-cuda/mmq-config-rdna3.cuh @@ -247,16 +247,27 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf // --------------------------------------------------------------------------------------------- CASE(GGML_TYPE_MXFP4, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 128, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 128, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 128, 1, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 128, 1, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 128, 1, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 128, 1, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 128, 1, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 128, 1, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 128, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 128, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 128, 2, 64, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 128, 2, 64, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 128, 2, 64, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 128, 2, 64, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_NVFP4, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, true); CASE(GGML_TYPE_NVFP4, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, true); diff --git a/ggml/src/ggml-cuda/mmq-config-rdna4.cuh b/ggml/src/ggml-cuda/mmq-config-rdna4.cuh index 9293d9d55885..8ef32188f690 100644 --- a/ggml/src/ggml-cuda/mmq-config-rdna4.cuh +++ b/ggml/src/ggml-cuda/mmq-config-rdna4.cuh @@ -261,17 +261,29 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf // --------------------------------------------------------------------------------------------- CASE(GGML_TYPE_MXFP4, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 128, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 128, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 128, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 128, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 128, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 128, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); + CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_NVFP4, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, true); CASE(GGML_TYPE_NVFP4, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, true); diff --git a/ggml/src/ggml-cuda/mmq-load-tiles.cuh b/ggml/src/ggml-cuda/mmq-load-tiles.cuh index 40c825405e26..7a37603fa22d 100644 --- a/ggml/src/ggml-cuda/mmq-load-tiles.cuh +++ b/ggml/src/ggml-cuda/mmq-load-tiles.cuh @@ -1862,3 +1862,73 @@ template static __device__ __forceinline_ x_u32_scale[i*sram_stride] = get_int_b4(bxi->d, 0); } } + +// Q4_0_ROCMFP4_FAST: 32 values / block, one UE4M3 scale per 32 values - same tile structure as MXFP4. +// Differences from the mxfp4 loader above: block_rocmfp4_fast layout, kvalues_rocmfp4 codebook +// (Codebook10), and the UE4M3 scale decoder without the e8m0 0.5 factor. +template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q4_0_rocmfp4_fast( + const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + constexpr int warp_size = ggml_cuda_get_physical_warp_size(); + constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; + constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); + constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); + +#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) + int * x_qs = (int *) x_tile; + float * x_df = (float *) (x_qs + MMQ_TILE_NE_K*2); +#else + constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_Q4_0_ROCMFP4_FAST, I); + int * x_qs = (int *) x_tile; + float * x_df = (float *) (x_qs + txs.qs); +#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) + + constexpr int threads_per_row = MMQ_ITER_K / (4 * QR_ROCMFP4); + constexpr int nrows = warp_size / threads_per_row; + const int txi = warp_size > threads_per_row ? threadIdx.x % threads_per_row : threadIdx.x; + const int kbx = txi / QI_ROCMFP4; + const int kqsx = txi % QI_ROCMFP4; + +#pragma unroll + for (int i0 = 0; i0 < I; i0 += nrows*nwarps) { + int i = i0 + (nrows == 1 ? threadIdx.y : threadIdx.y*nrows + threadIdx.x/threads_per_row); + + if (fallback) { + i = min(i, i_max); + } + + const block_rocmfp4_fast * bxi = (const block_rocmfp4_fast *) x + kbx0 + i*stride + kbx; + + const int aux_q4 = get_int_b1(bxi->qs, kqsx); + const int2 v = get_int_from_table_16(aux_q4, kvalues_rocmfp4); + const int k0 = kbx * (2 * QI_ROCMFP4) + kqsx; + +#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) + x_qs[i*sram_stride + k0 + 0] = v.x; + x_qs[i*sram_stride + k0 + QI_ROCMFP4] = v.y; +#else + x_qs[i*(2*MMQ_TILE_NE_K + 1) + k0 + 0] = v.x; + x_qs[i*(2*MMQ_TILE_NE_K + 1) + k0 + QI_ROCMFP4] = v.y; +#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) + } + + constexpr int blocks_per_tile_x_row = MMQ_TILE_NE_K / QI_ROCMFP4; + constexpr int rows_per_warp = warp_size / blocks_per_tile_x_row; + const int kbxd = threadIdx.x % blocks_per_tile_x_row; + +#pragma unroll + for (int i0 = 0; i0 < I; i0 += nwarps * rows_per_warp) { + int i = i0 + threadIdx.y * rows_per_warp + threadIdx.x / blocks_per_tile_x_row; + + if (fallback) { + i = min(i, i_max); + } + + const block_rocmfp4_fast * bxi = (const block_rocmfp4_fast *) x + kbx0 + i*stride + kbxd; + +#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) + x_df[i*sram_stride + kbxd] = rocmfp4_ue4m3_to_fp32_half_finite(bxi->e); +#else + x_df[i*(MMQ_TILE_NE_K/QI_ROCMFP4) + i/QI_ROCMFP4 + kbxd] = rocmfp4_ue4m3_to_fp32_half_finite(bxi->e); +#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) + } +} diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu index 0e5479dc8eea..31e5d3161e6c 100644 --- a/ggml/src/ggml-cuda/mmq.cu +++ b/ggml/src/ggml-cuda/mmq.cu @@ -104,6 +104,9 @@ static void ggml_cuda_mul_mat_q_switch_type(ggml_backend_cuda_context & ctx, con case GGML_TYPE_NVFP4: mul_mat_q_case(ctx, args, stream); break; + case GGML_TYPE_Q4_0_ROCMFP4_FAST: + mul_mat_q_case(ctx, args, stream); + break; default: GGML_ABORT("fatal error"); break; @@ -494,6 +497,7 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t // ------------------------------------------------- case GGML_TYPE_MXFP4: case GGML_TYPE_NVFP4: + case GGML_TYPE_Q4_0_ROCMFP4_FAST: mmq_supported = true; break; default: diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 333dac85dcfb..26df051abfc5 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -75,6 +75,8 @@ static mmq_q8_1_ds_layout mmq_get_q8_1_ds_layout(const ggml_type type_x) { return MMQ_Q8_1_DS_LAYOUT_D4; case GGML_TYPE_NVFP4: return MMQ_Q8_1_DS_LAYOUT_D4; + case GGML_TYPE_Q4_0_ROCMFP4_FAST: + return MMQ_Q8_1_DS_LAYOUT_D4; case GGML_TYPE_Q2_K: return MMQ_Q8_1_DS_LAYOUT_D2S6; case GGML_TYPE_Q3_K: @@ -406,6 +408,7 @@ static constexpr __host__ __device__ tile_x_sizes mmq_get_dp4a_tile_x_sizes(ggml case GGML_TYPE_Q8_0: return MMQ_DP4A_TXS_Q8_0; case GGML_TYPE_MXFP4: return MMQ_DP4A_TXS_Q8_1; case GGML_TYPE_NVFP4: return MMQ_DP4A_TXS_Q8_0_16; + case GGML_TYPE_Q4_0_ROCMFP4_FAST: return MMQ_DP4A_TXS_Q8_1; case GGML_TYPE_Q2_K: return MMQ_DP4A_TXS_Q2_K; case GGML_TYPE_Q3_K: return MMQ_DP4A_TXS_Q3_K; case GGML_TYPE_Q4_K: return MMQ_DP4A_TXS_Q4_K; @@ -686,6 +689,12 @@ static constexpr __device__ ggml_cuda_mmq_util_funcs ggml_cuda_mmq_get_util_func ggml_cuda_mmq_load_tiles_nvfp4, ggml_cuda_mmq_vec_dot_q8_0_16_q8_1_dp4a, ggml_cuda_mmq_write_back_dp4a); + case GGML_TYPE_Q4_0_ROCMFP4_FAST: + return ggml_cuda_mmq_util_funcs( + VDR_ROCMFP4_FAST_Q8_1_MMQ, + ggml_cuda_mmq_load_tiles_q4_0_rocmfp4_fast, + ggml_cuda_mmq_vec_dot_q8_0_q8_1_dp4a, + ggml_cuda_mmq_write_back_dp4a); default: return ggml_cuda_mmq_util_funcs(1, nullptr, nullptr, nullptr); } @@ -850,6 +859,12 @@ static constexpr __device__ ggml_cuda_mmq_util_funcs ggml_cuda_mmq_get_util_func ggml_cuda_mmq_load_tiles_nvfp4, ggml_cuda_mmq_vec_dot_q8_0_16_q8_1_mma, ggml_cuda_mmq_write_back_mma); + case GGML_TYPE_Q4_0_ROCMFP4_FAST: + return ggml_cuda_mmq_util_funcs( + -1, + ggml_cuda_mmq_load_tiles_q4_0_rocmfp4_fast, + ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma, + ggml_cuda_mmq_write_back_mma); default: return ggml_cuda_mmq_util_funcs(1, nullptr, nullptr, nullptr); } @@ -1930,6 +1945,7 @@ extern DECL_MMQ_CASE(GGML_TYPE_IQ4_XS); // ----------------------------------------- extern DECL_MMQ_CASE(GGML_TYPE_MXFP4); extern DECL_MMQ_CASE(GGML_TYPE_NVFP4); +extern DECL_MMQ_CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST); // ------------------------------------------------------------------------------------------------------------------------- diff --git a/ggml/src/ggml-cuda/template-instances/generate_cu_files.py b/ggml/src/ggml-cuda/template-instances/generate_cu_files.py index d7cd271675e0..d12cd8e98da4 100755 --- a/ggml/src/ggml-cuda/template-instances/generate_cu_files.py +++ b/ggml/src/ggml-cuda/template-instances/generate_cu_files.py @@ -40,7 +40,7 @@ "GGML_TYPE_Q4_0", "GGML_TYPE_Q4_1", "GGML_TYPE_Q5_0", "GGML_TYPE_Q5_1", "GGML_TYPE_Q8_0", "GGML_TYPE_Q2_K", "GGML_TYPE_Q3_K", "GGML_TYPE_Q4_K", "GGML_TYPE_Q5_K", "GGML_TYPE_Q6_K", "GGML_TYPE_IQ2_XXS", "GGML_TYPE_IQ2_XS", "GGML_TYPE_IQ2_S", "GGML_TYPE_IQ3_XXS", "GGML_TYPE_IQ3_S", - "GGML_TYPE_IQ1_S", "GGML_TYPE_IQ4_NL", "GGML_TYPE_IQ4_XS", "GGML_TYPE_MXFP4", "GGML_TYPE_NVFP4" + "GGML_TYPE_IQ1_S", "GGML_TYPE_IQ4_NL", "GGML_TYPE_IQ4_XS", "GGML_TYPE_MXFP4", "GGML_TYPE_NVFP4", "GGML_TYPE_Q4_0_ROCMFP4_FAST" ] SOURCE_MMQ = """// This file has been autogenerated by generate_cu_files.py, do not edit manually. diff --git a/ggml/src/ggml-cuda/template-instances/mmq-instance-q4_0_rocmfp4_fast.cu b/ggml/src/ggml-cuda/template-instances/mmq-instance-q4_0_rocmfp4_fast.cu new file mode 100644 index 000000000000..df63d09ae0c6 --- /dev/null +++ b/ggml/src/ggml-cuda/template-instances/mmq-instance-q4_0_rocmfp4_fast.cu @@ -0,0 +1,5 @@ +// This file has been autogenerated by generate_cu_files.py, do not edit manually. + +#include "../mmq.cuh" + +DECL_MMQ_CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST); diff --git a/ggml/src/ggml-cuda/vecdotq.cuh b/ggml/src/ggml-cuda/vecdotq.cuh index 7999f7011bc9..37a04cfd165f 100644 --- a/ggml/src/ggml-cuda/vecdotq.cuh +++ b/ggml/src/ggml-cuda/vecdotq.cuh @@ -306,6 +306,7 @@ template static __device__ __forceinline__ float vec_dot_q8_0_16_q8_1_ #define VDR_MXFP4_Q8_1_MMVQ 2 #define VDR_MXFP4_Q8_1_MMQ 4 +#define VDR_ROCMFP4_FAST_Q8_1_MMQ 4 static __device__ __forceinline__ float vec_dot_mxfp4_q8_1( const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) { From 3db7cc3b639ed960705981c097379444663d9c1b Mon Sep 17 00:00:00 2001 From: Baraxanx Date: Mon, 28 Sep 2026 11:54:47 -0400 Subject: [PATCH 3/5] cuda: drop unreferenced rocmfp4 HIP helper files Review finding (halo-box/strix-llama.cpp#40): three files added in the first commit are compiled by nothing and referenced nowhere: - ggml/rocmfp4/rocmfp4_hip.cu - two extern "C" dequant kernels; not in any CMake source list (ggml-hip globs only ../ggml-cuda/*.cu and the template instances) and no callers, so unreachable even from tests. - ggml/rocmfp4/rocmfp4_hip_codebook.cuh - included by nothing; live MMVQ/MMQ paths use get_int_from_table_16 + kvalues_rocmfp4 from ggml-common.h, which already lowers to permb on HIP. The single-nibble perm variant was for an FA K/V decode path not present in this PR. - ggml/rocmfpx/rocmfpx_hip_codebook.cuh - same, included by nothing. The live UE4M3 scale decoder is rocmfp4_hip_scale.cuh (included via common.cuh) and stays. No behavior change: verified zero references to these files/symbols in the rest of the tree. --- ggml/rocmfp4/rocmfp4_hip.cu | 93 --------------------------- ggml/rocmfp4/rocmfp4_hip_codebook.cuh | 87 ------------------------- ggml/rocmfpx/rocmfpx_hip_codebook.cuh | 53 --------------- 3 files changed, 233 deletions(-) delete mode 100644 ggml/rocmfp4/rocmfp4_hip.cu delete mode 100644 ggml/rocmfp4/rocmfp4_hip_codebook.cuh delete mode 100644 ggml/rocmfpx/rocmfpx_hip_codebook.cuh diff --git a/ggml/rocmfp4/rocmfp4_hip.cu b/ggml/rocmfp4/rocmfp4_hip.cu deleted file mode 100644 index 3174743cc41d..000000000000 --- a/ggml/rocmfp4/rocmfp4_hip.cu +++ /dev/null @@ -1,93 +0,0 @@ -// ROCmFPx quant formats, hand-ported into this fork. -// -// Origin: https://github.com/charlie12345/ROCmFPX - creator of the ROCmFP4 format. -// Ported from: https://github.com/ciru-ai/ROCmFPX - a fork of the above. -// -// Both upstream projects are MIT licensed and based on llama.cpp; upstream authors -// retain their authorship and MIT license credit. See LICENSE. - -#include "rocmfp4.h" - -#include - -#include "rocmfp4_hip_scale.cuh" - -// Standalone ROCm/HIP dequant kernel for integration tests and future fused -// paths. One lane owns one packed byte and writes the matching low/high -// half-block values, so each byte is read once. -extern "C" __global__ void rocmfp4_dequantize_q4_0_f32_kernel( - const block_rocmfp4 * __restrict__ x, - float * __restrict__ y, - int64_t k) { - const int64_t packed_idx = (int64_t) blockIdx.x*blockDim.x + threadIdx.x; - const int64_t nblocks = (k + QK_ROCMFP4 - 1) / QK_ROCMFP4; - const int64_t packed_count = nblocks * (QK_ROCMFP4/2); - - if (packed_idx >= packed_count) { - return; - } - - const int64_t ib = packed_idx / (QK_ROCMFP4/2); - const int tid = packed_idx - ib*(QK_ROCMFP4/2); - const int64_t base = ib*QK_ROCMFP4; - const uint8_t packed = x[ib].qs[tid]; - 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]); - - if (base + tid < k) { - y[base + tid] = (float) rocmfp4_decode_i8(packed & 0x0f) * d0; - } - if (base + tid + QK_ROCMFP4/2 < k) { - y[base + tid + QK_ROCMFP4/2] = (float) rocmfp4_decode_i8(packed >> 4) * d1; - } -} - -extern "C" __global__ void rocmfp4_dequantize_q4_0_fast_f32_kernel( - const block_rocmfp4_fast * __restrict__ x, - float * __restrict__ y, - int64_t k) { - const int64_t packed_idx = (int64_t) blockIdx.x*blockDim.x + threadIdx.x; - const int64_t nblocks = (k + QK_ROCMFP4 - 1) / QK_ROCMFP4; - const int64_t packed_count = nblocks * (QK_ROCMFP4/2); - - if (packed_idx >= packed_count) { - return; - } - - const int64_t ib = packed_idx / (QK_ROCMFP4/2); - const int tid = packed_idx - ib*(QK_ROCMFP4/2); - const int64_t base = ib*QK_ROCMFP4; - const uint8_t packed = x[ib].qs[tid]; - const float d = rocmfp4_ue4m3_to_fp32_half_finite(x[ib].e); - - if (base + tid < k) { - y[base + tid] = (float) rocmfp4_decode_i8(packed & 0x0f) * d; - } - if (base + tid + QK_ROCMFP4/2 < k) { - y[base + tid + QK_ROCMFP4/2] = (float) rocmfp4_decode_i8(packed >> 4) * d; - } -} - -extern "C" void rocmfp4_hip_dequantize_q4_0_to_f32( - const void * src, - float * dst, - int64_t k, - hipStream_t stream) { - const int64_t nblocks = (k + QK_ROCMFP4 - 1) / QK_ROCMFP4; - const int64_t packed_count = nblocks * (QK_ROCMFP4/2); - const dim3 block(256); - const dim3 grid((unsigned int) ((packed_count + block.x - 1) / block.x)); - rocmfp4_dequantize_q4_0_f32_kernel<<>>((const block_rocmfp4 *) src, dst, k); -} - -extern "C" void rocmfp4_hip_dequantize_q4_0_fast_to_f32( - const void * src, - float * dst, - int64_t k, - hipStream_t stream) { - const int64_t nblocks = (k + QK_ROCMFP4 - 1) / QK_ROCMFP4; - const int64_t packed_count = nblocks * (QK_ROCMFP4/2); - const dim3 block(256); - const dim3 grid((unsigned int) ((packed_count + block.x - 1) / block.x)); - rocmfp4_dequantize_q4_0_fast_f32_kernel<<>>((const block_rocmfp4_fast *) src, dst, k); -} diff --git a/ggml/rocmfp4/rocmfp4_hip_codebook.cuh b/ggml/rocmfp4/rocmfp4_hip_codebook.cuh deleted file mode 100644 index 7fa1357f8e8d..000000000000 --- a/ggml/rocmfp4/rocmfp4_hip_codebook.cuh +++ /dev/null @@ -1,87 +0,0 @@ -// ROCmFPx quant formats, hand-ported into this fork. -// -// Origin: https://github.com/charlie12345/ROCmFPX - creator of the ROCmFP4 format. -// Ported from: https://github.com/ciru-ai/ROCmFPX - a fork of the above. -// -// Both upstream projects are MIT licensed and based on llama.cpp; upstream authors -// retain their authorship and MIT license credit. See LICENSE. - -#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 - -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 get_int_from_table_16(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 get_int_from_table_16(q4, fallback_table).x; -#endif -} diff --git a/ggml/rocmfpx/rocmfpx_hip_codebook.cuh b/ggml/rocmfpx/rocmfpx_hip_codebook.cuh deleted file mode 100644 index b9e36528c2fc..000000000000 --- a/ggml/rocmfpx/rocmfpx_hip_codebook.cuh +++ /dev/null @@ -1,53 +0,0 @@ -// ROCmFPx quant formats, hand-ported into this fork. -// -// Origin: https://github.com/charlie12345/ROCmFPX - creator of the ROCmFP4 format. -// Ported from: https://github.com/ciru-ai/ROCmFPX - a fork of the above. -// -// Both upstream projects are MIT licensed and based on llama.cpp; upstream authors -// retain their authorship and MIT license credit. See LICENSE. - -#pragma once - -#include - -static __device__ __constant__ const int8_t rocmfpx_fp3_codebook[8] = { - 0, 1, 2, 4, 0, -1, -2, -4, -}; - -static __device__ __constant__ const int8_t rocmfpx_fp6_codebook[64] = { - 0, 1, 2, 3, 4, 5, 6, 7, - 8, 9, 10, 11, 12, 13, 14, 15, - 16, 17, 18, 19, 20, 21, 22, 23, - 24, 25, 26, 27, 28, 29, 30, 31, - 0, -1, -2, -3, -4, -5, -6, -7, - -8, -9, -10, -11, -12, -13, -14, -15, - -16, -17, -18, -19, -20, -21, -22, -23, - -24, -25, -26, -27, -28, -29, -30, -31, -}; - -static __device__ __forceinline__ int rocmfpx_pack4_fp3_codes(const uint32_t bits12) { - const uint32_t c0 = (bits12 >> 0) & 7u; - const uint32_t c1 = (bits12 >> 3) & 7u; - const uint32_t c2 = (bits12 >> 6) & 7u; - const uint32_t c3 = (bits12 >> 9) & 7u; - const char4 v = make_char4( - rocmfpx_fp3_codebook[c0], - rocmfpx_fp3_codebook[c1], - rocmfpx_fp3_codebook[c2], - rocmfpx_fp3_codebook[c3]); - return *((const int *) &v); -} - -static __device__ __forceinline__ int rocmfpx_pack4_fp6_codes(const uint32_t bits24) { - const uint32_t c0 = (bits24 >> 0) & 63u; - const uint32_t c1 = (bits24 >> 6) & 63u; - const uint32_t c2 = (bits24 >> 12) & 63u; - const uint32_t c3 = (bits24 >> 18) & 63u; - - const char4 v = make_char4( - rocmfpx_fp6_codebook[c0], - rocmfpx_fp6_codebook[c1], - rocmfpx_fp6_codebook[c2], - rocmfpx_fp6_codebook[c3]); - return *((const int *) &v); -} From d3a8ab54a2f8d9f8c551a68f60cad6463d059bf1 Mon Sep 17 00:00:00 2001 From: Baraxanx Date: Mon, 28 Sep 2026 13:01:18 -0400 Subject: [PATCH 4/5] cuda: gate Q4_0_ROCMFP4_FAST MMQ to validated architectures Review scope reduction (halo-box/strix-llama.cpp#40 review): keep the MMQ tile rows only where the path is actually validated - RDNA3.5 (gfx1151, test-backend-ops 78/0) and the Ampere config table (RTX 4090 sm_89). Drop the pascal-dp4a / pascal-older / rdna2 / rdna3 / cdna / rdna4 rows: no hardware to measure there, so those architectures keep today's dequant + BLAS fallback (zero behavior change) instead of an unvalidated tile policy. should_use_mmq now returns false for this type outside the gated set (including Blackwell, which has no rows - native FP4 MMQ is MXFP4/NVFP4-only there), so a missing config row can never reach ggml_cuda_mmq_get_config's abort. --- ggml/src/ggml-cuda/mmq-config-cdna.cuh | 7 ------- ggml/src/ggml-cuda/mmq-config-pascal-dp4a.cuh | 11 ----------- ggml/src/ggml-cuda/mmq-config-pascal-older.cuh | 11 ----------- ggml/src/ggml-cuda/mmq-config-rdna2.cuh | 11 ----------- ggml/src/ggml-cuda/mmq-config-rdna3.cuh | 11 ----------- ggml/src/ggml-cuda/mmq-config-rdna4.cuh | 12 ------------ ggml/src/ggml-cuda/mmq.cu | 10 +++++++++- 7 files changed, 9 insertions(+), 64 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq-config-cdna.cuh b/ggml/src/ggml-cuda/mmq-config-cdna.cuh index 18d9ac6dc698..4a8d89f72019 100644 --- a/ggml/src/ggml-cuda/mmq-config-cdna.cuh +++ b/ggml/src/ggml-cuda/mmq-config-cdna.cuh @@ -166,19 +166,12 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf // --------------------------------------------------------------------------------------------- CASE(GGML_TYPE_MXFP4, 512, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, true); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 512, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, true); CASE(GGML_TYPE_MXFP4, 512, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, true); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 512, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, true); CASE(GGML_TYPE_MXFP4, 512, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, true); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 512, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, true); CASE(GGML_TYPE_MXFP4, 512, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 512, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); CASE(GGML_TYPE_MXFP4, 512, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 512, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); CASE(GGML_TYPE_MXFP4, 512, 1, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 512, 1, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); CASE(GGML_TYPE_MXFP4, 512, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 512, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, true, false); CASE(GGML_TYPE_NVFP4, 512, 1, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, true, true); CASE(GGML_TYPE_NVFP4, 512, 1, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, true, true); diff --git a/ggml/src/ggml-cuda/mmq-config-pascal-dp4a.cuh b/ggml/src/ggml-cuda/mmq-config-pascal-dp4a.cuh index 87b76b6c8fc0..83eb7c146e11 100644 --- a/ggml/src/ggml-cuda/mmq-config-pascal-dp4a.cuh +++ b/ggml/src/ggml-cuda/mmq-config-pascal-dp4a.cuh @@ -246,27 +246,16 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf // --------------------------------------------------------------------------------------------- CASE(GGML_TYPE_MXFP4, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_NVFP4, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, true); CASE(GGML_TYPE_NVFP4, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, true); diff --git a/ggml/src/ggml-cuda/mmq-config-pascal-older.cuh b/ggml/src/ggml-cuda/mmq-config-pascal-older.cuh index 31078d1462b5..2a8dc9e1a93e 100644 --- a/ggml/src/ggml-cuda/mmq-config-pascal-older.cuh +++ b/ggml/src/ggml-cuda/mmq-config-pascal-older.cuh @@ -246,27 +246,16 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf // --------------------------------------------------------------------------------------------- CASE(GGML_TYPE_MXFP4, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_NVFP4, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, true); CASE(GGML_TYPE_NVFP4, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, true); diff --git a/ggml/src/ggml-cuda/mmq-config-rdna2.cuh b/ggml/src/ggml-cuda/mmq-config-rdna2.cuh index 6cd7ccc17c72..8324d9e1a830 100644 --- a/ggml/src/ggml-cuda/mmq-config-rdna2.cuh +++ b/ggml/src/ggml-cuda/mmq-config-rdna2.cuh @@ -246,27 +246,16 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf // --------------------------------------------------------------------------------------------- CASE(GGML_TYPE_MXFP4, 256, 2, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_NVFP4, 256, 2, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, true); CASE(GGML_TYPE_NVFP4, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, true); diff --git a/ggml/src/ggml-cuda/mmq-config-rdna3.cuh b/ggml/src/ggml-cuda/mmq-config-rdna3.cuh index 8af3423d7d68..3a3ef7bd9c09 100644 --- a/ggml/src/ggml-cuda/mmq-config-rdna3.cuh +++ b/ggml/src/ggml-cuda/mmq-config-rdna3.cuh @@ -247,27 +247,16 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf // --------------------------------------------------------------------------------------------- CASE(GGML_TYPE_MXFP4, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 128, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 128, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 128, 1, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 128, 1, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 128, 1, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 128, 1, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 128, 1, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 128, 1, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 128, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 128, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 128, 2, 64, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 128, 2, 64, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 128, 2, 64, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 128, 2, 64, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_NVFP4, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, true); CASE(GGML_TYPE_NVFP4, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, true); diff --git a/ggml/src/ggml-cuda/mmq-config-rdna4.cuh b/ggml/src/ggml-cuda/mmq-config-rdna4.cuh index 8ef32188f690..9293d9d55885 100644 --- a/ggml/src/ggml-cuda/mmq-config-rdna4.cuh +++ b/ggml/src/ggml-cuda/mmq-config-rdna4.cuh @@ -261,29 +261,17 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf // --------------------------------------------------------------------------------------------- CASE(GGML_TYPE_MXFP4, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 128, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 128, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true); CASE(GGML_TYPE_MXFP4, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 128, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 128, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 128, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 128, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_MXFP4, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); - CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false); CASE(GGML_TYPE_NVFP4, 128, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, true); CASE(GGML_TYPE_NVFP4, 128, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, true); diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu index 31e5d3161e6c..c191a4f92e82 100644 --- a/ggml/src/ggml-cuda/mmq.cu +++ b/ggml/src/ggml-cuda/mmq.cu @@ -497,9 +497,17 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t // ------------------------------------------------- case GGML_TYPE_MXFP4: case GGML_TYPE_NVFP4: - case GGML_TYPE_Q4_0_ROCMFP4_FAST: mmq_supported = true; break; + case GGML_TYPE_Q4_0_ROCMFP4_FAST: + // MMQ tile rows exist only for RDNA3.5 (gfx1151, validated) and the + // Ampere config table (validated on RTX 4090 / sm_89). Gate to exactly + // those: Blackwell/Rubin have no rows of their own (native FP4 MMQ is + // MXFP4/NVFP4-only there) and would otherwise fall through to the + // Ampere table unvalidated. Everywhere else falls back to dequant + BLAS. + mmq_supported = (ampere_mma_available(cc) && cc < GGML_CUDA_CC_BLACKWELL) || + GGML_CUDA_CC_IS_RDNA3_5(cc); + break; default: mmq_supported = false; break; From 40cc375d834ec81647969f7809d6481b98d409a7 Mon Sep 17 00:00:00 2001 From: Baraxanx Date: Tue, 29 Sep 2026 21:25:50 -0400 Subject: [PATCH 5/5] cuda: return 0 for invalid ROCmFP4 scale bytes in the HIP decoder The arithmetic UE4M3 half-scale decode (non-LUT branch) accepted bytes 0x7f-0xff and produced garbage scales, while the CPU decoder (rocmfp4.c) returns 0 for them. Loader validation of those bytes only runs with --check-tensors, so a crafted GGUF decoded differently on GPU and CPU. Guard x > 0x7e -> 0 in both branches to match CPU. Verified bit-identical to the CPU LUT decoder for all 256 scale byte values; nvcc sm_89 compile clean. No behavior change for valid bytes, so existing test-backend-ops results stand. Assisted-by: Hermes Agent (Qwen3.8-Flash-Next) --- ggml/rocmfp4/rocmfp4_hip_scale.cuh | 11 +++++++---- 1 file changed, 7 insertions(+), 4 deletions(-) diff --git a/ggml/rocmfp4/rocmfp4_hip_scale.cuh b/ggml/rocmfp4/rocmfp4_hip_scale.cuh index a783be2e9493..b69ac0d87f85 100644 --- a/ggml/rocmfp4/rocmfp4_hip_scale.cuh +++ b/ggml/rocmfp4/rocmfp4_hip_scale.cuh @@ -97,12 +97,15 @@ static __device__ __forceinline__ float rocmfp4_u32_as_f32(uint32_t bits) { #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. +// Decode invalid scale bytes (> 0x7e) as 0, matching the CPU decoder in +// rocmfp4.c. Loader validation of these bytes only runs with --check-tensors, +// so kernels must not trust them. static __device__ __forceinline__ float rocmfp4_ue4m3_to_fp32_half_finite(uint8_t x) { + if (x > 0x7e) { + return 0.0f; + } #if defined(GGML_USE_HIP) && GGML_ROCMFP4_USE_SCALE_LUT - return x <= 0x7e ? rocmfp4_scale_ue4m3_half_lut[x] : 0.0f; + return rocmfp4_scale_ue4m3_half_lut[x]; #else const int exp = (x >> 3) & 0xF; const int man = x & 0x7;