Skip to content

hip: ROCm/HIP-only perf (TOP_K wave32 hybrid, k-quant VDR) #153

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
raufaser wants to merge 2 commits into
base: v0.4.7
Choose a base branch
from
Open

hip: ROCm/HIP-only perf (TOP_K wave32 hybrid, k-quant VDR) #153

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
abe3026
hip: ROCm/HIP-only perf (TOP_K wave32 hybrid, k-quant VDR, APU back-out)
raufaser Sep 9, 2026
5344385
hip: address #153 review (sgemm revert, top-k seeding, APU note)
raufaser Sep 11, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
6 changes: 6 additions & 0 deletions docs/beellama-features.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -198,6 +198,12 @@ CDNA fast routing is compiled and selected by capability but remains
experimental until hardware parity and performance results are published.
MUSA explicitly remains on the portable route.

On HIP, Bee reports `integrated = false`, backing out the upstream APU
zero-copy host-buffer path after async-execution corruption was observed
(PPL 5.9243 -> 8.51+ without `HIP_LAUNCH_BLOCKING`). This changes APU
tensor placement off host-mapped memory and therefore VRAM headroom; the
trade-off is unmeasured (no APU hardware available).

Set `GGML_KVARN_DEBUG_ROUTES=1` to log the selected CUDA/HIP route, compute
capability, rotated/original domain, K/V bit widths, query and KV counts,
attached exact-tail rows and type for integrated entries, entry path, and
Expand Down
9 changes: 8 additions & 1 deletion ggml/src/ggml-cuda/ggml-cuda.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -307,7 +307,14 @@ static ggml_cuda_device_info ggml_cuda_init() {
info.default_tensor_split[id] = total_vram;
total_vram += device_vram;
#if defined(GGML_USE_HIP)
info.devices[id].integrated = prop.integrated;
// Fork divergence from PR #24233: integrated=true enables the CUDA
// host-buffer path (zero-copy UMA weights) on APUs, which corrupts
// full-model results under async execution on this box (PPL 5.9243
// -> 8.51+ without HIP_LAUNCH_BLOCKING). The back-out restores
// async-safe operation. Placement changes with it: supports_buft gates
// the CUDA-host buffer type on `integrated`, so APU tensors move off
// host-mapped memory and VRAM headroom changes (unmeasured, no APU HW here).
info.devices[id].integrated = false; // Temporarily disabled due to issues with corrupted output (e.g. #15034)
#else
info.devices[id].integrated = false; // Temporarily disabled due to issues with corrupted output (e.g. #15034)
#endif
Expand Down
51 changes: 45 additions & 6 deletions ggml/src/ggml-cuda/mmvq.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -27,9 +27,24 @@ static constexpr __device__ vec_dot_q_cuda_t get_vec_dot_q_cuda(ggml_type type)
case GGML_TYPE_NVFP4: return vec_dot_nvfp4_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;
case GGML_TYPE_Q5_K: return vec_dot_q5_K_q8_1;
case GGML_TYPE_Q6_K: return vec_dot_q6_K_q8_1;
case GGML_TYPE_Q4_K:
#if defined(RDNA3_0) || defined(RDNA3_5) || defined(RDNA4)
return vec_dot_q4_K_q8_1_vdr4;
#else
return vec_dot_q4_K_q8_1;
#endif
case GGML_TYPE_Q5_K:
#if defined(RDNA3_0) || defined(RDNA3_5) || defined(RDNA4)
return vec_dot_q5_K_q8_1_vdr4;
#else
return vec_dot_q5_K_q8_1;
#endif
case GGML_TYPE_Q6_K:
#if defined(RDNA3_0) || defined(RDNA3_5) || defined(RDNA4)
return vec_dot_q6_K_q8_1_vdr2;
#else
return vec_dot_q6_K_q8_1;
#endif
case GGML_TYPE_IQ2_XXS: return vec_dot_iq2_xxs_q8_1;
case GGML_TYPE_IQ2_XS: return vec_dot_iq2_xs_q8_1;
case GGML_TYPE_IQ2_S: return vec_dot_iq2_s_q8_1;
Expand Down Expand Up @@ -82,6 +97,7 @@ enum mmvq_parameter_table_id {
MMVQ_PARAMETERS_GCN,
MMVQ_PARAMETERS_RDNA2,
MMVQ_PARAMETERS_RDNA3_0,
MMVQ_PARAMETERS_RDNA3_5,
MMVQ_PARAMETERS_RDNA4,
MMVQ_PARAMETERS_GB10
};
Expand All @@ -91,7 +107,9 @@ static constexpr __device__ mmvq_parameter_table_id get_device_table_id() {
return MMVQ_PARAMETERS_RDNA4;
#elif defined(RDNA3_0)
return MMVQ_PARAMETERS_RDNA3_0;
#elif defined(RDNA2) || defined(RDNA3_5)
#elif defined(RDNA3_5)
return MMVQ_PARAMETERS_RDNA3_5;
#elif defined(RDNA2)
return MMVQ_PARAMETERS_RDNA2;
#elif defined(GCN) || defined(CDNA)
return MMVQ_PARAMETERS_GCN;
Expand All @@ -111,7 +129,10 @@ static __host__ mmvq_parameter_table_id get_device_table_id(int cc) {
if (GGML_CUDA_CC_IS_RDNA3_0(cc)) {
return MMVQ_PARAMETERS_RDNA3_0;
}
if (GGML_CUDA_CC_IS_RDNA2(cc) || GGML_CUDA_CC_IS_RDNA3_5(cc)) {
if (GGML_CUDA_CC_IS_RDNA3_5(cc)) {
return MMVQ_PARAMETERS_RDNA3_5;
}
if (GGML_CUDA_CC_IS_RDNA2(cc)) {
return MMVQ_PARAMETERS_RDNA2;
}
if (GGML_CUDA_CC_IS_GCN(cc) || GGML_CUDA_CC_IS_CDNA(cc)) {
Expand Down Expand Up @@ -260,7 +281,7 @@ static constexpr __host__ __device__ int get_mmvq_mmid_max_batch_rdna4(ggml_type
case GGML_TYPE_Q3_K: return 4;
case GGML_TYPE_Q4_0: return 7;
case GGML_TYPE_Q4_1: return 7;
case GGML_TYPE_Q4_K: return 4;
case GGML_TYPE_Q4_K: return 7;
case GGML_TYPE_Q5_0: return 7;
case GGML_TYPE_Q5_1: return 7;
case GGML_TYPE_Q6_0: return 6;
Expand Down Expand Up @@ -508,6 +529,24 @@ static constexpr __host__ __device__ int calc_nwarps(ggml_type type, int ncols_d
}
return 1;
}
if (table_id == MMVQ_PARAMETERS_RDNA3_5) {
// gfx1151 (Strix Halo iGPU): nwarps=1 (the RDNA2 table) underutilizes the
// wave32 datapath on the large-K decode matmuls; nwarps=8 (the RDNA3_0
// table) over-parallelizes the small ones. Swept 2025-08: nwarps=2 wins
// (~+0.6% decode on Qwen3.6-35B-A3B Q8_0), nwarps=4 regresses.
// Apply to the whole mmvq range (ncols_dst 1..8), not just decode: the
// speculative verify batch (n_draft+1 tokens) must use the same nwarps
// as decode so its per-row dot-product accumulation is bit-identical.
if (ncols_dst <= MMVQ_MAX_BATCH_SIZE) {
switch (type) {
case GGML_TYPE_Q8_0:
return 2;
default:
return 1;
}
}
return 1;
}
if (table_id == MMVQ_PARAMETERS_TURING) {
if (ncols_dst == 1) {
switch (type) {
Expand Down
Loading