diff --git a/rocm/ds4_rocm_attention.cuh b/rocm/ds4_rocm_attention.cuh
index 0d688326b..ae022238a 100644
--- a/rocm/ds4_rocm_attention.cuh
+++ b/rocm/ds4_rocm_attention.cuh
@@ -8,6 +8,13 @@
 #define DS4_ROCM_ATTENTION_INDEXED_SCORE_CAP \
     (256u + DS4_ROCM_ATTENTION_INDEXED_TOPK_CAP)
 
+// ---------------------------------------------------------------------------
+// Vectorized prefill kernels — use float4 loads to reduce bus transactions
+// on unified-memory APUs (Strix Halo).  The scalar inner dot-product loop is
+// replaced with 128-bit vector loads + dot4_f32 to cut memory controller
+// pressure by ~4x for the common head_dim=512 case.
+// ---------------------------------------------------------------------------
+
 __global__ static void attention_prefill_raw_kernel(
         float *heads,
         const float *sinks,
@@ -29,11 +36,20 @@ __global__ static void attention_prefill_raw_kernel(
     __shared__ float denom;
     float scale = rsqrtf((float)head_dim);
     float local_max = sinks[h];
+    const uint32_t n4 = head_dim >> 2u;
+    const uint32_t tail = n4 << 2u;
     __syncthreads();
     for (uint32_t r = threadIdx.x; r < raw_count; r += blockDim.x) {
-        const float *kv = raw_kv + (uint64_t)(raw_start + r) * head_dim;
+        const float4 *kv4 = (const float4 *)(raw_kv + (uint64_t)(raw_start + r) * head_dim);
+        const float4 *q4  = (const float4 *)qh;
         float dot = 0.0f;
-        for (uint32_t d = 0; d < head_dim; d++) dot += qh[d] * kv[d];
+        #pragma unroll 16
+        for (uint32_t i = 0; i < n4; i++) {
+            const float4 qv  = q4[i];
+            const float4 kvv = kv4[i];
+            dot += qv.x * kvv.x + qv.y * kvv.y + qv.z * kvv.z + qv.w * kvv.w;
+        }
+        for (uint32_t d = tail; d < head_dim; d++) dot += qh[d] * (raw_kv[(uint64_t)(raw_start + r) * head_dim + d]);
         scores[r] = dot * scale;
         local_max = fmaxf(local_max, scores[r]);
     }
@@ -94,11 +110,20 @@ __global__ static void attention_prefill_mixed_kernel(
     float local_max = sinks[h];
     uint32_t n_score = raw_count + visible_comp;
     if (n_score > DS4_ROCM_ATTENTION_PREFILL_MIXED_SCORE_CAP) return;
+    const uint32_t n4 = head_dim >> 2u;
+    const uint32_t tail = n4 << 2u;
 
     for (uint32_t r = threadIdx.x; r < raw_count; r += blockDim.x) {
-        const float *kvrow = raw_kv + (uint64_t)(raw_start + r) * head_dim;
+        const float4 *kv4 = (const float4 *)(raw_kv + (uint64_t)(raw_start + r) * head_dim);
+        const float4 *q4  = (const float4 *)qh;
         float dot = 0.0f;
-        for (uint32_t d = 0; d < head_dim; d++) dot += qh[d] * kvrow[d];
+        #pragma unroll 16
+        for (uint32_t i = 0; i < n4; i++) {
+            const float4 qv  = q4[i];
+            const float4 kvv = kv4[i];
+            dot += qv.x * kvv.x + qv.y * kvv.y + qv.z * kvv.z + qv.w * kvv.w;
+        }
+        for (uint32_t d = tail; d < head_dim; d++) dot += qh[d] * (raw_kv[(uint64_t)(raw_start + r) * head_dim + d]);
         scores[r] = dot * scale;
         local_max = fmaxf(local_max, scores[r]);
     }
@@ -106,9 +131,16 @@ __global__ static void attention_prefill_mixed_kernel(
         float add = use_comp_mask ? comp_mask[(uint64_t)t * n_comp + c] : 0.0f;
         float s = -INFINITY;
         if (add > -1.0e20f) {
-            const float *kvrow = comp_kv + (uint64_t)c * head_dim;
+            const float4 *kv4 = (const float4 *)(comp_kv + (uint64_t)c * head_dim);
+            const float4 *q4  = (const float4 *)qh;
             float dot = 0.0f;
-            for (uint32_t d = 0; d < head_dim; d++) dot += qh[d] * kvrow[d];
+            #pragma unroll 16
+            for (uint32_t i = 0; i < n4; i++) {
+                const float4 qv  = q4[i];
+                const float4 kvv = kv4[i];
+                dot += qv.x * kvv.x + qv.y * kvv.y + qv.z * kvv.z + qv.w * kvv.w;
+            }
+            for (uint32_t d = tail; d < head_dim; d++) dot += qh[d] * (comp_kv[(uint64_t)c * head_dim + d]);
             s = dot * scale + add;
         }
         scores[raw_count + c] = s;
@@ -154,9 +186,6 @@ __global__ static void attention_prefill_raw_softmax_kernel(
     uint32_t h = blockIdx.y;
     if (t >= n_tokens) return;
     float *row = scores + ((uint64_t)h * n_tokens + t) * n_keys;
-    __shared__ float partial[256];
-    __shared__ float max_s;
-    __shared__ float denom;
     float local_max = sinks[h];
     for (uint32_t k = threadIdx.x; k < n_keys; k += blockDim.x) {
         bool valid = k <= t && (window == 0 || t - k < window);
@@ -164,28 +193,14 @@ __global__ static void attention_prefill_raw_softmax_kernel(
         row[k] = s;
         local_max = fmaxf(local_max, s);
     }
-    partial[threadIdx.x] = local_max;
-    __syncthreads();
-    for (uint32_t stride = blockDim.x >> 1; stride > 0; stride >>= 1) {
-        if (threadIdx.x < stride) partial[threadIdx.x] = fmaxf(partial[threadIdx.x], partial[threadIdx.x + stride]);
-        __syncthreads();
-    }
-    if (threadIdx.x == 0) max_s = partial[0];
-    __syncthreads();
+    const float max_s = block_reduce_f32_max(local_max);
     float den_local = 0.0f;
     for (uint32_t k = threadIdx.x; k < n_keys; k += blockDim.x) {
         float p = isfinite(row[k]) ? expf(row[k] - max_s) : 0.0f;
         row[k] = p;
         den_local += p;
     }
-    partial[threadIdx.x] = den_local;
-    __syncthreads();
-    for (uint32_t stride = blockDim.x >> 1; stride > 0; stride >>= 1) {
-        if (threadIdx.x < stride) partial[threadIdx.x] += partial[threadIdx.x + stride];
-        __syncthreads();
-    }
-    if (threadIdx.x == 0) denom = partial[0] + expf(sinks[h] - max_s);
-    __syncthreads();
+    const float denom = block_reduce_f32_sum(den_local) + expf(sinks[h] - max_s);
     for (uint32_t k = threadIdx.x; k < n_keys; k += blockDim.x) row[k] /= denom;
 }
 
@@ -203,9 +218,6 @@ __global__ static void attention_prefill_mixed_softmax_kernel(
     uint32_t h = blockIdx.y;
     if (t >= n_tokens || ratio == 0) return;
     float *row = scores + ((uint64_t)h * n_tokens + t) * n_keys;
-    __shared__ float partial[256];
-    __shared__ float max_s;
-    __shared__ float denom;
     float local_max = sinks[h];
     const uint32_t visible_comp = (t + 1u) / ratio;
     for (uint32_t k = threadIdx.x; k < n_keys; k += blockDim.x) {
@@ -222,28 +234,14 @@ __global__ static void attention_prefill_mixed_softmax_kernel(
         row[k] = s;
         local_max = fmaxf(local_max, s);
     }
-    partial[threadIdx.x] = local_max;
-    __syncthreads();
-    for (uint32_t stride = blockDim.x >> 1; stride > 0; stride >>= 1) {
-        if (threadIdx.x < stride) partial[threadIdx.x] = fmaxf(partial[threadIdx.x], partial[threadIdx.x + stride]);
-        __syncthreads();
-    }
-    if (threadIdx.x == 0) max_s = partial[0];
-    __syncthreads();
+    const float max_s = block_reduce_f32_max(local_max);
     float den_local = 0.0f;
     for (uint32_t k = threadIdx.x; k < n_keys; k += blockDim.x) {
         float p = isfinite(row[k]) ? expf(row[k] - max_s) : 0.0f;
         row[k] = p;
         den_local += p;
     }
-    partial[threadIdx.x] = den_local;
-    __syncthreads();
-    for (uint32_t stride = blockDim.x >> 1; stride > 0; stride >>= 1) {
-        if (threadIdx.x < stride) partial[threadIdx.x] += partial[threadIdx.x + stride];
-        __syncthreads();
-    }
-    if (threadIdx.x == 0) denom = partial[0] + expf(sinks[h] - max_s);
-    __syncthreads();
+    const float denom = block_reduce_f32_sum(den_local) + expf(sinks[h] - max_s);
     for (uint32_t k = threadIdx.x; k < n_keys; k += blockDim.x) row[k] /= denom;
 }
 
@@ -264,9 +262,6 @@ __global__ static void attention_prefill_mixed_softmax_tile_kernel(
     if (t >= tile_tokens || ratio == 0) return;
     const uint32_t global_t = tile_start + t;
     float *row = scores + ((uint64_t)h * tile_tokens + t) * n_keys;
-    __shared__ float partial[256];
-    __shared__ float max_s;
-    __shared__ float denom;
     float local_max = sinks[h];
     const uint32_t visible_comp = (global_t + 1u) / ratio;
     for (uint32_t k = threadIdx.x; k < n_keys; k += blockDim.x) {
@@ -283,28 +278,14 @@ __global__ static void attention_prefill_mixed_softmax_tile_kernel(
         row[k] = s;
         local_max = fmaxf(local_max, s);
     }
-    partial[threadIdx.x] = local_max;
-    __syncthreads();
-    for (uint32_t stride = blockDim.x >> 1; stride > 0; stride >>= 1) {
-        if (threadIdx.x < stride) partial[threadIdx.x] = fmaxf(partial[threadIdx.x], partial[threadIdx.x + stride]);
-        __syncthreads();
-    }
-    if (threadIdx.x == 0) max_s = partial[0];
-    __syncthreads();
+    const float max_s = block_reduce_f32_max(local_max);
     float den_local = 0.0f;
     for (uint32_t k = threadIdx.x; k < n_keys; k += blockDim.x) {
         float p = isfinite(row[k]) ? expf(row[k] - max_s) : 0.0f;
         row[k] = p;
         den_local += p;
     }
-    partial[threadIdx.x] = den_local;
-    __syncthreads();
-    for (uint32_t stride = blockDim.x >> 1; stride > 0; stride >>= 1) {
-        if (threadIdx.x < stride) partial[threadIdx.x] += partial[threadIdx.x + stride];
-        __syncthreads();
-    }
-    if (threadIdx.x == 0) denom = partial[0] + expf(sinks[h] - max_s);
-    __syncthreads();
+    const float denom = block_reduce_f32_sum(den_local) + expf(sinks[h] - max_s);
     for (uint32_t k = threadIdx.x; k < n_keys; k += blockDim.x) row[k] /= denom;
 }
 
@@ -1484,3 +1465,279 @@ __global__ static void attention_decode_mixed_heads8_online_kernel(
         out4[lane + 96u] = o3;
     }
 }
+
+// ---------------------------------------------------------------------------
+// Tiled fused prefill kernel  —  tile KV rows with two-pass softmax per tile,
+// never write scores to global memory.  Replaces the cublas pack → SGEMM →
+// softmax → SGEMM → unpack pipeline with a single kernel for head_dim=512.
+// ---------------------------------------------------------------------------
+
+__global__ static void attention_prefill_tiled_raw_online_kernel(
+        float *heads,
+        const float *sinks,
+        const float *q,
+        const float *raw_kv,
+        uint32_t n_tokens,
+        uint32_t window,
+        uint32_t n_head,
+        uint32_t head_dim) {
+    uint32_t t = blockIdx.x;
+    uint32_t head_group = blockIdx.y;
+    if (t >= n_tokens || head_dim != 512u) return;
+    const uint32_t lane = threadIdx.x & 31u;
+    const uint32_t warp = threadIdx.x >> 5u;
+    const uint32_t head = head_group * 8u + warp;
+    const bool valid_head = head < n_head;
+
+    const uint32_t raw_count = window != 0u && t + 1u > window ? window : t + 1u;
+    const uint32_t raw_start = t + 1u - raw_count;
+    const uint32_t n_score = raw_count;
+    const float scale = rsqrtf((float)head_dim);
+
+    // Per-warp Q load (into registers)
+    const float4 *q4 = valid_head
+        ? (const float4 *)(q + ((uint64_t)t * n_head + head) * head_dim)
+        : NULL;
+    float4 qv0 = make_float4(0.0f, 0.0f, 0.0f, 0.0f);
+    float4 qv1 = qv0, qv2 = qv0, qv3 = qv0;
+    if (valid_head) {
+        qv0 = q4[lane +  0u];
+        qv1 = q4[lane + 32u];
+        qv2 = q4[lane + 64u];
+        qv3 = q4[lane + 96u];
+    }
+
+    // Running online-softmax state (per-warp)
+    float running_max  = valid_head ? sinks[head] : -INFINITY;
+    float running_denom = valid_head ? 1.0f : 0.0f;
+    float4 o0 = make_float4(0.0f, 0.0f, 0.0f, 0.0f);
+    float4 o1 = o0, o2 = o0, o3 = o0;
+
+    const uint32_t TILE = 384u;
+    __shared__ float4 kv_shared[4 * 128];
+    __shared__ float tile_scores[8 * TILE];  // 8 warps × TILE = 12 KiB
+
+    for (uint32_t tile0 = 0; tile0 < n_score; tile0 += TILE) {
+        const uint32_t tile_nr = n_score - tile0 < TILE ? n_score - tile0 : TILE;
+
+        // --- Pass 1: score computation for this tile ---
+        for (uint32_t row0 = 0; row0 < tile_nr; row0 += 4u) {
+            const uint32_t nr = tile_nr - row0 < 4u ? tile_nr - row0 : 4u;
+            for (uint32_t off = threadIdx.x; off < nr * 128u; off += blockDim.x) {
+                const uint32_t rr = off >> 7u;
+                const uint32_t c4 = off & 127u;
+                const uint32_t sr = tile0 + row0 + rr;
+                const float4 *src = (const float4 *)(raw_kv + (uint64_t)(raw_start + sr) * head_dim);
+                kv_shared[off] = src[c4];
+            }
+            __syncthreads();
+            if (valid_head) {
+                for (uint32_t rr = 0; rr < nr; rr++) {
+                    const float4 *kv4 = kv_shared + rr * 128u;
+                    float4 k0 = kv4[lane +  0u];
+                    float4 k1 = kv4[lane + 32u];
+                    float4 k2 = kv4[lane + 64u];
+                    float4 k3 = kv4[lane + 96u];
+                    float score = dot4_f32(qv0, k0) +
+                                  dot4_f32(qv1, k1) +
+                                  dot4_f32(qv2, k2) +
+                                  dot4_f32(qv3, k3);
+                    score = warp_sum_f32(score) * scale;
+                    score = __shfl_sync(FULL_WARP_MASK, score, 0);
+                    if (lane == 0u)
+                        tile_scores[warp * TILE + row0 + rr] = score;
+                }
+            }
+            __syncthreads();
+        }
+
+        // --- Reduce: tile max ---
+        float tile_max = -INFINITY;
+        if (valid_head) {
+            const float *ts = tile_scores + warp * TILE;
+            for (uint32_t i = lane; i < tile_nr; i += 32u)
+                tile_max = fmaxf(tile_max, ts[i]);
+            tile_max = warp_max_f32(tile_max);
+            tile_max = __shfl_sync(FULL_WARP_MASK, tile_max, 0);
+        }
+
+        // --- Rescale running accumulator ---
+        const float new_max = fmaxf(running_max, tile_max);
+        const float old_scale = expf(running_max - new_max);
+        const float sink_add  = expf(sinks[head] - new_max) - expf(sinks[head] - running_max) * old_scale;
+        running_denom = running_denom * old_scale;
+        o0.x *= old_scale; o0.y *= old_scale; o0.z *= old_scale; o0.w *= old_scale;
+        o1.x *= old_scale; o1.y *= old_scale; o1.z *= old_scale; o1.w *= old_scale;
+        o2.x *= old_scale; o2.y *= old_scale; o2.z *= old_scale; o2.w *= old_scale;
+        o3.x *= old_scale; o3.y *= old_scale; o3.z *= old_scale; o3.w *= old_scale;
+
+        // --- Softmax the tile scores, then accumulate weighted values ---
+        float tile_denom = 0.0f;
+        if (valid_head) {
+            float *ts = tile_scores + warp * TILE;
+            for (uint32_t i = lane; i < tile_nr; i += 32u) {
+                float p = expf(ts[i] - new_max);
+                ts[i] = p;
+                tile_denom += p;
+            }
+            tile_denom = warp_sum_f32(tile_denom);
+            tile_denom = __shfl_sync(FULL_WARP_MASK, tile_denom, 0);
+        }
+        running_denom += tile_denom + sink_add;
+        running_max = new_max;
+
+        // --- Pass 2: weighted value accumulation ---
+        for (uint32_t row0 = 0; row0 < tile_nr; row0 += 4u) {
+            const uint32_t nr = tile_nr - row0 < 4u ? tile_nr - row0 : 4u;
+            for (uint32_t off = threadIdx.x; off < nr * 128u; off += blockDim.x) {
+                const uint32_t rr = off >> 7u;
+                const uint32_t c4 = off & 127u;
+                const uint32_t sr = tile0 + row0 + rr;
+                const float4 *src = (const float4 *)(raw_kv + (uint64_t)(raw_start + sr) * head_dim);
+                kv_shared[off] = src[c4];
+            }
+            __syncthreads();
+            if (valid_head) {
+                const float *ts = tile_scores + warp * TILE;
+                for (uint32_t rr = 0; rr < nr; rr++) {
+                    const float p = ts[row0 + rr];
+                    const float4 *kv4 = kv_shared + rr * 128u;
+                    float4 k0 = kv4[lane +  0u];
+                    float4 k1 = kv4[lane + 32u];
+                    float4 k2 = kv4[lane + 64u];
+                    float4 k3 = kv4[lane + 96u];
+                    o0.x += k0.x * p; o0.y += k0.y * p; o0.z += k0.z * p; o0.w += k0.w * p;
+                    o1.x += k1.x * p; o1.y += k1.y * p; o1.z += k1.z * p; o1.w += k1.w * p;
+                    o2.x += k2.x * p; o2.y += k2.y * p; o2.z += k2.z * p; o2.w += k2.w * p;
+                    o3.x += k3.x * p; o3.y += k3.y * p; o3.z += k3.z * p; o3.w += k3.w * p;
+                }
+            }
+            __syncthreads();
+        }
+    }
+
+    // Final normalization and write-back
+    if (valid_head) {
+        const float inv_s = running_denom == 0.0f ? 0.0f : 1.0f / running_denom;
+        o0.x *= inv_s; o0.y *= inv_s; o0.z *= inv_s; o0.w *= inv_s;
+        o1.x *= inv_s; o1.y *= inv_s; o1.z *= inv_s; o1.w *= inv_s;
+        o2.x *= inv_s; o2.y *= inv_s; o2.z *= inv_s; o2.w *= inv_s;
+        o3.x *= inv_s; o3.y *= inv_s; o3.z *= inv_s; o3.w *= inv_s;
+        float4 *out4 = (float4 *)(heads + ((uint64_t)t * n_head + head) * head_dim);
+        out4[lane +  0u] = o0;
+        out4[lane + 32u] = o1;
+        out4[lane + 64u] = o2;
+        out4[lane + 96u] = o3;
+    }
+}
+
+
+// ---------------------------------------------------------------------------
+// Tiled fused single-pass prefill — online softmax with double-precision
+// correction.  Loads KV into shared memory ONCE per group of 4 rows for
+// both score and accumulation.  head_dim=512 only.
+// ---------------------------------------------------------------------------
+
+__global__ static void attention_prefill_sp_online_kernel(
+        float *heads,
+        const float *sinks,
+        const float *q,
+        const float *raw_kv,
+        uint32_t n_tokens,
+        uint32_t window,
+        uint32_t n_head,
+        uint32_t head_dim) {
+    uint32_t t = blockIdx.x;
+    uint32_t head_group = blockIdx.y;
+    if (t >= n_tokens || head_dim != 512u) return;
+    const uint32_t lane = threadIdx.x & 31u;
+    const uint32_t warp = threadIdx.x >> 5u;
+    const uint32_t head = head_group * 8u + warp;
+    const bool valid_head = head < n_head;
+
+    const uint32_t raw_count = window != 0u && t + 1u > window ? window : t + 1u;
+    const uint32_t raw_start = t + 1u - raw_count;
+    const float scale = rsqrtf((float)head_dim);
+
+    const float4 *q4 = valid_head
+        ? (const float4 *)(q + ((uint64_t)t * n_head + head) * head_dim)
+        : NULL;
+    float4 qv[4];
+    if (valid_head) {
+        qv[0] = q4[lane +  0u];
+        qv[1] = q4[lane + 32u];
+        qv[2] = q4[lane + 64u];
+        qv[3] = q4[lane + 96u];
+    }
+
+    double running_max  = valid_head ? (double)sinks[head] : -INFINITY;
+    double running_denom = valid_head ? 1.0 : 0.0;
+    float4 o0 = make_float4(0.0f, 0.0f, 0.0f, 0.0f);
+    float4 o1 = o0, o2 = o0, o3 = o0;
+
+    __shared__ float4 kv_shared[4 * 128];
+
+    for (uint32_t row0 = 0; row0 < raw_count; row0 += 4u) {
+        const uint32_t nr = raw_count - row0 < 4u ? raw_count - row0 : 4u;
+        for (uint32_t off = threadIdx.x; off < nr * 128u; off += blockDim.x) {
+            const uint32_t rr = off >> 7u;
+            const uint32_t c4 = off & 127u;
+            const uint32_t sr = raw_start + row0 + rr;
+            const float4 *src = (const float4 *)(raw_kv + (uint64_t)sr * head_dim);
+            kv_shared[off] = src[c4];
+        }
+        __syncthreads();
+        if (valid_head) {
+            for (uint32_t rr = 0; rr < nr; rr++) {
+                const float4 *kv4 = kv_shared + rr * 128u;
+                float4 k0 = kv4[lane +  0u];
+                float4 k1 = kv4[lane + 32u];
+                float4 k2 = kv4[lane + 64u];
+                float4 k3 = kv4[lane + 96u];
+                float score = (dot4_f32(qv[0], k0) + dot4_f32(qv[1], k1) +
+                               dot4_f32(qv[2], k2) + dot4_f32(qv[3], k3))
+                              * scale;
+                score = __shfl_sync(FULL_WARP_MASK, score, 0);
+
+                const double sd = (double)score;
+                const double new_max = fmax(running_max, sd);
+                const double rescale = exp(running_max - new_max);
+                const double row_scale = exp(sd - new_max);
+
+                running_denom = running_denom * rescale + row_scale;
+                o0.x = (float)((double)o0.x * rescale) + k0.x * (float)row_scale;
+                o0.y = (float)((double)o0.y * rescale) + k0.y * (float)row_scale;
+                o0.z = (float)((double)o0.z * rescale) + k0.z * (float)row_scale;
+                o0.w = (float)((double)o0.w * rescale) + k0.w * (float)row_scale;
+                o1.x = (float)((double)o1.x * rescale) + k1.x * (float)row_scale;
+                o1.y = (float)((double)o1.y * rescale) + k1.y * (float)row_scale;
+                o1.z = (float)((double)o1.z * rescale) + k1.z * (float)row_scale;
+                o1.w = (float)((double)o1.w * rescale) + k1.w * (float)row_scale;
+                o2.x = (float)((double)o2.x * rescale) + k2.x * (float)row_scale;
+                o2.y = (float)((double)o2.y * rescale) + k2.y * (float)row_scale;
+                o2.z = (float)((double)o2.z * rescale) + k2.z * (float)row_scale;
+                o2.w = (float)((double)o2.w * rescale) + k2.w * (float)row_scale;
+                o3.x = (float)((double)o3.x * rescale) + k3.x * (float)row_scale;
+                o3.y = (float)((double)o3.y * rescale) + k3.y * (float)row_scale;
+                o3.z = (float)((double)o3.z * rescale) + k3.z * (float)row_scale;
+                o3.w = (float)((double)o3.w * rescale) + k3.w * (float)row_scale;
+
+                running_max = new_max;
+            }
+        }
+        __syncthreads();
+    }
+
+    if (valid_head) {
+        const double inv_s = running_denom == 0.0 ? 0.0 : 1.0 / running_denom;
+        float4 *out4 = (float4 *)(heads + ((uint64_t)t * n_head + head) * head_dim);
+        out4[lane +  0u] = make_float4((float)(o0.x * inv_s), (float)(o0.y * inv_s),
+                                        (float)(o0.z * inv_s), (float)(o0.w * inv_s));
+        out4[lane + 32u] = make_float4((float)(o1.x * inv_s), (float)(o1.y * inv_s),
+                                        (float)(o1.z * inv_s), (float)(o1.w * inv_s));
+        out4[lane + 64u] = make_float4((float)(o2.x * inv_s), (float)(o2.y * inv_s),
+                                        (float)(o2.z * inv_s), (float)(o2.w * inv_s));
+        out4[lane + 96u] = make_float4((float)(o3.x * inv_s), (float)(o3.y * inv_s),
+                                        (float)(o3.z * inv_s), (float)(o3.w * inv_s));
+    }
+}
diff --git a/rocm/ds4_rocm_attention_launch.cuh b/rocm/ds4_rocm_attention_launch.cuh
index b9b43d958..32b10974b 100644
--- a/rocm/ds4_rocm_attention_launch.cuh
+++ b/rocm/ds4_rocm_attention_launch.cuh
@@ -139,6 +139,19 @@ extern "C" int ds4_gpu_attention_prefill_raw_heads_tensor(ds4_gpu_tensor *heads,
                                                                    head_dim);
         return cuda_ok(cudaGetLastError(), "attention raw window launch");
     }
+    /* Tiled fused online-prefill for larger windows (head_dim=512).
+     * Single-pass variant: online softmax with double-precision correction.
+     * Loads KV once, no separate score/softmax/value passes. */
+    if (n_tokens > 1 && head_dim == 512 && !g_quality_mode) {
+        dim3 grid(n_tokens, (n_head + 7u) / 8u, 1);
+        attention_prefill_sp_online_kernel<<<grid, 256>>>(
+                (float *)heads->ptr,
+                sinks,
+                (const float *)q->ptr,
+                (const float *)raw_kv->ptr,
+                n_tokens, window, n_head, head_dim);
+        return cuda_ok(cudaGetLastError(), "attention sp online launch");
+    }
     if (g_cublas_ready && n_tokens > 1 && head_dim == 512) {
         const uint32_t n_keys = n_tokens;
         const uint64_t score_count = (uint64_t)n_head * n_tokens * n_keys;
@@ -152,48 +165,76 @@ extern "C" int ds4_gpu_attention_prefill_raw_heads_tensor(ds4_gpu_tensor *heads,
         float *out_tmp = (float *)((char *)tmp + out_offset);
         const float alpha = 1.0f / sqrtf((float)head_dim);
         const float beta = 0.0f;
-        cublasStatus_t st = cublasSgemmStridedBatched(g_cublas,
-                                                      CUBLAS_OP_T,
-                                                      CUBLAS_OP_N,
-                                                      (int)n_keys,
-                                                      (int)n_tokens,
-                                                      (int)head_dim,
-                                                      &alpha,
-                                                      (const float *)raw_kv->ptr,
-                                                      (int)head_dim,
-                                                      0,
-                                                      (const float *)q->ptr,
-                                                      (int)(n_head * head_dim),
-                                                      (long long)head_dim,
-                                                      &beta,
-                                                      scores,
-                                                      (int)n_keys,
-                                                      (long long)n_keys * n_tokens,
-                                                      (int)n_head);
-        if (!cublas_ok(st, "attention raw score gemm")) return 0;
+        int gemm_ok = 0;
+#ifdef __HIP_PLATFORM_AMD__
+        if (g_hipblaslt_ready)
+            gemm_ok = hipblaslt_gemm_strided_batched_f32(
+                    scores,
+                    (const float *)raw_kv->ptr,
+                    (const float *)q->ptr,
+                    HIPBLAS_OP_T, HIPBLAS_OP_N,
+                    n_keys, n_tokens, head_dim,
+                    0, (int64_t)head_dim, (int64_t)n_keys * n_tokens,
+                    n_head, alpha, beta, "raw_score");
+#endif
+        if (!gemm_ok) {
+            cublasStatus_t st = cublasSgemmStridedBatched(g_cublas,
+                                                          CUBLAS_OP_T,
+                                                          CUBLAS_OP_N,
+                                                          (int)n_keys,
+                                                          (int)n_tokens,
+                                                          (int)head_dim,
+                                                          &alpha,
+                                                          (const float *)raw_kv->ptr,
+                                                          (int)head_dim,
+                                                          0,
+                                                          (const float *)q->ptr,
+                                                          (int)(n_head * head_dim),
+                                                          (long long)head_dim,
+                                                          &beta,
+                                                          scores,
+                                                          (int)n_keys,
+                                                          (long long)n_keys * n_tokens,
+                                                          (int)n_head);
+            if (!cublas_ok(st, "attention raw score gemm")) return 0;
+        }
         dim3 sgrid(n_tokens, n_head, 1);
         attention_prefill_raw_softmax_kernel<<<sgrid, 256>>>(scores, sinks, n_tokens, window, n_keys);
         if (!cuda_ok(cudaGetLastError(), "attention raw softmax launch")) return 0;
         const float one = 1.0f;
-        st = cublasSgemmStridedBatched(g_cublas,
-                                       CUBLAS_OP_N,
-                                       CUBLAS_OP_N,
-                                       (int)head_dim,
-                                       (int)n_tokens,
-                                       (int)n_keys,
-                                       &one,
-                                       (const float *)raw_kv->ptr,
-                                       (int)head_dim,
-                                       0,
-                                       scores,
-                                       (int)n_keys,
-                                       (long long)n_keys * n_tokens,
-                                       &beta,
-                                       out_tmp,
-                                       (int)head_dim,
-                                       (long long)head_dim * n_tokens,
-                                       (int)n_head);
-        if (!cublas_ok(st, "attention raw value gemm")) return 0;
+        gemm_ok = 0;
+#ifdef __HIP_PLATFORM_AMD__
+        if (g_hipblaslt_ready)
+            gemm_ok = hipblaslt_gemm_strided_batched_f32(
+                    out_tmp,
+                    (const float *)raw_kv->ptr,
+                    scores,
+                    HIPBLAS_OP_N, HIPBLAS_OP_N,
+                    head_dim, n_tokens, n_keys,
+                    0, (int64_t)n_keys * n_tokens, (int64_t)head_dim * n_tokens,
+                    n_head, one, beta, "raw_value");
+#endif
+        if (!gemm_ok) {
+            cublasStatus_t st = cublasSgemmStridedBatched(g_cublas,
+                                                           CUBLAS_OP_N,
+                                                           CUBLAS_OP_N,
+                                                           (int)head_dim,
+                                                           (int)n_tokens,
+                                                           (int)n_keys,
+                                                           &one,
+                                                           (const float *)raw_kv->ptr,
+                                                           (int)head_dim,
+                                                           0,
+                                                           scores,
+                                                           (int)n_keys,
+                                                           (long long)n_keys * n_tokens,
+                                                           &beta,
+                                                           out_tmp,
+                                                           (int)head_dim,
+                                                           (long long)head_dim * n_tokens,
+                                                           (int)n_head);
+            if (!cublas_ok(st, "attention raw value gemm")) return 0;
+        }
         uint64_t n = (uint64_t)n_tokens * n_head * head_dim;
         attention_prefill_unpack_heads_kernel<<<(n + 255) / 256, 256>>>((float *)heads->ptr,
                                                                         out_tmp,
@@ -717,25 +758,39 @@ static int attention_prefill_mixed_launch(
         if (!cuda_ok(cudaGetLastError(), "attention mixed kv pack launch")) return 0;
         const float alpha = 1.0f / sqrtf((float)head_dim);
         const float beta = 0.0f;
-        cublasStatus_t st = cublasSgemmStridedBatched(g_cublas,
-                                                      CUBLAS_OP_T,
-                                                      CUBLAS_OP_N,
-                                                      (int)n_keys,
-                                                      (int)n_tokens,
-                                                      (int)head_dim,
-                                                      &alpha,
-                                                      kv,
-                                                      (int)head_dim,
-                                                      0,
-                                                      (const float *)q->ptr,
-                                                      (int)(n_head * head_dim),
-                                                      (long long)head_dim,
-                                                      &beta,
-                                                      scores,
-                                                      (int)n_keys,
-                                                      (long long)n_keys * n_tokens,
-                                                      (int)n_head);
-        if (!cublas_ok(st, "attention mixed score gemm")) return 0;
+        int gemm_ok = 0;
+#ifdef __HIP_PLATFORM_AMD__
+        if (g_hipblaslt_ready)
+            gemm_ok = hipblaslt_gemm_strided_batched_f32(
+                    scores,
+                    kv,
+                    (const float *)q->ptr,
+                    HIPBLAS_OP_T, HIPBLAS_OP_N,
+                    n_keys, n_tokens, head_dim,
+                    0, (int64_t)head_dim, (int64_t)n_keys * n_tokens,
+                    n_head, alpha, beta, "mixed_score");
+#endif
+        if (!gemm_ok) {
+            cublasStatus_t st = cublasSgemmStridedBatched(g_cublas,
+                                                          CUBLAS_OP_T,
+                                                          CUBLAS_OP_N,
+                                                          (int)n_keys,
+                                                          (int)n_tokens,
+                                                          (int)head_dim,
+                                                          &alpha,
+                                                          kv,
+                                                          (int)head_dim,
+                                                          0,
+                                                          (const float *)q->ptr,
+                                                          (int)(n_head * head_dim),
+                                                          (long long)head_dim,
+                                                          &beta,
+                                                          scores,
+                                                          (int)n_keys,
+                                                          (long long)n_keys * n_tokens,
+                                                          (int)n_head);
+            if (!cublas_ok(st, "attention mixed score gemm")) return 0;
+        }
         dim3 sgrid(n_tokens, n_head, 1);
         attention_prefill_mixed_softmax_kernel<<<sgrid, 256>>>(
                 scores,
@@ -749,25 +804,39 @@ static int attention_prefill_mixed_launch(
                 n_keys);
         if (!cuda_ok(cudaGetLastError(), "attention mixed softmax launch")) return 0;
         const float one = 1.0f;
-        st = cublasSgemmStridedBatched(g_cublas,
-                                       CUBLAS_OP_N,
-                                       CUBLAS_OP_N,
-                                       (int)head_dim,
-                                       (int)n_tokens,
-                                       (int)n_keys,
-                                       &one,
-                                       kv,
-                                       (int)head_dim,
-                                       0,
-                                       scores,
-                                       (int)n_keys,
-                                       (long long)n_keys * n_tokens,
-                                       &beta,
-                                       out_tmp,
-                                       (int)head_dim,
-                                       (long long)head_dim * n_tokens,
-                                       (int)n_head);
-        if (!cublas_ok(st, "attention mixed value gemm")) return 0;
+        gemm_ok = 0;
+#ifdef __HIP_PLATFORM_AMD__
+        if (g_hipblaslt_ready)
+            gemm_ok = hipblaslt_gemm_strided_batched_f32(
+                    out_tmp,
+                    kv,
+                    scores,
+                    HIPBLAS_OP_N, HIPBLAS_OP_N,
+                    head_dim, n_tokens, n_keys,
+                    0, (int64_t)n_keys * n_tokens, (int64_t)head_dim * n_tokens,
+                    n_head, one, beta, "mixed_value");
+#endif
+        if (!gemm_ok) {
+            cublasStatus_t st = cublasSgemmStridedBatched(g_cublas,
+                                                           CUBLAS_OP_N,
+                                                           CUBLAS_OP_N,
+                                                           (int)head_dim,
+                                                           (int)n_tokens,
+                                                           (int)n_keys,
+                                                           &one,
+                                                           kv,
+                                                           (int)head_dim,
+                                                           0,
+                                                           scores,
+                                                           (int)n_keys,
+                                                           (long long)n_keys * n_tokens,
+                                                           &beta,
+                                                           out_tmp,
+                                                           (int)head_dim,
+                                                           (long long)head_dim * n_tokens,
+                                                           (int)n_head);
+            if (!cublas_ok(st, "attention mixed value gemm")) return 0;
+        }
         uint64_t n = (uint64_t)n_tokens * n_head * head_dim;
         attention_prefill_unpack_heads_kernel<<<(n + 255) / 256, 256>>>((float *)heads->ptr,
                                                                         out_tmp,
diff --git a/rocm/ds4_rocm_hipblaslt.cuh b/rocm/ds4_rocm_hipblaslt.cuh
index ce5c15a86..8343bae76 100644
--- a/rocm/ds4_rocm_hipblaslt.cuh
+++ b/rocm/ds4_rocm_hipblaslt.cuh
@@ -135,3 +135,279 @@ static int hipblaslt_gemm_tn_f16_out_f16(
                                         NULL, 0, 0),
                         label ? label : "gemm");
 }
+
+// ---------------------------------------------------------------------------
+// FP32 strided-batched GEMM via hipBLASLt  (for attention prefill SGEMMs).
+// Keyed on (m, n, k, batch_count, stride_a, stride_b, stride_c, trans).
+// ---------------------------------------------------------------------------
+
+struct cuda_hipblaslt_gemm_plan_sb_f32 {
+    uint32_t m, n, k;
+    uint32_t batch_count;
+    int64_t stride_a, stride_b, stride_c;
+    hipblasOperation_t trans_a, trans_b;
+    hipblasLtMatmulDesc_t desc;
+    hipblasLtMatrixLayout_t a_desc, b_desc, c_desc, d_desc;
+    hipblasLtMatmulAlgo_t algo;
+};
+
+static std::vector<cuda_hipblaslt_gemm_plan_sb_f32> g_hipblaslt_sb_plans;
+
+static void hipblaslt_sb_plan_clear(void) {
+    for (size_t i = 0; i < g_hipblaslt_sb_plans.size(); i++) {
+        auto &p = g_hipblaslt_sb_plans[i];
+        if (p.d_desc) (void)hipblasLtMatrixLayoutDestroy(p.d_desc);
+        if (p.c_desc) (void)hipblasLtMatrixLayoutDestroy(p.c_desc);
+        if (p.b_desc) (void)hipblasLtMatrixLayoutDestroy(p.b_desc);
+        if (p.a_desc) (void)hipblasLtMatrixLayoutDestroy(p.a_desc);
+        if (p.desc) (void)hipblasLtMatmulDescDestroy(p.desc);
+    }
+    g_hipblaslt_sb_plans.clear();
+}
+
+static cuda_hipblaslt_gemm_plan_sb_f32 *hipblaslt_gemm_plan_get_sb_f32(
+        hipblasOperation_t trans_a, hipblasOperation_t trans_b,
+        uint32_t m, uint32_t n, uint32_t k,
+        int64_t stride_a, int64_t stride_b, int64_t stride_c,
+        uint32_t batch_count, const char *label) {
+    for (size_t i = 0; i < g_hipblaslt_sb_plans.size(); i++) {
+        auto &p = g_hipblaslt_sb_plans[i];
+        if (p.m == m && p.n == n && p.k == k &&
+            p.batch_count == batch_count &&
+            p.stride_a == stride_a && p.stride_b == stride_b && p.stride_c == stride_c &&
+            p.trans_a == trans_a && p.trans_b == trans_b)
+            return &p;
+    }
+
+    hipblasLtMatmulDesc_t desc = NULL;
+    hipblasLtMatrixLayout_t a_desc = NULL, b_desc = NULL, c_desc = NULL, d_desc = NULL;
+    hipblasLtMatmulPreference_t pref = NULL;
+    hipblasLtMatmulHeuristicResult_t heur[8];
+    int returned = 0, ok = 0;
+    do {
+        if (!hipblaslt_ok(hipblasLtMatmulDescCreate(&desc, HIPBLAS_COMPUTE_32F, HIP_R_32F),
+                          "sb desc create")) break;
+        if (!hipblaslt_ok(hipblasLtMatmulDescSetAttribute(desc, HIPBLASLT_MATMUL_DESC_TRANSA,
+                                                          &trans_a, sizeof(trans_a)),
+                          "sb set transA")) break;
+        if (!hipblaslt_ok(hipblasLtMatmulDescSetAttribute(desc, HIPBLASLT_MATMUL_DESC_TRANSB,
+                                                          &trans_b, sizeof(trans_b)),
+                          "sb set transB")) break;
+
+        // Dimensions: for op(A) × op(B), A is (m × k), B is (k × n), C/D is (m × n)
+        uint32_t a_rows = (trans_a == HIPBLAS_OP_N) ? m : k;
+        uint32_t a_cols = (trans_a == HIPBLAS_OP_N) ? k : m;
+        uint32_t a_ld   = a_rows;
+
+        uint32_t b_rows = (trans_b == HIPBLAS_OP_N) ? k : n;
+        uint32_t b_cols = (trans_b == HIPBLAS_OP_N) ? n : k;
+        uint32_t b_ld   = b_rows;
+
+        if (!hipblaslt_ok(hipblasLtMatrixLayoutCreate(&a_desc, HIP_R_32F, a_rows, a_cols, a_ld),
+                          "sb A layout")) break;
+        if (!hipblaslt_ok(hipblasLtMatrixLayoutCreate(&b_desc, HIP_R_32F, b_rows, b_cols, b_ld),
+                          "sb B layout")) break;
+        if (!hipblaslt_ok(hipblasLtMatrixLayoutCreate(&c_desc, HIP_R_32F, m, n, m),
+                          "sb C layout")) break;
+        if (!hipblaslt_ok(hipblasLtMatrixLayoutCreate(&d_desc, HIP_R_32F, m, n, m),
+                          "sb D layout")) break;
+
+        // Set batch count and strides per-layout (hipBLASLt 7.2 API)
+        (void)hipblasLtMatrixLayoutSetAttribute(a_desc, HIPBLASLT_MATRIX_LAYOUT_BATCH_COUNT,
+                                                 &batch_count, sizeof(batch_count));
+        (void)hipblasLtMatrixLayoutSetAttribute(b_desc, HIPBLASLT_MATRIX_LAYOUT_BATCH_COUNT,
+                                                 &batch_count, sizeof(batch_count));
+        (void)hipblasLtMatrixLayoutSetAttribute(c_desc, HIPBLASLT_MATRIX_LAYOUT_BATCH_COUNT,
+                                                 &batch_count, sizeof(batch_count));
+        (void)hipblasLtMatrixLayoutSetAttribute(d_desc, HIPBLASLT_MATRIX_LAYOUT_BATCH_COUNT,
+                                                 &batch_count, sizeof(batch_count));
+
+        if (stride_a != 0)
+            (void)hipblasLtMatrixLayoutSetAttribute(a_desc, HIPBLASLT_MATRIX_LAYOUT_STRIDED_BATCH_OFFSET,
+                                                     &stride_a, sizeof(stride_a));
+        if (stride_b != 0)
+            (void)hipblasLtMatrixLayoutSetAttribute(b_desc, HIPBLASLT_MATRIX_LAYOUT_STRIDED_BATCH_OFFSET,
+                                                     &stride_b, sizeof(stride_b));
+        if (stride_c != 0)
+            (void)hipblasLtMatrixLayoutSetAttribute(c_desc, HIPBLASLT_MATRIX_LAYOUT_STRIDED_BATCH_OFFSET,
+                                                     &stride_c, sizeof(stride_c));
+        (void)hipblasLtMatrixLayoutSetAttribute(d_desc, HIPBLASLT_MATRIX_LAYOUT_STRIDED_BATCH_OFFSET,
+                                                 &stride_c, sizeof(stride_c));
+
+        if (!hipblaslt_ok(hipblasLtMatmulPreferenceCreate(&pref), "sb pref create")) break;
+        const size_t max_workspace = 0;
+        if (!hipblaslt_ok(hipblasLtMatmulPreferenceSetAttribute(
+                pref, HIPBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES,
+                &max_workspace, sizeof(max_workspace)),
+                "sb set max workspace")) break;
+        if (!hipblaslt_ok(hipblasLtMatmulAlgoGetHeuristic(g_hipblaslt, desc,
+                                                          a_desc, b_desc, c_desc, d_desc,
+                                                          pref, 8, heur, &returned),
+                          "sb algo heuristic")) break;
+        if (returned <= 0 || heur[0].state != HIPBLAS_STATUS_SUCCESS) {
+            fprintf(stderr, "ds4: hipBLASLt no sb algo for %s m=%u n=%u k=%u b=%u\n",
+                    label ? label : "sb_gemm", m, n, k, batch_count);
+            break;
+        }
+        ok = 1;
+    } while (0);
+    if (pref) (void)hipblasLtMatmulPreferenceDestroy(pref);
+    if (!ok) {
+        if (d_desc) (void)hipblasLtMatrixLayoutDestroy(d_desc);
+        if (c_desc) (void)hipblasLtMatrixLayoutDestroy(c_desc);
+        if (b_desc) (void)hipblasLtMatrixLayoutDestroy(b_desc);
+        if (a_desc) (void)hipblasLtMatrixLayoutDestroy(a_desc);
+        if (desc) (void)hipblasLtMatmulDescDestroy(desc);
+        return NULL;
+    }
+
+    cuda_hipblaslt_gemm_plan_sb_f32 p;
+    p.m = m; p.n = n; p.k = k;
+    p.batch_count = batch_count;
+    p.stride_a = stride_a; p.stride_b = stride_b; p.stride_c = stride_c;
+    p.trans_a = trans_a; p.trans_b = trans_b;
+    p.desc = desc;
+    p.a_desc = a_desc; p.b_desc = b_desc; p.c_desc = c_desc; p.d_desc = d_desc;
+    p.algo = heur[0].algo;
+    g_hipblaslt_sb_plans.push_back(p);
+    return &g_hipblaslt_sb_plans.back();
+}
+
+static int hipblaslt_gemm_strided_batched_f32(
+        float *out,
+        const float *a,
+        const float *b,
+        hipblasOperation_t trans_a, hipblasOperation_t trans_b,
+        uint32_t m, uint32_t n, uint32_t k,
+        int64_t stride_a, int64_t stride_b, int64_t stride_c,
+        uint32_t batch_count,
+        float alpha, float beta,
+        const char *label) {
+    if (!g_hipblaslt_ready || !out || !a || !b ||
+        m == 0 || n == 0 || k == 0 || batch_count == 0) return 0;
+    cuda_hipblaslt_gemm_plan_sb_f32 *p = hipblaslt_gemm_plan_get_sb_f32(
+            trans_a, trans_b, m, n, k, stride_a, stride_b, stride_c, batch_count, label);
+    if (!p) return 0;
+    return hipblaslt_ok(hipblasLtMatmul(g_hipblaslt, p->desc, &alpha,
+                                        a, p->a_desc,
+                                        b, p->b_desc,
+                                        &beta,
+                                        out, p->c_desc,
+                                        out, p->d_desc,
+                                        &p->algo,
+                                        NULL, 0, 0),
+                        label ? label : "sb_gemm");
+}
+
+// ---------------------------------------------------------------------------
+// FP16→FP32 GEMM via hipBLASLt  (for MoE prefill matmuls: FP16 weights +
+// FP16 activations → FP32 output).  Keyed on (out_dim, n_tok, in_dim).
+// ---------------------------------------------------------------------------
+
+struct cuda_hipblaslt_gemm_plan_f16_out_f32 {
+    uint32_t out_dim, n_tok, in_dim;
+    hipblasLtMatmulDesc_t desc;
+    hipblasLtMatrixLayout_t a_desc, b_desc, c_desc, d_desc;
+    hipblasLtMatmulAlgo_t algo;
+};
+
+static std::vector<cuda_hipblaslt_gemm_plan_f16_out_f32> g_hipblaslt_f16_f32_plans;
+
+static void hipblaslt_f16_f32_plan_clear(void) {
+    for (auto &p : g_hipblaslt_f16_f32_plans) {
+        if (p.d_desc) (void)hipblasLtMatrixLayoutDestroy(p.d_desc);
+        if (p.c_desc) (void)hipblasLtMatrixLayoutDestroy(p.c_desc);
+        if (p.b_desc) (void)hipblasLtMatrixLayoutDestroy(p.b_desc);
+        if (p.a_desc) (void)hipblasLtMatrixLayoutDestroy(p.a_desc);
+        if (p.desc) (void)hipblasLtMatmulDescDestroy(p.desc);
+    }
+    g_hipblaslt_f16_f32_plans.clear();
+}
+
+static cuda_hipblaslt_gemm_plan_f16_out_f32 *hipblaslt_gemm_plan_get_f16_f32(
+        uint32_t out_dim, uint32_t n_tok, uint32_t in_dim, const char *label) {
+    for (auto &p : g_hipblaslt_f16_f32_plans)
+        if (p.out_dim == out_dim && p.n_tok == n_tok && p.in_dim == in_dim)
+            return &p;
+
+    hipblasLtMatmulDesc_t desc = NULL;
+    hipblasLtMatrixLayout_t a_desc = NULL, b_desc = NULL, c_desc = NULL, d_desc = NULL;
+    hipblasLtMatmulPreference_t pref = NULL;
+    hipblasLtMatmulHeuristicResult_t heur[8];
+    int returned = 0, ok = 0;
+    do {
+        if (!hipblaslt_ok(hipblasLtMatmulDescCreate(&desc, HIPBLAS_COMPUTE_32F, HIP_R_32F),
+                          "f16_f32 desc")) break;
+        hipblasOperation_t op_a = HIPBLAS_OP_T;
+        hipblasOperation_t op_b = HIPBLAS_OP_N;
+        if (!hipblaslt_ok(hipblasLtMatmulDescSetAttribute(desc, HIPBLASLT_MATMUL_DESC_TRANSA,
+                                                          &op_a, sizeof(op_a)), "set transA")) break;
+        if (!hipblaslt_ok(hipblasLtMatmulDescSetAttribute(desc, HIPBLASLT_MATMUL_DESC_TRANSB,
+                                                          &op_b, sizeof(op_b)), "set transB")) break;
+        // A = FP16 weight (in_dim × out_dim, transposed)
+        if (!hipblaslt_ok(hipblasLtMatrixLayoutCreate(&a_desc, HIP_R_16F, in_dim, out_dim, in_dim),
+                          "f16_f32 A")) break;
+        // B = FP16 activation (in_dim × n_tok)
+        if (!hipblaslt_ok(hipblasLtMatrixLayoutCreate(&b_desc, HIP_R_16F, in_dim, n_tok, in_dim),
+                          "f16_f32 B")) break;
+        // C = FP32 output (out_dim × n_tok)
+        if (!hipblaslt_ok(hipblasLtMatrixLayoutCreate(&c_desc, HIP_R_32F, out_dim, n_tok, out_dim),
+                          "f16_f32 C")) break;
+        if (!hipblaslt_ok(hipblasLtMatrixLayoutCreate(&d_desc, HIP_R_32F, out_dim, n_tok, out_dim),
+                          "f16_f32 D")) break;
+        if (!hipblaslt_ok(hipblasLtMatmulPreferenceCreate(&pref), "f16_f32 pref")) break;
+        const size_t max_workspace = 0;
+        if (!hipblaslt_ok(hipblasLtMatmulPreferenceSetAttribute(
+                pref, HIPBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES,
+                &max_workspace, sizeof(max_workspace)), "f16_f32 ws")) break;
+        if (!hipblaslt_ok(hipblasLtMatmulAlgoGetHeuristic(g_hipblaslt, desc,
+                    a_desc, b_desc, c_desc, d_desc, pref, 8, heur, &returned),
+                    "f16_f32 heuristic")) break;
+        if (returned <= 0 || heur[0].state != HIPBLAS_STATUS_SUCCESS) {
+            fprintf(stderr, "ds4: hipBLASLt no f16_f32 algo for %s m=%u n=%u k=%u\n",
+                    label ? label : "gemm", out_dim, n_tok, in_dim);
+            break;
+        }
+        ok = 1;
+    } while (0);
+    if (pref) (void)hipblasLtMatmulPreferenceDestroy(pref);
+    if (!ok) {
+        if (d_desc) (void)hipblasLtMatrixLayoutDestroy(d_desc);
+        if (c_desc) (void)hipblasLtMatrixLayoutDestroy(c_desc);
+        if (b_desc) (void)hipblasLtMatrixLayoutDestroy(b_desc);
+        if (a_desc) (void)hipblasLtMatrixLayoutDestroy(a_desc);
+        if (desc) (void)hipblasLtMatmulDescDestroy(desc);
+        return NULL;
+    }
+    cuda_hipblaslt_gemm_plan_f16_out_f32 p;
+    p.out_dim = out_dim; p.n_tok = n_tok; p.in_dim = in_dim;
+    p.desc = desc;
+    p.a_desc = a_desc; p.b_desc = b_desc; p.c_desc = c_desc; p.d_desc = d_desc;
+    p.algo = heur[0].algo;
+    g_hipblaslt_f16_f32_plans.push_back(p);
+    return &g_hipblaslt_f16_f32_plans.back();
+}
+
+static int hipblaslt_gemm_tn_f16_out_f32(
+        float *out,
+        const __half *w_rowmajor_out_in,
+        const __half *x_rowmajor_tok_in,
+        uint32_t out_dim,
+        uint32_t n_tok,
+        uint32_t in_dim,
+        const char *label) {
+    if (!g_hipblaslt_ready || !out || !w_rowmajor_out_in || !x_rowmajor_tok_in ||
+        out_dim == 0 || n_tok == 0 || in_dim == 0) return 0;
+    cuda_hipblaslt_gemm_plan_f16_out_f32 *p = hipblaslt_gemm_plan_get_f16_f32(
+            out_dim, n_tok, in_dim, label);
+    if (!p) return 0;
+    const float alpha = 1.0f, beta = 0.0f;
+    return hipblaslt_ok(hipblasLtMatmul(g_hipblaslt, p->desc, &alpha,
+                    w_rowmajor_out_in, p->a_desc,
+                    x_rowmajor_tok_in, p->b_desc,
+                    &beta,
+                    out, p->c_desc,
+                    out, p->d_desc,
+                    &p->algo, NULL, 0, 0),
+            label ? label : "gemm_f16_f32");
+}
diff --git a/rocm/ds4_rocm_matmul.cuh b/rocm/ds4_rocm_matmul.cuh
index 0f9e06625..e0bbced68 100644
--- a/rocm/ds4_rocm_matmul.cuh
+++ b/rocm/ds4_rocm_matmul.cuh
@@ -128,6 +128,12 @@ static int cuda_matmul_q8_0_tensor_f16_gemm(
     if (!xh) return 0;
     f32_to_f16_kernel<<<(xh_count + 255u) / 256u, 256>>>(xh, (const float *)x->ptr, xh_count);
     if (!cuda_ok(cudaGetLastError(), "q8 f16 activation convert launch")) return 0;
+#ifdef __HIP_PLATFORM_AMD__
+    if (g_hipblaslt_ready &&
+        hipblaslt_gemm_tn_f16_out_f32((float *)out->ptr, w_f16, xh,
+                                      (uint32_t)out_dim, (uint32_t)n_tok, (uint32_t)in_dim, "q8_matmulf16"))
+        return 1;
+#endif
     const float alpha = 1.0f;
     const float beta = 0.0f;
     cublasStatus_t st = cublasGemmEx(g_cublas,
@@ -645,6 +651,12 @@ extern "C" int ds4_gpu_matmul_f16_tensor(ds4_gpu_tensor *out, const void *model_
         if (!xh) return 0;
         f32_to_f16_kernel<<<(xh_count + 255) / 256, 256>>>(xh, (const float *)x->ptr, xh_count);
         if (!cuda_ok(cudaGetLastError(), "f16 activation convert launch")) return 0;
+#ifdef __HIP_PLATFORM_AMD__
+        if (g_hipblaslt_ready &&
+            hipblaslt_gemm_tn_f16_out_f32((float *)out->ptr, w, xh,
+                                          (uint32_t)out_dim, (uint32_t)n_tok, (uint32_t)in_dim, "f16_matmulf16"))
+            return 1;
+#endif
         const float alpha = 1.0f;
         const float beta = 0.0f;
         cublasStatus_t st = cublasGemmEx(g_cublas,
diff --git a/rocm/ds4_rocm_runtime.cuh b/rocm/ds4_rocm_runtime.cuh
index 3bd786f8e..93bd72308 100644
--- a/rocm/ds4_rocm_runtime.cuh
+++ b/rocm/ds4_rocm_runtime.cuh
@@ -4407,6 +4407,8 @@ extern "C" void ds4_gpu_cleanup(void) {
     cuda_shared_gate_up_async_cleanup();
 #ifdef __HIP_PLATFORM_AMD__
     hipblaslt_gemm_plan_clear();
+    hipblaslt_sb_plan_clear();
+    hipblaslt_f16_f32_plan_clear();
 #endif
     if (g_cublas_ready) {
         (void)cublasDestroy(g_cublas);