diff --git a/.gitignore b/.gitignore index 8d86412f5..6e4dfd850 100644 --- a/.gitignore +++ b/.gitignore @@ -21,3 +21,9 @@ __pycache__/ /misc/ .*.swp .DS_Store +tests/test_qwen35_rocm +tests/test_qwen_* +tests/test_q4k_top8 +tests/test_ssd_residency +*.orig +labestia_server.log diff --git a/Makefile b/Makefile index f9e4f7c22..fed2d47f4 100644 --- a/Makefile +++ b/Makefile @@ -18,7 +18,10 @@ DEPFLAGS ?= -MMD -MP BUILD_GIT_SHA ?= $(shell git rev-parse --short=12 HEAD 2>/dev/null || echo unknown) BUILD_GIT_SUFFIX ?= $(shell test -z "$$(git status --porcelain --untracked-files=normal 2>/dev/null)" || printf '%s' -dirty) -CFLAGS += -DDS4_BUILD_GIT_SHA=\"$(BUILD_GIT_SHA)$(BUILD_GIT_SUFFIX)\" +# Pass the sha as a bare token and stringify it in ds4_build.c: a quoted +# -D value does not survive the extra shell evaluation in recursive targets +# (strix-halo/cuda pass CFLAGS="$(CFLAGS) ..." through $(MAKE)). +CFLAGS += -DDS4_BUILD_GIT_SHA=$(BUILD_GIT_SHA)$(BUILD_GIT_SUFFIX) LDLIBS ?= -lm -pthread METAL_SRCS := $(wildcard metal/*.metal) diff --git a/RESULT_QWEN_ROCM.md b/RESULT_QWEN_ROCM.md new file mode 100644 index 000000000..2906a8b80 --- /dev/null +++ b/RESULT_QWEN_ROCM.md @@ -0,0 +1,82 @@ +# Qwen3.6-35B-A3B on ROCm (Radeon 780M) — result + +Date: 2026-07-16. Hardware: Minisforum UM790 Pro, Ryzen 9 7940HS, **iGPU Radeon +780M (gfx1103, run as gfx1100 via HSA_OVERRIDE_GFX_VERSION=11.0.0)**, 64 GB DDR5, +Pop!_OS, ROCm 6.3.1. Branch `rocm-qwen` (fork of andreaborio/ds4). + +## Status: **Qwen3.6-35B-A3B runs end-to-end on the 780M GPU.** + +Resident mode (19.37 GiB payload wired in GTT, no SSD streaming). Coherent +generation; numerically matches the scalar CPU reference. + +### Measured (greedy, temp 0) + +| path | prefill t/s | generation t/s | notes | +|---|---|---|---| +| ds4 **ROCm GPU** (this port) | 11.8–14.3 | **13.4–14.9** | resident, 780M | +| ds4 CPU reference | ~10 | ~10–11.5 | scalar oracle | +| llama.cpp CPU (same artifact) | ~100 (pp64) | 16.1 (tg64) | 8 threads | + +The GPU path already beats the ds4 CPU reference. It does not yet beat +llama.cpp — expected for a v1 correctness-first port (no HIP-graph batching, F32 +throughout, top-8 MoE run as two top-4 halves). Andrea's Metal numbers (58–66 +t/s on M5 Pro) are on ~2× the memory bandwidth and a mature kernel set. + +### Correctness (M2 parity gate) + +CPU-vs-GPU top-logprobs, prompt "The sea is", 5 steps: +- **top-1 agreement: 5/5** (GPU selects the same token as the CPU reference at + every step) +- **cosine 0.9994–0.9999** on the shared top-k logit vectors (Andrea's own + Metal-vs-llama.cpp gate is 0.99994) +- max logit diff 0.5–1.7 out of ~26 (sub-1%) — pure f32 accumulation-order + difference between the scalar CPU path and the GPU kernels. + +Greedy 96-token text is token-identical to the CPU golden for the first ~40 +tokens, then diverges at a near-tie — the documented upstream behavior, not a +bug. + +Per-op kernel parity (`tests/test_qwen35_rocm.c`, all 22 checks pass on the +780M vs `ds4_qwen_ref.c`): split_q_gate, sigmoid_mul (elements/rows), rope, +router (exact top-8 + weights), gqa decode/prefill, gated_delta_step (kd=128 +fast path + kd=64 generic), gated_delta_sequence_128, causal_conv step/sequence, +rmsnorm_gated, gated_delta_controls, dequant_embedding_q8_0. The GDN recurrence +(the novel, highest-risk kernel) matches to ~1e-8 relative. + +## What was built + +- **`rocm/ds4_rocm_qwen.cuh`**: 19 HIP kernels (1:1 port of `metal/qwen35.metal`; + Metal simdgroup → RDNA3 wave32, simd_sum/max/min → `__shfl_xor` butterfly) + + the 20 `ds4_gpu_qwen35_*` entry points. Byte-stride args mirror `ds4_metal.m`. +- **Top-8 MoE**: the shared ROCm Q4_K MoE core is specialized for DeepSeek's + top-6 (fused `sum6` kernels). Qwen's top-8 runs as **two top-4 halves + add** + (`moe_split_selected_kernel` + `routed_moe_batch_two_half`), leaving the + DeepSeek path untouched. +- **Host wiring (`ds4.c`)**: Qwen GPU forward/session/dispatch guards widened + from `__APPLE__` to `DS4_QWEN_GPU_BUILD`. No `DS4_BACKEND_ROCM` enum exists — + ROCm runs under `DS4_BACKEND_CUDA` — so `ds4_backend_is_qwen_gpu()` replaces + the `backend==METAL` checks. Env gate `DS4_QWEN_EXPERIMENTAL_ROCM=1`; resident + forced (SSD streaming for Qwen not ported). +- **Artifact**: `tests/qwen/normalize_qwen36_gguf.py` reproduces the blessed + `Qwen3.6-35B-A3B-ds4-Q4_K_S.gguf` from the Unsloth UD-Q4_K_S source (Q6_K + ffn_down_exps → Q4_K, output → Q8_0, pad-id + chat template). ds4 accepts it; + inventory **f32=361 / q8_0=252 / q4_k=120 == the prereg histogram exactly**. +- **Forward-ported La Bestia ROCm hardening** (multi-model range cache, readback + event ring + host-side waits, adaptive free-floor) onto the fork; build-sha + bug fixed. + +## Run + +```sh +DS4_QWEN_EXPERIMENTAL_ROCM=1 HSA_OVERRIDE_GFX_VERSION=11.0.0 \ + ./ds4 -m gguf/Qwen3.6-35B-A3B-ds4-Q4_K_S.gguf -c 4096 -p "…" +``` + +## Not done / next + +- Soak at 8k and 32k context; formal `ds4-bench` CSV. +- Throughput: HIP-graph batching, fused gate/up+SwiGLU for the two-half MoE, + parallel GQA decode variant (only the serial GQA kernel is wired). +- Server path on ROCm (chat.html) — engine wiring is done; needs a server smoke. +- **Caveat**: the Qwen server is text-only (no tool-call), so it cannot replace + Ollama for JARVIS function-calling until upstream adds tool parsing. diff --git a/ds4.c b/ds4.c index e388b3686..79de649ac 100644 --- a/ds4.c +++ b/ds4.c @@ -78,6 +78,25 @@ static bool ds4_backend_uses_graph(ds4_backend backend) { return backend == DS4_BACKEND_METAL || backend == DS4_BACKEND_CUDA; } +/* The Qwen3.6 GPU runtime is Apple Metal today; the ROCm build reuses the + * legacy metal_graph_* graph backend, which presents at runtime as the CUDA + * backend enum (ds4_build_backend() still reports "rocm"). This helper marks + * the graph backend that carries the Qwen GPU path for the current build so + * the many backend equality checks stay in one place. */ +#if defined(__APPLE__) || (defined(DS4_ROCM_BUILD) && !defined(DS4_NO_GPU)) +#define DS4_QWEN_GPU_BUILD 1 +#endif +static bool ds4_backend_is_qwen_gpu(ds4_backend backend) { +#if defined(__APPLE__) + return backend == DS4_BACKEND_METAL; +#elif defined(DS4_ROCM_BUILD) && !defined(DS4_NO_GPU) + return backend == DS4_BACKEND_CUDA; +#else + (void)backend; + return false; +#endif +} + static bool ds4_backend_supports_ssd_streaming(ds4_backend backend) { if (backend == DS4_BACKEND_METAL) return true; if (backend == DS4_BACKEND_CUDA) { @@ -12665,7 +12684,7 @@ void ds4_internal_qwen35_gpu_graph_free( qwen35_gpu_graph_free(storage); } -#if defined(__APPLE__) +#if defined(__APPLE__) || defined(DS4_ROCM_BUILD) /* ========================================================================= * Qwen3.6 Metal one-token forward. * ========================================================================= @@ -14018,7 +14037,7 @@ static DS4_MAYBE_UNUSED bool qwen35_gpu_forward_token( return qwen35_gpu_forward_token_commands( logits, model, weights, graph, token, position, false, true); } -#endif /* __APPLE__ */ +#endif /* __APPLE__ || DS4_ROCM_BUILD */ /* ========================================================================= * Metal Release Graph State. @@ -28503,7 +28522,7 @@ struct ds4_session { ds4_dist_session *distributed; #ifndef DS4_NO_GPU ds4_gpu_graph graph; -#if defined(__APPLE__) +#ifdef DS4_QWEN_GPU_BUILD ds4_qwen35_gpu_graph qwen35_gpu_graph; #endif #endif @@ -28839,7 +28858,7 @@ static bool ds4_session_is_qwen35_cpu(const ds4_session *s) { static DS4_MAYBE_UNUSED bool ds4_session_is_qwen35_metal( const ds4_session *s) { return ds4_session_is_qwen35(s) && s->engine->qwen_metal_runtime && - s->engine->backend == DS4_BACKEND_METAL; + ds4_backend_is_qwen_gpu(s->engine->backend); } static uint32_t ds4_session_vocab_size(const ds4_session *s) { @@ -28888,7 +28907,7 @@ static bool ds4_session_qwen35_timeline_valid(const ds4_session *s) { if (ds4_session_is_qwen35_cpu(s)) { return s->qwen35_cpu_cache.n_tokens == checkpoint_len; } -#if defined(__APPLE__) && !defined(DS4_NO_GPU) +#ifdef DS4_QWEN_GPU_BUILD if (ds4_session_is_qwen35_metal(s)) { return s->qwen35_gpu_graph.state_valid && s->qwen35_gpu_graph.n_tokens == checkpoint_len; @@ -28902,7 +28921,7 @@ static bool ds4_session_qwen35_reset_timeline(ds4_session *s) { bool ok = true; if (ds4_session_is_qwen35_cpu(s)) { ds4_qwen35_cpu_cache_reset(&s->qwen35_cpu_cache); -#if defined(__APPLE__) && !defined(DS4_NO_GPU) +#ifdef DS4_QWEN_GPU_BUILD } else if (ds4_session_is_qwen35_metal(s)) { /* qwen35_gpu_forward_token already performs a full reset after any * late, state-mutating failure. Avoid zeroing the entire context a @@ -31037,11 +31056,18 @@ static bool ds4_engine_configure_streaming_auto_cache(ds4_engine *e) { return true; } - if (e->backend != DS4_BACKEND_METAL) { + /* The adaptive planner is capability-gated, not backend-gated: any + * backend that can produce a live host-memory snapshot gets the adaptive + * budget (ROCm UMA APUs need it as much as Metal does — without a + * host-aware floor the kernel TTM can evict the GPU queue buffers under + * peak cache pressure and freeze generation). Backends without a + * snapshot keep the legacy fixed-fraction budget. */ + ds4_ssd_host_memory memory; + if (e->backend != DS4_BACKEND_METAL && + !ds4_gpu_host_memory_snapshot(&memory)) { return ds4_engine_configure_streaming_legacy_auto_cache(e); } - ds4_ssd_host_memory memory; if (!ds4_gpu_host_memory_snapshot(&memory)) { fprintf(stderr, "ds4: SSD streaming auto cache: host memory snapshot unavailable; " @@ -31672,9 +31698,9 @@ static bool ds4_engine_qwen35_metal_options_valid( const ds4_engine_options *opt, bool load_slice) { if (!e || !opt) return false; - if (e->backend != DS4_BACKEND_METAL) { + if (!ds4_backend_is_qwen_gpu(e->backend)) { fprintf(stderr, - "ds4: experimental Qwen Metal inference requires --metal\n"); + "ds4: experimental Qwen GPU inference requires the graph backend\n"); return false; } if (e->quality) { @@ -31743,7 +31769,7 @@ static bool ds4_engine_qwen35_metal_options_valid( return true; } -#if defined(__APPLE__) && !defined(DS4_NO_GPU) +#ifdef DS4_QWEN_GPU_BUILD static bool ds4_engine_configure_qwen35_metal_streaming(ds4_engine *e) { if (!e || e->backend != DS4_BACKEND_METAL || e->residency != DS4_RESIDENCY_SSD || !e->ssd_streaming) { @@ -31941,7 +31967,7 @@ static bool ds4_engine_configure_qwen35_metal_streaming(ds4_engine *e) { } static bool ds4_engine_configure_qwen35_metal_resident(ds4_engine *e) { - if (!e || e->backend != DS4_BACKEND_METAL || + if (!e || !ds4_backend_is_qwen_gpu(e->backend) || e->residency != DS4_RESIDENCY_RESIDENT || e->ssd_streaming || !e->model.map || e->model.tensor_data_pos >= e->model.size) { return false; @@ -31949,7 +31975,7 @@ static bool ds4_engine_configure_qwen35_metal_resident(ds4_engine *e) { if (!ds4_gpu_init()) { fprintf(stderr, - "ds4: Metal backend unavailable while initializing Qwen resident mode\n"); + "ds4: GPU backend unavailable while initializing Qwen resident mode\n"); return false; } e->metal_ready = true; @@ -32150,12 +32176,17 @@ int ds4_engine_open(ds4_engine **out, const ds4_engine_options *opt) { return 0; } - const char *experimental_metal = - getenv("DS4_QWEN_EXPERIMENTAL_METAL"); - if (!experimental_metal || strcmp(experimental_metal, "1") != 0) { +#ifdef DS4_ROCM_BUILD + const char *experimental_gate = getenv("DS4_QWEN_EXPERIMENTAL_ROCM"); + const char *experimental_gate_name = "DS4_QWEN_EXPERIMENTAL_ROCM"; +#else + const char *experimental_gate = getenv("DS4_QWEN_EXPERIMENTAL_METAL"); + const char *experimental_gate_name = "DS4_QWEN_EXPERIMENTAL_METAL"; +#endif + if (!experimental_gate || strcmp(experimental_gate, "1") != 0) { fprintf(stderr, - "ds4: Qwen Metal inference requires " - "DS4_QWEN_EXPERIMENTAL_METAL=1\n"); + "ds4: Qwen GPU inference requires %s=1\n", + experimental_gate_name); ds4_engine_close(e); return 1; } @@ -32164,7 +32195,7 @@ int ds4_engine_open(ds4_engine **out, const ds4_engine_options *opt) { ds4_engine_close(e); return 1; } -#if defined(__APPLE__) && !defined(DS4_NO_GPU) +#ifdef DS4_QWEN_GPU_BUILD if (!ds4_engine_resolve_residency( e, opt, false, 0, 0, false)) { ds4_engine_close(e); @@ -32172,6 +32203,14 @@ int ds4_engine_open(ds4_engine **out, const ds4_engine_options *opt) { } vocab_load(&e->vocab, &e->model); bool configured = false; +#ifdef DS4_ROCM_BUILD + /* Qwen SSD streaming is not ported to ROCm; the ~19.4 GiB payload is + * resident on the 64 GB box, the regime this port targets. */ + e->residency = DS4_RESIDENCY_RESIDENT; + e->ssd_streaming = false; + if (opt->warm_weights) model_warm_weights(&e->model); + configured = ds4_engine_configure_qwen35_metal_resident(e); +#else if (e->residency == DS4_RESIDENCY_SSD && e->ssd_streaming) { if (opt->warm_weights) { fprintf(stderr, @@ -32185,9 +32224,10 @@ int ds4_engine_open(ds4_engine **out, const ds4_engine_options *opt) { configured = ds4_engine_configure_qwen35_metal_resident(e); } +#endif if (!configured) { fprintf(stderr, - "ds4: Qwen Metal residency configuration failed for %s mode\n", + "ds4: Qwen GPU residency configuration failed for %s mode\n", ds4_residency_mode_name(e->residency)); ds4_engine_close(e); return 1; @@ -32195,7 +32235,7 @@ int ds4_engine_open(ds4_engine **out, const ds4_engine_options *opt) { e->qwen_raw_runtime = true; e->qwen_metal_runtime = true; fprintf(stderr, - "ds4: experimental Qwen Metal %s runtime enabled\n", + "ds4: experimental Qwen GPU %s runtime enabled\n", ds4_residency_mode_name(e->residency)); *out = e; return 0; @@ -32866,7 +32906,7 @@ int ds4_session_create(ds4_session **out, ds4_engine *e, int ctx_size) { return 1; } if (qwen35 && e->backend != DS4_BACKEND_CPU && - !(e->backend == DS4_BACKEND_METAL && e->qwen_metal_runtime)) { + !(ds4_backend_is_qwen_gpu(e->backend) && e->qwen_metal_runtime)) { fprintf(stderr, "ds4: Qwen session backend does not match its enabled runtime\n"); return 1; @@ -32940,7 +32980,7 @@ int ds4_session_create(ds4_session **out, ds4_engine *e, int ctx_size) { ds4_session *s = xcalloc(1, sizeof(*s)); s->engine = e; s->ctx_size = ctx_size; -#if defined(__APPLE__) +#ifdef DS4_QWEN_GPU_BUILD if (qwen35) { s->prefill_cap = 1u; if (!qwen35_gpu_graph_alloc( @@ -33037,7 +33077,7 @@ void ds4_session_free(ds4_session *s) { } #ifndef DS4_NO_GPU else if (ds4_session_is_qwen35_metal(s)) { -#if defined(__APPLE__) +#ifdef DS4_QWEN_GPU_BUILD if (getenv("DS4_METAL_MEMORY_REPORT") != NULL) { ds4_gpu_print_memory_report("before Qwen graph free"); } @@ -33772,7 +33812,7 @@ static int ds4_session_sync_qwen35_with_forward( return 0; } -#if defined(__APPLE__) && !defined(DS4_NO_GPU) +#ifdef DS4_QWEN_GPU_BUILD static bool ds4_session_qwen35_metal_forward_commit( ds4_session *s, int token, @@ -34099,7 +34139,7 @@ int ds4_session_sync(ds4_session *s, const ds4_tokens *prompt, char *err, size_t return ds4_session_sync_qwen35_with_forward( s, prompt, err, errlen, qwen35_cpu_forward_token); } -#if defined(__APPLE__) && !defined(DS4_NO_GPU) +#ifdef DS4_QWEN_GPU_BUILD if (ds4_session_is_qwen35_metal(s)) { return ds4_session_sync_qwen35_metal( s, prompt, err, errlen); @@ -34519,7 +34559,7 @@ static int ds4_session_eval_qwen35_with_forward( return 0; } -#if defined(__APPLE__) && !defined(DS4_NO_GPU) +#ifdef DS4_QWEN_GPU_BUILD static int ds4_session_eval_qwen35_metal( ds4_session *s, int token, @@ -34584,7 +34624,7 @@ static int ds4_session_eval_internal(ds4_session *s, int token, bool probe_mtp, return ds4_session_eval_qwen35_with_forward( s, token, err, errlen, qwen35_cpu_forward_token); } -#if defined(__APPLE__) && !defined(DS4_NO_GPU) +#ifdef DS4_QWEN_GPU_BUILD if (ds4_session_is_qwen35_metal(s)) { return ds4_session_eval_qwen35_metal( s, token, err, errlen); diff --git a/ds4_build.c b/ds4_build.c index bda77ef30..042479594 100644 --- a/ds4_build.c +++ b/ds4_build.c @@ -2,9 +2,14 @@ #include +/* The Makefile passes DS4_BUILD_GIT_SHA as a bare token (no embedded quotes): + * quoted -D values lose their quotes in recursive targets, where CFLAGS goes + * through one extra shell evaluation. Stringify it here instead. */ #ifndef DS4_BUILD_GIT_SHA -#define DS4_BUILD_GIT_SHA "unknown" +#define DS4_BUILD_GIT_SHA unknown #endif +#define DS4_BUILD_STR2(x) #x +#define DS4_BUILD_STR(x) DS4_BUILD_STR2(x) const char *ds4_build_backend(void) { #ifdef DS4_NO_GPU @@ -29,7 +34,7 @@ const char *ds4_build_arch(void) { } const char *ds4_build_git_sha(void) { - return DS4_BUILD_GIT_SHA; + return DS4_BUILD_STR(DS4_BUILD_GIT_SHA); } void ds4_build_info_print(FILE *fp) { diff --git a/ds4_rocm.cu b/ds4_rocm.cu index d5ea9b53a..5a14b587e 100644 --- a/ds4_rocm.cu +++ b/ds4_rocm.cu @@ -128,4 +128,6 @@ typedef struct { #include "rocm/ds4_rocm_hc_output_launch.cuh" +#include "rocm/ds4_rocm_qwen.cuh" + #include "rocm/ds4_rocm_current_api_compat.cuh" diff --git a/ds4_rocm.h b/ds4_rocm.h index 50481b8d5..c05b11259 100644 --- a/ds4_rocm.h +++ b/ds4_rocm.h @@ -1,6 +1,15 @@ #pragma once +/* Opt in to the __shfl_*_sync / __ballot_sync warp builtins, which ROCm 6.3 + * gates behind this macro (hip/amd_detail/amd_warp_sync_functions.h). The + * backend calls __shfl_sync/__shfl_down_sync with an 8-byte MASK_T mask. */ +#define HIP_ENABLE_WARP_SYNC_BUILTINS 1 + #include +/* hipBLAS 2.x hides the hipDataType-based GemmEx/GemmStridedBatchedEx behind + * HIPBLAS_V2; without it only the deprecated hipblasDatatype_t overloads are + * declared and the CUDA_R_* (now hipDataType) arguments don't match. */ +#define HIPBLAS_V2 #include #include #include @@ -79,8 +88,11 @@ #define CUBLAS_DEFAULT_MATH HIPBLAS_DEFAULT_MATH #define CUBLAS_COMPUTE_32F HIPBLAS_COMPUTE_32F #define CUBLAS_TF32_TENSOR_OP_MATH HIPBLAS_TF32_TENSOR_OP_MATH -#define CUDA_R_16F HIPBLAS_R_16F -#define CUDA_R_32F HIPBLAS_R_32F +/* hipBLAS 2.x (ROCm 6.3) hipblasGemmEx/GemmStridedBatchedEx take hipDataType + * (HIP_R_*), not the deprecated hipblasDatatype_t (HIPBLAS_R_*). The compute + * type stays hipblasComputeType_t (HIPBLAS_COMPUTE_32F). */ +#define CUDA_R_16F HIP_R_16F +#define CUDA_R_32F HIP_R_32F #define cublasCreate hipblasCreate #define cublasDestroy hipblasDestroy @@ -92,6 +104,13 @@ namespace cub = hipcub; +/* CUDA's __syncwarp has no HIP equivalent (not provided even with + * HIP_ENABLE_WARP_SYNC_BUILTINS on ROCm 6.3). On RDNA3 wave32 the lanes of a + * wave execute in lockstep, so warp reconvergence is implicit; a wavefront- + * scoped fence supplies the compiler ordering the call is relied on for. + * Variadic so both __syncwarp() and __syncwarp(mask) map through. */ +#define __syncwarp(...) __builtin_amdgcn_fence(__ATOMIC_ACQ_REL, "wavefront") + static __device__ __forceinline__ int32_t __vcmpne4(uint32_t a, uint32_t b) { // For each byte: 0xFF if a != b, 0x00 if a == b uint32_t diff = a ^ b; diff --git a/rocm/ds4_rocm_attention_launch.cuh b/rocm/ds4_rocm_attention_launch.cuh index b9b43d958..0a194eec4 100644 --- a/rocm/ds4_rocm_attention_launch.cuh +++ b/rocm/ds4_rocm_attention_launch.cuh @@ -990,11 +990,29 @@ extern "C" int ds4_gpu_attention_output_q8_batch_tensor( cuda_model_range_ptr(model_map, out_b_offset, out_b_bytes, "attn_out_b")); if (!out_a || !out_b) return 0; + /* NOTE: routing small streaming batches (MTP verify) through the cublas + * f16 branch was tried (DS4_ROCM_ATTN_OUT_FAST_BATCH) and reverted: it + * halved this projection's cost but reintroduced the intermittent HSA + * queue wedge under multi-stream load. The q8 small-n ntok kernels give + * most of the win with none of the fragility. */ const int attn_output_cublas = cuda_runtime_config()->attention_output_cublas_all && (n_tokens == 1u || !g_ssd_streaming_mode); if (!attn_output_cublas) { - if ((group_dim & 31u) == 0u && rank <= UINT32_MAX && n_tokens <= UINT32_MAX) { + const int outproj_profile = getenv("DS4_ROCM_ATTN_OUT_PROFILE") != NULL; + double outproj_t0 = 0.0; + if (outproj_profile) { + (void)cudaDeviceSynchronize(); + outproj_t0 = cuda_wall_sec(); + } + /* Same small-batch dispatch rationale as cuda_matmul_q8_0_tensor_labeled: + * the grouped sharedx kernel tiles 16 tokens per block and wastes most + * of the tile on MTP verify suffixes (n_tokens 2-5). */ + const int smalln_warp8 = + n_tokens > 1u && n_tokens < 8u && + getenv("DS4_ROCM_SMALLN_SHAREDX") == NULL; + if (!smalln_warp8 && + (group_dim & 31u) == 0u && rank <= UINT32_MAX && n_tokens <= UINT32_MAX) { const uint32_t rows_per_block = 32u; const uint32_t tile = 32u; const uint32_t block_tile = 16u; @@ -1022,6 +1040,26 @@ extern "C" int ds4_gpu_attention_output_q8_batch_tensor( blocks_a); } if (!cuda_ok(cudaGetLastError(), "attention_output_q8_a f32 batch launch")) return 0; + if (outproj_profile) { + (void)cudaDeviceSynchronize(); + const double t_a = cuda_wall_sec(); + const int rc = cuda_matmul_q8_0_tensor_labeled(out, model_map, model_size, + out_b_offset, low_dim, out_dim, + low, n_tokens, "attn_output_b"); + (void)cudaDeviceSynchronize(); + const double t_b = cuda_wall_sec(); + fprintf(stderr, + DS4_GPU_LOG_PREFIX "attn_out profile tokens=%u a=%.3f ms b=%.3f ms " + "(low_dim=%llu out_dim=%llu groups=%u gdim=%llu)\n", + n_tokens, + (t_a - outproj_t0) * 1000.0, + (t_b - t_a) * 1000.0, + (unsigned long long)low_dim, + (unsigned long long)out_dim, + n_groups, + (unsigned long long)group_dim); + return rc; + } return cuda_matmul_q8_0_tensor_labeled(out, model_map, model_size, diff --git a/rocm/ds4_rocm_current_api_compat.cuh b/rocm/ds4_rocm_current_api_compat.cuh index 574f6e61b..120165974 100644 --- a/rocm/ds4_rocm_current_api_compat.cuh +++ b/rocm/ds4_rocm_current_api_compat.cuh @@ -1,10 +1,17 @@ +static cudaEvent_t *selected_readback_event_slot(uint64_t event_value) { + if (event_value == 0) return NULL; + return &g_selected_readback_events[event_value % + DS4_SELECTED_READBACK_EVENT_RING]; +} + extern "C" int ds4_gpu_signal_selected_readback_ready(uint64_t *event_value) { if (!event_value) return 0; *event_value = 0; - if (!g_selected_readback_event) { + const uint64_t value = g_selected_readback_event_value + 1; + cudaEvent_t *slot = selected_readback_event_slot(value); + if (!*slot) { cudaError_t err = - cudaEventCreateWithFlags(&g_selected_readback_event, - cudaEventDisableTiming); + cudaEventCreateWithFlags(slot, cudaEventDisableTiming); if (err != cudaSuccess) { fprintf(stderr, DS4_GPU_LOG_PREFIX "selected readback event creation failed: %s\n", @@ -13,7 +20,7 @@ extern "C" int ds4_gpu_signal_selected_readback_ready(uint64_t *event_value) { return 0; } } - cudaError_t err = cudaEventRecord(g_selected_readback_event, 0); + cudaError_t err = cudaEventRecord(*slot, 0); if (err != cudaSuccess) { fprintf(stderr, DS4_GPU_LOG_PREFIX "selected readback event record failed: %s\n", @@ -21,13 +28,15 @@ extern "C" int ds4_gpu_signal_selected_readback_ready(uint64_t *event_value) { (void)cudaGetLastError(); return 0; } - *event_value = ++g_selected_readback_event_value; + g_selected_readback_event_value = value; + *event_value = value; return 1; } extern "C" int ds4_gpu_commit_and_wait_selected_readback(uint64_t event_value, const char *label) { - if (event_value == 0 || !g_selected_readback_event) return 0; - cudaError_t err = cudaEventSynchronize(g_selected_readback_event); + cudaEvent_t *slot = selected_readback_event_slot(event_value); + if (!slot || !*slot) return 0; + cudaError_t err = cudaEventSynchronize(*slot); if (err != cudaSuccess) { fprintf(stderr, DS4_GPU_LOG_PREFIX "selected readback wait failed for %s: %s\n", @@ -50,10 +59,10 @@ extern "C" int ds4_gpu_tensor_read_after_selected_event( uint64_t bytes, uint64_t event_value, const char *label) { + cudaEvent_t *slot = selected_readback_event_slot(event_value); if (!tensor || !data || offset > tensor->bytes || bytes > tensor->bytes - offset || - event_value == 0 || - !g_selected_readback_event) { + !slot || !*slot) { return 0; } if (!g_selected_readback_stream) { @@ -68,18 +77,21 @@ extern "C" int ds4_gpu_tensor_read_after_selected_event( return 0; } } -#ifdef __HIP_PLATFORM_AMD__ - cudaError_t err = hipStreamWaitEvent(g_selected_readback_stream, - g_selected_readback_event, - 0); -#else - cudaError_t err = cudaStreamWaitEvent(g_selected_readback_stream, - g_selected_readback_event, - 0); -#endif + /* + * Wait for the router's event on the HOST, not with a device-side + * hipStreamWaitEvent barrier. On this ROCm/APU stack a pending + * stream-wait barrier occasionally misses its signal and wedges the HSA + * hardware queue it occupies; every stream multiplexed onto that queue + * (including the pinned-staging upload streams) then stalls behind it and + * the process deadlocks in cudaStreamSynchronize/cudaDeviceSynchronize + * with an idle GPU. A host-side event wait provides the same ordering + * guarantee (the copy below starts only after the router output is + * complete) without enqueueing any device-side dependency. + */ + cudaError_t err = cudaEventSynchronize(*slot); if (err != cudaSuccess) { fprintf(stderr, - DS4_GPU_LOG_PREFIX "selected readback stream wait failed for %s: %s\n", + DS4_GPU_LOG_PREFIX "selected readback event wait failed for %s: %s\n", label ? label : "selected-id readback", cudaGetErrorString(err)); (void)cudaGetLastError(); @@ -174,9 +186,23 @@ extern "C" void ds4_gpu_set_streaming_expert_cache_budget(uint32_t experts) { g_stream_expert_cache_budget = experts; } +/* Fail-closed correctness floor from the adaptive planner (ds4_ssd.c). The + * ROCm cache does not yet shrink under live pressure the way Metal does, so + * the floor currently guards only the startup budget; storing it keeps the + * model contract visible to shared engine code and future enforcement. */ +static uint32_t g_stream_expert_cache_required_floor; + extern "C" void ds4_gpu_set_streaming_expert_cache_required_floor( uint32_t experts) { - (void)experts; + g_stream_expert_cache_required_floor = experts; + if (experts != 0 && g_stream_expert_cache_budget != 0 && + g_stream_expert_cache_budget < experts) { + fprintf(stderr, + DS4_GPU_LOG_PREFIX "streaming expert cache budget %u is below " + "the model's correctness floor %u\n", + g_stream_expert_cache_budget, + experts); + } } extern "C" void ds4_gpu_set_streaming_expert_cache_expert_bytes(uint64_t bytes) { @@ -199,9 +225,93 @@ extern "C" uint64_t ds4_gpu_recommended_working_set_size(void) { return (uint64_t)total_b; } +/* + * Linux/ROCm host memory snapshot for the adaptive SSD expert-cache planner. + * + * Field mapping from the Darwin collector (ds4_metal.m) to Linux sources: + * physical_bytes <- /proc/meminfo MemTotal + * recommended_bytes <- hipMemGetInfo() total: on UMA APUs the GTT pool is + * the platform's working-set envelope, the closest + * analogue of recommendedMaxWorkingSetSize + * task_footprint <- /proc/self/status VmRSS + * free_bytes <- MemFree (MemAvailable would double-count the + * file-backed credit the planner adds separately) + * purgeable_bytes <- 0 (no Linux analogue) + * inactive_bytes <- Inactive(file): reclaimable page-cache, the + * conservative counterpart of Mach inactive_count + * file_backed_bytes <- Cached + * pressure_normal <- PSI /proc/pressure/memory, "some avg10" < 5.0 + * + * Motivating failure (Radeon 780M, 64 GB, GTT 48 GiB): with the legacy fixed + * pool/8 floor the expert cache peaked after ~150-250 generated tokens and + * the kernel TTM evicted the GPU queue buffers themselves — dmesg shows + * "Not enough memory for command submission!" and "Freeing queue vital + * buffer, queue evicted" — freezing generation with an idle GPU. A + * host-aware budget prevents it; this snapshot is what the planner needs to + * compute one. + */ extern "C" int ds4_gpu_host_memory_snapshot(ds4_ssd_host_memory *out) { - if (out) memset(out, 0, sizeof(*out)); + if (!out) return 0; + memset(out, 0, sizeof(*out)); +#if !defined(__linux__) return 0; +#else + size_t gtt_free = 0; + size_t gtt_total = 0; + if (cudaMemGetInfo(>t_free, >t_total) != cudaSuccess || gtt_total == 0) { + (void)cudaGetLastError(); + return 0; + } + + uint64_t mem_total = 0, mem_free = 0, inactive_file = 0, cached = 0; + FILE *mi = fopen("/proc/meminfo", "r"); + if (!mi) return 0; + char line[160]; + while (fgets(line, sizeof(line), mi)) { + unsigned long long kb = 0; + if (sscanf(line, "MemTotal: %llu kB", &kb) == 1) mem_total = (uint64_t)kb << 10; + else if (sscanf(line, "MemFree: %llu kB", &kb) == 1) mem_free = (uint64_t)kb << 10; + else if (sscanf(line, "Inactive(file): %llu kB", &kb) == 1) inactive_file = (uint64_t)kb << 10; + else if (sscanf(line, "Cached: %llu kB", &kb) == 1) cached = (uint64_t)kb << 10; + } + fclose(mi); + if (mem_total == 0) return 0; + + uint64_t vm_rss = 0; + FILE *st = fopen("/proc/self/status", "r"); + if (st) { + while (fgets(line, sizeof(line), st)) { + unsigned long long kb = 0; + if (sscanf(line, "VmRSS: %llu kB", &kb) == 1) { + vm_rss = (uint64_t)kb << 10; + break; + } + } + fclose(st); + } + + out->physical_bytes = mem_total; + out->recommended_bytes = (uint64_t)gtt_total; + out->task_footprint_bytes = vm_rss; + out->free_bytes = mem_free; + out->purgeable_bytes = 0; + out->inactive_bytes = inactive_file; + out->file_backed_bytes = cached; + + FILE *psi = fopen("/proc/pressure/memory", "r"); + if (psi) { + while (fgets(line, sizeof(line), psi)) { + float avg10 = 0.0f; + if (sscanf(line, "some avg10=%f", &avg10) == 1) { + out->pressure_status_available = true; + out->pressure_normal = avg10 < 5.0f; + break; + } + } + fclose(psi); + } + return 1; +#endif } extern "C" uint32_t ds4_gpu_stream_expert_cache_configured_count(void) { @@ -359,3 +469,16 @@ extern "C" int ds4_gpu_routed_moe_set_selected_override( g_routed_moe_selected_override_n = n_selected; return 1; } + +/* Documented contract (ds4_gpu.h): zero on non-Metal backends. The ROCm + * runtime keeps its own optional counters behind DS4_ROCM_STREAM_CACHE_STATS; + * wiring them here is future work. */ +extern "C" void ds4_gpu_stream_expert_cache_stats( + uint64_t *hits, uint64_t *misses, uint64_t *pread_bytes, + double *pread_ms, double *split_resident_wait_ms) { + if (hits) *hits = 0; + if (misses) *misses = 0; + if (pread_bytes) *pread_bytes = 0; + if (pread_ms) *pread_ms = 0.0; + if (split_resident_wait_ms) *split_resident_wait_ms = 0.0; +} diff --git a/rocm/ds4_rocm_matmul.cuh b/rocm/ds4_rocm_matmul.cuh index 0f9e06625..5b6878ef9 100644 --- a/rocm/ds4_rocm_matmul.cuh +++ b/rocm/ds4_rocm_matmul.cuh @@ -293,6 +293,47 @@ static int cuda_matmul_q8_0_tensor_labeled(ds4_gpu_tensor *out, const void *mode return cuda_ok(cudaGetLastError(), "matmul_q8_0 f32 batch wmma 4w launch"); } #endif + /* + * Small batches (MTP verify suffixes, 2-4 tokens): both per-token grid + * kernels (weights re-read once PER TOKEN) and the sharedx tiles + * (sized for 16-token tiles) waste most of their bandwidth here. + * Quantize the activations once and run the ntok kernel, which reads + * each weight row once for all tokens — the same trick that makes the + * single-token decode path fast. DS4_ROCM_SMALLN_SHAREDX restores the + * old dispatch for comparison. + */ + const int smalln_preq = + n_tok <= 4u && getenv("DS4_ROCM_SMALLN_SHAREDX") == NULL; + if (smalln_preq) { + const uint64_t xq_bytes = n_tok * blocks * 32u; + const uint64_t scale_off = (xq_bytes + 15u) & ~15ull; + const uint64_t tmp_bytes = scale_off + n_tok * blocks * sizeof(float); + void *tmp = cuda_tmp_alloc(tmp_bytes, "q8_0 smalln prequant"); + if (tmp) { + int8_t *xq = (int8_t *)tmp; + float *xscale = (float *)((char *)tmp + scale_off); + dim3 qgrid((unsigned)blocks, (unsigned)n_tok, 1); + quantize_q8_0_f32_kernel<<>>(xq, xscale, (const float *)x->ptr, in_dim, blocks); + if (cuda_ok(cudaGetLastError(), "matmul_q8_0 smalln quantize launch")) { + const uint32_t rpb = 8u; + matmul_q8_0_preq_rows_w32_ntok_kernel<<< + (unsigned)((out_dim + rpb - 1u) / rpb), + rpb * 32u>>>( + (float *)out->ptr, + reinterpret_cast(wptr), + xq, + xscale, + in_dim, + out_dim, + blocks, + rpb, + (uint32_t)n_tok, + 1); + return cuda_ok(cudaGetLastError(), "matmul_q8_0 smalln ntok launch"); + } + } + /* fall through to the generic paths on alloc/launch failure */ + } if ((in_dim & 31u) == 0u && out_dim <= UINT32_MAX && n_tok <= UINT32_MAX) { const uint32_t rows_per_block = 32u; const uint32_t tile = 32u; diff --git a/rocm/ds4_rocm_moe_launch.cuh b/rocm/ds4_rocm_moe_launch.cuh index 6eafc377b..2d70dec73 100644 --- a/rocm/ds4_rocm_moe_launch.cuh +++ b/rocm/ds4_rocm_moe_launch.cuh @@ -1735,8 +1735,81 @@ extern "C" int ds4_gpu_routed_moe_one_tensor(ds4_gpu_tensor *out, ds4_gpu_tensor expert_in_dim, expert_mid_dim, out_dim, selected, weights, n_total_expert, n_expert, clamp, x, layer_index, 1); } +/* Split a [n_token][n_expert] routed selection into two [n_token][half] halves. + * Used to run Qwen's top-8 route as two top-4 passes on the top-6-specialized + * ROCm MoE core. */ +__global__ static void moe_split_selected_kernel( + int32_t *h0s, int32_t *h1s, float *h0w, float *h1w, + const int32_t *sel, const float *w, + uint32_t n_token, uint32_t n_expert, uint32_t half) { + const uint32_t t = blockIdx.x * blockDim.x + threadIdx.x; + if (t >= n_token) return; + for (uint32_t j = 0; j < half; j++) { + h0s[t * half + j] = sel[t * n_expert + j]; + h0w[t * half + j] = w[t * n_expert + j]; + h1s[t * half + j] = sel[t * n_expert + half + j]; + h1w[t * half + j] = w[t * n_expert + half + j]; + } +} + +/* Qwen top-8 as two top-4 halves + add, entirely on device. gate/up/mid/down + * scratch (sized for n_expert by the caller) is reused across the two 4-expert + * launches; the two partial outputs are separate temporaries. */ +static int routed_moe_batch_two_half( + ds4_gpu_tensor *out, ds4_gpu_tensor *gate, ds4_gpu_tensor *up, + ds4_gpu_tensor *mid, ds4_gpu_tensor *down, const void *model_map, + uint64_t model_size, uint64_t gate_offset, uint64_t up_offset, + uint64_t down_offset, uint32_t gate_type, uint32_t down_type, + uint64_t gate_expert_bytes, uint64_t gate_row_bytes, + uint64_t down_expert_bytes, uint64_t down_row_bytes, + uint32_t expert_in_dim, uint32_t expert_mid_dim, uint32_t out_dim, + const ds4_gpu_tensor *selected, const ds4_gpu_tensor *weights, + uint32_t n_total_expert, uint32_t n_expert, float clamp, + const ds4_gpu_tensor *x, uint32_t layer_index, uint32_t n_tokens) { + const uint32_t half = n_expert / 2u; + if (n_expert != 8u || (n_expert & 1u) != 0u) return 0; + void *buf = NULL; + const size_t sel_bytes = (size_t)n_tokens * half * sizeof(int32_t); + const size_t w_bytes = (size_t)n_tokens * half * sizeof(float); + const size_t part_bytes = (size_t)n_tokens * out_dim * sizeof(float); + const size_t total = 2 * sel_bytes + 2 * w_bytes + 2 * part_bytes; + if (cudaMalloc(&buf, total) != cudaSuccess) return 0; + char *p = (char *)buf; + ds4_gpu_tensor h0s{ p, sel_bytes, 0 }; p += sel_bytes; + ds4_gpu_tensor h1s{ p, sel_bytes, 0 }; p += sel_bytes; + ds4_gpu_tensor h0w{ p, w_bytes, 0 }; p += w_bytes; + ds4_gpu_tensor h1w{ p, w_bytes, 0 }; p += w_bytes; + ds4_gpu_tensor part0{ p, part_bytes, 0 }; p += part_bytes; + ds4_gpu_tensor part1{ p, part_bytes, 0 }; + moe_split_selected_kernel<<<(n_tokens + 255u) / 256u, 256>>>( + (int32_t *)h0s.ptr, (int32_t *)h1s.ptr, (float *)h0w.ptr, (float *)h1w.ptr, + (const int32_t *)selected->ptr, (const float *)weights->ptr, n_tokens, n_expert, half); + int ok = cuda_ok(cudaGetLastError(), "qwen moe split"); + if (ok) ok = routed_moe_launch(&part0, gate, up, mid, down, model_map, model_size, + gate_offset, up_offset, down_offset, gate_type, down_type, + gate_expert_bytes, gate_row_bytes, down_expert_bytes, down_row_bytes, + expert_in_dim, expert_mid_dim, out_dim, &h0s, &h0w, n_total_expert, half, clamp, x, layer_index, n_tokens); + if (ok) ok = routed_moe_launch(&part1, gate, up, mid, down, model_map, model_size, + gate_offset, up_offset, down_offset, gate_type, down_type, + gate_expert_bytes, gate_row_bytes, down_expert_bytes, down_row_bytes, + expert_in_dim, expert_mid_dim, out_dim, &h1s, &h1w, n_total_expert, half, clamp, x, layer_index, n_tokens); + if (ok) ok = ds4_gpu_add_tensor(out, &part0, &part1, n_tokens * out_dim); + cudaDeviceSynchronize(); + cudaFree(buf); + return ok; +} + extern "C" int ds4_gpu_routed_moe_batch_tensor(ds4_gpu_tensor *out, ds4_gpu_tensor *gate, ds4_gpu_tensor *up, ds4_gpu_tensor *mid, ds4_gpu_tensor *down, const void *model_map, uint64_t model_size, uint64_t gate_offset, uint64_t up_offset, uint64_t down_offset, uint32_t gate_type, uint32_t down_type, uint64_t gate_expert_bytes, uint64_t gate_row_bytes, uint64_t down_expert_bytes, uint64_t down_row_bytes, uint32_t expert_in_dim, uint32_t expert_mid_dim, uint32_t out_dim, const ds4_gpu_tensor *selected, const ds4_gpu_tensor *weights, uint32_t n_total_expert, uint32_t n_expert, float clamp, const ds4_gpu_tensor *x, uint32_t layer_index, uint32_t n_tokens, bool *mid_is_f16) { if (mid_is_f16) *mid_is_f16 = false; + /* Qwen routes top-8; the shared core is specialized for DeepSeek's top-6. + * Run 8 as two top-4 halves + add without touching the sum6 kernels. */ + if (n_expert > DS4_ROCM_N_EXPERT_USED) { + return routed_moe_batch_two_half(out, gate, up, mid, down, model_map, model_size, + gate_offset, up_offset, down_offset, gate_type, down_type, + gate_expert_bytes, gate_row_bytes, down_expert_bytes, down_row_bytes, + expert_in_dim, expert_mid_dim, out_dim, selected, weights, + n_total_expert, n_expert, clamp, x, layer_index, n_tokens); + } return routed_moe_launch(out, gate, up, mid, down, model_map, model_size, gate_offset, up_offset, down_offset, gate_type, down_type, diff --git a/rocm/ds4_rocm_q8.cuh b/rocm/ds4_rocm_q8.cuh index 5b2423de3..310db13aa 100644 --- a/rocm/ds4_rocm_q8.cuh +++ b/rocm/ds4_rocm_q8.cuh @@ -381,6 +381,47 @@ __global__ static void matmul_q8_0_preq_batch_warp8_kernel( if (lane == 0) out[tok * out_dim + row] = acc; } +/* + * Small-batch (2-4 token) prequantized matmul: one warp per output row, the + * row's weight blocks are read ONCE and applied to every token's quantized + * activations. The per-token grid kernels re-read the full weight matrix for + * each token, which is what made the MTP verify's attention/Q projections + * n-times more expensive than a single-token decode. + */ +__global__ static void matmul_q8_0_preq_rows_w32_ntok_kernel( + float *out, + const unsigned char *w, + const int8_t *xq, + const float *xscale, + uint64_t in_dim, + uint64_t out_dim, + uint64_t blocks, + uint32_t rows_per_block, + uint32_t n_tok, + int use_dp4a) { + const uint64_t row = (uint64_t)blockIdx.x * rows_per_block + (threadIdx.x >> 5u); + const uint32_t lane = threadIdx.x & 31u; + if (row >= out_dim || n_tok > 4u) return; + const unsigned char *wr = w + row * blocks * 34u; + float acc[4] = {0.0f, 0.0f, 0.0f, 0.0f}; + for (uint64_t b = lane; b < blocks; b += 32u) { + const uint64_t i0 = b * 32u; + const uint64_t bn = in_dim - i0 < 32u ? in_dim - i0 : 32u; + const __half *scale_h = (const __half *)(wr + b * 34u); + const int8_t *qs = (const int8_t *)(wr + b * 34u + 2u); + const float ws = __half2float(*scale_h); + for (uint32_t t = 0; t < n_tok; t++) { + const int8_t *xqb = xq + ((uint64_t)t * blocks + b) * 32u; + const int dot = dot_i8_block(qs, xqb, bn, use_dp4a); + acc[t] += ws * xscale[(uint64_t)t * blocks + b] * (float)dot; + } + } + for (uint32_t t = 0; t < n_tok; t++) { + const float v = warp_sum_f32(acc[t]); + if (lane == 0u) out[(uint64_t)t * out_dim + row] = v; + } +} + __device__ static float q8_0_scale_scalar(const unsigned char *blk) { const uint16_t bits = (uint16_t)blk[0] | ((uint16_t)blk[1] << 8); return __half2float(__ushort_as_half((unsigned short)bits)); diff --git a/rocm/ds4_rocm_qwen.cuh b/rocm/ds4_rocm_qwen.cuh new file mode 100644 index 000000000..2dfaffcd2 --- /dev/null +++ b/rocm/ds4_rocm_qwen.cuh @@ -0,0 +1,1320 @@ +// Qwen3.6-35B-A3B decode/prefill primitives for DS4's ROCm backend. +// +// Direct HIP translation of metal/qwen35.metal. Buffers contain F32 values and +// every stride is expressed in bytes, so the host binds packed rows or larger +// workspaces exactly as the Metal host does. The numeric contract is identical +// to ds4_qwen_ref.c / ds4_qwen.c: arithmetic order is preserved so the ROCm +// kernels reproduce the scalar reference (and the Metal kernels) bit-comparably +// within the tolerances used by tests/qwen. +// +// simdgroup mapping: RDNA3 runs wave32 (via HSA_OVERRIDE_GFX_VERSION=11.0.0), +// so one Metal simdgroup == one 32-lane wavefront. lane == threadIdx.x & 31, +// wave index == threadIdx.x >> 5, threads_per_simdgroup == 32. Metal +// simd_sum/simd_max/simd_min become butterfly (__shfl_xor) reductions that +// broadcast the result to every lane, matching Metal's SIMD-wide semantics. + +#ifndef DS4_ROCM_QWEN_CUH +#define DS4_ROCM_QWEN_CUH + +// ---- args structs (byte-stride layout mirrors metal/qwen35.metal) ---------- + +struct qrocm_split_q_gate { + uint32_t n_token, n_query_head, head_dim, reserved; + uint64_t projection_token_stride, projection_head_stride, projection_dim_stride; + uint64_t query_token_stride, query_head_stride, query_dim_stride; + uint64_t gate_token_stride, gate_head_stride, gate_dim_stride; +}; + +struct qrocm_sigmoid_mul { + uint64_t n_value, input_stride, gate_stride, output_stride; +}; + +struct qrocm_sigmoid_mul_rows { + uint32_t n_row, row_width; + uint64_t input_row_stride, input_dim_stride, gate_row_stride; + uint64_t output_row_stride, output_dim_stride; +}; + +struct qrocm_rope { + uint32_t n_token, n_head, head_dim, n_rot; + float theta; uint32_t reserved; + uint64_t source_token_stride, source_head_stride, source_dim_stride; + uint64_t output_token_stride, output_head_stride, output_dim_stride; + uint64_t position_stride; +}; + +struct qrocm_conv_step { + uint32_t n_channel, kernel_size; + uint64_t input_channel_stride, weight_channel_stride, weight_tap_stride; + uint64_t state_channel_stride, state_tap_stride, output_channel_stride; +}; + +struct qrocm_conv_sequence { + uint32_t n_token, n_channel, kernel_size, reserved; + uint64_t input_token_stride, input_channel_stride; + uint64_t weight_channel_stride, weight_tap_stride; + uint64_t state_channel_stride, state_tap_stride; + uint64_t output_token_stride, output_channel_stride; +}; + +struct qrocm_gated_delta_step { + uint32_t n_key_head, n_value_head, key_dim, value_dim; + uint64_t query_head_stride, query_dim_stride; + uint64_t key_head_stride, key_dim_stride; + uint64_t value_head_stride, value_dim_stride; + uint64_t log_decay_head_stride, beta_head_stride; + uint64_t state_head_stride, state_value_stride, state_key_stride; + uint64_t output_head_stride, output_dim_stride; +}; + +struct qrocm_gated_delta_sequence { + uint32_t n_token, n_key_head, n_value_head, key_dim, value_dim, reserved; + uint64_t projection_token_stride, query_offset, key_offset, value_offset; + uint64_t query_head_stride, query_dim_stride; + uint64_t key_head_stride, key_dim_stride; + uint64_t value_head_stride, value_dim_stride; + uint64_t log_decay_token_stride, log_decay_head_stride; + uint64_t beta_token_stride, beta_head_stride; + uint64_t state_head_stride, state_value_stride, state_key_stride; + uint64_t output_token_stride, output_head_stride, output_dim_stride; +}; + +struct qrocm_rmsnorm_gated { + uint32_t n_vector, dim; float epsilon; uint32_t reserved; + uint64_t input_vector_stride, input_dim_stride; + uint64_t gate_vector_stride, gate_dim_stride, weight_dim_stride; + uint64_t output_vector_stride, output_dim_stride; +}; + +struct qrocm_embedding_q8_0 { + uint32_t row_index, n_embd, block_size, reserved; + uint64_t source_row_stride, source_block_stride; + uint64_t source_scale_offset, source_quant_offset, source_quant_stride; + uint64_t output_dim_stride; +}; + +struct qrocm_embedding_q8_0_batch { + uint32_t n_token, n_row, n_embd, block_size; + uint64_t source_row_stride, source_block_stride; + uint64_t source_scale_offset, source_quant_offset, source_quant_stride; + uint64_t token_id_stride, output_token_stride, output_dim_stride; +}; + +struct qrocm_gated_delta_controls { + uint32_t n_token, n_value_head; + uint64_t alpha_logit_token_stride, alpha_logit_head_stride; + uint64_t beta_logit_token_stride, beta_logit_head_stride; + uint64_t ssm_a_head_stride, dt_bias_head_stride; + uint64_t log_decay_token_stride, log_decay_head_stride; + uint64_t beta_token_stride, beta_head_stride; +}; + +struct qrocm_gqa_decode { + uint32_t n_kv, n_query_head, n_kv_head, head_dim; + uint64_t query_head_stride, query_dim_stride; + uint64_t key_token_stride, key_head_stride, key_dim_stride; + uint64_t value_token_stride, value_head_stride, value_dim_stride; + uint64_t output_head_stride, output_dim_stride; +}; + +struct qrocm_gqa_prefill { + uint32_t position0, n_token, n_query_head, n_kv_head, head_dim, reserved; + uint64_t query_token_stride, query_head_stride, query_dim_stride; + uint64_t key_token_stride, key_head_stride, key_dim_stride; + uint64_t value_token_stride, value_head_stride, value_dim_stride; + uint64_t output_token_stride, output_head_stride, output_dim_stride; +}; + +struct qrocm_router_top8 { + uint32_t n_token, reserved; + uint64_t logits_token_stride, logits_stride; + uint64_t selected_token_stride, selected_stride; + uint64_t selected_weight_token_stride, selected_weight_stride; +}; + +// ---- device helpers -------------------------------------------------------- + +__device__ __forceinline__ float qrocm_sigmoid(float x) { + if (x >= 0.0f) return 1.0f / (1.0f + expf(-x)); + const float e = expf(x); + return e / (1.0f + e); +} +__device__ __forceinline__ float qrocm_silu(float x) { return x * qrocm_sigmoid(x); } +__device__ __forceinline__ float qrocm_softplus(float x) { + if (x > 20.0f) return x; + if (x < -20.0f) return expf(x); + if (x < -10.0f) { + const float e = expf(x); + return e - 0.5f * e * e + (e * e * e) / 3.0f; + } + if (x > 10.0f) return x + logf(1.0f + expf(-x)); + return logf(1.0f + expf(x)); +} +__device__ __forceinline__ float qrocm_ld(const char *b, uint64_t o) { + return *((const float *)(b + o)); +} +__device__ __forceinline__ void qrocm_st(char *b, uint64_t o, float v) { + *((float *)(b + o)) = v; +} +// All-lane (broadcast) 32-wide reductions == Metal simd_sum/max/min. +__device__ __forceinline__ float qrocm_wsum(float v) { +#pragma unroll + for (int o = 16; o > 0; o >>= 1) v += __shfl_xor(v, o, 32); + return v; +} +__device__ __forceinline__ float qrocm_wmax(float v) { +#pragma unroll + for (int o = 16; o > 0; o >>= 1) v = fmaxf(v, __shfl_xor(v, o, 32)); + return v; +} +__device__ __forceinline__ uint32_t qrocm_wmin_u(uint32_t v) { +#pragma unroll + for (int o = 16; o > 0; o >>= 1) { + uint32_t t = (uint32_t)__shfl_xor((int)v, o, 32); + if (t < v) v = t; + } + return v; +} + +// ---- router: stable softmax + deterministic top-8 (serial reference) ------- + +__global__ void qrocm_k_router_serial( + qrocm_router_top8 args, const char *logits, + char *selected, char *selected_weight) { + constexpr uint32_t n_expert = 256u, n_selected = 8u; + __shared__ float probability[256]; + const uint32_t tid = threadIdx.x, token = blockIdx.x; + if (token >= args.n_token || blockDim.x < n_expert || + args.logits_stride < sizeof(float) || + args.selected_stride < sizeof(int32_t) || + args.selected_weight_stride < sizeof(float)) return; + + const uint64_t logits_base = (uint64_t)token * args.logits_token_stride; + const uint64_t selected_base = (uint64_t)token * args.selected_token_stride; + const uint64_t sw_base = (uint64_t)token * args.selected_weight_token_stride; + + __shared__ float ctl[2]; + if (tid == 0u) { + float maximum = qrocm_ld(logits, logits_base); + bool finite = isfinite(maximum); + for (uint32_t e = 1u; e < n_expert; e++) { + const float v = qrocm_ld(logits, logits_base + (uint64_t)e * args.logits_stride); + finite = finite && isfinite(v); + if (v > maximum) maximum = v; + } + ctl[0] = maximum; + ctl[1] = finite ? 1.0f : 0.0f; + } + __syncthreads(); + + const float maximum = ctl[0]; + const bool finite = ctl[1] != 0.0f; + if (!finite) { + if (tid < n_selected) { + *((int32_t *)(selected + selected_base + (uint64_t)tid * args.selected_stride)) = -1; + qrocm_st(selected_weight, sw_base + (uint64_t)tid * args.selected_weight_stride, 0.0f); + } + return; + } + if (tid < n_expert) + probability[tid] = expf(qrocm_ld(logits, logits_base + (uint64_t)tid * args.logits_stride) - maximum); + __syncthreads(); + if (tid != 0u) return; + + float total = 0.0f; + for (uint32_t e = 0u; e < n_expert; e++) total += probability[e]; + for (uint32_t e = 0u; e < n_expert; e++) probability[e] /= total; + + int32_t chosen[n_selected]; + float chosen_weight[n_selected]; + for (uint32_t slot = 0u; slot < n_selected; slot++) { + uint32_t best = n_expert; + for (uint32_t e = 0u; e < n_expert; e++) { + bool used = false; + for (uint32_t p = 0u; p < slot; p++) if (chosen[p] == (int32_t)e) { used = true; break; } + if (used) continue; + if (best == n_expert || probability[e] > probability[best] || + (probability[e] == probability[best] && e < best)) best = e; + } + chosen[slot] = (int32_t)best; + chosen_weight[slot] = probability[best]; + } + float selected_total = 0.0f; + for (uint32_t slot = 0u; slot < n_selected; slot++) selected_total += chosen_weight[slot]; + for (uint32_t slot = 0u; slot < n_selected; slot++) { + *((int32_t *)(selected + selected_base + (uint64_t)slot * args.selected_stride)) = chosen[slot]; + qrocm_st(selected_weight, sw_base + (uint64_t)slot * args.selected_weight_stride, + chosen_weight[slot] / selected_total); + } +} + +// ---- embedding dequant (Q8_0) --------------------------------------------- + +__global__ void qrocm_k_embed_q8_0( + qrocm_embedding_q8_0 args, const char *embedding, char *output) { + const uint32_t dim = blockIdx.x * blockDim.x + threadIdx.x; + if (dim >= args.n_embd || args.block_size != 32u) return; + const uint32_t block = dim / args.block_size; + const uint32_t within = dim - block * args.block_size; + const uint64_t bo = (uint64_t)args.row_index * args.source_row_stride + + (uint64_t)block * args.source_block_stride; + const __half scale = *((const __half *)(embedding + bo + args.source_scale_offset)); + const int8_t q = *((const int8_t *)(embedding + bo + args.source_quant_offset + + (uint64_t)within * args.source_quant_stride)); + qrocm_st(output, (uint64_t)dim * args.output_dim_stride, (float)scale * (float)q); +} + +__global__ void qrocm_k_embed_q8_0_batch( + qrocm_embedding_q8_0_batch args, const char *embedding, + const char *token_ids, char *output) { + const uint64_t gid = (uint64_t)blockIdx.x * blockDim.x + threadIdx.x; + const uint64_t total = (uint64_t)args.n_token * args.n_embd; + if (gid >= total || args.n_embd == 0u || args.block_size != 32u) return; + const uint64_t token = gid / args.n_embd; + const uint32_t dim = (uint32_t)(gid - token * args.n_embd); + const int32_t row = *((const int32_t *)(token_ids + token * args.token_id_stride)); + if (row < 0 || (uint32_t)row >= args.n_row) return; + const uint32_t block = dim / args.block_size; + const uint32_t within = dim - block * args.block_size; + const uint64_t bo = (uint64_t)(uint32_t)row * args.source_row_stride + + (uint64_t)block * args.source_block_stride; + const __half scale = *((const __half *)(embedding + bo + args.source_scale_offset)); + const int8_t q = *((const int8_t *)(embedding + bo + args.source_quant_offset + + (uint64_t)within * args.source_quant_stride)); + qrocm_st(output, token * args.output_token_stride + (uint64_t)dim * args.output_dim_stride, + (float)scale * (float)q); +} + +// ---- gated-delta controls -------------------------------------------------- + +__global__ void qrocm_k_gd_controls( + qrocm_gated_delta_controls args, const char *alpha_logit, + const char *beta_logit, const char *ssm_a, const char *dt_bias, + char *log_decay, char *beta) { + const uint64_t total = (uint64_t)args.n_token * args.n_value_head; + const uint64_t gid = (uint64_t)blockIdx.x * blockDim.x + threadIdx.x; + if (gid >= total || args.n_value_head == 0u) return; + const uint32_t token = (uint32_t)(gid / args.n_value_head); + const uint32_t head = (uint32_t)(gid - (uint64_t)token * args.n_value_head); + const float alpha = qrocm_ld(alpha_logit, + (uint64_t)token * args.alpha_logit_token_stride + (uint64_t)head * args.alpha_logit_head_stride); + const float beta_v = qrocm_sigmoid(qrocm_ld(beta_logit, + (uint64_t)token * args.beta_logit_token_stride + (uint64_t)head * args.beta_logit_head_stride)); + const float a = qrocm_ld(ssm_a, (uint64_t)head * args.ssm_a_head_stride); + const float bias = qrocm_ld(dt_bias, (uint64_t)head * args.dt_bias_head_stride); + qrocm_st(log_decay, + (uint64_t)token * args.log_decay_token_stride + (uint64_t)head * args.log_decay_head_stride, + a * qrocm_softplus(alpha + bias)); + qrocm_st(beta, + (uint64_t)token * args.beta_token_stride + (uint64_t)head * args.beta_head_stride, beta_v); +} + +// ---- split q/gate ---------------------------------------------------------- + +__global__ void qrocm_k_split_q_gate( + qrocm_split_q_gate args, const char *projection, char *query, char *gate) { + const uint64_t head_values = (uint64_t)args.n_query_head * args.head_dim; + const uint64_t total = (uint64_t)args.n_token * head_values; + const uint64_t gid = (uint64_t)blockIdx.x * blockDim.x + threadIdx.x; + if (gid >= total || args.head_dim == 0 || args.n_query_head == 0) return; + const uint64_t token = gid / head_values; + const uint64_t wt = gid - token * head_values; + const uint64_t head = wt / args.head_dim; + const uint64_t dim = wt - head * args.head_dim; + const uint64_t pbase = token * args.projection_token_stride + head * args.projection_head_stride; + const float q = qrocm_ld(projection, pbase + dim * args.projection_dim_stride); + const float g = qrocm_ld(projection, pbase + ((uint64_t)args.head_dim + dim) * args.projection_dim_stride); + qrocm_st(query, token * args.query_token_stride + head * args.query_head_stride + dim * args.query_dim_stride, q); + qrocm_st(gate, token * args.gate_token_stride + head * args.gate_head_stride + dim * args.gate_dim_stride, g); +} + +// ---- sigmoid-mul (elementwise + row-wise) ---------------------------------- + +__global__ void qrocm_k_sigmoid_mul( + qrocm_sigmoid_mul args, const char *input, const char *gate_logit, char *output) { + const uint64_t gid = (uint64_t)blockIdx.x * blockDim.x + threadIdx.x; + if (gid >= args.n_value) return; + const float x = qrocm_ld(input, gid * args.input_stride); + const float z = qrocm_ld(gate_logit, gid * args.gate_stride); + qrocm_st(output, gid * args.output_stride, x * qrocm_sigmoid(z)); +} + +__global__ void qrocm_k_sigmoid_mul_rows( + qrocm_sigmoid_mul_rows args, const char *input, const char *gate_logit, char *output) { + const uint64_t total = (uint64_t)args.n_row * args.row_width; + const uint64_t gid = (uint64_t)blockIdx.x * blockDim.x + threadIdx.x; + if (gid >= total || args.row_width == 0u) return; + const uint64_t row = gid / args.row_width; + const uint64_t dim = gid - row * args.row_width; + const float x = qrocm_ld(input, row * args.input_row_stride + dim * args.input_dim_stride); + const float z = qrocm_ld(gate_logit, row * args.gate_row_stride); + qrocm_st(output, row * args.output_row_stride + dim * args.output_dim_stride, x * qrocm_sigmoid(z)); +} + +// ---- RoPE (split-half NeoX over first n_rot dims) -------------------------- + +__global__ void qrocm_k_rope( + qrocm_rope args, const char *source, const uint8_t *position, char *output) { + const uint64_t head_values = (uint64_t)args.n_head * args.head_dim; + const uint64_t total = (uint64_t)args.n_token * head_values; + const uint64_t gid = (uint64_t)blockIdx.x * blockDim.x + threadIdx.x; + if (gid >= total || args.n_head == 0 || args.head_dim == 0 || + args.n_rot == 0 || args.n_rot > args.head_dim || + (args.n_rot & 1u) != 0 || !(args.theta > 0.0f)) return; + const uint64_t token = gid / head_values; + const uint64_t wt = gid - token * head_values; + const uint64_t head = wt / args.head_dim; + const uint64_t dim = wt - head * args.head_dim; + const uint64_t sbase = token * args.source_token_stride + head * args.source_head_stride; + const uint64_t obase = token * args.output_token_stride + head * args.output_head_stride; + if (dim >= (uint64_t)args.n_rot) { + qrocm_st(output, obase + dim * args.output_dim_stride, + qrocm_ld(source, sbase + dim * args.source_dim_stride)); + return; + } + const uint64_t half_rot = (uint64_t)args.n_rot / 2u; + if (dim >= half_rot) return; + const float exponent = (2.0f * (float)dim) / (float)args.n_rot; + const uint32_t pos = *((const uint32_t *)(position + token * args.position_stride)); + const float angle = (float)pos / powf(args.theta, exponent); + const float c = cosf(angle), s = sinf(angle); + const float a = qrocm_ld(source, sbase + dim * args.source_dim_stride); + const float b = qrocm_ld(source, sbase + (dim + half_rot) * args.source_dim_stride); + qrocm_st(output, obase + dim * args.output_dim_stride, a * c - b * s); + qrocm_st(output, obase + (dim + half_rot) * args.output_dim_stride, b * c + a * s); +} + +// ---- causal depthwise conv1d (k=4) + SiLU ---------------------------------- + +__global__ void qrocm_k_conv_step( + qrocm_conv_step args, const char *input, const char *weight, + char *state, char *output) { + const uint64_t channel = (uint64_t)blockIdx.x * blockDim.x + threadIdx.x; + if (channel >= (uint64_t)args.n_channel || args.kernel_size < 2u) return; + const uint64_t sbase = channel * args.state_channel_stride; + const uint64_t wbase = channel * args.weight_channel_stride; + const float current = qrocm_ld(input, channel * args.input_channel_stride); + float total = current * qrocm_ld(weight, wbase + ((uint64_t)args.kernel_size - 1u) * args.weight_tap_stride); + const uint64_t hist = (uint64_t)args.kernel_size - 1u; + for (uint64_t tap = 0; tap < hist; tap++) + total += qrocm_ld(state, sbase + tap * args.state_tap_stride) * + qrocm_ld(weight, wbase + tap * args.weight_tap_stride); + qrocm_st(output, channel * args.output_channel_stride, qrocm_silu(total)); + for (uint64_t tap = 0; tap + 1u < hist; tap++) + qrocm_st(state, sbase + tap * args.state_tap_stride, + qrocm_ld(state, sbase + (tap + 1u) * args.state_tap_stride)); + qrocm_st(state, sbase + (hist - 1u) * args.state_tap_stride, current); +} + +__global__ void qrocm_k_conv_sequence( + qrocm_conv_sequence args, const char *input, const char *weight, + char *state, char *output) { + const uint32_t channel = blockIdx.x * blockDim.x + threadIdx.x; + if (channel >= args.n_channel || args.n_token == 0u || args.kernel_size != 4u) return; + const uint64_t sbase = (uint64_t)channel * args.state_channel_stride; + const uint64_t wbase = (uint64_t)channel * args.weight_channel_stride; + float history[3]; + for (uint32_t tap = 0u; tap < 3u; tap++) + history[tap] = qrocm_ld(state, sbase + (uint64_t)tap * args.state_tap_stride); + for (uint32_t token = 0u; token < args.n_token; token++) { + const uint64_t io = (uint64_t)token * args.input_token_stride + + (uint64_t)channel * args.input_channel_stride; + const float current = qrocm_ld(input, io); + float total = current * qrocm_ld(weight, wbase + 3u * args.weight_tap_stride); + for (uint32_t tap = 0u; tap < 3u; tap++) + total += history[tap] * qrocm_ld(weight, wbase + (uint64_t)tap * args.weight_tap_stride); + qrocm_st(output, (uint64_t)token * args.output_token_stride + + (uint64_t)channel * args.output_channel_stride, qrocm_silu(total)); + history[0] = history[1]; history[1] = history[2]; history[2] = current; + } + for (uint32_t tap = 0u; tap < 3u; tap++) + qrocm_st(state, sbase + (uint64_t)tap * args.state_tap_stride, history[tap]); +} + +// ---- gated DeltaNet recurrent step (generic) ------------------------------- +// One block owns one value head for one decode token. blockDim.x is a multiple +// of 32; the block cooperatively reduces query/key norms then each lane strides +// over value rows. scratch holds 2*n_wave floats. +__global__ void qrocm_k_gd_step( + qrocm_gated_delta_step args, const char *query, const char *key, + const char *value, const char *log_decay, const char *beta, + char *state, char *output) { + extern __shared__ float s_gd[]; + const uint32_t value_head = blockIdx.x, tid = threadIdx.x; + const uint32_t lane = tid & 31u, wave = tid >> 5; + const uint32_t n_wave = (blockDim.x + 31u) / 32u; + if (value_head >= args.n_value_head || args.n_key_head == 0 || + args.key_dim == 0 || args.value_dim == 0 || + args.n_value_head % args.n_key_head != 0) return; + float *query_partial = s_gd; + float *key_partial = s_gd + n_wave; + const uint32_t key_head = value_head % args.n_key_head; + const uint64_t query_base = (uint64_t)key_head * args.query_head_stride; + const uint64_t key_base = (uint64_t)key_head * args.key_head_stride; + + float qsq = 0.0f, ksq = 0.0f; + for (uint32_t dim = tid; dim < args.key_dim; dim += blockDim.x) { + const float q = qrocm_ld(query, query_base + (uint64_t)dim * args.query_dim_stride); + const float k = qrocm_ld(key, key_base + (uint64_t)dim * args.key_dim_stride); + qsq += q * q; ksq += k * k; + } + qsq = qrocm_wsum(qsq); ksq = qrocm_wsum(ksq); + if (lane == 0u) { query_partial[wave] = qsq; key_partial[wave] = ksq; } + __syncthreads(); + if (wave == 0u) { + float q = lane < n_wave ? query_partial[lane] : 0.0f; + float k = lane < n_wave ? key_partial[lane] : 0.0f; + q = qrocm_wsum(q); k = qrocm_wsum(k); + if (lane == 0u) { + query_partial[0] = (1.0f / sqrtf((float)args.key_dim)) / sqrtf(q + 1.0e-6f); + key_partial[0] = 1.0f / sqrtf(k + 1.0e-6f); + } + } + __syncthreads(); + const float query_inverse = query_partial[0]; + const float key_inverse = key_partial[0]; + const float decay = expf(qrocm_ld(log_decay, (uint64_t)value_head * args.log_decay_head_stride)); + const float stepv = qrocm_ld(beta, (uint64_t)value_head * args.beta_head_stride); + const uint64_t value_base = (uint64_t)value_head * args.value_head_stride; + const uint64_t state_head_base = (uint64_t)value_head * args.state_head_stride; + const uint64_t output_base = (uint64_t)value_head * args.output_head_stride; + + for (uint32_t vd = tid; vd < args.value_dim; vd += blockDim.x) { + const uint64_t state_row = state_head_base + (uint64_t)vd * args.state_value_stride; + float memory = 0.0f; + for (uint32_t kd = 0; kd < args.key_dim; kd++) { + const uint64_t so = state_row + (uint64_t)kd * args.state_key_stride; + const float decayed = qrocm_ld(state, so) * decay; + qrocm_st(state, so, decayed); + const float nk = qrocm_ld(key, key_base + (uint64_t)kd * args.key_dim_stride) * key_inverse; + memory += decayed * nk; + } + const float target = qrocm_ld(value, value_base + (uint64_t)vd * args.value_dim_stride); + const float delta = (target - memory) * stepv; + float result = 0.0f; + for (uint32_t kd = 0; kd < args.key_dim; kd++) { + const uint64_t so = state_row + (uint64_t)kd * args.state_key_stride; + const float nk = qrocm_ld(key, key_base + (uint64_t)kd * args.key_dim_stride) * key_inverse; + const float updated = qrocm_ld(state, so) + nk * delta; + qrocm_st(state, so, updated); + const float nq = qrocm_ld(query, query_base + (uint64_t)kd * args.query_dim_stride) * query_inverse; + result += updated * nq; + } + qrocm_st(output, output_base + (uint64_t)vd * args.output_dim_stride, result); + } +} + +// gated DeltaNet decode specialized for key_dim=128: one wave per value row, +// four state cells per lane in registers, 32x4 block advances four rows. +// grid = (ceil(value_dim/4), n_value_head); block = (32, 4). +__global__ void qrocm_k_gd_step_128( + qrocm_gated_delta_step args, const char *query, const char *key, + const char *value, const char *log_decay, const char *beta, + char *state, char *output) { + constexpr uint32_t dims_per_lane = 4u, rows_per_group = 4u; + __shared__ float norm[2]; + const uint32_t value_head = blockIdx.y, row_in_group = threadIdx.y; + const uint32_t lane = threadIdx.x; + if (value_head >= args.n_value_head || args.n_key_head == 0u || + args.key_dim != 128u || args.value_dim == 0u || + args.n_value_head % args.n_key_head != 0u) return; + const uint32_t key_head = value_head % args.n_key_head; + const uint64_t query_base = (uint64_t)key_head * args.query_head_stride; + const uint64_t key_base = (uint64_t)key_head * args.key_head_stride; + if (row_in_group == 0u) { + float qsq = 0.0f, ksq = 0.0f; + for (uint32_t j = 0u; j < dims_per_lane; j++) { + const uint32_t dim = lane * dims_per_lane + j; + const float q = qrocm_ld(query, query_base + (uint64_t)dim * args.query_dim_stride); + const float k = qrocm_ld(key, key_base + (uint64_t)dim * args.key_dim_stride); + qsq += q * q; ksq += k * k; + } + qsq = qrocm_wsum(qsq); ksq = qrocm_wsum(ksq); + if (lane == 0u) { + norm[0] = (1.0f / sqrtf((float)args.key_dim)) / sqrtf(qsq + 1.0e-6f); + norm[1] = 1.0f / sqrtf(ksq + 1.0e-6f); + } + } + __syncthreads(); + const uint32_t value_dim = blockIdx.x * rows_per_group + row_in_group; + if (value_dim >= args.value_dim) return; + const float query_inverse = norm[0], key_inverse = norm[1]; + const float decay = expf(qrocm_ld(log_decay, (uint64_t)value_head * args.log_decay_head_stride)); + const float stepv = qrocm_ld(beta, (uint64_t)value_head * args.beta_head_stride); + const uint64_t state_row = (uint64_t)value_head * args.state_head_stride + + (uint64_t)value_dim * args.state_value_stride; + float sv[dims_per_lane], nk[dims_per_lane], nq[dims_per_lane]; + float memory_partial = 0.0f; + for (uint32_t j = 0u; j < dims_per_lane; j++) { + const uint32_t dim = lane * dims_per_lane + j; + const uint64_t so = state_row + (uint64_t)dim * args.state_key_stride; + nk[j] = qrocm_ld(key, key_base + (uint64_t)dim * args.key_dim_stride) * key_inverse; + nq[j] = qrocm_ld(query, query_base + (uint64_t)dim * args.query_dim_stride) * query_inverse; + sv[j] = qrocm_ld(state, so) * decay; + memory_partial += sv[j] * nk[j]; + } + const float memory = qrocm_wsum(memory_partial); + const uint64_t value_offset = (uint64_t)value_head * args.value_head_stride + + (uint64_t)value_dim * args.value_dim_stride; + const float target = qrocm_ld(value, value_offset); + const float delta = (target - memory) * stepv; + float result_partial = 0.0f; + for (uint32_t j = 0u; j < dims_per_lane; j++) { + const uint32_t dim = lane * dims_per_lane + j; + const uint64_t so = state_row + (uint64_t)dim * args.state_key_stride; + sv[j] += nk[j] * delta; + qrocm_st(state, so, sv[j]); + result_partial += sv[j] * nq[j]; + } + const float result = qrocm_wsum(result_partial); + if (lane == 0u) { + qrocm_st(output, (uint64_t)value_head * args.output_head_stride + + (uint64_t)value_dim * args.output_dim_stride, result); + } +} + +// gated DeltaNet prefill, key_dim=value_dim=128; one wave per value row keeps +// its 128 state cells in registers across the whole token chunk (serial tokens, +// parallel rows/heads). grid = (ceil(value_dim/4), n_value_head); block=(32,4). +__global__ void qrocm_k_gd_sequence_128( + qrocm_gated_delta_sequence args, const char *projection, + const char *log_decay, const char *beta, char *state, char *output) { + constexpr uint32_t dims_per_lane = 4u, rows_per_group = 4u; + __shared__ float norm[2]; + const uint32_t value_head = blockIdx.y, row_in_group = threadIdx.y; + const uint32_t lane = threadIdx.x; + if (value_head >= args.n_value_head || args.n_token == 0u || + args.n_key_head == 0u || args.key_dim != 128u || + args.value_dim != 128u || args.n_value_head % args.n_key_head != 0u) return; + const uint32_t value_dim = blockIdx.x * rows_per_group + row_in_group; + const bool active_row = value_dim < args.value_dim; + const uint32_t key_head = value_head % args.n_key_head; + const uint64_t state_row = (uint64_t)value_head * args.state_head_stride + + (uint64_t)value_dim * args.state_value_stride; + float sv[dims_per_lane] = { 0.0f, 0.0f, 0.0f, 0.0f }; + if (active_row) { + for (uint32_t j = 0u; j < dims_per_lane; j++) { + const uint32_t dim = lane * dims_per_lane + j; + sv[j] = qrocm_ld(state, state_row + (uint64_t)dim * args.state_key_stride); + } + } + for (uint32_t token = 0u; token < args.n_token; token++) { + const uint64_t pbase = (uint64_t)token * args.projection_token_stride; + const uint64_t query_base = pbase + args.query_offset + (uint64_t)key_head * args.query_head_stride; + const uint64_t key_base = pbase + args.key_offset + (uint64_t)key_head * args.key_head_stride; + if (row_in_group == 0u) { + float qsq = 0.0f, ksq = 0.0f; + for (uint32_t j = 0u; j < dims_per_lane; j++) { + const uint32_t dim = lane * dims_per_lane + j; + const float q = qrocm_ld(projection, query_base + (uint64_t)dim * args.query_dim_stride); + const float k = qrocm_ld(projection, key_base + (uint64_t)dim * args.key_dim_stride); + qsq += q * q; ksq += k * k; + } + qsq = qrocm_wsum(qsq); ksq = qrocm_wsum(ksq); + if (lane == 0u) { + norm[0] = (1.0f / sqrtf((float)args.key_dim)) / sqrtf(qsq + 1.0e-6f); + norm[1] = 1.0f / sqrtf(ksq + 1.0e-6f); + } + } + __syncthreads(); + if (active_row) { + const float query_inverse = norm[0], key_inverse = norm[1]; + const float decay = expf(qrocm_ld(log_decay, + (uint64_t)token * args.log_decay_token_stride + + (uint64_t)value_head * args.log_decay_head_stride)); + const float stepv = qrocm_ld(beta, + (uint64_t)token * args.beta_token_stride + + (uint64_t)value_head * args.beta_head_stride); + float nk[dims_per_lane], nq[dims_per_lane]; + float memory_partial = 0.0f; + for (uint32_t j = 0u; j < dims_per_lane; j++) { + const uint32_t dim = lane * dims_per_lane + j; + nk[j] = qrocm_ld(projection, key_base + (uint64_t)dim * args.key_dim_stride) * key_inverse; + nq[j] = qrocm_ld(projection, query_base + (uint64_t)dim * args.query_dim_stride) * query_inverse; + sv[j] *= decay; + memory_partial += sv[j] * nk[j]; + } + const float memory = qrocm_wsum(memory_partial); + const float target = qrocm_ld(projection, pbase + args.value_offset + + (uint64_t)value_head * args.value_head_stride + + (uint64_t)value_dim * args.value_dim_stride); + const float delta = (target - memory) * stepv; + float result_partial = 0.0f; + for (uint32_t j = 0u; j < dims_per_lane; j++) { + sv[j] += nk[j] * delta; + result_partial += sv[j] * nq[j]; + } + const float result = qrocm_wsum(result_partial); + if (lane == 0u) { + qrocm_st(output, (uint64_t)token * args.output_token_stride + + (uint64_t)value_head * args.output_head_stride + + (uint64_t)value_dim * args.output_dim_stride, result); + } + } + __syncthreads(); + } + if (active_row) { + for (uint32_t j = 0u; j < dims_per_lane; j++) { + const uint32_t dim = lane * dims_per_lane + j; + qrocm_st(state, state_row + (uint64_t)dim * args.state_key_stride, sv[j]); + } + } +} + +// ---- GQA decode / prefill (online softmax) --------------------------------- +// One block owns one query head; blockDim.x >= head_dim (multiple of 32). +// scratch = n_wave + 4 floats. +__global__ void qrocm_k_gqa_decode( + qrocm_gqa_decode args, const char *query, const char *key_cache, + const char *value_cache, char *output) { + extern __shared__ float s_gqa[]; + const uint32_t query_head = blockIdx.x, tid = threadIdx.x; + const uint32_t lane = tid & 31u, wave = tid >> 5; + const uint32_t n_wave = (blockDim.x + 31u) / 32u; + if (query_head >= args.n_query_head || args.n_kv == 0u || + args.n_kv_head == 0u || args.head_dim == 0u || + args.n_query_head % args.n_kv_head != 0u || blockDim.x < args.head_dim) return; + if (n_wave > 32u) return; + const uint32_t control = n_wave; + const uint32_t query_per_kv = args.n_query_head / args.n_kv_head; + const uint32_t kv_head = query_head / query_per_kv; + const uint64_t query_base = (uint64_t)query_head * args.query_head_stride; + const uint64_t output_base = (uint64_t)query_head * args.output_head_stride; + const float scale = 1.0f / sqrtf((float)args.head_dim); + float accumulator = 0.0f; + if (tid == 0u) { + s_gqa[control + 0u] = -INFINITY; s_gqa[control + 1u] = 0.0f; + s_gqa[control + 2u] = 0.0f; s_gqa[control + 3u] = 0.0f; + } + __syncthreads(); + for (uint32_t token = 0; token < args.n_kv; token++) { + const uint64_t key_base = (uint64_t)token * args.key_token_stride + + (uint64_t)kv_head * args.key_head_stride; + float dot = 0.0f; + for (uint32_t dim = tid; dim < args.head_dim; dim += blockDim.x) + dot += qrocm_ld(query, query_base + (uint64_t)dim * args.query_dim_stride) * + qrocm_ld(key_cache, key_base + (uint64_t)dim * args.key_dim_stride); + dot = qrocm_wsum(dot); + if (lane == 0u) s_gqa[wave] = dot; + __syncthreads(); + if (wave == 0u) { + float d = lane < n_wave ? s_gqa[lane] : 0.0f; + d = qrocm_wsum(d); + if (lane == 0u) { + const float score = d * scale; + const float pm = s_gqa[control + 0u]; + const float nm = fmaxf(pm, score); + const float pf = expf(pm - nm), cf = expf(score - nm); + s_gqa[control + 0u] = nm; + s_gqa[control + 1u] = s_gqa[control + 1u] * pf + cf; + s_gqa[control + 2u] = pf; s_gqa[control + 3u] = cf; + } + } + __syncthreads(); + if (tid < args.head_dim) { + const uint64_t vo = (uint64_t)token * args.value_token_stride + + (uint64_t)kv_head * args.value_head_stride + + (uint64_t)tid * args.value_dim_stride; + accumulator = accumulator * s_gqa[control + 2u] + + qrocm_ld(value_cache, vo) * s_gqa[control + 3u]; + } + __syncthreads(); + } + if (tid < args.head_dim) + qrocm_st(output, output_base + (uint64_t)tid * args.output_dim_stride, + accumulator / s_gqa[control + 1u]); +} + +__global__ void qrocm_k_gqa_prefill( + qrocm_gqa_prefill args, const char *query, const char *key_cache, + const char *value_cache, char *output) { + extern __shared__ float s_gqp[]; + const uint32_t query_head = blockIdx.x, query_token = blockIdx.y, tid = threadIdx.x; + const uint32_t lane = tid & 31u, wave = tid >> 5; + const uint32_t n_wave = (blockDim.x + 31u) / 32u; + if (query_head >= args.n_query_head || query_token >= args.n_token || + args.n_kv_head == 0u || args.head_dim == 0u || + args.n_query_head % args.n_kv_head != 0u || blockDim.x < args.head_dim) return; + if (n_wave > 32u) return; + const uint32_t control = n_wave; + const uint32_t query_per_kv = args.n_query_head / args.n_kv_head; + const uint32_t kv_head = query_head / query_per_kv; + const uint32_t n_kv = args.position0 + query_token + 1u; + const uint64_t query_base = (uint64_t)query_token * args.query_token_stride + + (uint64_t)query_head * args.query_head_stride; + const uint64_t output_base = (uint64_t)query_token * args.output_token_stride + + (uint64_t)query_head * args.output_head_stride; + const float scale = 1.0f / sqrtf((float)args.head_dim); + float accumulator = 0.0f; + if (tid == 0u) { + s_gqp[control + 0u] = -INFINITY; s_gqp[control + 1u] = 0.0f; + s_gqp[control + 2u] = 0.0f; s_gqp[control + 3u] = 0.0f; + } + __syncthreads(); + for (uint32_t token = 0u; token < n_kv; token++) { + const uint64_t key_base = (uint64_t)token * args.key_token_stride + + (uint64_t)kv_head * args.key_head_stride; + float dot = 0.0f; + for (uint32_t dim = tid; dim < args.head_dim; dim += blockDim.x) + dot += qrocm_ld(query, query_base + (uint64_t)dim * args.query_dim_stride) * + qrocm_ld(key_cache, key_base + (uint64_t)dim * args.key_dim_stride); + dot = qrocm_wsum(dot); + if (lane == 0u) s_gqp[wave] = dot; + __syncthreads(); + if (wave == 0u) { + float d = lane < n_wave ? s_gqp[lane] : 0.0f; + d = qrocm_wsum(d); + if (lane == 0u) { + const float score = d * scale; + const float pm = s_gqp[control + 0u]; + const float nm = fmaxf(pm, score); + const float pf = expf(pm - nm), cf = expf(score - nm); + s_gqp[control + 0u] = nm; + s_gqp[control + 1u] = s_gqp[control + 1u] * pf + cf; + s_gqp[control + 2u] = pf; s_gqp[control + 3u] = cf; + } + } + __syncthreads(); + if (tid < args.head_dim) { + const uint64_t vo = (uint64_t)token * args.value_token_stride + + (uint64_t)kv_head * args.value_head_stride + + (uint64_t)tid * args.value_dim_stride; + accumulator = accumulator * s_gqp[control + 2u] + + qrocm_ld(value_cache, vo) * s_gqp[control + 3u]; + } + __syncthreads(); + } + if (tid < args.head_dim) + qrocm_st(output, output_base + (uint64_t)tid * args.output_dim_stride, + accumulator / s_gqp[control + 1u]); +} + +// ---- gated RMSNorm --------------------------------------------------------- +// One block per row; scratch = n_wave floats. +__global__ void qrocm_k_rmsnorm_gated( + qrocm_rmsnorm_gated args, const char *input, const char *gate, + const char *weight, char *output) { + extern __shared__ float s_rms[]; + const uint32_t row = blockIdx.x, tid = threadIdx.x; + const uint32_t lane = tid & 31u, wave = tid >> 5; + const uint32_t n_wave = (blockDim.x + 31u) / 32u; + if (row >= args.n_vector || args.dim == 0 || !(args.epsilon > 0.0f)) return; + const uint64_t input_base = (uint64_t)row * args.input_vector_stride; + float square = 0.0f; + for (uint32_t dim = tid; dim < args.dim; dim += blockDim.x) { + const float x = qrocm_ld(input, input_base + (uint64_t)dim * args.input_dim_stride); + square += x * x; + } + square = qrocm_wsum(square); + if (lane == 0u) s_rms[wave] = square; + __syncthreads(); + if (wave == 0u) { + float total = lane < n_wave ? s_rms[lane] : 0.0f; + total = qrocm_wsum(total); + if (lane == 0u) s_rms[0] = 1.0f / sqrtf(total / (float)args.dim + args.epsilon); + } + __syncthreads(); + const float inverse_rms = s_rms[0]; + const uint64_t gate_base = (uint64_t)row * args.gate_vector_stride; + const uint64_t output_base = (uint64_t)row * args.output_vector_stride; + for (uint32_t dim = tid; dim < args.dim; dim += blockDim.x) { + const float x = qrocm_ld(input, input_base + (uint64_t)dim * args.input_dim_stride); + const float z = qrocm_ld(gate, gate_base + (uint64_t)dim * args.gate_dim_stride); + const float w = qrocm_ld(weight, (uint64_t)dim * args.weight_dim_stride); + qrocm_st(output, output_base + (uint64_t)dim * args.output_dim_stride, + x * inverse_rms * w * qrocm_silu(z)); + } +} + +// =========================================================================== +// Host entry points. Stride expressions mirror ds4_metal.m exactly (the +// porting contract); geometry maps Metal threadgroups/simdgroups to HIP +// blocks/waves. Resident-only: weights are resolved through the model range +// registry (the whole payload is registered in resident mode). +// =========================================================================== + +static uint32_t qrocm_elem_threads(uint64_t count) { + uint32_t n = count < 256u ? (uint32_t)count : 256u; + return n == 0u ? 1u : n; +} +// clamp(preferred,32,1024) floored to a multiple of 32; *nsg = nth/32. +static uint32_t qrocm_reduction_threads(uint32_t preferred, uint32_t *nsg) { + uint32_t nth = preferred < 32u ? 32u : (preferred > 1024u ? 1024u : preferred); + nth -= nth % 32u; + if (nth == 0u) nth = 32u; + *nsg = nth / 32u; + return nth; +} + +extern "C" int ds4_gpu_qwen35_split_q_gate_batch_tensor( + ds4_gpu_tensor *query, ds4_gpu_tensor *gate, const ds4_gpu_tensor *projection, + uint32_t n_token, uint32_t n_query_head, uint32_t head_dim) { + if (!query || !gate || !projection || !n_token || !n_query_head || !head_dim) return 0; + const uint64_t row_values = (uint64_t)n_query_head * head_dim; + const uint64_t values = row_values * n_token; + if (values > UINT32_MAX) return 0; + const uint64_t out_bytes = values * 4u; + if (!cuda_tensor_has_bytes(projection, out_bytes * 2u) || + !cuda_tensor_has_bytes(query, out_bytes) || !cuda_tensor_has_bytes(gate, out_bytes)) return 0; + qrocm_split_q_gate a{}; + a.n_token = n_token; a.n_query_head = n_query_head; a.head_dim = head_dim; + a.projection_token_stride = 2u * row_values * 4u; + a.projection_head_stride = 2ull * head_dim * 4u; + a.projection_dim_stride = 4u; + a.query_token_stride = row_values * 4u; a.query_head_stride = (uint64_t)head_dim * 4u; a.query_dim_stride = 4u; + a.gate_token_stride = row_values * 4u; a.gate_head_stride = (uint64_t)head_dim * 4u; a.gate_dim_stride = 4u; + const uint32_t nth = qrocm_elem_threads(values); + qrocm_k_split_q_gate<<<(unsigned)((values + nth - 1) / nth), nth>>>( + a, (const char *)projection->ptr, (char *)query->ptr, (char *)gate->ptr); + return cuda_ok(cudaGetLastError(), "qwen35 split_q_gate"); +} +extern "C" int ds4_gpu_qwen35_split_q_gate_tensor( + ds4_gpu_tensor *query, ds4_gpu_tensor *gate, const ds4_gpu_tensor *projection, + uint32_t n_query_head, uint32_t head_dim) { + return ds4_gpu_qwen35_split_q_gate_batch_tensor(query, gate, projection, 1u, n_query_head, head_dim); +} + +extern "C" int ds4_gpu_qwen35_sigmoid_mul_tensor( + ds4_gpu_tensor *out, const ds4_gpu_tensor *input, const ds4_gpu_tensor *gate, + uint32_t n_value, bool broadcast_gate) { + if (!out || !input || !gate || !n_value) return 0; + const uint64_t bytes = (uint64_t)n_value * 4u; + if (!cuda_tensor_has_bytes(input, bytes) || !cuda_tensor_has_bytes(out, bytes) || + !cuda_tensor_has_bytes(gate, broadcast_gate ? 4u : bytes)) return 0; + qrocm_sigmoid_mul a{}; + a.n_value = n_value; a.input_stride = 4u; a.gate_stride = broadcast_gate ? 0u : 4u; a.output_stride = 4u; + const uint32_t nth = qrocm_elem_threads(n_value); + qrocm_k_sigmoid_mul<<<(unsigned)(((uint64_t)n_value + nth - 1) / nth), nth>>>( + a, (const char *)input->ptr, (const char *)gate->ptr, (char *)out->ptr); + return cuda_ok(cudaGetLastError(), "qwen35 sigmoid_mul"); +} + +extern "C" int ds4_gpu_qwen35_sigmoid_mul_rows_tensor( + ds4_gpu_tensor *out, const ds4_gpu_tensor *input, const ds4_gpu_tensor *gate, + uint32_t n_row, uint32_t row_width) { + if (!out || !input || !gate || !n_row || !row_width) return 0; + const uint64_t values = (uint64_t)n_row * row_width; + if (values > UINT32_MAX) return 0; + const uint64_t bytes = values * 4u; + if (!cuda_tensor_has_bytes(input, bytes) || !cuda_tensor_has_bytes(out, bytes) || + !cuda_tensor_has_bytes(gate, (uint64_t)n_row * 4u)) return 0; + qrocm_sigmoid_mul_rows a{}; + a.n_row = n_row; a.row_width = row_width; + a.input_row_stride = (uint64_t)row_width * 4u; a.input_dim_stride = 4u; a.gate_row_stride = 4u; + a.output_row_stride = (uint64_t)row_width * 4u; a.output_dim_stride = 4u; + const uint32_t nth = qrocm_elem_threads(values); + qrocm_k_sigmoid_mul_rows<<<(unsigned)((values + nth - 1) / nth), nth>>>( + a, (const char *)input->ptr, (const char *)gate->ptr, (char *)out->ptr); + return cuda_ok(cudaGetLastError(), "qwen35 sigmoid_mul_rows"); +} + +static int qrocm_rope_impl( + ds4_gpu_tensor *values, const ds4_gpu_tensor *positions, const uint32_t *single_position, + uint32_t n_token, uint32_t n_head, uint32_t head_dim, uint32_t n_rot, float theta) { + if (!values || (positions == NULL) == (single_position == NULL)) return 0; + if (!n_token || !n_head || !head_dim) return 0; + if (n_rot == 0u) return 1; + if (n_rot > head_dim || (n_rot & 1u) != 0u || !(theta > 0.0f) || !isfinite(theta)) return 0; + const uint64_t row_values = (uint64_t)n_head * head_dim; + const uint64_t count = row_values * n_token; + if (count > UINT32_MAX) return 0; + if (!cuda_tensor_has_bytes(values, count * 4u)) return 0; + if (positions && !cuda_tensor_has_bytes(positions, (uint64_t)n_token * 4u)) return 0; + qrocm_rope a{}; + a.n_token = n_token; a.n_head = n_head; a.head_dim = head_dim; a.n_rot = n_rot; a.theta = theta; + a.source_token_stride = row_values * 4u; a.source_head_stride = (uint64_t)head_dim * 4u; a.source_dim_stride = 4u; + a.output_token_stride = row_values * 4u; a.output_head_stride = (uint64_t)head_dim * 4u; a.output_dim_stride = 4u; + a.position_stride = 4u; + // Positions live in device memory; for the scalar form stage the single + // value into a tiny device buffer so the kernel reads it uniformly. + const uint8_t *pos_ptr = NULL; + uint32_t *staged = NULL; + if (positions) { + pos_ptr = (const uint8_t *)positions->ptr; + } else { + if (cudaMalloc((void **)&staged, sizeof(uint32_t)) != cudaSuccess) return 0; + if (cudaMemcpy(staged, single_position, sizeof(uint32_t), cudaMemcpyHostToDevice) != cudaSuccess) { + cudaFree(staged); return 0; + } + a.position_stride = 0u; + pos_ptr = (const uint8_t *)staged; + } + const uint32_t nth = qrocm_elem_threads(count); + qrocm_k_rope<<<(unsigned)((count + nth - 1) / nth), nth>>>( + a, (const char *)values->ptr, pos_ptr, (char *)values->ptr); + const int ok = cuda_ok(cudaGetLastError(), "qwen35 rope"); + if (staged) { cudaDeviceSynchronize(); cudaFree(staged); } + return ok; +} +extern "C" int ds4_gpu_qwen35_rope_prefix_tensor( + ds4_gpu_tensor *values, uint32_t n_head, uint32_t head_dim, uint32_t n_rot, + uint32_t position, float theta) { + return qrocm_rope_impl(values, NULL, &position, 1u, n_head, head_dim, n_rot, theta); +} +extern "C" int ds4_gpu_qwen35_rope_prefix_batch_tensor( + ds4_gpu_tensor *values, const ds4_gpu_tensor *positions, uint32_t n_token, + uint32_t n_head, uint32_t head_dim, uint32_t n_rot, float theta) { + return qrocm_rope_impl(values, positions, NULL, n_token, n_head, head_dim, n_rot, theta); +} + +extern "C" int ds4_gpu_qwen35_causal_conv_step_tensor( + ds4_gpu_tensor *out, ds4_gpu_tensor *state, const ds4_gpu_tensor *input, + const void *model_map, uint64_t model_size, uint64_t weight_offset, + uint32_t n_channel, uint32_t kernel_size) { + if (!out || !state || !input || !model_map || !n_channel || kernel_size < 2u) return 0; + const uint64_t weight_bytes = (uint64_t)n_channel * kernel_size * 4u; + const uint64_t state_bytes = (uint64_t)n_channel * (kernel_size - 1u) * 4u; + if (!cuda_model_range_fits(model_size, weight_offset, weight_bytes)) return 0; + if (!cuda_tensor_has_bytes(input, (uint64_t)n_channel * 4u) || + !cuda_tensor_has_bytes(state, state_bytes) || + !cuda_tensor_has_bytes(out, (uint64_t)n_channel * 4u)) return 0; + const char *w = cuda_model_range_ptr(model_map, weight_offset, weight_bytes, "qwen35_conv_step_w"); + if (!w) return 0; + qrocm_conv_step a{}; + a.n_channel = n_channel; a.kernel_size = kernel_size; + a.input_channel_stride = 4u; a.weight_channel_stride = (uint64_t)kernel_size * 4u; a.weight_tap_stride = 4u; + a.state_channel_stride = (uint64_t)(kernel_size - 1u) * 4u; a.state_tap_stride = 4u; a.output_channel_stride = 4u; + const uint32_t nth = qrocm_elem_threads(n_channel); + qrocm_k_conv_step<<<(unsigned)(((uint64_t)n_channel + nth - 1) / nth), nth>>>( + a, (const char *)input->ptr, w, (char *)state->ptr, (char *)out->ptr); + return cuda_ok(cudaGetLastError(), "qwen35 conv_step"); +} + +extern "C" int ds4_gpu_qwen35_causal_conv_sequence_tensor( + ds4_gpu_tensor *out, ds4_gpu_tensor *state, const ds4_gpu_tensor *input, + const void *model_map, uint64_t model_size, uint64_t weight_offset, + uint32_t n_token, uint32_t n_channel, uint32_t kernel_size) { + if (!out || !state || !input || !model_map || !n_token || !n_channel || kernel_size != 4u) return 0; + const uint64_t activation_bytes = (uint64_t)n_channel * n_token * 4u; + const uint64_t state_bytes = (uint64_t)n_channel * 3u * 4u; + const uint64_t weight_bytes = (uint64_t)n_channel * 4u * 4u; + if (!cuda_model_range_fits(model_size, weight_offset, weight_bytes)) return 0; + if (!cuda_tensor_has_bytes(input, activation_bytes) || + !cuda_tensor_has_bytes(state, state_bytes) || + !cuda_tensor_has_bytes(out, activation_bytes)) return 0; + const char *w = cuda_model_range_ptr(model_map, weight_offset, weight_bytes, "qwen35_conv_seq_w"); + if (!w) return 0; + qrocm_conv_sequence a{}; + a.n_token = n_token; a.n_channel = n_channel; a.kernel_size = 4u; + a.input_token_stride = (uint64_t)n_channel * 4u; a.input_channel_stride = 4u; + a.weight_channel_stride = 16u; a.weight_tap_stride = 4u; + a.state_channel_stride = 12u; a.state_tap_stride = 4u; + a.output_token_stride = (uint64_t)n_channel * 4u; a.output_channel_stride = 4u; + const uint32_t nth = qrocm_elem_threads(n_channel); + qrocm_k_conv_sequence<<<(unsigned)(((uint64_t)n_channel + nth - 1) / nth), nth>>>( + a, (const char *)input->ptr, w, (char *)state->ptr, (char *)out->ptr); + return cuda_ok(cudaGetLastError(), "qwen35 conv_sequence"); +} + +extern "C" int ds4_gpu_qwen35_gated_delta_step_tensor( + ds4_gpu_tensor *out, ds4_gpu_tensor *state, const ds4_gpu_tensor *query, + const ds4_gpu_tensor *key, const ds4_gpu_tensor *value, const ds4_gpu_tensor *log_decay, + const ds4_gpu_tensor *beta, uint32_t n_key_head, uint32_t n_value_head, + uint32_t key_dim, uint32_t value_dim) { + if (!out || !state || !query || !key || !value || !log_decay || !beta) return 0; + if (!n_key_head || !n_value_head || !key_dim || !value_dim || + n_value_head % n_key_head != 0u) return 0; + const uint64_t key_bytes = (uint64_t)n_key_head * key_dim * 4u; + const uint64_t value_bytes = (uint64_t)n_value_head * value_dim * 4u; + const uint64_t control_bytes = (uint64_t)n_value_head * 4u; + const uint64_t state_bytes = (uint64_t)n_value_head * value_dim * key_dim * 4u; + if (!cuda_tensor_has_bytes(query, key_bytes) || !cuda_tensor_has_bytes(key, key_bytes) || + !cuda_tensor_has_bytes(value, value_bytes) || !cuda_tensor_has_bytes(log_decay, control_bytes) || + !cuda_tensor_has_bytes(beta, control_bytes) || !cuda_tensor_has_bytes(state, state_bytes) || + !cuda_tensor_has_bytes(out, value_bytes)) return 0; + qrocm_gated_delta_step a{}; + a.n_key_head = n_key_head; a.n_value_head = n_value_head; a.key_dim = key_dim; a.value_dim = value_dim; + a.query_head_stride = (uint64_t)key_dim * 4u; a.query_dim_stride = 4u; + a.key_head_stride = (uint64_t)key_dim * 4u; a.key_dim_stride = 4u; + a.value_head_stride = (uint64_t)value_dim * 4u; a.value_dim_stride = 4u; + a.log_decay_head_stride = 4u; a.beta_head_stride = 4u; + a.state_head_stride = (uint64_t)value_dim * key_dim * 4u; a.state_value_stride = (uint64_t)key_dim * 4u; a.state_key_stride = 4u; + a.output_head_stride = (uint64_t)value_dim * 4u; a.output_dim_stride = 4u; + if (key_dim == 128u) { + dim3 grid((value_dim + 3u) / 4u, n_value_head, 1u); + dim3 block(32u, 4u, 1u); + qrocm_k_gd_step_128<<>>( + a, (const char *)query->ptr, (const char *)key->ptr, (const char *)value->ptr, + (const char *)log_decay->ptr, (const char *)beta->ptr, (char *)state->ptr, (char *)out->ptr); + return cuda_ok(cudaGetLastError(), "qwen35 gd_step_128"); + } + uint32_t nsg = 0; const uint32_t nth = qrocm_reduction_threads(128u, &nsg); + qrocm_k_gd_step<<>>( + a, (const char *)query->ptr, (const char *)key->ptr, (const char *)value->ptr, + (const char *)log_decay->ptr, (const char *)beta->ptr, (char *)state->ptr, (char *)out->ptr); + return cuda_ok(cudaGetLastError(), "qwen35 gd_step"); +} + +extern "C" int ds4_gpu_qwen35_gated_delta_sequence_128_tensor( + ds4_gpu_tensor *out, ds4_gpu_tensor *state, const ds4_gpu_tensor *projection, + const ds4_gpu_tensor *log_decay, const ds4_gpu_tensor *beta, + uint32_t n_token, uint32_t n_key_head, uint32_t n_value_head) { + if (!out || !state || !projection || !log_decay || !beta) return 0; + if (!n_token || !n_key_head || !n_value_head || n_value_head % n_key_head != 0u) return 0; + const uint64_t query_values = (uint64_t)n_key_head * 128u; + const uint64_t value_values = (uint64_t)n_value_head * 128u; + const uint64_t proj_row = query_values + query_values + value_values; + const uint64_t proj_bytes = proj_row * n_token * 4u; + const uint64_t out_bytes = value_values * n_token * 4u; + const uint64_t control_bytes = (uint64_t)n_value_head * n_token * 4u; + const uint64_t state_bytes = (uint64_t)n_value_head * 128u * 128u * 4u; + if (!cuda_tensor_has_bytes(projection, proj_bytes) || !cuda_tensor_has_bytes(log_decay, control_bytes) || + !cuda_tensor_has_bytes(beta, control_bytes) || !cuda_tensor_has_bytes(state, state_bytes) || + !cuda_tensor_has_bytes(out, out_bytes)) return 0; + qrocm_gated_delta_sequence a{}; + a.n_token = n_token; a.n_key_head = n_key_head; a.n_value_head = n_value_head; + a.key_dim = 128u; a.value_dim = 128u; + a.projection_token_stride = proj_row * 4u; + a.query_offset = 0u; a.key_offset = query_values * 4u; a.value_offset = (query_values + query_values) * 4u; + a.query_head_stride = 512u; a.query_dim_stride = 4u; + a.key_head_stride = 512u; a.key_dim_stride = 4u; + a.value_head_stride = 512u; a.value_dim_stride = 4u; + a.log_decay_token_stride = (uint64_t)n_value_head * 4u; a.log_decay_head_stride = 4u; + a.beta_token_stride = (uint64_t)n_value_head * 4u; a.beta_head_stride = 4u; + a.state_head_stride = 65536u; a.state_value_stride = 512u; a.state_key_stride = 4u; + a.output_token_stride = value_values * 4u; a.output_head_stride = 512u; a.output_dim_stride = 4u; + dim3 grid(32u, n_value_head, 1u); + dim3 block(32u, 4u, 1u); + qrocm_k_gd_sequence_128<<>>( + a, (const char *)projection->ptr, (const char *)log_decay->ptr, + (const char *)beta->ptr, (char *)state->ptr, (char *)out->ptr); + return cuda_ok(cudaGetLastError(), "qwen35 gd_sequence_128"); +} + +static int qrocm_gd_controls_impl( + ds4_gpu_tensor *log_decay, ds4_gpu_tensor *beta, + const ds4_gpu_tensor *alpha_logit, const ds4_gpu_tensor *beta_logit, + const void *model_map, uint64_t model_size, uint64_t ssm_a_offset, uint64_t dt_bias_offset, + uint32_t n_token, uint32_t n_value_head) { + if (!log_decay || !beta || !alpha_logit || !beta_logit || !model_map || !n_token || !n_value_head) return 0; + const uint64_t bytes = (uint64_t)n_token * n_value_head * 4u; + const uint64_t weight_bytes = (uint64_t)n_value_head * 4u; + if (!cuda_model_range_fits(model_size, ssm_a_offset, weight_bytes) || + !cuda_model_range_fits(model_size, dt_bias_offset, weight_bytes)) return 0; + if (!cuda_tensor_has_bytes(alpha_logit, bytes) || !cuda_tensor_has_bytes(beta_logit, bytes) || + !cuda_tensor_has_bytes(log_decay, bytes) || !cuda_tensor_has_bytes(beta, bytes)) return 0; + const char *ssm_a = cuda_model_range_ptr(model_map, ssm_a_offset, weight_bytes, "qwen35_ssm_a"); + const char *dt_bias = cuda_model_range_ptr(model_map, dt_bias_offset, weight_bytes, "qwen35_dt_bias"); + if (!ssm_a || !dt_bias) return 0; + qrocm_gated_delta_controls a{}; + a.n_token = n_token; a.n_value_head = n_value_head; + a.alpha_logit_token_stride = (uint64_t)n_value_head * 4u; a.alpha_logit_head_stride = 4u; + a.beta_logit_token_stride = (uint64_t)n_value_head * 4u; a.beta_logit_head_stride = 4u; + a.ssm_a_head_stride = 4u; a.dt_bias_head_stride = 4u; + a.log_decay_token_stride = (uint64_t)n_value_head * 4u; a.log_decay_head_stride = 4u; + a.beta_token_stride = (uint64_t)n_value_head * 4u; a.beta_head_stride = 4u; + const uint64_t count = (uint64_t)n_token * n_value_head; + const uint32_t nth = qrocm_elem_threads(count); + qrocm_k_gd_controls<<<(unsigned)((count + nth - 1) / nth), nth>>>( + a, (const char *)alpha_logit->ptr, (const char *)beta_logit->ptr, ssm_a, dt_bias, + (char *)log_decay->ptr, (char *)beta->ptr); + return cuda_ok(cudaGetLastError(), "qwen35 gd_controls"); +} +extern "C" int ds4_gpu_qwen35_gated_delta_controls_tensor( + ds4_gpu_tensor *log_decay, ds4_gpu_tensor *beta, + const ds4_gpu_tensor *alpha_logit, const ds4_gpu_tensor *beta_logit, + const void *model_map, uint64_t model_size, uint64_t ssm_a_offset, uint64_t dt_bias_offset, + uint32_t n_value_head) { + return qrocm_gd_controls_impl(log_decay, beta, alpha_logit, beta_logit, model_map, model_size, + ssm_a_offset, dt_bias_offset, 1u, n_value_head); +} +extern "C" int ds4_gpu_qwen35_gated_delta_controls_batch_tensor( + ds4_gpu_tensor *log_decay, ds4_gpu_tensor *beta, + const ds4_gpu_tensor *alpha_logit, const ds4_gpu_tensor *beta_logit, + const void *model_map, uint64_t model_size, uint64_t ssm_a_offset, uint64_t dt_bias_offset, + uint32_t n_token, uint32_t n_value_head) { + return qrocm_gd_controls_impl(log_decay, beta, alpha_logit, beta_logit, model_map, model_size, + ssm_a_offset, dt_bias_offset, n_token, n_value_head); +} + +extern "C" int ds4_gpu_qwen35_rmsnorm_gated_tensor( + ds4_gpu_tensor *out, const ds4_gpu_tensor *input, const ds4_gpu_tensor *gate, + const void *model_map, uint64_t model_size, uint64_t weight_offset, + uint32_t n_vector, uint32_t dim, float epsilon) { + if (!out || !input || !gate || !model_map || !n_vector || !dim || + !(epsilon > 0.0f) || !isfinite(epsilon)) return 0; + const uint64_t values = (uint64_t)n_vector * dim; + const uint64_t bytes = values * 4u; + const uint64_t weight_bytes = (uint64_t)dim * 4u; + if (!cuda_model_range_fits(model_size, weight_offset, weight_bytes)) return 0; + if (!cuda_tensor_has_bytes(input, bytes) || !cuda_tensor_has_bytes(gate, bytes) || + !cuda_tensor_has_bytes(out, bytes)) return 0; + const char *w = cuda_model_range_ptr(model_map, weight_offset, weight_bytes, "qwen35_rmsnorm_w"); + if (!w) return 0; + qrocm_rmsnorm_gated a{}; + a.n_vector = n_vector; a.dim = dim; a.epsilon = epsilon; + a.input_vector_stride = (uint64_t)dim * 4u; a.input_dim_stride = 4u; + a.gate_vector_stride = (uint64_t)dim * 4u; a.gate_dim_stride = 4u; a.weight_dim_stride = 4u; + a.output_vector_stride = (uint64_t)dim * 4u; a.output_dim_stride = 4u; + uint32_t nsg = 0; const uint32_t nth = qrocm_reduction_threads(128u, &nsg); + qrocm_k_rmsnorm_gated<<>>( + a, (const char *)input->ptr, (const char *)gate->ptr, w, (char *)out->ptr); + return cuda_ok(cudaGetLastError(), "qwen35 rmsnorm_gated"); +} + +extern "C" int ds4_gpu_qwen35_dequant_embedding_q8_0_tensor( + ds4_gpu_tensor *out, const void *model_map, uint64_t model_size, + uint64_t embedding_offset, uint32_t row_index, uint32_t n_embd) { + if (!out || !model_map || !n_embd || (n_embd % 32u) != 0u) return 0; + const uint64_t row_bytes = (uint64_t)(n_embd / 32u) * 34u; + const uint64_t row_offset = embedding_offset + (uint64_t)row_index * row_bytes; + if (!cuda_model_range_fits(model_size, row_offset, row_bytes)) return 0; + if (!cuda_tensor_has_bytes(out, (uint64_t)n_embd * 4u)) return 0; + const char *row = cuda_model_range_ptr(model_map, row_offset, row_bytes, "qwen35_embed_row"); + if (!row) return 0; + qrocm_embedding_q8_0 a{}; + a.row_index = 0u; a.n_embd = n_embd; a.block_size = 32u; + a.source_row_stride = row_bytes; a.source_block_stride = 34u; + a.source_scale_offset = 0u; a.source_quant_offset = 2u; a.source_quant_stride = 1u; a.output_dim_stride = 4u; + const uint32_t nth = qrocm_elem_threads(n_embd); + qrocm_k_embed_q8_0<<<(unsigned)(((uint64_t)n_embd + nth - 1) / nth), nth>>>( + a, row, (char *)out->ptr); + return cuda_ok(cudaGetLastError(), "qwen35 embed_q8_0"); +} + +extern "C" int ds4_gpu_qwen35_dequant_embedding_q8_0_batch_tensor( + ds4_gpu_tensor *out, const ds4_gpu_tensor *token_ids, const void *model_map, + uint64_t model_size, uint64_t embedding_offset, uint32_t n_token, uint32_t n_row, uint32_t n_embd) { + if (!out || !token_ids || !model_map || !n_token || !n_row || !n_embd || (n_embd % 32u) != 0u) return 0; + const uint64_t row_bytes = (uint64_t)(n_embd / 32u) * 34u; + const uint64_t embedding_bytes = (uint64_t)n_row * row_bytes; + const uint64_t values = (uint64_t)n_token * n_embd; + if (values > UINT32_MAX) return 0; + if (!cuda_model_range_fits(model_size, embedding_offset, embedding_bytes)) return 0; + if (!cuda_tensor_has_bytes(token_ids, (uint64_t)n_token * 4u) || + !cuda_tensor_has_bytes(out, values * 4u)) return 0; + const char *table = cuda_model_range_ptr(model_map, embedding_offset, embedding_bytes, "qwen35_embed_table"); + if (!table) return 0; + qrocm_embedding_q8_0_batch a{}; + a.n_token = n_token; a.n_row = n_row; a.n_embd = n_embd; a.block_size = 32u; + a.source_row_stride = row_bytes; a.source_block_stride = 34u; + a.source_scale_offset = 0u; a.source_quant_offset = 2u; a.source_quant_stride = 1u; + a.token_id_stride = 4u; a.output_token_stride = (uint64_t)n_embd * 4u; a.output_dim_stride = 4u; + const uint32_t nth = qrocm_elem_threads(values); + qrocm_k_embed_q8_0_batch<<<(unsigned)((values + nth - 1) / nth), nth>>>( + a, table, (const char *)token_ids->ptr, (char *)out->ptr); + return cuda_ok(cudaGetLastError(), "qwen35 embed_q8_0_batch"); +} + +extern "C" int ds4_gpu_qwen35_gqa_decode_tensor( + ds4_gpu_tensor *out, const ds4_gpu_tensor *query, const ds4_gpu_tensor *key_cache, + const ds4_gpu_tensor *value_cache, uint32_t n_kv, uint32_t n_query_head, + uint32_t n_kv_head, uint32_t head_dim) { + if (!out || !query || !key_cache || !value_cache) return 0; + if (!n_kv || !n_query_head || !n_kv_head || !head_dim || n_query_head % n_kv_head != 0u) return 0; + const uint64_t query_bytes = (uint64_t)n_query_head * head_dim * 4u; + const uint64_t cache_bytes = (uint64_t)n_kv * n_kv_head * head_dim * 4u; + if (!cuda_tensor_has_bytes(query, query_bytes) || !cuda_tensor_has_bytes(key_cache, cache_bytes) || + !cuda_tensor_has_bytes(value_cache, cache_bytes) || !cuda_tensor_has_bytes(out, query_bytes)) return 0; + qrocm_gqa_decode a{}; + a.n_kv = n_kv; a.n_query_head = n_query_head; a.n_kv_head = n_kv_head; a.head_dim = head_dim; + a.query_head_stride = (uint64_t)head_dim * 4u; a.query_dim_stride = 4u; + a.key_token_stride = (uint64_t)n_kv_head * head_dim * 4u; a.key_head_stride = (uint64_t)head_dim * 4u; a.key_dim_stride = 4u; + a.value_token_stride = (uint64_t)n_kv_head * head_dim * 4u; a.value_head_stride = (uint64_t)head_dim * 4u; a.value_dim_stride = 4u; + a.output_head_stride = (uint64_t)head_dim * 4u; a.output_dim_stride = 4u; + const uint32_t preferred = ((head_dim + 31u) / 32u) * 32u; + uint32_t nsg = 0; const uint32_t nth = qrocm_reduction_threads(preferred, &nsg); + if (nth < head_dim || nth != preferred) return 0; + qrocm_k_gqa_decode<<>>( + a, (const char *)query->ptr, (const char *)key_cache->ptr, + (const char *)value_cache->ptr, (char *)out->ptr); + return cuda_ok(cudaGetLastError(), "qwen35 gqa_decode"); +} + +extern "C" int ds4_gpu_qwen35_gqa_prefill_tensor( + ds4_gpu_tensor *out, const ds4_gpu_tensor *query, const ds4_gpu_tensor *key_cache, + const ds4_gpu_tensor *value_cache, uint32_t position0, uint32_t n_token, + uint32_t n_query_head, uint32_t n_kv_head, uint32_t head_dim) { + if (!out || !query || !key_cache || !value_cache) return 0; + if (!n_token || !n_query_head || !n_kv_head || !head_dim || n_query_head % n_kv_head != 0u) return 0; + if (position0 > UINT32_MAX - n_token) return 0; + const uint32_t n_kv = position0 + n_token; + const uint64_t query_bytes = (uint64_t)n_query_head * head_dim * n_token * 4u; + const uint64_t cache_bytes = (uint64_t)n_kv_head * head_dim * n_kv * 4u; + if (!cuda_tensor_has_bytes(query, query_bytes) || !cuda_tensor_has_bytes(key_cache, cache_bytes) || + !cuda_tensor_has_bytes(value_cache, cache_bytes) || !cuda_tensor_has_bytes(out, query_bytes)) return 0; + qrocm_gqa_prefill a{}; + a.position0 = position0; a.n_token = n_token; a.n_query_head = n_query_head; + a.n_kv_head = n_kv_head; a.head_dim = head_dim; + a.query_token_stride = (uint64_t)n_query_head * head_dim * 4u; a.query_head_stride = (uint64_t)head_dim * 4u; a.query_dim_stride = 4u; + a.key_token_stride = (uint64_t)n_kv_head * head_dim * 4u; a.key_head_stride = (uint64_t)head_dim * 4u; a.key_dim_stride = 4u; + a.value_token_stride = (uint64_t)n_kv_head * head_dim * 4u; a.value_head_stride = (uint64_t)head_dim * 4u; a.value_dim_stride = 4u; + a.output_token_stride = (uint64_t)n_query_head * head_dim * 4u; a.output_head_stride = (uint64_t)head_dim * 4u; a.output_dim_stride = 4u; + const uint32_t preferred = ((head_dim + 31u) / 32u) * 32u; + uint32_t nsg = 0; const uint32_t nth = qrocm_reduction_threads(preferred, &nsg); + if (nth < head_dim || nth != preferred) return 0; + dim3 grid(n_query_head, n_token, 1u); + qrocm_k_gqa_prefill<<>>( + a, (const char *)query->ptr, (const char *)key_cache->ptr, + (const char *)value_cache->ptr, (char *)out->ptr); + return cuda_ok(cudaGetLastError(), "qwen35 gqa_prefill"); +} + +extern "C" int ds4_gpu_qwen35_router_softmax_top8_batch_tensor( + ds4_gpu_tensor *selected, ds4_gpu_tensor *selected_weight, + const ds4_gpu_tensor *logits, uint32_t n_token) { + if (!selected || !selected_weight || !logits || !n_token) return 0; + const uint64_t logits_bytes = (uint64_t)n_token * 256u * 4u; + const uint64_t selected_bytes = (uint64_t)n_token * 8u * 4u; + if (!cuda_tensor_has_bytes(logits, logits_bytes) || !cuda_tensor_has_bytes(selected, selected_bytes) || + !cuda_tensor_has_bytes(selected_weight, selected_bytes)) return 0; + qrocm_router_top8 a{}; + a.n_token = n_token; + a.logits_token_stride = 1024u; a.logits_stride = 4u; + a.selected_token_stride = 32u; a.selected_stride = 4u; + a.selected_weight_token_stride = 32u; a.selected_weight_stride = 4u; + qrocm_k_router_serial<<>>( + a, (const char *)logits->ptr, (char *)selected->ptr, (char *)selected_weight->ptr); + return cuda_ok(cudaGetLastError(), "qwen35 router_top8"); +} +extern "C" int ds4_gpu_qwen35_router_softmax_top8_tensor( + ds4_gpu_tensor *selected, ds4_gpu_tensor *selected_weight, const ds4_gpu_tensor *logits) { + return ds4_gpu_qwen35_router_softmax_top8_batch_tensor(selected, selected_weight, logits, 1u); +} + +// Resident Qwen top-8 routed MoE. The shared ROCm Q4_K MoE core is specialized +// for DeepSeek's top-6 (fused sum6 kernels); Qwen routes top-8. Run it as two +// independent 4-expert halves and add the partial outputs — the same split the +// Metal resident/streaming paths use. selected_half0/half1 are offset-0/16 +// views into selected_top8, so they already point at experts [0..3] and [4..7]. +extern "C" int ds4_gpu_qwen35_routed_moe_top8_tensor( + ds4_gpu_tensor *out, ds4_gpu_tensor *partial0, ds4_gpu_tensor *partial1, + ds4_gpu_tensor *gate, ds4_gpu_tensor *up, ds4_gpu_tensor *mid, ds4_gpu_tensor *experts, + const void *model_map, uint64_t model_size, uint64_t gate_offset, uint64_t up_offset, + uint64_t down_offset, uint32_t gate_type, uint32_t down_type, + uint64_t gate_expert_bytes, uint64_t gate_row_bytes, uint64_t down_expert_bytes, + uint64_t down_row_bytes, uint32_t expert_in_dim, uint32_t expert_mid_dim, uint32_t out_dim, + const ds4_gpu_tensor *selected_top8, const ds4_gpu_tensor *weights_top8, + const ds4_gpu_tensor *selected_half0, const ds4_gpu_tensor *selected_half1, + const ds4_gpu_tensor *weights_half0, const ds4_gpu_tensor *weights_half1, + uint32_t n_total_expert, float clamp, const ds4_gpu_tensor *x, + uint32_t layer_index, bool trusted_gpu_route) { + (void)selected_top8; (void)weights_top8; (void)trusted_gpu_route; + if (!partial0 || !partial1 || !selected_half0 || !selected_half1 || + !weights_half0 || !weights_half1) return 0; + int ok = ds4_gpu_routed_moe_one_tensor( + partial0, gate, up, mid, experts, model_map, model_size, gate_offset, up_offset, down_offset, + gate_type, down_type, gate_expert_bytes, gate_row_bytes, down_expert_bytes, down_row_bytes, + expert_in_dim, expert_mid_dim, out_dim, selected_half0, weights_half0, + n_total_expert, 4u, clamp, x, layer_index); + if (!ok) return 0; + ok = ds4_gpu_routed_moe_one_tensor( + partial1, gate, up, mid, experts, model_map, model_size, gate_offset, up_offset, down_offset, + gate_type, down_type, gate_expert_bytes, gate_row_bytes, down_expert_bytes, down_row_bytes, + expert_in_dim, expert_mid_dim, out_dim, selected_half1, weights_half1, + n_total_expert, 4u, clamp, x, layer_index); + if (!ok) return 0; + return ds4_gpu_add_tensor(out, partial0, partial1, out_dim); +} + +// Profiling counters (stubs — resident route always takes the GPU path). +static uint64_t g_qwen35_resident_gpu_route_calls; +static uint64_t g_qwen35_gdn128_parallel_calls; +extern "C" void ds4_gpu_internal_qwen35_resident_route_stats_reset(void) { g_qwen35_resident_gpu_route_calls = 0; } +extern "C" void ds4_gpu_internal_qwen35_resident_route_stats_add(uint64_t calls) { g_qwen35_resident_gpu_route_calls += calls; } +extern "C" uint64_t ds4_gpu_internal_qwen35_resident_gpu_route_calls(void) { return g_qwen35_resident_gpu_route_calls; } +extern "C" uint64_t ds4_gpu_internal_qwen35_resident_host_readbacks(void) { return 0; } +extern "C" void ds4_gpu_internal_qwen35_gdn128_stats_reset(void) { g_qwen35_gdn128_parallel_calls = 0; } +extern "C" uint64_t ds4_gpu_internal_qwen35_gdn128_parallel_calls(void) { return g_qwen35_gdn128_parallel_calls; } + +#endif // DS4_ROCM_QWEN_CUH diff --git a/rocm/ds4_rocm_runtime.cuh b/rocm/ds4_rocm_runtime.cuh index 3bd786f8e..2734a08a4 100644 --- a/rocm/ds4_rocm_runtime.cuh +++ b/rocm/ds4_rocm_runtime.cuh @@ -12,7 +12,19 @@ static int g_model_cache_full; static int g_ssd_streaming_mode; static cudaStream_t g_model_upload_stream; static cudaStream_t g_selected_readback_stream; -static cudaEvent_t g_selected_readback_event; +/* + * Selected-id readback events form a small ring indexed by the monotonically + * increasing event value. A single shared event object is not safe here: + * three threads (decode encoder, single-token async load worker, batch async + * load worker) record and wait concurrently under MTP speculative decoding, + * and on HIP re-recording an event while a hipStreamWaitEvent from an earlier + * record is still pending resets the underlying HSA signal — the pending + * barrier then never fires and the readback stream wedges (observed as a + * permanent hipDeviceSynchronize hang with an idle GPU). Per-value slots keep + * every pending wait bound to its own record. + */ +#define DS4_SELECTED_READBACK_EVENT_RING 8u +static cudaEvent_t g_selected_readback_events[DS4_SELECTED_READBACK_EVENT_RING]; static uint64_t g_selected_readback_event_value; static cublasHandle_t g_cublas; static int g_cublas_ready; @@ -25,6 +37,9 @@ enum { DS4_ROCM_N_EXPERT = 256u, DS4_ROCM_MAX_N_EXPERT = 384u, DS4_ROCM_N_EXPERT_USED = 6u, + /* 36 workers were tried for MTP verify unions (no measured gain — the + * load is latency-bound, not queue-depth-bound) and reverted: more + * concurrent GPU-touching threads widen the HSA-wedge race window. */ DS4_ROCM_STREAM_READ_WORKERS = DS4_ROCM_N_EXPERT_USED * 3u, DS4_ROCM_STREAM_READ_MAX_JOBS = DS4_ROCM_MAX_N_EXPERT * 3u, DS4_ROCM_COMPRESSOR_MAX_RATIO = 128u @@ -314,8 +329,215 @@ static void cuda_stream_cache_stats_note_resident(void) { } } +/* ---------------------------------------------------------------------- + * Host-RAM LFU expert-slice cache (Phase 3) + * + * The per-expert streaming read path (cuda_stream_read_job_run) issues a + * buffered pread and then drops the OS page cache for that range + * (cuda_model_drop_file_pages / FADV_DONTNEED), so the kernel keeps NO RAM + * copy of routed experts. This cache is the only RAM tier between the + * 16-GiB-class GTT resident cache and the ~5 GB/s NVMe: it holds pageable + * copies of raw file bytes keyed by their file offset (which uniquely + * identifies a gate/up/down expert slice, model map is read-only/immutable), + * so a repeat demand-load is served from RAM (~90 GB/s) instead of the SSD. + * + * Policy: LFU with a last-used tiebreak (popular experts survive regardless + * of recency, which suits MoE routing skew). Thread-safe: the read-worker + * pool calls get()/insert() from up to DS4_ROCM_STREAM_READ_WORKERS threads, + * all serialized on g_host_cache_mutex. Budget: DS4_ROCM_HOST_CACHE_GB + * (0 disables; default scales to available RAM, capped at 12 GiB), clamped + * to 60%% of MemAvailable so it never fights the OS/JARVIS for RAM. + * -------------------------------------------------------------------- */ +struct ds4_host_cache_entry { + void *data; + uint64_t bytes; + uint64_t freq; + uint64_t last; +}; +static std::unordered_map g_host_cache; +static pthread_mutex_t g_host_cache_mutex = PTHREAD_MUTEX_INITIALIZER; +static uint64_t g_host_cache_budget = 0; /* 0 until initialized */ +static uint64_t g_host_cache_bytes = 0; +static uint64_t g_host_cache_clock = 0; +static int g_host_cache_ready = 0; +static uint64_t g_host_cache_hits = 0; +static uint64_t g_host_cache_misses = 0; +static uint64_t g_host_cache_inserts = 0; +static uint64_t g_host_cache_evictions = 0; + +static uint64_t cuda_host_cache_mem_available_bytes(void) { + FILE *f = fopen("/proc/meminfo", "r"); + if (!f) return 0; + char key[64]; + unsigned long long val = 0; + char unit[16]; + uint64_t avail = 0; + while (fscanf(f, "%63s %llu %15s", key, &val, unit) == 3) { + if (strcmp(key, "MemAvailable:") == 0) { avail = (uint64_t)val * 1024ull; break; } + } + fclose(f); + return avail; +} + +/* Must be called with g_host_cache_mutex held. */ +static void cuda_host_cache_init_locked(void) { + if (g_host_cache_ready) return; + g_host_cache_ready = 1; + /* Default OFF: on DeepSeek V4 Flash the per-token routed-expert working set + * (~72 GiB) dwarfs any RAM cache that fits alongside JARVIS, so a sub-model + * cache yields a low hit rate and its bookkeeping can cost more than it + * saves. Opt in with DS4_ROCM_HOST_CACHE_GB when the cache can hold a large + * fraction of the working set (more RAM) or the workload has higher expert + * reuse. See LABESTIA_DESIGN.md Sec.Q3. */ + uint64_t budget = 0; + const char *env = getenv("DS4_ROCM_HOST_CACHE_GB"); + if (env && *env) { + char *end = NULL; + const double gb = strtod(env, &end); + if (end && *end == '\0' && gb >= 0.0) { + budget = (uint64_t)(gb * 1073741824.0); + } else { + fprintf(stderr, DS4_GPU_LOG_PREFIX + "ignoring invalid DS4_ROCM_HOST_CACHE_GB=\"%s\"\n", env); + } + } + /* Never take more than 40% of currently-available RAM: the buffered-pread + * fallback and the OS still need page cache, and on this 47 GiB box shared + * with JARVIS a larger cache pushes the system into reclaim/swap and + * *regresses* throughput hard (measured: a ~19 GiB resident cache dropped + * generation from ~1.9 to ~0.5 t/s). Keep the cache well clear of that. */ + const uint64_t avail = cuda_host_cache_mem_available_bytes(); + if (avail) { + const uint64_t cap = avail / 10ull * 4ull; + if (budget > cap) budget = cap; + } + g_host_cache_budget = budget; + if (budget) { + fprintf(stderr, DS4_GPU_LOG_PREFIX + "host expert cache enabled: %.2f GiB budget\n", + (double)budget / 1073741824.0); + } +} + +/* On hit, copies bytes into dst and returns 1. Thread-safe. */ +static int cuda_host_cache_get(uint64_t offset, uint64_t bytes, void *dst) { + int hit = 0; + pthread_mutex_lock(&g_host_cache_mutex); + cuda_host_cache_init_locked(); + if (g_host_cache_budget) { + std::unordered_map::iterator it = + g_host_cache.find(offset); + if (it != g_host_cache.end() && it->second.bytes == bytes) { + memcpy(dst, it->second.data, (size_t)bytes); + it->second.freq++; + it->second.last = ++g_host_cache_clock; + g_host_cache_hits++; + hit = 1; + } else { + g_host_cache_misses++; + } + } + pthread_mutex_unlock(&g_host_cache_mutex); + return hit; +} + +/* Inserts a private copy of src. Evicts LFU victims to stay under budget. */ +static void cuda_host_cache_insert(uint64_t offset, uint64_t bytes, const void *src) { + if (bytes == 0) return; + pthread_mutex_lock(&g_host_cache_mutex); + cuda_host_cache_init_locked(); + if (!g_host_cache_budget || bytes > g_host_cache_budget || + g_host_cache.find(offset) != g_host_cache.end()) { + pthread_mutex_unlock(&g_host_cache_mutex); + return; + } + while (g_host_cache_bytes + bytes > g_host_cache_budget && !g_host_cache.empty()) { + /* Sampled LFU: evicting the true global minimum is O(N) per eviction, + * which serializes the 18 read workers under heavy thrash and erases + * the cache's benefit. Instead sample a bounded window of buckets + * (rotating start so coverage spreads) and evict the min-freq entry + * found — O(window), a good LFU approximation. */ + const size_t nbuckets = g_host_cache.bucket_count(); + const size_t start = nbuckets ? (size_t)(g_host_cache_clock % nbuckets) : 0; + uint64_t victim_key = 0; + int have_victim = 0; + uint64_t best_freq = UINT64_MAX; + uint64_t best_last = UINT64_MAX; + int sampled = 0; + for (size_t b = 0; b < nbuckets && sampled < 32; b++) { + const size_t bi = (start + b) % nbuckets; + for (std::unordered_map::local_iterator lit = + g_host_cache.begin(bi); lit != g_host_cache.end(bi); ++lit) { + if (!have_victim || lit->second.freq < best_freq || + (lit->second.freq == best_freq && lit->second.last < best_last)) { + best_freq = lit->second.freq; + best_last = lit->second.last; + victim_key = lit->first; + have_victim = 1; + } + if (++sampled >= 32) break; + } + } + if (!have_victim) break; + std::unordered_map::iterator v = + g_host_cache.find(victim_key); + if (v == g_host_cache.end()) break; + free(v->second.data); + g_host_cache_bytes -= v->second.bytes; + g_host_cache.erase(v); + g_host_cache_evictions++; + } + void *copy = malloc((size_t)bytes); + if (copy) { + memcpy(copy, src, (size_t)bytes); + ds4_host_cache_entry e; + e.data = copy; + e.bytes = bytes; + e.freq = 1; + e.last = ++g_host_cache_clock; + g_host_cache[offset] = e; + g_host_cache_bytes += bytes; + g_host_cache_inserts++; + /* Periodic frequency decay so a burst-popular expert cannot squat the + * cache forever (LFU aging), mirroring the reference implementation. */ + if ((g_host_cache_inserts & 4095ull) == 0ull) { + for (std::unordered_map::iterator it = + g_host_cache.begin(); it != g_host_cache.end(); ++it) { + it->second.freq >>= 1; + } + } + } + pthread_mutex_unlock(&g_host_cache_mutex); +} + +static void cuda_host_cache_release(void) { + pthread_mutex_lock(&g_host_cache_mutex); + for (std::unordered_map::iterator it = + g_host_cache.begin(); it != g_host_cache.end(); ++it) { + free(it->second.data); + } + g_host_cache.clear(); + g_host_cache_bytes = 0; + pthread_mutex_unlock(&g_host_cache_mutex); +} + static void cuda_stream_cache_stats_print(const char *label) { if (!cuda_stream_cache_stats_on()) return; + if (g_host_cache_budget) { + const uint64_t total = g_host_cache_hits + g_host_cache_misses; + fprintf(stderr, + DS4_GPU_LOG_PREFIX "host cache stats %s: " + "hits=%llu misses=%llu hit_rate=%.1f%% inserts=%llu " + "evictions=%llu resident=%.2f/%.2f GiB\n", + label ? label : "", + (unsigned long long)g_host_cache_hits, + (unsigned long long)g_host_cache_misses, + total ? 100.0 * (double)g_host_cache_hits / (double)total : 0.0, + (unsigned long long)g_host_cache_inserts, + (unsigned long long)g_host_cache_evictions, + (double)g_host_cache_bytes / 1073741824.0, + (double)g_host_cache_budget / 1073741824.0); + } fprintf(stderr, DS4_GPU_LOG_PREFIX "stream cache stats %s: " "selected calls=%llu slots=%llu hits=%llu misses=%llu; " @@ -506,6 +728,7 @@ static void cuda_stream_read_stage_release(void) { g_stream_read_upload_streams[i] = NULL; } } + cuda_host_cache_release(); } static void cuda_stream_batch_selected_cache_release(void) { @@ -969,8 +1192,111 @@ static int cuda_stream_resident_evict_one( return 1; } +/* Device bytes kept free for the KV cache, BLAS workspaces, transient graph + * buffers, and other processes sharing the GPU pool. The reserve was + * historically a flat 16 GiB, sized for machines whose hipMemGetInfo() total + * is a ~110 GiB GTT aperture (Strix Halo class). On devices with a small + * pool (e.g. Phoenix APUs whose ROCm pool is the ~24 GiB GTT) a flat 16 GiB + * starves the streaming expert cache to zero, so the default scales with the + * pool and DS4_ROCM_STREAM_RESERVE_MB overrides it explicitly. */ static uint64_t cuda_stream_resident_free_reserve_bytes(void) { - return 16ull * 1024ull * 1024ull * 1024ull; + static uint64_t cached_reserve = UINT64_MAX; + if (cached_reserve != UINT64_MAX) return cached_reserve; + + const char *env = getenv("DS4_ROCM_STREAM_RESERVE_MB"); + if (env && *env) { + char *end = NULL; + const unsigned long long mb = strtoull(env, &end, 10); + if (end && *end == '\0' && mb <= (UINT64_MAX >> 20)) { + cached_reserve = (uint64_t)mb << 20; + return cached_reserve; + } + fprintf(stderr, + DS4_GPU_LOG_PREFIX "ignoring invalid DS4_ROCM_STREAM_RESERVE_MB=\"%s\"\n", + env); + } + + size_t free_b = 0; + size_t total_b = 0; + if (cudaMemGetInfo(&free_b, &total_b) == cudaSuccess && total_b > 0) { + (void)free_b; + /* + * Adaptive default. On UMA APUs the GTT pool is ordinary system RAM: + * if the expert cache grows until only a thin slice of the pool is + * free while the rest of the system is also tight, the kernel's TTM + * starts evicting GPU queue buffers ("Freeing queue vital buffer, + * queue evicted" + "Not enough memory for command submission!" in + * dmesg) and every later HIP sync blocks forever with an idle GPU. + * A fixed pool fraction cannot see any of that, so derive the floor + * from the host memory that is actually available at startup: + * + * headroom = max(4 GiB, MemTotal/12) (co-tenant + kernel margin) + * reserve = pool_total - (MemAvailable - headroom) + * + * i.e. the cache may consume at most what the host can spare beyond + * the headroom; everything else stays free inside the pool. On the + * 64 GB / 48 GiB-GTT reference box this reproduces the empirically + * validated 14 GiB floor (250/300-token soaks clean vs reliable + * mid-generation freezes with the old pool/8 = 6 GiB default). + */ + uint64_t mem_total = 0; + uint64_t mem_available = 0; +#if defined(__linux__) + FILE *mi = fopen("/proc/meminfo", "r"); + if (mi) { + char line[128]; + while (fgets(line, sizeof(line), mi)) { + unsigned long long kb = 0; + if (sscanf(line, "MemTotal: %llu kB", &kb) == 1) { + mem_total = (uint64_t)kb << 10; + } else if (sscanf(line, "MemAvailable: %llu kB", &kb) == 1) { + mem_available = (uint64_t)kb << 10; + } + if (mem_total && mem_available) break; + } + fclose(mi); + } +#endif + /* + * The floor is also a THROUGHPUT optimum, not only a safety line: on + * UMA boxes every GiB the expert cache wires is a GiB of page cache + * the 80 GB model file loses, and streaming reads slow down more than + * the extra residency gains (measured: a 6.6 GiB floor grew the cache + * and LOST ~10%% generation vs the 14 GiB floor on the same box). + * pool/4 keeps the fresh-state floor near the measured optimum while + * the host-aware term above still raises it when RAM is tight. + */ + const uint64_t min_reserve = + (uint64_t)total_b / 4u > (2ull << 30) ? (uint64_t)total_b / 4u + : (2ull << 30); + const uint64_t max_reserve = (uint64_t)total_b / 2u; + uint64_t reserve; + if (mem_total && mem_available) { + uint64_t headroom = mem_total / 12u; + if (headroom < (4ull << 30)) headroom = 4ull << 30; + const uint64_t allowed = + mem_available > headroom ? mem_available - headroom : 0; + reserve = (uint64_t)total_b > allowed ? (uint64_t)total_b - allowed + : min_reserve; + } else { + /* No /proc/meminfo (or non-Linux): legacy pool/8 default. */ + reserve = (uint64_t)total_b / 8u; + } + if (reserve < min_reserve) reserve = min_reserve; + if (reserve > max_reserve) reserve = max_reserve; + fprintf(stderr, + DS4_GPU_LOG_PREFIX "streaming cache free-floor: %.2f GiB " + "(pool %.2f GiB, host available %.2f GiB; " + "DS4_ROCM_STREAM_RESERVE_MB overrides)\n", + (double)reserve / 1073741824.0, + (double)total_b / 1073741824.0, + (double)mem_available / 1073741824.0); + cached_reserve = reserve; + return cached_reserve; + } + (void)cudaGetLastError(); + cached_reserve = 16ull << 30; + return cached_reserve; } static int cuda_stream_resident_make_room( @@ -1138,13 +1464,31 @@ static void cuda_stream_read_job_run(cuda_stream_read_job *job) { if (job) job->errnum = EINVAL; return; } + /* (a) Serve from the host RAM cache if present (copies under the cache + * lock, so no eviction race), skipping the NVMe read entirely. */ + if (cuda_host_cache_get(job->offset, job->bytes, job->host_buf)) { + job->ok = 1; + return; + } if (cuda_pread_full(g_model_fd, job->host_buf, job->bytes, job->offset)) { job->ok = 1; + /* (b) Populate the host cache with a private copy (host_buf is the + * shared per-worker pinned staging buffer, so we must not donate it). */ + cuda_host_cache_insert(job->offset, job->bytes, job->host_buf); } else { job->errnum = errno ? errno : EIO; } } +/* GPU uploads from the read workers are serialized: concurrent + * cudaMemcpyAsync+cudaStreamSynchronize pairs from many pool threads + * (while the encoder thread issues its own device syncs) intermittently + * wedge the HSA queues on this ROCm/APU stack — observed as generation + * freezing mid-request with one thread spinning and zero disk traffic, + * with worker stacks parked in this stream sync. The SSD preads stay + * fully parallel; the serialized H2D copies amount to ~1 ms per layer. */ +static pthread_mutex_t g_stream_read_upload_mutex = PTHREAD_MUTEX_INITIALIZER; + static int cuda_stream_read_job_upload( cuda_stream_read_job *job, cudaStream_t stream) { @@ -1152,12 +1496,14 @@ static int cuda_stream_read_job_upload( if (job) job->errnum = EINVAL; return 0; } + pthread_mutex_lock(&g_stream_read_upload_mutex); cudaError_t err = cudaMemcpyAsync(job->dst, job->host_buf, (size_t)job->bytes, cudaMemcpyHostToDevice, stream); if (err == cudaSuccess) err = cudaStreamSynchronize(stream); + pthread_mutex_unlock(&g_stream_read_upload_mutex); if (err != cudaSuccess) { fprintf(stderr, DS4_GPU_LOG_PREFIX "streaming read-worker upload failed: %s\n", @@ -3383,10 +3729,28 @@ static const char *cuda_model_range_ptr(const void *model_map, uint64_t offset, auto exact = g_model_range_by_offset.find(offset); if (exact != g_model_range_by_offset.end()) { const cuda_model_range &r = g_model_ranges[exact->second]; - if (r.host_base == model_map && end >= offset && bytes <= r.bytes) return r.device_ptr; + if (r.host_base == model_map && end >= offset && bytes <= r.bytes && + !r.host_registered) { + return r.device_ptr; + } } + /* + * Prefer device-resident copies over host-registered mmap windows: the + * GPU reads registered host pages at ~7-8 GB/s on this APU (measured on + * the attention output projections), an order of magnitude slower than + * GTT-resident copies. Host-registered ranges created early (prefill + * runs before the static decode map install) must not shadow the device + * copies installed later. + */ for (const cuda_model_range &r : g_model_ranges) { - if (r.host_base == model_map && offset >= r.offset && end >= offset && end <= r.offset + r.bytes) { + if (r.host_base == model_map && !r.host_registered && + offset >= r.offset && end >= offset && end <= r.offset + r.bytes) { + return r.device_ptr + (offset - r.offset); + } + } + for (const cuda_model_range &r : g_model_ranges) { + if (r.host_base == model_map && r.host_registered && + offset >= r.offset && end >= offset && end <= r.offset + r.bytes) { return r.device_ptr + (offset - r.offset); } if (r.host_base == model_map && r.host_registered && r.registered_base && r.registered_device_base) { @@ -3466,6 +3830,77 @@ static const char *cuda_model_range_ptr(const void *model_map, uint64_t offset, return (const char *)dev; } +/* + * Materialize a device-resident copy of a model range, bypassing (and then + * retiring) any host-registered mmap window that covers it. Used by the + * streaming span installs: weights that stay hot on the GPU (attention + * projections, MTP draft tables) must not be served from registered host + * pages (~7-8 GB/s reads on this APU vs GTT-resident copies). + */ +static const char *cuda_model_range_ptr_device(const void *model_map, uint64_t offset, uint64_t bytes, const char *what) { + if (bytes == 0) return cuda_model_ptr(model_map, offset); + if (cuda_model_image_owned(model_map)) return cuda_model_ptr(model_map, offset); + const uint64_t end = offset + bytes; + if (end < offset) return NULL; + for (const cuda_model_range &r : g_model_ranges) { + if (r.host_base == model_map && !r.host_registered && + offset >= r.offset && end <= r.offset + r.bytes) { + return r.device_ptr + (offset - r.offset); + } + } + + void *dev = NULL; + cudaError_t err = cudaMalloc(&dev, (size_t)bytes); + if (err != cudaSuccess) { + (void)cudaGetLastError(); + fprintf(stderr, DS4_GPU_LOG_PREFIX "device model range alloc failed for %s (%.2f MiB): %s\n", + what ? what : "weights", (double)bytes / 1048576.0, cudaGetErrorString(err)); + return NULL; + } + const char *src = (const char *)model_map + offset; + const uint64_t chunk = 64ull * 1024ull * 1024ull; + for (uint64_t done = 0; done < bytes; done += chunk) { + uint64_t n = bytes - done < chunk ? bytes - done : chunk; + err = cudaMemcpy((char *)dev + done, src + done, (size_t)n, cudaMemcpyHostToDevice); + if (err != cudaSuccess) { + fprintf(stderr, DS4_GPU_LOG_PREFIX "device model range copy failed for %s: %s\n", + what ? what : "weights", cudaGetErrorString(err)); + (void)cudaFree(dev); + (void)cudaGetLastError(); + return NULL; + } + } + + /* Retire host-registered windows now shadowed by the device copy: they + * would only pin page-cache RAM. Entries fully inside [offset,end) go + * away; partial overlaps are kept (still needed for their tails). */ + size_t w = 0; + for (size_t i = 0; i < g_model_ranges.size(); i++) { + cuda_model_range &r = g_model_ranges[i]; + const int retire = + r.host_base == model_map && r.host_registered && + r.offset >= offset && r.offset + r.bytes <= end; + if (retire) { + if (r.registered_base) (void)cudaHostUnregister(r.registered_base); + continue; + } + if (w != i) g_model_ranges[w] = g_model_ranges[i]; + w++; + } + if (w != g_model_ranges.size()) { + g_model_ranges.resize(w); + g_model_range_by_offset.clear(); + for (size_t i = 0; i < g_model_ranges.size(); i++) { + g_model_range_by_offset[g_model_ranges[i].offset] = i; + } + } + + g_model_ranges.push_back({model_map, offset, bytes, (char *)dev, NULL, NULL, 0, 0, 0}); + g_model_range_by_offset[offset] = g_model_ranges.size() - 1u; + g_model_range_bytes += bytes; + return (const char *)dev; +} + static int cuda_model_range_is_cached(const void *model_map, uint64_t offset, uint64_t bytes) { if (bytes == 0) return 1; if (cuda_model_image_owned(model_map)) return 1; @@ -4003,6 +4438,14 @@ static void cuda_model_discard_source_pages(const void *model_map, uint64_t mode } static void cuda_model_drop_file_pages(uint64_t offset, uint64_t bytes) { + /* By default we drop the page cache for each streamed range so the 81 GB + * model does not evict everything else on a RAM-tight box. When there is + * spare RAM (e.g. after shrinking the BIOS VRAM carve), keeping the pages + * lets the OS page cache act as a free, self-evicting (never-OOM) host + * cache for hot experts at ~RAM speed. DS4_ROCM_KEEP_PAGE_CACHE=1 opts in. */ + static int keep = -1; + if (keep < 0) keep = getenv("DS4_ROCM_KEEP_PAGE_CACHE") != NULL ? 1 : 0; + if (keep) return; #if defined(POSIX_FADV_DONTNEED) if (g_model_fd < 0 || bytes == 0) return; (void)posix_fadvise(g_model_fd, (off_t)offset, (off_t)bytes, POSIX_FADV_DONTNEED); @@ -4451,9 +4894,11 @@ extern "C" void ds4_gpu_cleanup(void) { (void)cudaStreamDestroy(g_selected_readback_stream); g_selected_readback_stream = NULL; } - if (g_selected_readback_event) { - (void)cudaEventDestroy(g_selected_readback_event); - g_selected_readback_event = NULL; + for (uint32_t i = 0; i < DS4_SELECTED_READBACK_EVENT_RING; i++) { + if (g_selected_readback_events[i]) { + (void)cudaEventDestroy(g_selected_readback_events[i]); + g_selected_readback_events[i] = NULL; + } } g_selected_readback_event_value = 0; cuda_model_image_release_all(); @@ -4607,7 +5052,22 @@ extern "C" int ds4_gpu_set_model_map(const void *model_map, uint64_t model_size) const int multi_model = g_model_host_base != NULL && (g_model_host_base != model_map || g_model_registered_size != model_size); - cuda_model_range_release_all(); + const int same_map_resized = + g_model_host_base == model_map && g_model_registered_size != model_size; + if (!multi_model || same_map_resized) { + /* First registration, or the same mapping re-registered with a new + * size (a genuine remap): cached ranges for it are stale. */ + cuda_model_range_release_all(); + } + /* + * Cached model ranges are keyed by (host_base, offset), so mappings for + * several models (main + MTP support model) can coexist in the registry. + * Do NOT release them when switching between live mappings: the streaming + * session re-registers the main model spans after the MTP model loads, + * and releasing here would evict the MTP draft's device-resident expert + * tables (routed MoE then fails per token at decode time). The full + * release still happens in ds4_gpu_cleanup() at engine close. + */ cuda_q8_f16_cache_release_all(); g_q8_f16_disabled_after_oom = 0; g_q8_f16_budget_notice_printed = 0; diff --git a/tests/qwen/normalize_qwen36_gguf.py b/tests/qwen/normalize_qwen36_gguf.py new file mode 100644 index 000000000..fdf53a9a7 --- /dev/null +++ b/tests/qwen/normalize_qwen36_gguf.py @@ -0,0 +1,154 @@ +#!/usr/bin/env python3 +"""Normalize the Unsloth Qwen3.6-35B-A3B UD-Q4_K_S GGUF into the DS4 artifact. + +Reproduces the qualified-artifact provenance documented in tests/qwen/README.md: + - converts the Q6_K `ffn_down_exps` banks to the uniform Q4_K cache layout + - converts `output.weight` from Q6_K to Q8_0 + - restores the official padding token ID (248044) + - restores the canonical chat template (qwen36_chat_template.jinja) +Every other tensor payload is copied byte-identical (the container is +rewritten, so absolute offsets may differ; DS4 validates inventory and +payload sizes, not file offsets). + +Requantization uses ggml's reference quantizers (ggml_quantize_chunk) via +ctypes — the same code path llama.cpp uses — so the produced Q4_K/Q8_0 +blocks follow the reference layout DS4 validates. + +Usage: + normalize_qwen36_gguf.py SRC.gguf DST.gguf \ + --ggml-lib llama.cpp/build-cpu/bin/libggml-base.so \ + --gguf-py llama.cpp/gguf-py \ + --chat-template tests/qwen/qwen36_chat_template.jinja +""" + +import argparse +import ctypes +import sys +from pathlib import Path + +EXPECTED_PAD_ID = 248044 +GGML_TYPE_Q8_0 = 8 +GGML_TYPE_Q4_K = 12 + + +def main() -> int: + ap = argparse.ArgumentParser() + ap.add_argument("src") + ap.add_argument("dst") + ap.add_argument("--ggml-lib", required=True) + ap.add_argument("--gguf-py", required=True) + ap.add_argument("--chat-template", required=True) + args = ap.parse_args() + + sys.path.insert(0, args.gguf_py) + import gguf # noqa: E402 + from gguf import GGUFReader, GGUFWriter, GGMLQuantizationType # noqa: E402 + from gguf.constants import GGUFValueType # noqa: E402 + from gguf.quants import quant_shape_from_byte_shape # noqa: E402 + import numpy as np # noqa: E402 + + ggml = ctypes.CDLL(args.ggml_lib) + ggml.ggml_quantize_chunk.restype = ctypes.c_size_t + ggml.ggml_quantize_chunk.argtypes = [ + ctypes.c_int, ctypes.POINTER(ctypes.c_float), ctypes.c_void_p, + ctypes.c_int64, ctypes.c_int64, ctypes.c_int64, + ctypes.POINTER(ctypes.c_float), + ] + + template = Path(args.chat_template).read_text(encoding="utf-8") + reader = GGUFReader(args.src) + + arch = str(reader.get_field("general.architecture").contents()) + if arch != "qwen35moe": + print(f"unexpected architecture {arch!r}", file=sys.stderr) + return 1 + + writer = GGUFWriter(args.dst, arch, endianess=reader.endianess) + + # --- metadata: mirror gguf_new_metadata.py's copy loop --- + pad_key = "tokenizer.ggml.padding_token_id" + saw_pad = False + for field in reader.fields.values(): + if (field.name == "general.architecture" + or field.name.startswith("GGUF.")): + continue + if field.name.startswith("tokenizer.chat_template"): + continue # replaced below with the canonical template + val_type = field.types[0] + sub_type = (field.types[-1] + if val_type == GGUFValueType.ARRAY else None) + value = field.contents() + if field.name == pad_key: + print(f"padding_token_id: {value} -> {EXPECTED_PAD_ID}") + value = EXPECTED_PAD_ID + saw_pad = True + writer.add_key_value(field.name, value, val_type, sub_type=sub_type) + if not saw_pad: + print(f"padding_token_id: absent -> {EXPECTED_PAD_ID}") + writer.add_key_value(pad_key, EXPECTED_PAD_ID, GGUFValueType.UINT32) + writer.add_chat_template(template) + + # --- identify the Q6_K set to convert (fail closed) --- + q6_names = sorted(t.name for t in reader.tensors + if t.tensor_type == GGMLQuantizationType.Q6_K) + bad = [n for n in q6_names + if not (n.endswith("ffn_down_exps.weight") or n == "output.weight")] + if bad or "output.weight" not in q6_names: + print(f"unexpected Q6_K tensor set: {q6_names}", file=sys.stderr) + return 1 + print(f"converting {len(q6_names)} Q6_K tensors: {q6_names}") + + def requantize(tensor, target: int): + ttype = GGMLQuantizationType(target) + block, tsize = gguf.constants.GGML_QUANT_SIZES[ttype] + logical = quant_shape_from_byte_shape(tensor.data.shape, + tensor.tensor_type) + n_per_row = int(logical[-1]) + nrows = int(np.prod(logical[:-1], dtype=np.int64)) if len(logical) > 1 else 1 + f32 = gguf.quants.dequantize(tensor.data, tensor.tensor_type) + f32 = np.ascontiguousarray(f32, dtype=np.float32).reshape(nrows, + n_per_row) + row_bytes = (n_per_row // block) * tsize + out = np.empty((*logical[:-1], row_bytes), dtype=np.uint8) + written = ggml.ggml_quantize_chunk( + target, + f32.ctypes.data_as(ctypes.POINTER(ctypes.c_float)), + out.ctypes.data_as(ctypes.c_void_p), + 0, nrows, n_per_row, None) + assert written == out.nbytes, (tensor.name, written, out.nbytes) + return out, ttype + + # --- tensor info pass, then data pass (same order) --- + converted = {} + total = 0 + for t in reader.tensors: + if t.tensor_type == GGMLQuantizationType.Q6_K: + target = (GGML_TYPE_Q8_0 if t.name == "output.weight" + else GGML_TYPE_Q4_K) + data, dtype = requantize(t, target) + converted[t.name] = data + print(f" {t.name}: Q6_K -> {dtype.name} " + f"({t.n_bytes} -> {data.nbytes} bytes)") + writer.add_tensor_info(t.name, data.shape, data.dtype, + data.nbytes, dtype) + total += data.nbytes + else: + writer.add_tensor_info(t.name, t.data.shape, t.data.dtype, + t.data.nbytes, t.tensor_type) + total += t.data.nbytes + + print(f"tensor payload total: {total} bytes over {len(reader.tensors)} tensors") + writer.write_header_to_file() + writer.write_kv_data_to_file() + writer.write_ti_data_to_file() + for t in reader.tensors: + data = converted.get(t.name) + writer.write_tensor_data(data if data is not None else t.data, + tensor_endianess=reader.endianess) + writer.close() + print("done") + return 0 + + +if __name__ == "__main__": + sys.exit(main()) diff --git a/tests/qwen/qwen36_rocm_cpu_parity_cpu.json b/tests/qwen/qwen36_rocm_cpu_parity_cpu.json new file mode 100644 index 000000000..373ccaed9 --- /dev/null +++ b/tests/qwen/qwen36_rocm_cpu_parity_cpu.json @@ -0,0 +1,13 @@ +{ + "source":"ds4", + "prompt_tokens":23, + "ctx":2048, + "top_k":20, + "steps":[ + {"step":0,"selected":{"id":8160,"text":"Here","bytes":[72,101,114,101]},"top_logprobs":[{"token":{"id":8160,"text":"Here","bytes":[72,101,114,101]},"logit":26.3645287,"logprob":-0.00633561937},{"token":{"id":90700,"text":"Thinking","bytes":[84,104,105,110,107,105,110,103]},"logit":20.9081535,"logprob":-5.46271086},{"token":{"id":760,"text":"The","bytes":[84,104,101]},"logit":20.025013,"logprob":-6.34585142},{"token":{"id":77264,"text":"Hmm","bytes":[72,109,109]},"logit":16.7915306,"logprob":-9.57933331},{"token":{"id":97237,"text":"用户","bytes":[231,148,168,230,136,183]},"logit":16.2901764,"logprob":-10.0806875},{"token":{"id":6527,"text":"here","bytes":[104,101,114,101]},"logit":15.9703617,"logprob":-10.4005022},{"token":{"id":9764,"text":"Let","bytes":[76,101,116]},"logit":15.659955,"logprob":-10.7109089},{"token":{"id":31248,"text":"Okay","bytes":[79,107,97,121]},"logit":15.6425457,"logprob":-10.7283182},{"token":{"id":1421,"text":"User","bytes":[85,115,101,114]},"logit":15.3505135,"logprob":-11.0203505},{"token":{"id":107680,"text":"嗯","bytes":[229,151,175]},"logit":14.8491106,"logprob":-11.5217533},{"token":{"id":37405,"text":"Think","bytes":[84,104,105,110,107]},"logit":14.7588625,"logprob":-11.6120014},{"token":{"id":1596,"text":"We","bytes":[87,101]},"logit":14.6933584,"logprob":-11.6775055},{"token":{"id":97490,"text":"这里","bytes":[232,191,153,233,135,140]},"logit":14.6026001,"logprob":-11.7682638},{"token":{"id":248069,"text":"","bytes":[60,47,116,104,105,110,107,62]},"logit":14.5207281,"logprob":-11.8501358},{"token":{"id":13784,"text":"Wait","bytes":[87,97,105,116]},"logit":14.200161,"logprob":-12.1707029},{"token":{"id":2381,"text":"For","bytes":[70,111,114]},"logit":14.1980619,"logprob":-12.172802},{"token":{"id":5514,"text":" Here","bytes":[32,72,101,114,101]},"logit":13.9533691,"logprob":-12.4174948},{"token":{"id":1919,"text":"This","bytes":[84,104,105,115]},"logit":13.9133081,"logprob":-12.4575558},{"token":{"id":58523,"text":"thought","bytes":[116,104,111,117,103,104,116]},"logit":13.9035168,"logprob":-12.4673471},{"token":{"id":79420,"text":"thinking","bytes":[116,104,105,110,107,105,110,103]},"logit":13.6566639,"logprob":-12.7142}]}, + {"step":1,"selected":{"id":579,"text":"'s","bytes":[39,115]},"top_logprobs":[{"token":{"id":579,"text":"'s","bytes":[39,115]},"logit":28.9169903,"logprob":-1.60773889e-05},{"token":{"id":513,"text":" are","bytes":[32,97,114,101]},"logit":16.6726284,"logprob":-12.2443781},{"token":{"id":725,"text":"’s","bytes":[226,128,153,115]},"logit":16.0319366,"logprob":-12.8850698},{"token":{"id":369,"text":" is","bytes":[32,105,115]},"logit":15.9129763,"logprob":-13.0040302},{"token":{"id":2224,"text":"'re","bytes":[39,114,101]},"logit":15.8135166,"logprob":-13.1034899},{"token":{"id":11,"text":",","bytes":[44]},"logit":15.326745,"logprob":-13.5902615},{"token":{"id":4035,"text":"'d","bytes":[39,100]},"logit":14.1342802,"logprob":-14.7827263},{"token":{"id":2688,"text":"'m","bytes":[39,109]},"logit":13.9989977,"logprob":-14.9180088},{"token":{"id":1144,"text":" need","bytes":[32,110,101,101,100]},"logit":13.6401329,"logprob":-15.2768736},{"token":{"id":12887,"text":"'S","bytes":[39,83]},"logit":13.6071291,"logprob":-15.3098774},{"token":{"id":303,"text":" in","bytes":[32,105,110]},"logit":13.2463465,"logprob":-15.67066},{"token":{"id":728,"text":" me","bytes":[32,109,101]},"logit":12.9936876,"logprob":-15.9233189},{"token":{"id":2400,"text":" full","bytes":[32,102,117,108,108]},"logit":12.9626446,"logprob":-15.9543619},{"token":{"id":364,"text":" for","bytes":[32,102,111,114]},"logit":12.8884306,"logprob":-16.0285759},{"token":{"id":3172,"text":"'ll","bytes":[39,108,108]},"logit":12.4917727,"logprob":-16.4252338},{"token":{"id":2908,"text":"'ve","bytes":[39,118,101]},"logit":12.304842,"logprob":-16.6121635},{"token":{"id":328,"text":" \"","bytes":[32,34]},"logit":12.261776,"logprob":-16.6552296},{"token":{"id":1156,"text":" user","bytes":[32,117,115,101,114]},"logit":12.1131907,"logprob":-16.8038158},{"token":{"id":82,"text":"s","bytes":[115]},"logit":12.1107435,"logprob":-16.806263},{"token":{"id":264,"text":" a","bytes":[32,97]},"logit":11.8687143,"logprob":-17.0482922}]}, + {"step":2,"selected":{"id":264,"text":" a","bytes":[32,97]},"top_logprobs":[{"token":{"id":264,"text":" a","bytes":[32,97]},"logit":28.8784103,"logprob":-0.000102516737},{"token":{"id":7047,"text":" thinking","bytes":[32,116,104,105,110,107,105,110,103]},"logit":18.9177227,"logprob":-9.96078968},{"token":{"id":821,"text":" my","bytes":[32,109,121]},"logit":18.8193455,"logprob":-10.0591669},{"token":{"id":1817,"text":" process","bytes":[32,112,114,111,99,101,115,115]},"logit":16.2643642,"logprob":-12.6141481},{"token":{"id":449,"text":" an","bytes":[32,97,110]},"logit":15.5150623,"logprob":-13.3634501},{"token":{"id":1156,"text":" user","bytes":[32,117,115,101,114]},"logit":15.448514,"logprob":-13.4299984},{"token":{"id":279,"text":" the","bytes":[32,116,104,101]},"logit":15.0898628,"logprob":-13.7886496},{"token":{"id":303,"text":" in","bytes":[32,105,110]},"logit":14.6273785,"logprob":-14.2511339},{"token":{"id":296,"text":" o","bytes":[32,111]},"logit":13.7528677,"logprob":-15.1256447},{"token":{"id":353,"text":" I","bytes":[32,73]},"logit":13.7449436,"logprob":-15.1335688},{"token":{"id":310,"text":" to","bytes":[32,116,111]},"logit":13.6215048,"logprob":-15.2570076},{"token":{"id":328,"text":" \"","bytes":[32,34]},"logit":13.3874712,"logprob":-15.4910412},{"token":{"id":1683,"text":" think","bytes":[32,116,104,105,110,107]},"logit":13.356246,"logprob":-15.5222664},{"token":{"id":4821,"text":" unit","bytes":[32,117,110,105,116]},"logit":13.1761332,"logprob":-15.7023792},{"token":{"id":1602,"text":" being","bytes":[32,98,101,105,110,103]},"logit":13.1016912,"logprob":-15.7768211},{"token":{"id":295,"text":" m","bytes":[32,109]},"logit":13.0432177,"logprob":-15.8352947},{"token":{"id":913,"text":" type","bytes":[32,116,121,112,101]},"logit":12.7224064,"logprob":-16.1561069},{"token":{"id":974,"text":" off","bytes":[32,111,102,102]},"logit":12.6834116,"logprob":-16.1951008},{"token":{"id":4732,"text":" outside","bytes":[32,111,117,116,115,105,100,101]},"logit":12.574173,"logprob":-16.3043404},{"token":{"id":685,"text":" up","bytes":[32,117,112]},"logit":12.5073423,"logprob":-16.37117}]}, + {"step":3,"selected":{"id":7047,"text":" thinking","bytes":[32,116,104,105,110,107,105,110,103]},"top_logprobs":[{"token":{"id":7047,"text":" thinking","bytes":[32,116,104,105,110,107,105,110,103]},"logit":31.5377865,"logprob":-2.20038655e-05},{"token":{"id":9117,"text":" sea","bytes":[32,115,101,97]},"logit":20.5045261,"logprob":-11.0332823},{"token":{"id":99496,"text":"思考","bytes":[230,128,157,232,128,131]},"logit":18.4133682,"logprob":-13.1244402},{"token":{"id":50534,"text":" Thinking","bytes":[32,84,104,105,110,107,105,110,103]},"logit":17.6258259,"logprob":-13.9119825},{"token":{"id":90700,"text":"Thinking","bytes":[84,104,105,110,107,105,110,103]},"logit":17.1496849,"logprob":-14.3881235},{"token":{"id":8008,"text":" feeling","bytes":[32,102,101,101,108,105,110,103]},"logit":16.8346825,"logprob":-14.703126},{"token":{"id":3874,"text":" description","bytes":[32,100,101,115,99,114,105,112,116,105,111,110]},"logit":16.8215923,"logprob":-14.7162161},{"token":{"id":79420,"text":"thinking","bytes":[116,104,105,110,107,105,110,103]},"logit":16.2964134,"logprob":-15.241395},{"token":{"id":1683,"text":" think","bytes":[32,116,104,105,110,107]},"logit":16.2397499,"logprob":-15.2980585},{"token":{"id":3272,"text":" thought","bytes":[32,116,104,111,117,103,104,116]},"logit":15.4231186,"logprob":-16.1146908},{"token":{"id":14593,"text":" Sea","bytes":[32,83,101,97]},"logit":15.2089996,"logprob":-16.3288097},{"token":{"id":12463,"text":" considering","bytes":[32,99,111,110,115,105,100,101,114,105,110,103]},"logit":15.0542517,"logprob":-16.4835567},{"token":{"id":31626,"text":" reasoning","bytes":[32,114,101,97,115,111,110,105,110,103]},"logit":14.7922468,"logprob":-16.7455616},{"token":{"id":26884,"text":" descriptions","bytes":[32,100,101,115,99,114,105,112,116,105,111,110,115]},"logit":14.3499689,"logprob":-17.1878395},{"token":{"id":21988,"text":" describing","bytes":[32,100,101,115,99,114,105,98,105,110,103]},"logit":14.1731405,"logprob":-17.3646679},{"token":{"id":6236,"text":" meeting","bytes":[32,109,101,101,116,105,110,103]},"logit":14.0771875,"logprob":-17.4606209},{"token":{"id":99150,"text":"想到","bytes":[230,131,179,229,136,176]},"logit":14.066309,"logprob":-17.4715004},{"token":{"id":5732,"text":" starting","bytes":[32,115,116,97,114,116,105,110,103]},"logit":13.9979496,"logprob":-17.5398598},{"token":{"id":1156,"text":" user","bytes":[32,117,115,101,114]},"logit":13.847106,"logprob":-17.6907024},{"token":{"id":709,"text":" function","bytes":[32,102,117,110,99,116,105,111,110]},"logit":13.8331747,"logprob":-17.7046337}]}, + {"step":4,"selected":{"id":1817,"text":" process","bytes":[32,112,114,111,99,101,115,115]},"top_logprobs":[{"token":{"id":1817,"text":" process","bytes":[32,112,114,111,99,101,115,115]},"logit":28.0503674,"logprob":-4.14358328e-05},{"token":{"id":8240,"text":" sequence","bytes":[32,115,101,113,117,101,110,99,101]},"logit":16.2345161,"logprob":-11.8158922},{"token":{"id":7047,"text":" thinking","bytes":[32,116,104,105,110,107,105,110,103]},"logit":15.4725637,"logprob":-12.5778446},{"token":{"id":97424,"text":"过程","bytes":[232,191,135,231,168,139]},"logit":15.3419323,"logprob":-12.7084761},{"token":{"id":2193,"text":" context","bytes":[32,99,111,110,116,101,120,116]},"logit":14.9792595,"logprob":-13.0711489},{"token":{"id":27375,"text":" proces","bytes":[32,112,114,111,99,101,115]},"logit":14.8572407,"logprob":-13.1931677},{"token":{"id":1156,"text":" user","bytes":[32,117,115,101,114]},"logit":14.6157913,"logprob":-13.434617},{"token":{"id":10978,"text":"_process","bytes":[95,112,114,111,99,101,115,115]},"logit":14.5034294,"logprob":-13.546979},{"token":{"id":17128,"text":" processor","bytes":[32,112,114,111,99,101,115,115,111,114]},"logit":14.2893429,"logprob":-13.7610655},{"token":{"id":80828,"text":" procession","bytes":[32,112,114,111,99,101,115,115,105,111,110]},"logit":13.947381,"logprob":-14.1030273},{"token":{"id":3618,"text":" box","bytes":[32,98,111,120]},"logit":13.5775976,"logprob":-14.4728107},{"token":{"id":709,"text":" function","bytes":[32,102,117,110,99,116,105,111,110]},"logit":13.5747595,"logprob":-14.4756489},{"token":{"id":4480,"text":"process","bytes":[112,114,111,99,101,115,115]},"logit":13.5438938,"logprob":-14.5065145},{"token":{"id":9640,"text":" prompt","bytes":[32,112,114,111,109,112,116]},"logit":13.4515285,"logprob":-14.5988798},{"token":{"id":11036,"text":" processes","bytes":[32,112,114,111,99,101,115,115,101,115]},"logit":13.4487524,"logprob":-14.601656},{"token":{"id":8427,"text":" processing","bytes":[32,112,114,111,99,101,115,115,105,110,103]},"logit":13.444418,"logprob":-14.6059904},{"token":{"id":1965,"text":" response","bytes":[32,114,101,115,112,111,110,115,101]},"logit":13.3923044,"logprob":-14.6581039},{"token":{"id":3296,"text":" question","bytes":[32,113,117,101,115,116,105,111,110]},"logit":13.3486319,"logprob":-14.7017765},{"token":{"id":176460,"text":" processus","bytes":[32,112,114,111,99,101,115,115,117,115]},"logit":13.3467045,"logprob":-14.7037039},{"token":{"id":1785,"text":" system","bytes":[32,115,121,115,116,101,109]},"logit":13.2427607,"logprob":-14.8076477}]} + ] +} diff --git a/tests/qwen/qwen36_rocm_cpu_parity_gpu.json b/tests/qwen/qwen36_rocm_cpu_parity_gpu.json new file mode 100644 index 000000000..3e2055f10 --- /dev/null +++ b/tests/qwen/qwen36_rocm_cpu_parity_gpu.json @@ -0,0 +1,13 @@ +{ + "source":"ds4", + "prompt_tokens":23, + "ctx":2048, + "top_k":20, + "steps":[ + {"step":0,"selected":{"id":8160,"text":"Here","bytes":[72,101,114,101]},"top_logprobs":[{"token":{"id":8160,"text":"Here","bytes":[72,101,114,101]},"logit":26.4651203,"logprob":-0.00880505145},{"token":{"id":90700,"text":"Thinking","bytes":[84,104,105,110,107,105,110,103]},"logit":21.4824028,"logprob":-4.99152279},{"token":{"id":760,"text":"The","bytes":[84,104,101]},"logit":20.0973778,"logprob":-6.37654781},{"token":{"id":77264,"text":"Hmm","bytes":[72,109,109]},"logit":16.70047,"logprob":-9.77345562},{"token":{"id":97237,"text":"用户","bytes":[231,148,168,230,136,183]},"logit":16.380703,"logprob":-10.0932226},{"token":{"id":6527,"text":"here","bytes":[104,101,114,101]},"logit":15.8159809,"logprob":-10.6579447},{"token":{"id":1421,"text":"User","bytes":[85,115,101,114]},"logit":15.5212784,"logprob":-10.9526472},{"token":{"id":31248,"text":"Okay","bytes":[79,107,97,121]},"logit":15.3949118,"logprob":-11.0790138},{"token":{"id":9764,"text":"Let","bytes":[76,101,116]},"logit":15.3086329,"logprob":-11.1652927},{"token":{"id":37405,"text":"Think","bytes":[84,104,105,110,107]},"logit":14.8650551,"logprob":-11.6088705},{"token":{"id":97490,"text":"这里","bytes":[232,191,153,233,135,140]},"logit":14.6166658,"logprob":-11.8572598},{"token":{"id":107680,"text":"嗯","bytes":[229,151,175]},"logit":14.5236959,"logprob":-11.9502296},{"token":{"id":1596,"text":"We","bytes":[87,101]},"logit":14.4066362,"logprob":-12.0672894},{"token":{"id":248069,"text":"","bytes":[60,47,116,104,105,110,107,62]},"logit":14.3791637,"logprob":-12.0947618},{"token":{"id":2381,"text":"For","bytes":[70,111,114]},"logit":14.3306446,"logprob":-12.143281},{"token":{"id":79420,"text":"thinking","bytes":[116,104,105,110,107,105,110,103]},"logit":14.0683899,"logprob":-12.4055357},{"token":{"id":58523,"text":"thought","bytes":[116,104,111,117,103,104,116]},"logit":13.9908848,"logprob":-12.4830408},{"token":{"id":5514,"text":" Here","bytes":[32,72,101,114,101]},"logit":13.9333534,"logprob":-12.5405722},{"token":{"id":99496,"text":"思考","bytes":[230,128,157,232,128,131]},"logit":13.7437401,"logprob":-12.7301855},{"token":{"id":13784,"text":"Wait","bytes":[87,97,105,116]},"logit":13.6402884,"logprob":-12.8336372}]}, + {"step":1,"selected":{"id":579,"text":"'s","bytes":[39,115]},"top_logprobs":[{"token":{"id":579,"text":"'s","bytes":[39,115]},"logit":28.8340626,"logprob":-1.50857331e-05},{"token":{"id":513,"text":" are","bytes":[32,97,114,101]},"logit":16.3498611,"logprob":-12.4842167},{"token":{"id":369,"text":" is","bytes":[32,105,115]},"logit":15.9522533,"logprob":-12.8818245},{"token":{"id":725,"text":"’s","bytes":[226,128,153,115]},"logit":15.9463663,"logprob":-12.8877115},{"token":{"id":2224,"text":"'re","bytes":[39,114,101]},"logit":15.5324249,"logprob":-13.3016529},{"token":{"id":11,"text":",","bytes":[44]},"logit":15.2031202,"logprob":-13.6309576},{"token":{"id":4035,"text":"'d","bytes":[39,100]},"logit":14.3280487,"logprob":-14.5060291},{"token":{"id":2688,"text":"'m","bytes":[39,109]},"logit":13.9793921,"logprob":-14.8546858},{"token":{"id":12887,"text":"'S","bytes":[39,83]},"logit":13.8920355,"logprob":-14.9420424},{"token":{"id":1144,"text":" need","bytes":[32,110,101,101,100]},"logit":13.4665651,"logprob":-15.3675127},{"token":{"id":303,"text":" in","bytes":[32,105,110]},"logit":12.7628098,"logprob":-16.0712681},{"token":{"id":364,"text":" for","bytes":[32,102,111,114]},"logit":12.6958141,"logprob":-16.1382637},{"token":{"id":2400,"text":" full","bytes":[32,102,117,108,108]},"logit":12.6271,"logprob":-16.2069778},{"token":{"id":728,"text":" me","bytes":[32,109,101]},"logit":12.4628849,"logprob":-16.3711929},{"token":{"id":2908,"text":"'ve","bytes":[39,118,101]},"logit":12.3610888,"logprob":-16.4729881},{"token":{"id":3172,"text":"'ll","bytes":[39,108,108]},"logit":12.2786512,"logprob":-16.5554256},{"token":{"id":328,"text":" \"","bytes":[32,34]},"logit":12.2781944,"logprob":-16.5558834},{"token":{"id":1156,"text":" user","bytes":[32,117,115,101,114]},"logit":12.2362633,"logprob":-16.5978146},{"token":{"id":82,"text":"s","bytes":[115]},"logit":12.2327423,"logprob":-16.6013355},{"token":{"id":914,"text":"'t","bytes":[39,116]},"logit":12.0121708,"logprob":-16.821907}]}, + {"step":2,"selected":{"id":264,"text":" a","bytes":[32,97]},"top_logprobs":[{"token":{"id":264,"text":" a","bytes":[32,97]},"logit":29.0423965,"logprob":-7.35700232e-05},{"token":{"id":821,"text":" my","bytes":[32,109,121]},"logit":18.8938465,"logprob":-10.1486235},{"token":{"id":7047,"text":" thinking","bytes":[32,116,104,105,110,107,105,110,103]},"logit":18.3613281,"logprob":-10.6811419},{"token":{"id":1817,"text":" process","bytes":[32,112,114,111,99,101,115,115]},"logit":16.5793457,"logprob":-12.4631243},{"token":{"id":1156,"text":" user","bytes":[32,117,115,101,114]},"logit":15.5150213,"logprob":-13.5274487},{"token":{"id":279,"text":" the","bytes":[32,116,104,101]},"logit":15.1303501,"logprob":-13.9121199},{"token":{"id":449,"text":" an","bytes":[32,97,110]},"logit":15.0499487,"logprob":-13.9925213},{"token":{"id":303,"text":" in","bytes":[32,105,110]},"logit":14.6810513,"logprob":-14.3614187},{"token":{"id":353,"text":" I","bytes":[32,73]},"logit":13.9311485,"logprob":-15.1113214},{"token":{"id":296,"text":" o","bytes":[32,111]},"logit":13.869812,"logprob":-15.172658},{"token":{"id":4821,"text":" unit","bytes":[32,117,110,105,116]},"logit":13.710084,"logprob":-15.332386},{"token":{"id":328,"text":" \"","bytes":[32,34]},"logit":13.4147472,"logprob":-15.6277227},{"token":{"id":295,"text":" m","bytes":[32,109]},"logit":13.3375998,"logprob":-15.7048702},{"token":{"id":913,"text":" type","bytes":[32,116,121,112,101]},"logit":13.2652874,"logprob":-15.7771826},{"token":{"id":1683,"text":" think","bytes":[32,116,104,105,110,107]},"logit":13.2039471,"logprob":-15.8385229},{"token":{"id":310,"text":" to","bytes":[32,116,111]},"logit":13.1517487,"logprob":-15.8907213},{"token":{"id":1602,"text":" being","bytes":[32,98,101,105,110,103]},"logit":12.9114323,"logprob":-16.1310387},{"token":{"id":974,"text":" off","bytes":[32,111,102,102]},"logit":12.563343,"logprob":-16.4791279},{"token":{"id":4732,"text":" outside","bytes":[32,111,117,116,115,105,100,101]},"logit":12.4369907,"logprob":-16.6054802},{"token":{"id":293,"text":" d","bytes":[32,100]},"logit":12.4093399,"logprob":-16.633131}]}, + {"step":3,"selected":{"id":7047,"text":" thinking","bytes":[32,116,104,105,110,107,105,110,103]},"top_logprobs":[{"token":{"id":7047,"text":" thinking","bytes":[32,116,104,105,110,107,105,110,103]},"logit":31.9806862,"logprob":-1.84522323e-05},{"token":{"id":9117,"text":" sea","bytes":[32,115,101,97]},"logit":20.7030392,"logprob":-11.2776651},{"token":{"id":3874,"text":" description","bytes":[32,100,101,115,99,114,105,112,116,105,111,110]},"logit":18.5572586,"logprob":-13.4234457},{"token":{"id":99496,"text":"思考","bytes":[230,128,157,232,128,131]},"logit":18.3761635,"logprob":-13.6045408},{"token":{"id":8008,"text":" feeling","bytes":[32,102,101,101,108,105,110,103]},"logit":18.111742,"logprob":-13.8689623},{"token":{"id":50534,"text":" Thinking","bytes":[32,84,104,105,110,107,105,110,103]},"logit":17.2601357,"logprob":-14.7205687},{"token":{"id":90700,"text":"Thinking","bytes":[84,104,105,110,107,105,110,103]},"logit":16.8952141,"logprob":-15.0854902},{"token":{"id":79420,"text":"thinking","bytes":[116,104,105,110,107,105,110,103]},"logit":16.4698772,"logprob":-15.5108271},{"token":{"id":3272,"text":" thought","bytes":[32,116,104,111,117,103,104,116]},"logit":16.1708317,"logprob":-15.8098726},{"token":{"id":1683,"text":" think","bytes":[32,116,104,105,110,107]},"logit":16.1689663,"logprob":-15.811738},{"token":{"id":31626,"text":" reasoning","bytes":[32,114,101,97,115,111,110,105,110,103]},"logit":15.435915,"logprob":-16.5447903},{"token":{"id":1156,"text":" user","bytes":[32,117,115,101,114]},"logit":15.384572,"logprob":-16.5961323},{"token":{"id":26884,"text":" descriptions","bytes":[32,100,101,115,99,114,105,112,116,105,111,110,115]},"logit":15.2972393,"logprob":-16.683466},{"token":{"id":1817,"text":" process","bytes":[32,112,114,111,99,101,115,115]},"logit":15.1948195,"logprob":-16.7858849},{"token":{"id":12463,"text":" considering","bytes":[32,99,111,110,115,105,100,101,114,105,110,103]},"logit":15.1519098,"logprob":-16.8287945},{"token":{"id":14593,"text":" Sea","bytes":[32,83,101,97]},"logit":15.0595312,"logprob":-16.9211731},{"token":{"id":21988,"text":" describing","bytes":[32,100,101,115,99,114,105,98,105,110,103]},"logit":14.9148207,"logprob":-17.0658836},{"token":{"id":6236,"text":" meeting","bytes":[32,109,101,101,116,105,110,103]},"logit":14.8479872,"logprob":-17.1327171},{"token":{"id":10976,"text":" thoughts","bytes":[32,116,104,111,117,103,104,116,115]},"logit":14.612422,"logprob":-17.3682823},{"token":{"id":5732,"text":" starting","bytes":[32,115,116,97,114,116,105,110,103]},"logit":14.258853,"logprob":-17.7218513}]}, + {"step":4,"selected":{"id":1817,"text":" process","bytes":[32,112,114,111,99,101,115,115]},"top_logprobs":[{"token":{"id":1817,"text":" process","bytes":[32,112,114,111,99,101,115,115]},"logit":29.0582008,"logprob":-2.38481298e-05},{"token":{"id":8240,"text":" sequence","bytes":[32,115,101,113,117,101,110,99,101]},"logit":16.6968384,"logprob":-12.3613863},{"token":{"id":97424,"text":"过程","bytes":[232,191,135,231,168,139]},"logit":16.3850746,"logprob":-12.6731501},{"token":{"id":7047,"text":" thinking","bytes":[32,116,104,105,110,107,105,110,103]},"logit":15.8587818,"logprob":-13.1994429},{"token":{"id":2193,"text":" context","bytes":[32,99,111,110,116,101,120,116]},"logit":15.4854555,"logprob":-13.5727692},{"token":{"id":27375,"text":" proces","bytes":[32,112,114,111,99,101,115]},"logit":15.3771286,"logprob":-13.6810961},{"token":{"id":10978,"text":"_process","bytes":[95,112,114,111,99,101,115,115]},"logit":15.1318502,"logprob":-13.9263744},{"token":{"id":1156,"text":" user","bytes":[32,117,115,101,114]},"logit":15.0163784,"logprob":-14.0418463},{"token":{"id":17128,"text":" processor","bytes":[32,112,114,111,99,101,115,115,111,114]},"logit":14.6207247,"logprob":-14.4375},{"token":{"id":9640,"text":" prompt","bytes":[32,112,114,111,109,112,116]},"logit":14.2864037,"logprob":-14.771821},{"token":{"id":4480,"text":"process","bytes":[112,114,111,99,101,115,115]},"logit":14.2633858,"logprob":-14.7948389},{"token":{"id":80828,"text":" procession","bytes":[32,112,114,111,99,101,115,115,105,111,110]},"logit":14.2171955,"logprob":-14.8410292},{"token":{"id":709,"text":" function","bytes":[32,102,117,110,99,116,105,111,110]},"logit":13.9554691,"logprob":-15.1027555},{"token":{"id":3296,"text":" question","bytes":[32,113,117,101,115,116,105,111,110]},"logit":13.9503088,"logprob":-15.1079159},{"token":{"id":176460,"text":" processus","bytes":[32,112,114,111,99,101,115,115,117,115]},"logit":13.9207258,"logprob":-15.1374989},{"token":{"id":25,"text":":","bytes":[58]},"logit":13.8134136,"logprob":-15.2448111},{"token":{"id":11036,"text":" processes","bytes":[32,112,114,111,99,101,115,115,101,115]},"logit":13.7603607,"logprob":-15.297864},{"token":{"id":11870,"text":" logic","bytes":[32,108,111,103,105,99]},"logit":13.6129761,"logprob":-15.4452486},{"token":{"id":1965,"text":" response","bytes":[32,114,101,115,112,111,110,115,101]},"logit":13.6086102,"logprob":-15.4496145},{"token":{"id":7188,"text":"Process","bytes":[80,114,111,99,101,115,115]},"logit":13.608387,"logprob":-15.4498377}]} + ] +} diff --git a/tests/test_qwen35_rocm.c b/tests/test_qwen35_rocm.c new file mode 100644 index 000000000..5ede3165b --- /dev/null +++ b/tests/test_qwen35_rocm.c @@ -0,0 +1,377 @@ +// Per-op parity harness: ROCm Qwen3.6 kernels vs the scalar reference +// (ds4_qwen_ref.c). Random inputs, host reference, device kernel, compare. +// Covers the ops that do not read model-map weights (the highest-risk GDN and +// GQA kernels among them). Weight-based ops (conv, rmsnorm_gated, controls, +// embedding) are validated separately once a model map is registered. + +#include "ds4_gpu.h" +#include "ds4_qwen_ref.h" + +#include +#include +#include +#include +#include + +static uint32_t g_seed = 0x1234567u; +static float frand(void) { + g_seed = g_seed * 1664525u + 1013904223u; + return ((float)(g_seed >> 8) / (float)(1u << 24)) * 2.0f - 1.0f; // [-1,1) +} + +static int compare(const char *name, const float *got, const float *want, + size_t n, float tol) { + double worst = 0.0; size_t at = 0; + for (size_t i = 0; i < n; i++) { + const float scale = fmaxf(1.0f, fabsf(want[i])); + const double e = fabs((double)got[i] - (double)want[i]) / scale; + if (!isfinite(got[i])) { fprintf(stderr, "%s[%zu] not finite\n", name, i); return 1; } + if (e > worst) { worst = e; at = i; } + } + if (worst > tol) { + fprintf(stderr, "FAIL %s: worst rel %.3g at %zu (got %.9g want %.9g)\n", + name, worst, at, (double)got[at], (double)want[at]); + return 1; + } + fprintf(stderr, "ok %s (worst rel %.3g over %zu)\n", name, worst, n); + return 0; +} + +// Helper: alloc a device tensor, upload host floats. +static ds4_gpu_tensor *up_f32(const float *h, size_t n) { + ds4_gpu_tensor *t = ds4_gpu_tensor_alloc(n * sizeof(float)); + if (!t) return NULL; + ds4_gpu_tensor_write(t, 0, h, n * sizeof(float)); + return t; +} +static void down_f32(const ds4_gpu_tensor *t, float *h, size_t n) { + ds4_gpu_tensor_read(t, 0, h, n * sizeof(float)); +} + +static int fails = 0; + +static void test_split_q_gate(void) { + enum { NT = 3, NH = 4, HD = 8 }; + float proj[NT*NH*2*HD], q_ref[NT*NH*HD], g_ref[NT*NH*HD]; + for (int i = 0; i < NT*NH*2*HD; i++) proj[i] = frand(); + ds4_qwen_ref_split_q_gate_f32(q_ref, g_ref, proj, NT, NH, HD); + ds4_gpu_tensor *dp = up_f32(proj, NT*NH*2*HD); + ds4_gpu_tensor *dq = ds4_gpu_tensor_alloc(NT*NH*HD*sizeof(float)); + ds4_gpu_tensor *dg = ds4_gpu_tensor_alloc(NT*NH*HD*sizeof(float)); + if (!ds4_gpu_qwen35_split_q_gate_batch_tensor(dq, dg, dp, NT, NH, HD)) { fprintf(stderr,"split dispatch failed\n"); fails++; return; } + ds4_gpu_synchronize(); + float q_got[NT*NH*HD], g_got[NT*NH*HD]; + down_f32(dq, q_got, NT*NH*HD); down_f32(dg, g_got, NT*NH*HD); + fails += compare("split_q_gate.query", q_got, q_ref, NT*NH*HD, 1e-6f); + fails += compare("split_q_gate.gate", g_got, g_ref, NT*NH*HD, 1e-6f); + ds4_gpu_tensor_free(dp); ds4_gpu_tensor_free(dq); ds4_gpu_tensor_free(dg); +} + +static void test_sigmoid_mul_elements(void) { + enum { N = 257 }; + float in[N], gate[N], ref[N]; + for (int i = 0; i < N; i++) { in[i] = frand(); gate[i] = frand()*4.0f; } + ds4_qwen_ref_sigmoid_gate_elements_f32(ref, in, gate, N); + ds4_gpu_tensor *di = up_f32(in, N), *dg = up_f32(gate, N); + ds4_gpu_tensor *dout = ds4_gpu_tensor_alloc(N*sizeof(float)); + if (!ds4_gpu_qwen35_sigmoid_mul_tensor(dout, di, dg, N, false)) { fprintf(stderr,"sigmoid_mul dispatch failed\n"); fails++; return; } + ds4_gpu_synchronize(); + float got[N]; down_f32(dout, got, N); + fails += compare("sigmoid_mul.elements", got, ref, N, 2e-6f); + ds4_gpu_tensor_free(di); ds4_gpu_tensor_free(dg); ds4_gpu_tensor_free(dout); +} + +static void test_sigmoid_mul_rows(void) { + enum { R = 5, W = 40 }; + float in[R*W], gate[R], ref[R*W]; + for (int i = 0; i < R*W; i++) in[i] = frand(); + for (int i = 0; i < R; i++) gate[i] = frand()*4.0f; + // reference: broadcast scalar gate over each row + for (int r = 0; r < R; r++) { + const float g = 1.0f / (1.0f + expf(-gate[r])); + for (int c = 0; c < W; c++) ref[r*W+c] = in[r*W+c] * g; + } + ds4_gpu_tensor *di = up_f32(in, R*W), *dg = up_f32(gate, R); + ds4_gpu_tensor *dout = ds4_gpu_tensor_alloc(R*W*sizeof(float)); + if (!ds4_gpu_qwen35_sigmoid_mul_rows_tensor(dout, di, dg, R, W)) { fprintf(stderr,"sigmoid_rows dispatch failed\n"); fails++; return; } + ds4_gpu_synchronize(); + float got[R*W]; down_f32(dout, got, R*W); + fails += compare("sigmoid_mul.rows", got, ref, R*W, 2e-6f); + ds4_gpu_tensor_free(di); ds4_gpu_tensor_free(dg); ds4_gpu_tensor_free(dout); +} + +static void test_rope(void) { + enum { NT = 2, NH = 3, HD = 16, NROT = 8 }; + float vals[NT*NH*HD], ref[NT*NH*HD]; + uint32_t pos[NT] = { 5, 37 }; + for (int i = 0; i < NT*NH*HD; i++) { vals[i] = frand(); ref[i] = vals[i]; } + ds4_qwen_ref_text_rope_f32(ref, pos, NT, NH, HD, NROT, 1000000.0f); + ds4_gpu_tensor *dv = up_f32(vals, NT*NH*HD); + ds4_gpu_tensor *dpos = ds4_gpu_tensor_alloc(NT*sizeof(uint32_t)); + ds4_gpu_tensor_write(dpos, 0, pos, NT*sizeof(uint32_t)); + if (!ds4_gpu_qwen35_rope_prefix_batch_tensor(dv, dpos, NT, NH, HD, NROT, 1000000.0f)) { fprintf(stderr,"rope dispatch failed\n"); fails++; return; } + ds4_gpu_synchronize(); + float got[NT*NH*HD]; down_f32(dv, got, NT*NH*HD); + // 5e-6: RoPE differs only in transcendental rounding (powf/cosf/sinf vs the + // reference libm), never in structure — the same tolerance the GQA online + // softmax uses for its exp() chain. + fails += compare("rope.batch", got, ref, NT*NH*HD, 5e-6f); + ds4_gpu_tensor_free(dv); ds4_gpu_tensor_free(dpos); +} + +static void test_gqa_prefill(void) { + // Reference GQA is over a whole sequence; prefill kernel writes per-token + // causal outputs into the same [token][head][dim] layout. + enum { NT = 6, QH = 8, KVH = 2, HD = 16 }; + float q[NT*QH*HD], k[NT*KVH*HD], v[NT*KVH*HD], ref[NT*QH*HD]; + for (int i = 0; i < NT*QH*HD; i++) q[i] = frand(); + for (int i = 0; i < NT*KVH*HD; i++) { k[i] = frand(); v[i] = frand(); } + ds4_qwen_ref_causal_gqa_f32(ref, q, k, v, NT, QH, KVH, HD); + ds4_gpu_tensor *dq = up_f32(q, NT*QH*HD), *dk = up_f32(k, NT*KVH*HD), *dv = up_f32(v, NT*KVH*HD); + ds4_gpu_tensor *dout = ds4_gpu_tensor_alloc(NT*QH*HD*sizeof(float)); + if (!ds4_gpu_qwen35_gqa_prefill_tensor(dout, dq, dk, dv, 0, NT, QH, KVH, HD)) { fprintf(stderr,"gqa_prefill dispatch failed\n"); fails++; return; } + ds4_gpu_synchronize(); + float got[NT*QH*HD]; down_f32(dout, got, NT*QH*HD); + fails += compare("gqa.prefill", got, ref, NT*QH*HD, 5e-6f); + ds4_gpu_tensor_free(dq); ds4_gpu_tensor_free(dk); ds4_gpu_tensor_free(dv); ds4_gpu_tensor_free(dout); +} + +static void test_gqa_decode(void) { + // Decode = the last query row of a causal sequence vs the full K/V cache. + enum { NKV = 20, QH = 8, KVH = 2, HD = 16 }; + float q[QH*HD], k[NKV*KVH*HD], v[NKV*KVH*HD], ref[QH*HD]; + for (int i = 0; i < QH*HD; i++) q[i] = frand(); + for (int i = 0; i < NKV*KVH*HD; i++) { k[i] = frand(); v[i] = frand(); } + // reference: single-query causal attention over NKV cached tokens + const float scale = 1.0f / sqrtf((float)HD); + const int qpk = QH / KVH; + for (int h = 0; h < QH; h++) { + const int kvh = h / qpk; + float mx = -INFINITY; + for (int t = 0; t < NKV; t++) { + float d = 0; for (int c = 0; c < HD; c++) d += q[h*HD+c]*k[(t*KVH+kvh)*HD+c]; + d *= scale; if (d > mx) mx = d; + } + float sum = 0, acc[HD]; for (int c=0;cfloat exact + ds4_gpu_tensor_free(dout); + (void)row_bytes; +} + +int main(void) { + if (!ds4_gpu_init()) { fprintf(stderr, "ds4_gpu_init failed\n"); return 2; } + test_split_q_gate(); + test_sigmoid_mul_elements(); + test_sigmoid_mul_rows(); + test_rope(); + test_router(); + test_gqa_prefill(); + test_gqa_decode(); + test_gd_step(128); // step_128 fast path + test_gd_step(64); // generic reduction path + test_gd_sequence_128(); + test_conv_step(); + test_conv_sequence(); + test_rmsnorm_gated(); + test_gd_controls(); + test_embedding_q8_0(); + ds4_gpu_cleanup(); + if (fails) { fprintf(stderr, "\n%d qwen35 ROCm parity check(s) FAILED\n", fails); return 1; } + fprintf(stderr, "\nAll qwen35 ROCm parity checks passed\n"); + return 0; +}