From 58f09db0b27e8180bc4a953e35f2bec3410eaef0 Mon Sep 17 00:00:00 2001
From: Bartlomiej Kocot <barkocot@amd.com>
Date: Wed, 21 Oct 2026 14:34:59 -0500
Subject: [PATCH 1/6] perf(rocke): depthwise 7x7 and grouped cpg4/8 wgrad on
 batched mfma 4x4x4

---
 .../common/conv_direct_dw_wgrad_mfma.py       | 603 +++++++++++
 .../common/conv_direct_grouped_wgrad_mfma.py  | 432 ++++++++
 .../library/kernels/conv/direct_aot_jobs.py   | 103 +-
 .../library/kernels/conv/direct_variants.py   |  38 +-
 .../tests/parity/conv_direct_grouped_emit.c   | 110 +-
 .../tests/parity/conv_direct_grouped_emit.py  | 102 +-
 .../tests/test_builder_signature_contract.py  |   2 +
 .../tests/test_direct_conv_correctness.py     |  48 +-
 .../library/tests/test_direct_kernel_cache.py |  13 +-
 .../rocke/platform/cpp/core/arch/data.cpp     |   1 +
 .../rocke/platform/cpp/core/ir/ir_tile.cpp    |   1 +
 .../platform/cpp/core/lower_llvm/data.cpp     |   3 +
 .../platform/cpp/core/lower_llvm/mma.cpp      |   6 +
 .../instance_conv_direct_dw_wgrad_mfma.h      | 101 ++
 .../instance_conv_direct_grouped_wgrad_mfma.h | 106 ++
 .../rocke/platform/cpp/include/rocke/ir.h     |   4 +
 .../common/conv_direct_dw_wgrad_mfma.cpp      | 939 ++++++++++++++++++
 .../common/conv_direct_grouped_wgrad_mfma.cpp | 699 +++++++++++++
 .../dsl_docs/architecture/kernel_taxonomy.md  |   2 +
 .../dsl_docs/instances/convolution.md         |  38 +-
 .../platform/dsl_docs/instances/index.md      |   2 +
 .../platform/dsl_docs/ir_lowering/ir_model.md |   1 +
 .../dsl_docs/reference/op_vocabulary.md       |   1 +
 .../platform/python/rocke/core/arch/target.py |   1 +
 .../rocke/platform/python/rocke/core/ir.py    |  10 +
 .../platform/python/rocke/core/lower_hip.py   |   7 +
 .../platform/python/rocke/core/lower_llvm.py  |  24 +
 .../python/rocke/instances/SUPPORT_MATRIX.md  |   1 +
 .../rocke_representative_ir_sha256.json       | 129 ++-
 .../instances/rocke_ir_parity_harness.py      |  38 +-
 30 files changed, 3524 insertions(+), 41 deletions(-)
 create mode 100644 dnn-providers/hip-kernel-provider/rocke/library/kernels/common/conv_direct_dw_wgrad_mfma.py
 create mode 100644 dnn-providers/hip-kernel-provider/rocke/library/kernels/common/conv_direct_grouped_wgrad_mfma.py
 create mode 100644 dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_conv_direct_dw_wgrad_mfma.h
 create mode 100644 dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_conv_direct_grouped_wgrad_mfma.h
 create mode 100644 dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/conv_direct_dw_wgrad_mfma.cpp
 create mode 100644 dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/conv_direct_grouped_wgrad_mfma.cpp

diff --git a/dnn-providers/hip-kernel-provider/rocke/library/kernels/common/conv_direct_dw_wgrad_mfma.py b/dnn-providers/hip-kernel-provider/rocke/library/kernels/common/conv_direct_dw_wgrad_mfma.py
new file mode 100644
index 000000000000..1daf8e5aa2a8
--- /dev/null
+++ b/dnn-providers/hip-kernel-provider/rocke/library/kernels/common/conv_direct_dw_wgrad_mfma.py
@@ -0,0 +1,603 @@
+# Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
+# SPDX-License-Identifier: MIT
+
+"""Depthwise wgrad on the batched 4x4x4 MFMA (stride 1, any filter size).
+
+A depthwise wgrad has no channel reduction, so the MFMA's batch axis is the
+only place the channels fit: ``mfma_f32_4x4x4`` runs 16 independent 4x4x4
+products per wave and each takes one channel. Per input row ``y`` and per
+4-column group ``q0 .. q0 + 3`` of the output, one block computes
+
+    A[s][k] = X[y][q0 + k + s - PAD]          (Hankel: shift down the rows)
+    B[k][r] = dY[y + PAD - r][q0 + k]
+    C[s][r] += sum_k A[s][k] B[k][r]          = dW[c][r][s] over those columns
+
+with the filter tiled into 4x4 tap groups (one accumulator each, lanes past
+the filter clamped and dropped by the epilogue). Lane ``l`` of a wave serves
+channel ``l // 4`` of the wave's 16 and supplies row ``l % 4`` of A and column
+``l % 4`` of B; its accumulator holds column ``r`` and items ``s``.
+
+Both operands want consecutive columns of one channel while NHWC keeps the
+channel contiguous, so rows reach them in two steps: an LDS DMA brings the
+row in NHWC order, and ``ds_read_tr16_b64`` hands lane ``L`` four consecutive
+columns of channel ``L``, which it stores channel-major (X into ``main_in``,
+the newest dY row into ``dnew``).
+
+* A: the four lanes of a channel each read a quarter of its ``main_in`` row
+  as aligned 16-byte words and share them with quad broadcasts; every A
+  fragment is then two ``v_perm_b32`` of that window (a lane's tap shift is
+  its byte selector), so no LDS read is unaligned.
+* B: the KH dY rows a row pairs with stay in registers. Lane ``(g, r)``
+  holds dY row ``y + PAD - 4g - r``; each step rotates the quad by one
+  (``quad_perm``), carries lane 3 of group ``g - 1`` into lane 0 of ``g``,
+  and takes the newest row into lane 0 of group 0.
+
+A block owns 64 channels and a ``4 * q_tiles`` column tile (or
+``unfold_n`` small images packed side by side), and walks the (images, row
+chunk) units ``z, z + gridDim.z, ...``; every unit accumulates into the same
+registers and the block stores one partial slot ``ws[part][tap][c]`` that
+:class:`DwWgradFoldSpec` sums into dW. When a unit is a whole image the
+prologue stages only the PAD dY rows below its top edge.
+"""
+
+from __future__ import annotations
+
+from dataclasses import dataclass
+
+from rocke.core.ir import I32, I64, IRBuilder, KernelDef, Value, VectorType
+from rocke.helpers.spec import ceil_div
+
+from kernels.common.conv_direct_grouped import (
+    DirectConvProblem,
+    _io_type,
+    _mfma,
+    dilation_reason,
+    emit_direct_params,
+)
+
+_WAVE = 64
+_NW = 4  # waves per block
+_BC = 64  # channels per block: 16 per wave
+_TAPS = 4  # M = N of the 4x4x4 atom
+_CH = _BC // 8  # 16-byte chunks per staged pixel
+
+
+@dataclass(frozen=True)
+class DwWgradMfmaSpec:
+    problem: DirectConvProblem
+    name: str = "dw_wgrad_mfma"
+    q_tiles: int = 8  # output columns per block / 4
+    rows_per_chunk: int = 16  # input rows per unit
+    prefetch: int = 2  # staging buffers
+    target_blocks: int = 2048
+    # Images packed side by side in one column tile (small images), at the
+    # pitch of image_pitch(): neighbours share their zero padding.
+    unfold_n: int = 1
+
+    @property
+    def block_q(self) -> int:
+        return 4 * self.q_tiles
+
+    @property
+    def rg(self) -> int:
+        return ceil_div(self.problem.KH, _TAPS)
+
+    @property
+    def sg(self) -> int:
+        return ceil_div(self.problem.KW, _TAPS)
+
+    @property
+    def preshift(self) -> bool:
+        # A lane's tap shift reaches past one v_perm (3 halves) from KW 4 on.
+        return min(_TAPS, self.problem.KW) - 1 > 2
+
+    def win_words(self) -> int:
+        return 2 * (self.q_tiles - 1) + 2 * (self.sg - 1) + 3
+
+    def win_reads(self) -> int:
+        return ceil_div(self.win_words() + (1 if self.preshift else 0), 4)
+
+    def in_pitch(self) -> int:
+        """main_in elements per channel: the lane's whole aligned window."""
+        return 8 * self.win_reads()
+
+    def in_groups(self) -> int:
+        """4-column groups of X staged and transposed per row."""
+        cols = 2 * (self.win_words() + (1 if self.preshift else 0))
+        return ceil_div(cols, 4)
+
+    def d_pitch(self) -> int:
+        return self.block_q + 4
+
+    def dnew_elems(self) -> int:
+        return _BC * self.d_pitch()
+
+    def ix(self) -> int:
+        """DMA instructions per wave per staged X row."""
+        return ceil_div(ceil_div(4 * self.in_groups() * _CH, _WAVE), _NW)
+
+    def iy(self) -> int:
+        return ceil_div(ceil_div(self.block_q * _CH, _WAVE), _NW)
+
+    def kernel_name(self) -> str:
+        from rocke.helpers.spec import kernel_name_join
+
+        p = self.problem
+        flags = {"bf16": True} if p.dtype == "bf16" else {}
+        return kernel_name_join(
+            self.name,
+            f"f{p.KH}x{p.KW}p{p.PAD}s{p.stride}",
+            f"qt{self.q_tiles}",
+            f"rpc{self.rows_per_chunk}",
+            f"pf{self.prefetch}",
+            f"un{self.unfold_n}",
+            f"tb{self.target_blocks}",
+            flags=flags,
+        )
+
+    def n_channel_tiles(self) -> int:
+        return ceil_div(self.problem.groups, _BC)
+
+    def n_q(self) -> int:
+        return ceil_div(self.problem.Wo, self.block_q)
+
+    def image_pitch(self) -> int:
+        """Columns between packed images: an output's KW taps stay inside
+        its own image and padding, and the dY rows do not overlap."""
+        p = self.problem
+        return max(p.Wo, p.Wo + p.KW - 1 - p.PAD, p.PAD + p.W)
+
+    def n_units(self) -> int:
+        p = self.problem
+        return ceil_div(p.N, self.unfold_n) * ceil_div(p.H, self.rows_per_chunk)
+
+    def n_z(self) -> int:
+        planes = self.n_channel_tiles() * self.n_q()
+        return min(self.n_units(), max(1, ceil_div(self.target_blocks, planes)))
+
+    def n_partitions(self) -> int:
+        return self.n_q() * self.n_z()
+
+    def workspace_bytes(self) -> int:
+        p = self.problem
+        return self.n_partitions() * p.KH * p.KW * p.groups * 4
+
+    def launch_geometry(self):
+        return (self.n_channel_tiles(), self.n_q(), self.n_z()), (_NW * _WAVE, 1, 1)
+
+    def lds_bytes(self) -> int:
+        stage = self.prefetch * _NW * (self.ix() + self.iy()) * _WAVE * 16
+        return stage + (_BC * self.in_pitch() + self.dnew_elems()) * 2
+
+
+def is_valid_dw_wgrad_mfma_spec(spec: DwWgradMfmaSpec, arch: str = "gfx950"):
+    p = spec.problem
+    why = dilation_reason(p)
+    if why is not None:
+        return False, why
+    if p.cpg != 1 or p.kpg != 1:
+        return False, "depthwise only"
+    if p.dtype not in ("fp16", "bf16"):
+        return False, "fp16/bf16 only"
+    if p.stride != 1:
+        return False, "stride 1 only"
+    if p.PAD > p.KH - 1 or p.PAD > p.KW - 1:
+        return False, "PAD <= K - 1"
+    if p.groups % 8:
+        return False, "C must be a multiple of 8"
+    if not 1 <= spec.q_tiles <= 8:
+        return False, "q_tiles in 1..8"
+    if spec.prefetch not in (1, 2, 3, 4):
+        return False, "prefetch in 1..4"
+    if spec.unfold_n not in (1, 2, 3, 4):
+        return False, "unfold_n in 1..4"
+    if spec.unfold_n > 1 and spec.unfold_n * spec.image_pitch() > spec.block_q:
+        return False, (
+            f"{spec.unfold_n} images at pitch {spec.image_pitch()} exceed the "
+            f"{spec.block_q}-column tile"
+        )
+    if spec.rows_per_chunk < 1:
+        return False, "rows_per_chunk >= 1"
+    if spec.lds_bytes() > 160 * 1024:
+        return False, f"LDS {spec.lds_bytes()} B > 160 KiB"
+    if arch not in ("gfx950",):
+        return False, "ds_read_tr16_b64 needs gfx950"
+    return True, "ok"
+
+
+def build_dw_wgrad_mfma(spec: DwWgradMfmaSpec, arch: str = "gfx950") -> KernelDef:
+    ok, why = is_valid_dw_wgrad_mfma_spec(spec, arch)
+    if not ok:
+        raise ValueError(f"invalid DwWgradMfmaSpec: {why}")
+    p = spec.problem
+    KH, KW, PAD = p.KH, p.KW, p.PAD
+    TAPS = KH * KW
+    QT, BQ = spec.q_tiles, spec.block_q
+    RG, SG = spec.rg, spec.sg
+    D = spec.prefetch
+    IX, IY = spec.ix(), spec.iy()
+    GI = spec.in_groups()
+    XC = 4 * GI  # staged X columns
+    REAL_COLS = BQ + KW - 1
+    HALF = spec.in_pitch()
+    DP = spec.d_pitch()
+    WW = spec.win_words()
+    io_type = _io_type(p.dtype)
+    T = _NW * _WAVE
+
+    b = IRBuilder(spec.kernel_name())
+    b.kernel.attrs["max_workgroup_size"] = T
+    b.kernel.attrs["lds_alias_scopes"] = True
+    params = emit_direct_params(b, direction="wgrad", io_type=io_type)
+
+    def ci(v: int) -> Value:
+        return b.const_i32(v)
+
+    c0, c1, c2 = ci(0), ci(1), ci(2)
+    oob = ci((1 << 31) - 1)
+    p_N = params["p_N"]
+    p_C = params["p_groups"]
+    p_H, p_W = params["p_Hi"], params["p_Wi"]
+    p_Ho, p_Wo = params["p_Ho"], params["p_Wo"]
+    sdy_n, sdy_h, sdy_w = (
+        params["p_dY_stride_n"],
+        params["p_dY_stride_ho"],
+        params["p_dY_stride_wo"],
+    )
+    sx_n, sx_h, sx_w = (
+        params["p_X_stride_n"],
+        params["p_X_stride_hi"],
+        params["p_X_stride_wi"],
+    )
+
+    tid = b.thread_id_x()
+    lane = b.mod(tid, ci(_WAVE))
+    wave = b.readfirstlane(b.div(tid, ci(_WAVE)))
+    ctile, qtile, zslot = b.block_id_x(), b.block_id_y(), b.block_id_z()
+    c_blk0 = b.mul(ctile, ci(_BC))
+    q0 = b.mul(qtile, ci(BQ))
+    n_q = b.div(b.add(p_Wo, ci(BQ - 1)), ci(BQ))
+    UN = spec.unfold_n
+    n_chunks = b.div(b.add(p_H, ci(spec.rows_per_chunk - 1)), ci(spec.rows_per_chunk))
+    n_ngrp = b.div(b.add(p_N, ci(UN - 1)), ci(UN))
+    n_units = b.mul(n_ngrp, n_chunks)
+    # Same formula as DwWgradMfmaSpec.n_z, which sizes the grid.
+    n_ct = b.div(b.add(p_C, ci(_BC - 1)), ci(_BC))
+    planes = b.mul(n_ct, n_q)
+    n_z = b.smin(
+        n_units,
+        b.smax(b.div(b.add(ci(spec.target_blocks), b.sub(planes, c1)), planes), c1),
+    )
+
+    a_rsrc = b.buffer_rsrc(params["A"], params["A_bytes"])  # dY
+    b_rsrc = b.buffer_rsrc(params["B"], params["B_bytes"])  # X
+
+    # ---- LDS: one allocation per staging buffer (one alias scope each) -----
+    X_ELEMS = _NW * IX * _WAVE * 8
+    Y_ELEMS = _NW * IY * _WAVE * 8
+    x_stage = [
+        b.smem_alloc(io_type, [1, X_ELEMS], name_hint=f"xs{i}") for i in range(D)
+    ]
+    y_stage = [
+        b.smem_alloc(io_type, [1, Y_ELEMS], name_hint=f"ys{i}") for i in range(D)
+    ]
+    x_addr = [b.smem_addr_of(s) for s in x_stage]
+    y_addr = [b.smem_addr_of(s) for s in y_stage]
+    MIN = _BC * HALF
+    DNEW = spec.dnew_elems()
+    main_in = b.smem_alloc(io_type, [1, MIN], name_hint="min")
+    # The newest dY row, channel-major; older rows live in registers.
+    dnew = b.smem_alloc(io_type, [1, DNEW], name_hint="dnew")
+
+    # ---- DMA lane geometry: slot L = (wave * I + i) * 64 + lane -> (col, chunk)
+    def dma_lanes(n_instr: int, cols: int):
+        out = []
+        for i in range(n_instr):
+            slot = b.add(b.mul(b.add(b.mul(wave, ci(n_instr)), ci(i)), ci(_WAVE)), lane)
+            col = b.div(slot, ci(_CH))
+            cc = b.add(c_blk0, b.mul(b.mod(slot, ci(_CH)), ci(8)))
+            out.append((col, b.land(b.cmp_lt(col, ci(cols)), b.cmp_lt(cc, p_C)), cc))
+        return out
+
+    # Lane offsets: (offset without the image and row, ok, packed image j).
+    if UN > 1:
+        # Image j owns tile columns [j * P, (j + 1) * P); with one column
+        # tile (q0 = 0) X column x sits at j * P + PAD + x, dY column q at
+        # j * P + q.
+        pitch = b.smax(b.smax(p_Wo, b.add(p_Wo, ci(KW - 1 - PAD))), b.add(p_W, ci(PAD)))
+
+        def split(col):
+            j = b.div(col, pitch)
+            return j, b.sub(col, b.mul(j, pitch))
+
+    else:
+
+        def split(col):
+            return c0, b.add(q0, col)
+
+    x_offs = []
+    for col, okl, cc in dma_lanes(IX, min(XC, REAL_COLS)):
+        j, u = split(col)
+        wi = b.sub(u, ci(PAD))
+        okc = b.land(okl, b.land(b.cmp_ge(wi, c0), b.cmp_lt(wi, p_W)))
+        okc = b.land(okc, b.cmp_lt(j, ci(UN)))
+        off = b.add(b.add(b.mul(wi, sx_w), cc), b.mul(j, sx_n))
+        x_offs.append((off, okc, j))
+    y_offs = []
+    for col, okl, cc in dma_lanes(IY, BQ):
+        j, wo = split(col)
+        ok = b.land(b.land(okl, b.cmp_lt(wo, p_Wo)), b.cmp_lt(j, ci(UN)))
+        off = b.add(b.add(b.mul(wo, sdy_w), cc), b.mul(j, sdy_n))
+        y_offs.append((off, ok, j))
+
+    def dma(rsrc, addr, n_instr, i, voff):
+        lds = b.smem_ptr_add(
+            addr,
+            b.zext(b.mul(b.add(b.mul(wave, ci(n_instr)), ci(i)), ci(_WAVE * 16)), I64),
+        )
+        b.async_buffer_load_lds(rsrc, lds, voff, c0, 4)
+
+    # ---- transpose geometry (lane L ends up with channel L) ---------------
+    # ds_read_tr16_b64: out lane (b, i', j') item x = in lane (b, x, i') item
+    # j' (lanes as 16b + 4i + j), so in lane (b, i, j) reads column 4g + i,
+    # channels 4 * (4b + j) .. + 3.
+    tr_row = b.mod(b.div(lane, ci(4)), ci(4))
+    tr_cg = b.add(b.mul(b.div(lane, ci(16)), ci(4)), b.mod(lane, ci(4)))
+    tr_base = b.add(b.mul(tr_row, ci(_BC)), b.mul(tr_cg, ci(4)))
+
+    def transpose_rows(stage_x, stage_y, with_x: bool):
+        """Staged NHWC rows -> main_in / dnew, channel-major."""
+        reads = []
+        if with_x:
+            for g in range(ceil_div(GI, _NW)):
+                gg = b.add(wave, ci(g * _NW))
+                reads.append((g, gg, True))
+        for g in range(ceil_div(QT, _NW)):
+            gg = b.add(wave, ci(g * _NW))
+            reads.append((g, gg, False))
+
+        def one(gg, is_x):
+            if is_x:
+                v = b.ds_read_tr16_b64(
+                    stage_x, c0, b.add(tr_base, b.mul(gg, ci(4 * _BC))), dtype=io_type
+                )
+                dst = b.add(b.mul(lane, ci(HALF)), b.mul(gg, ci(4)))
+                b.smem_store_vN(main_in, [c0, dst], v, 4)
+            else:
+                v = b.ds_read_tr16_b64(
+                    stage_y, c0, b.add(tr_base, b.mul(gg, ci(4 * _BC))), dtype=io_type
+                )
+                dst = b.add(b.mul(lane, ci(DP)), b.mul(gg, ci(4)))
+                b.smem_store_vN(dnew, [c0, dst], v, 4)
+
+        for g, gg, is_x in reads:
+            n_groups = GI if is_x else QT
+            if (g + 1) * _NW > n_groups:
+                with b.scf_if(b.cmp_lt(gg, ci(n_groups))):
+                    one(gg, is_x)
+            else:
+                one(gg, is_x)
+
+    # ---- per-lane MFMA operand geometry ------------------------------------
+    lane4 = b.mod(lane, ci(4))
+    ch_w = b.add(b.mul(wave, ci(16)), b.div(lane, ci(4)))  # channel in the tile
+    lane_s = b.smin(lane4, ci(min(_TAPS, KW) - 1))
+    if spec.preshift:
+        lane_word = b.div(lane_s, c2)
+    else:
+        lane_word = c0
+    lane_res = b.sub(lane_s, b.mul(lane_word, c2))
+    sel0 = ci(0x03020100)
+    sel1 = ci(0x01010101)
+    pre_sel = b.add(sel0, b.mul(sel1, b.mul(lane_word, ci(4))))
+    res_sel = b.add(sel0, b.mul(sel1, b.mul(lane_res, c2)))
+    frag_t = VectorType(io_type, 4)
+    # The four lanes of a channel each read a quarter of its main_in window
+    # (16 bytes per read) and share them with quad broadcasts.
+    NWR = spec.win_reads()
+    a_reads = [
+        b.add(
+            b.mul(ch_w, ci(HALF)),
+            b.mul(b.smin(b.add(lane4, ci(4 * k)), ci(NWR - 1)), ci(8)),
+        )
+        for k in range(ceil_div(NWR, 4))
+    ]
+    # Likewise the newest dY row: lane j of a channel reads q tiles j, j + 4.
+    d_reads = [
+        b.add(
+            b.mul(ch_w, ci(DP)),
+            b.mul(b.smin(b.add(lane4, ci(4 * k)), ci(QT - 1)), ci(4)),
+        )
+        for k in range(ceil_div(QT, 4))
+    ]
+    head = b.cmp_eq(lane4, c0)
+
+    def bcast(v: Value, src: int) -> Value:
+        return b.quad_perm(v, [src] * 4)
+
+    def words(v: Value, n: int):
+        w = b.vec_bitcast(v, VectorType(I32, n))
+        return [b.vec_extract(w, e) for e in range(n)]
+
+    def shift_ring(ring):
+        """Advance the dY ring one row: lane (g, r) takes (g, r - 1)'s row,
+        lane (g, 0) takes (g - 1, 3)'s, and lane (0, 0) the newest row."""
+        fresh = []
+        for k, idx in enumerate(d_reads):
+            v = b.smem_load_vN(dnew, c0, idx, dtype=io_type, n=4)
+            fresh.append(words(v, 2))
+        rot = [
+            [[b.quad_perm(d, [3, 0, 1, 2]) for d in ring[g][qt]] for qt in range(QT)]
+            for g in range(RG)
+        ]
+        out = []
+        for g in range(RG):
+            row = []
+            for qt in range(QT):
+                if g == 0:
+                    src = [bcast(d, qt % 4) for d in fresh[qt // 4]]
+                else:
+                    src = rot[g - 1][qt]
+                row.append([b.select(head, src[e], rot[g][qt][e]) for e in range(2)])
+            out.append(row)
+        return out
+
+    def compute(ring, accs):
+        """MFMA the X row in main_in against the dY rows in ``ring``."""
+        part = [
+            words(b.smem_load_vN(main_in, c0, idx, dtype=io_type, n=8), 4)
+            for idx in a_reads
+        ]
+        win = [bcast(part[m // 4][e], m % 4) for m in range(NWR) for e in range(4)]
+        if spec.preshift:
+            win = [b.perm_b32(win[j + 1], win[j], pre_sel) for j in range(WW)]
+        bfr = [
+            [b.vec_bitcast(b.vec_pack(ring[g][qt], I32), frag_t) for qt in range(QT)]
+            for g in range(RG)
+        ]
+        accs = list(accs)
+        for sgi in range(SG):
+            for qt in range(QT):
+                k = 2 * qt + 2 * sgi
+                a0 = b.perm_b32(win[k + 1], win[k], res_sel)
+                a1 = b.perm_b32(win[k + 2], win[k + 1], res_sel)
+                af = b.vec_bitcast(b.vec_pack([a0, a1], I32), frag_t)
+                for rgi in range(RG):
+                    i = rgi * SG + sgi
+                    accs[i] = _mfma(b, p.dtype, "4x4x4", af, bfr[rgi][qt], accs[i])
+        return accs
+
+    zero_acc = b.zero_vec_f32(4)
+    n_acc = RG * SG
+    PER_STEP = IX + IY
+
+    def pack(accs, ring):
+        return list(accs) + [d for g in ring for qt in g for d in qt]
+
+    def unpack(vals):
+        accs = list(vals[:n_acc])
+        flat = list(vals[n_acc:])
+        ring = [
+            [flat[(g * QT + qt) * 2 : (g * QT + qt) * 2 + 2] for qt in range(QT)]
+            for g in range(RG)
+        ]
+        return accs, ring
+
+    def emit_body(PRE: int):
+        """Unit loop and epilogue, PRE dY rows staged before a unit's first
+        X row: KH - 1 in general, PAD when a unit is a whole image (the
+        rows above it are padding, which the zeroed ring already holds)."""
+        unit_loop = b.scf_for_iter(
+            zslot,
+            n_units,
+            n_z,
+            [(f"acc{PRE}_{i}", zero_acc) for i in range(n_acc)],
+            iv_name=f"unit{PRE}",
+            elide_trailing_barrier=False,
+        )
+        with unit_loop as (unit, carried):
+            # A fresh ring per unit: its zeros stand in for the dY rows above
+            # the image, which the whole-image prologue does not stage.
+            accs = list(carried)
+            ring = [[[c0, c0] for _ in range(QT)] for _ in range(RG)]
+            n_img = b.mul(b.mod(unit, n_ngrp), ci(UN))
+            chunk = b.div(unit, n_ngrp)
+            y0 = b.mul(chunk, ci(spec.rows_per_chunk))
+            n_rows = b.smin(ci(spec.rows_per_chunk), b.sub(p_H, y0))
+            n_steps = b.add(n_rows, ci(PRE))
+            x_img = b.mul(n_img, sx_n)
+            y_img = b.mul(n_img, sdy_n)
+
+            def issue(buf: int, t: Value):
+                """DMA step t: X row y0 - PRE + t (steps >= PRE) and dY row y + PAD."""
+                y = b.add(b.sub(y0, ci(PRE)), t)
+                x_ok = b.land(b.cmp_ge(t, ci(PRE)), b.cmp_lt(t, n_steps))
+                x_row = b.add(x_img, b.mul(y, sx_h))
+                for i, (off, okl, j) in enumerate(x_offs):
+                    ok = b.land(okl, x_ok)
+                    if UN > 1:
+                        ok = b.land(ok, b.cmp_lt(b.add(n_img, j), p_N))
+                    voff = b.select(ok, b.mul(b.add(off, x_row), c2), oob)
+                    dma(b_rsrc, x_addr[buf], IX, i, voff)
+                pr = b.add(y, ci(PAD))
+                y_ok = b.land(
+                    b.land(b.cmp_ge(pr, c0), b.cmp_lt(pr, p_Ho)), b.cmp_lt(t, n_steps)
+                )
+                y_row = b.add(y_img, b.mul(pr, sdy_h))
+                for i, (off, okl, j) in enumerate(y_offs):
+                    ok = b.land(okl, y_ok)
+                    if UN > 1:
+                        ok = b.land(ok, b.cmp_lt(b.add(n_img, j), p_N))
+                    voff = b.select(ok, b.mul(b.add(off, y_row), c2), oob)
+                    dma(a_rsrc, y_addr[buf], IY, i, voff)
+
+            def step(t: Value, buf: int, accs, ring, do_compute: bool):
+                # Issue the row D - 1 ahead into the buffer the last step drained.
+                issue((buf + D - 1) % D, b.add(t, ci(D - 1)))
+                b.s_waitcnt(vmcnt=(D - 1) * PER_STEP, lgkmcnt=0)
+                b.s_barrier_bare()
+                transpose_rows(x_stage[buf], y_stage[buf], do_compute)
+                b.s_waitcnt(lgkmcnt=0)
+                b.s_barrier_bare()
+                ring = shift_ring(ring)
+                if do_compute:
+                    accs = compute(ring, accs)
+                return accs, ring
+
+            for t in range(D - 1):
+                issue(t % D, ci(t))
+            # Prologue: the PRE dY rows the first input row pairs with.
+            for t in range(PRE):
+                accs, ring = step(ci(t), t % D, accs, ring, False)
+            # Main: one input row per step, unrolled by D so buffers are constant.
+            trips = b.smax(b.div(b.add(n_rows, ci(D - 1)), ci(D)), c1)
+            row_loop = b.scf_for_iter(
+                c0,
+                trips,
+                c1,
+                [(f"a{PRE}_{i}", v) for i, v in enumerate(pack(accs, ring))],
+                iv_name=f"rb{PRE}",
+                elide_trailing_barrier=False,
+            )
+            with row_loop as (rb, la):
+                la_acc, la_ring = unpack(list(la))
+                for u in range(D):
+                    t = b.add(b.add(b.mul(rb, ci(D)), ci(PRE)), ci(u))
+                    la_acc, la_ring = step(t, (PRE + u) % D, la_acc, la_ring, True)
+                b.scf_yield(*pack(la_acc, la_ring))
+            # Drain the tail DMAs before the next unit reuses the buffers.
+            b.s_waitcnt(vmcnt=0, lgkmcnt=0)
+            b.s_barrier_bare()
+            b.scf_yield(*unpack(list(row_loop.results))[0])
+        accs = list(unit_loop.results)
+
+        # ---- epilogue: ws[part][tap][c], part = zslot * n_q + qtile -------------
+        part = b.add(b.mul(zslot, n_q), qtile)
+        c_glb = b.add(c_blk0, ch_w)
+        c_ok = b.cmp_lt(c_glb, p_C)
+        for rgi in range(RG):
+            r = b.add(lane4, ci(rgi * _TAPS))
+            row0 = b.add(
+                b.mul(b.add(b.mul(part, ci(TAPS)), b.mul(r, ci(KW))), p_C), c_glb
+            )
+            with b.scf_if(b.land(c_ok, b.cmp_lt(r, ci(KH)))):
+                for sgi in range(SG):
+                    acc = accs[rgi * SG + sgi]
+                    for i in range(_TAPS):
+                        s = sgi * _TAPS + i
+                        if s < KW:
+                            off = b.add(row0, b.mul(ci(s), p_C))
+                            b.global_store(
+                                params["D"], off, b.vec_extract(acc, i), align=4
+                            )
+
+    if PAD == KH - 1:
+        emit_body(PAD)
+    else:
+        # Whole images per unit (one row chunk) skip the KH - 1 - PAD zero rows.
+        with b.scf_if_else(b.cmp_eq(n_chunks, c1)) as (whole, chunked):
+            with whole:
+                emit_body(PAD)
+            with chunked:
+                emit_body(KH - 1)
+    return b.kernel
diff --git a/dnn-providers/hip-kernel-provider/rocke/library/kernels/common/conv_direct_grouped_wgrad_mfma.py b/dnn-providers/hip-kernel-provider/rocke/library/kernels/common/conv_direct_grouped_wgrad_mfma.py
new file mode 100644
index 000000000000..ff382f08e5ce
--- /dev/null
+++ b/dnn-providers/hip-kernel-provider/rocke/library/kernels/common/conv_direct_grouped_wgrad_mfma.py
@@ -0,0 +1,432 @@
+# Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
+# SPDX-License-Identifier: MIT
+
+"""Grouped wgrad for 4 and 8 channels per group on the batched 4x4x4 MFMA.
+
+With ``cpg = kpg = 4`` a group's weight gradient per tap is one 4x4 block,
+which is exactly one batch of ``mfma_f32_4x4x4``: a wave runs 16 groups at
+once. Per input row ``y``, tap ``(r, s)`` and 4-column group ``q0``:
+
+    A[c][k] = X[y][q0 + k + s - PAD][c]      (lane = channel c, 4 columns)
+    B[k][o] = dY[y + PAD - r][q0 + k][o]
+    C[c][o] += sum_k A[c][k] B[k][o]          = dW[o][r][s][c] over those columns
+
+Both operands come straight out of the NHWC row staged in LDS:
+``ds_read_tr16_b64`` hands lane ``L`` four consecutive columns of channel
+``L``, and the tap shift ``s`` is just a row offset of the per-lane read
+addresses, so every read is aligned and nothing is re-laid out.
+
+``cpg = kpg = 8``: a group is four 4x4 blocks. Lanes ``4b .. 4b + 3`` hold
+channels ``4b ..`` of the A read; a second dY read with the two 4-channel
+halves of each group swapped (``channel ^ 4``) gives the off-diagonal blocks,
+so each tap carries two accumulators.
+
+A block of ``waves_q`` waves owns 64 channels and ``4 * wave_q4 * waves_q``
+output columns (each wave its own columns) and walks (image, row chunk)
+units; dY rows stay in a register ring of KH rows. The waves reduce through
+LDS at the end and the block stores one partial slot ``ws[part][tap][o]``
+(``tap = (r * KW + s) * cpg + c``) that ``DwWgradFoldSpec`` folds into dW.
+"""
+
+from __future__ import annotations
+
+from dataclasses import dataclass
+
+from rocke.core.ir import F32, I64, IRBuilder, KernelDef, Value
+from rocke.helpers.spec import ceil_div
+
+from kernels.common.conv_direct_grouped import (
+    DirectConvProblem,
+    _io_type,
+    _mfma,
+    dilation_reason,
+    emit_direct_params,
+)
+
+_WAVE = 64
+_BC = 64  # channels per block
+_CH = _BC // 8  # 16-byte chunks per staged pixel
+
+
+@dataclass(frozen=True)
+class GroupedWgradMfmaSpec:
+    problem: DirectConvProblem
+    name: str = "grouped_wgrad_mfma"
+    waves_q: int = 4  # waves along the output columns
+    wave_q4: int = 1  # 4-column groups per wave
+    rows_per_chunk: int = 56  # input rows per unit
+    prefetch: int = 2  # staging buffers
+    target_blocks: int = 1024
+
+    @property
+    def threads(self) -> int:
+        return self.waves_q * _WAVE
+
+    @property
+    def block_q(self) -> int:
+        return 4 * self.wave_q4 * self.waves_q
+
+    @property
+    def halves(self) -> int:
+        """4x4 blocks along k per group: 1 for cpg 4, 2 for cpg 8."""
+        return self.problem.cpg // 4
+
+    def x_cols(self) -> int:
+        return self.block_q + self.problem.KW - 1
+
+    def ix(self) -> int:
+        return ceil_div(ceil_div(self.x_cols() * _CH, _WAVE), self.waves_q)
+
+    def iy(self) -> int:
+        return ceil_div(ceil_div(self.block_q * _CH, _WAVE), self.waves_q)
+
+    def kernel_name(self) -> str:
+        from rocke.helpers.spec import kernel_name_join
+
+        p = self.problem
+        flags = {"bf16": True} if p.dtype == "bf16" else {}
+        return kernel_name_join(
+            self.name,
+            f"f{p.KH}x{p.KW}p{p.PAD}s{p.stride}c{p.cpg}",
+            f"wq{self.waves_q}",
+            f"q4{self.wave_q4}",
+            f"rpc{self.rows_per_chunk}",
+            f"pf{self.prefetch}",
+            f"tb{self.target_blocks}",
+            flags=flags,
+        )
+
+    def n_channel_tiles(self) -> int:
+        return ceil_div(self.problem.total_c, _BC)
+
+    def n_q(self) -> int:
+        return ceil_div(self.problem.Wo, self.block_q)
+
+    def n_units(self) -> int:
+        p = self.problem
+        return p.N * ceil_div(p.H, self.rows_per_chunk)
+
+    def n_z(self) -> int:
+        planes = self.n_channel_tiles() * self.n_q()
+        return min(self.n_units(), max(1, ceil_div(self.target_blocks, planes)))
+
+    def n_partitions(self) -> int:
+        return self.n_q() * self.n_z()
+
+    def fold_taps(self) -> int:
+        p = self.problem
+        return p.KH * p.KW * p.cpg
+
+    def fold_channels(self) -> int:
+        return self.problem.total_k
+
+    def workspace_bytes(self) -> int:
+        return self.n_partitions() * self.fold_taps() * self.problem.total_k * 4
+
+    def launch_geometry(self):
+        return (self.n_channel_tiles(), self.n_q(), self.n_z()), (self.threads, 1, 1)
+
+    def n_acc(self) -> int:
+        p = self.problem
+        return p.KH * p.KW * self.halves
+
+    def lds_bytes(self) -> int:
+        stage = self.prefetch * self.waves_q * (self.ix() + self.iy()) * _WAVE * 16
+        red = self.waves_q * _WAVE * 4 * 4
+        return stage + red
+
+
+def is_valid_grouped_wgrad_mfma_spec(spec: GroupedWgradMfmaSpec, arch: str = "gfx950"):
+    p = spec.problem
+    why = dilation_reason(p)
+    if why is not None:
+        return False, why
+    if p.cpg != p.kpg or p.cpg not in (4, 8):
+        return False, "cpg = kpg in (4, 8)"
+    if p.dtype not in ("fp16", "bf16"):
+        return False, "fp16/bf16 only"
+    if p.stride != 1:
+        return False, "stride 1 only"
+    if p.PAD > p.KH - 1 or p.PAD > p.KW - 1:
+        return False, "PAD <= K - 1"
+    if p.total_c % 8:
+        return False, "C must be a multiple of 8"
+    if spec.waves_q not in (1, 2, 4, 8):
+        return False, "waves_q in (1, 2, 4, 8)"
+    if spec.wave_q4 not in (1, 2, 4):
+        return False, "wave_q4 in (1, 2, 4)"
+    if spec.prefetch not in (2, 3, 4):
+        return False, "prefetch in 2..4"
+    if spec.rows_per_chunk < 1:
+        return False, "rows_per_chunk >= 1"
+    if spec.n_acc() * 4 > 128:
+        return False, "too many accumulators"
+    if spec.lds_bytes() > 160 * 1024:
+        return False, f"LDS {spec.lds_bytes()} B > 160 KiB"
+    if arch not in ("gfx950",):
+        return False, "ds_read_tr16_b64 needs gfx950"
+    return True, "ok"
+
+
+def build_grouped_wgrad_mfma(
+    spec: GroupedWgradMfmaSpec, arch: str = "gfx950"
+) -> KernelDef:
+    ok, why = is_valid_grouped_wgrad_mfma_spec(spec, arch)
+    if not ok:
+        raise ValueError(f"invalid GroupedWgradMfmaSpec: {why}")
+    p = spec.problem
+    KH, KW, PAD, CPG = p.KH, p.KW, p.PAD, p.cpg
+    HV = spec.halves
+    TAPS = spec.fold_taps()
+    WQ, Q4 = spec.waves_q, spec.wave_q4
+    BQ = spec.block_q
+    XC = spec.x_cols()
+    D = spec.prefetch
+    IX, IY = spec.ix(), spec.iy()
+    T = spec.threads
+    PRE = KH - 1  # dY rows a unit needs before its first input row
+    io_type = _io_type(p.dtype)
+
+    b = IRBuilder(spec.kernel_name())
+    b.kernel.attrs["max_workgroup_size"] = T
+    b.kernel.attrs["lds_alias_scopes"] = True
+    params = emit_direct_params(b, direction="wgrad", io_type=io_type)
+
+    def ci(v: int) -> Value:
+        return b.const_i32(v)
+
+    c0, c1, c2 = ci(0), ci(1), ci(2)
+    oob = ci((1 << 31) - 1)
+    p_N = params["p_N"]
+    p_C = params["p_total_c"]
+    p_K = params["p_total_k"]
+    p_H, p_W = params["p_Hi"], params["p_Wi"]
+    p_Ho, p_Wo = params["p_Ho"], params["p_Wo"]
+    sdy_n, sdy_h, sdy_w = (
+        params["p_dY_stride_n"],
+        params["p_dY_stride_ho"],
+        params["p_dY_stride_wo"],
+    )
+    sx_n, sx_h, sx_w = (
+        params["p_X_stride_n"],
+        params["p_X_stride_hi"],
+        params["p_X_stride_wi"],
+    )
+
+    tid = b.thread_id_x()
+    lane = b.mod(tid, ci(_WAVE))
+    wave = b.readfirstlane(b.div(tid, ci(_WAVE)))
+    ctile, qtile, zslot = b.block_id_x(), b.block_id_y(), b.block_id_z()
+    c_blk0 = b.mul(ctile, ci(_BC))
+    q0 = b.mul(qtile, ci(BQ))
+    n_q = b.div(b.add(p_Wo, ci(BQ - 1)), ci(BQ))
+    n_chunks = b.div(b.add(p_H, ci(spec.rows_per_chunk - 1)), ci(spec.rows_per_chunk))
+    n_units = b.mul(p_N, n_chunks)
+    # Same formula as GroupedWgradMfmaSpec.n_z, which sizes the grid.
+    n_ct = b.div(b.add(p_C, ci(_BC - 1)), ci(_BC))
+    planes = b.mul(n_ct, n_q)
+    n_z = b.smin(
+        n_units,
+        b.smax(b.div(b.add(ci(spec.target_blocks), b.sub(planes, c1)), planes), c1),
+    )
+
+    a_rsrc = b.buffer_rsrc(params["A"], params["A_bytes"])  # dY
+    b_rsrc = b.buffer_rsrc(params["B"], params["B_bytes"])  # X
+
+    # ---- LDS: one allocation per staging buffer (one alias scope each) -----
+    X_ELEMS = WQ * IX * _WAVE * 8
+    Y_ELEMS = WQ * IY * _WAVE * 8
+    x_stage = [
+        b.smem_alloc(io_type, [1, X_ELEMS], name_hint=f"xs{i}") for i in range(D)
+    ]
+    y_stage = [
+        b.smem_alloc(io_type, [1, Y_ELEMS], name_hint=f"ys{i}") for i in range(D)
+    ]
+    x_addr = [b.smem_addr_of(s) for s in x_stage]
+    y_addr = [b.smem_addr_of(s) for s in y_stage]
+
+    # ---- DMA lane geometry: slot L = (wave * I + i) * 64 + lane -> (col, chunk)
+    def dma_lanes(n_instr: int, cols: int):
+        out = []
+        for i in range(n_instr):
+            slot = b.add(b.mul(b.add(b.mul(wave, ci(n_instr)), ci(i)), ci(_WAVE)), lane)
+            col = b.div(slot, ci(_CH))
+            cc = b.add(c_blk0, b.mul(b.mod(slot, ci(_CH)), ci(8)))
+            out.append((col, b.land(b.cmp_lt(col, ci(cols)), b.cmp_lt(cc, p_C)), cc))
+        return out
+
+    x_offs = []
+    for col, okl, cc in dma_lanes(IX, XC):
+        wi = b.add(b.sub(q0, ci(PAD)), col)
+        okc = b.land(okl, b.land(b.cmp_ge(wi, c0), b.cmp_lt(wi, p_W)))
+        x_offs.append((b.add(b.mul(wi, sx_w), cc), okc))
+    y_offs = []
+    for col, okl, cc in dma_lanes(IY, BQ):
+        wo = b.add(q0, col)
+        y_offs.append((b.add(b.mul(wo, sdy_w), cc), b.land(okl, b.cmp_lt(wo, p_Wo))))
+
+    def dma(rsrc, addr, n_instr, i, voff):
+        lds = b.smem_ptr_add(
+            addr,
+            b.zext(b.mul(b.add(b.mul(wave, ci(n_instr)), ci(i)), ci(_WAVE * 16)), I64),
+        )
+        b.async_buffer_load_lds(rsrc, lds, voff, c0, 4)
+
+    # ---- operand reads ------------------------------------------------------
+    # ds_read_tr16_b64: out lane (b, i', j') item x = in lane (b, x, i') item
+    # j' (lanes as 16b + 4i + j), so in lane (b, i, j) reads column col0 + i,
+    # channels 4 * (4b + j) .. + 3, and out lane L gets channel L.
+    tr_row = b.mod(b.div(lane, ci(4)), ci(4))
+    tr_cg = b.add(b.mul(b.div(lane, ci(16)), ci(4)), b.mod(lane, ci(4)))
+    tr_base = b.add(b.mul(tr_row, ci(_BC)), b.mul(tr_cg, ci(4)))
+    # The partner half of an 8-channel group (channel ^ 4).
+    tr_swap = b.add(b.mul(tr_row, ci(_BC)), b.mul(b.xor(tr_cg, c1), ci(4)))
+    col_w = b.mul(wave, ci(4 * Q4))  # this wave's first column in the tile
+
+    def read(stage, base: Value, col: int) -> Value:
+        idx = b.add(base, b.mul(b.add(col_w, ci(col)), ci(_BC)))
+        return b.ds_read_tr16_b64(stage, c0, idx, dtype=io_type)
+
+    zero_acc = b.zero_vec_f32(4)
+    zero_frag = b.zero_vec(io_type, 4)
+    n_acc = KH * KW * HV
+    # The carried ring holds the dY fragments of the KH - 1 rows before the
+    # current one, newest first: [row][q4][half].
+    n_ring = (KH - 1) * Q4 * HV
+
+    def acc_i(r, s, h):
+        return (r * KW + s) * HV + h
+
+    def pack(accs, ring):
+        return list(accs) + [v for row in ring for q in row for v in q]
+
+    def unpack(vals):
+        accs = list(vals[:n_acc])
+        flat = list(vals[n_acc:])
+        ring = []
+        for j in range(KH - 1):
+            ring.append(
+                [flat[(j * Q4 + q) * HV : (j * Q4 + q) * HV + HV] for q in range(Q4)]
+            )
+        return accs, ring
+
+    unit_loop = b.scf_for_iter(
+        zslot,
+        n_units,
+        n_z,
+        [(f"acc{i}", zero_acc) for i in range(n_acc)]
+        + [(f"ring{i}", zero_frag) for i in range(n_ring)],
+        iv_name="unit",
+        elide_trailing_barrier=False,
+    )
+    with unit_loop as (unit, carried):
+        accs, ring = unpack(list(carried))
+        n_img = b.mod(unit, p_N)
+        chunk = b.div(unit, p_N)
+        y0 = b.mul(chunk, ci(spec.rows_per_chunk))
+        n_rows = b.smin(ci(spec.rows_per_chunk), b.sub(p_H, y0))
+        n_steps = b.add(n_rows, ci(PRE))
+        x_img = b.mul(n_img, sx_n)
+        y_img = b.mul(n_img, sdy_n)
+        PER_STEP = IX + IY
+
+        def issue(buf: int, t: Value):
+            """DMA step t: X row y0 - PRE + t (steps >= PRE) and dY row y + PAD."""
+            y = b.add(b.sub(y0, ci(PRE)), t)
+            x_ok = b.land(b.cmp_ge(t, ci(PRE)), b.cmp_lt(t, n_steps))
+            x_row = b.add(x_img, b.mul(y, sx_h))
+            for i, (off, okl) in enumerate(x_offs):
+                voff = b.select(b.land(okl, x_ok), b.mul(b.add(off, x_row), c2), oob)
+                dma(b_rsrc, x_addr[buf], IX, i, voff)
+            pr = b.add(y, ci(PAD))
+            y_ok = b.land(
+                b.land(b.cmp_ge(pr, c0), b.cmp_lt(pr, p_Ho)), b.cmp_lt(t, n_steps)
+            )
+            y_row = b.add(y_img, b.mul(pr, sdy_h))
+            for i, (off, okl) in enumerate(y_offs):
+                voff = b.select(b.land(okl, y_ok), b.mul(b.add(off, y_row), c2), oob)
+                dma(a_rsrc, y_addr[buf], IY, i, voff)
+
+        def step(t: Value, buf: int, accs, ring, do_compute: bool):
+            """Barrier, DMA step t + D - 1, read and MFMA step t."""
+            # Step t's DMA has landed (steps up to t + D - 2 may still fly)
+            # and every wave is done reading the buffer refilled below.
+            b.s_waitcnt(vmcnt=(D - 2) * PER_STEP, lgkmcnt=0)
+            b.s_barrier_bare()
+            issue((buf + D - 1) % D, b.add(t, ci(D - 1)))
+            newest = []
+            for q in range(Q4):
+                h = [read(y_stage[buf], tr_base, 4 * q)]
+                if HV == 2:
+                    h.append(read(y_stage[buf], tr_swap, 4 * q))
+                newest.append(h)
+            full = [newest] + list(ring)  # full[r] = the dY row of tap r
+            if do_compute:
+                accs = list(accs)
+                for s in range(KW):
+                    for q in range(Q4):
+                        a = read(x_stage[buf], tr_base, 4 * q + s)
+                        for r in range(KH):
+                            for h in range(HV):
+                                i = acc_i(r, s, h)
+                                accs[i] = _mfma(
+                                    b, p.dtype, "4x4x4", a, full[r][q][h], accs[i]
+                                )
+            return accs, full[: KH - 1]
+
+        for t in range(D - 1):
+            issue(t, ci(t))
+        # Prologue: the PRE dY rows the first input row pairs with.
+        for t in range(PRE):
+            accs, ring = step(ci(t), t % D, accs, ring, False)
+        # Main: one input row per step, unrolled by D so buffers are constant.
+        trips = b.smax(b.div(b.add(n_rows, ci(D - 1)), ci(D)), c1)
+        row_loop = b.scf_for_iter(
+            c0,
+            trips,
+            c1,
+            [(f"a{i}", v) for i, v in enumerate(pack(accs, ring))],
+            iv_name="rb",
+            elide_trailing_barrier=False,
+        )
+        with row_loop as (rb, la):
+            la_acc, la_ring = unpack(list(la))
+            for u in range(D):
+                t = b.add(b.add(b.mul(rb, ci(D)), ci(PRE)), ci(u))
+                la_acc, la_ring = step(t, (PRE + u) % D, la_acc, la_ring, True)
+            b.scf_yield(*pack(la_acc, la_ring))
+        # Drain the tail DMAs before the next unit reuses the buffers.
+        b.s_waitcnt(vmcnt=0, lgkmcnt=0)
+        b.s_barrier_bare()
+        b.scf_yield(*list(row_loop.results))
+    accs, _ = unpack(list(unit_loop.results))
+
+    # ---- epilogue: sum the waves in LDS, store ws[part][tap][o] -------------
+    # Lane L, item i of an accumulator of half h: X channel c_L = 4 * (L // 4)
+    # + i, dY channel o_L = (4 * (L // 4)) ^ (4 * h) + L % 4, both in the tile.
+    red = b.smem_alloc(F32, [1, WQ * _WAVE * 4], name_hint="red")
+    part = b.add(b.mul(zslot, n_q), qtile)
+    lane4 = b.mod(lane, ci(4))
+    quad = b.mul(b.div(lane, ci(4)), ci(4))
+    for i_acc in range(n_acc):
+        r, s, h = i_acc // (KW * HV), (i_acc // HV) % KW, i_acc % HV
+        b.smem_store_vN_f32(red, [c0, b.mul(tid, ci(4))], accs[i_acc], 4)
+        b.sync()
+        with b.scf_if(b.cmp_lt(tid, ci(_WAVE))):
+            tot = b.smem_load_vN_f32(red, c0, b.mul(lane, ci(4)), n=4)
+            for w in range(1, WQ):
+                v = b.smem_load_vN_f32(
+                    red, c0, b.mul(b.add(lane, ci(w * _WAVE)), ci(4)), n=4
+                )
+                tot = b.vector_add(tot, v)
+            o_loc = b.add(b.xor(quad, ci(4 * h)), lane4)
+            o_glb = b.add(c_blk0, o_loc)
+            # Channel c of the group: (4 * (L // 4) + i) % cpg.
+            with b.scf_if(b.cmp_lt(o_glb, p_K)):
+                for e in range(4):
+                    c_in = b.mod(b.add(quad, ci(e)), ci(CPG))
+                    tap = b.add(ci((r * KW + s) * CPG), c_in)
+                    off = b.add(b.mul(b.add(b.mul(part, ci(TAPS)), tap), p_K), o_glb)
+                    b.global_store(params["D"], off, b.vec_extract(tot, e), align=4)
+        b.sync()
+    return b.kernel
diff --git a/dnn-providers/hip-kernel-provider/rocke/library/kernels/conv/direct_aot_jobs.py b/dnn-providers/hip-kernel-provider/rocke/library/kernels/conv/direct_aot_jobs.py
index 20b172860e6c..381754effe22 100644
--- a/dnn-providers/hip-kernel-provider/rocke/library/kernels/conv/direct_aot_jobs.py
+++ b/dnn-providers/hip-kernel-provider/rocke/library/kernels/conv/direct_aot_jobs.py
@@ -171,6 +171,22 @@
 WGRAD_GATHER_PREFETCH = (1, 2)
 WGRAD_GATHER_RPS = (8, 16)
 WGRAD_GATHER_BLOCKS = (256, 1024, 2048)
+# Depthwise wgrad on the batched 4x4x4 MFMA (DwWgradMfmaSpec): (4-column
+# tiles per block, images packed per tile), input rows per unit, staging
+# buffers, workgroups. Measured on gfx950 bf16 7x7 layers (ConvNeXt):
+# 16-wide tiles lead below 32 columns, 32-wide on 56x56; 7x7 images pack 2 to
+# a 20-wide or 3 to a 32-wide tile.
+WGRAD_DW_MFMA_TILES = ((4, 1), (8, 1), (5, 2), (8, 3))
+WGRAD_DW_MFMA_RPC = (7, 14, 28)
+WGRAD_DW_MFMA_PREFETCH = (1, 2)
+WGRAD_DW_MFMA_BLOCKS = (1024, 2048)
+# Grouped wgrad, 4 or 8 channels per group, on the batched 4x4x4 MFMA
+# (GroupedWgradMfmaSpec): (waves_q, wave_q4), rows per unit, staging
+# buffers, workgroups.
+WGRAD_GROUPED_MFMA_WAVES = ((4, 2), (4, 1), (2, 1), (2, 2))
+WGRAD_GROUPED_MFMA_RPC = (28, 56)
+WGRAD_GROUPED_MFMA_PREFETCH = (2, 3)
+WGRAD_GROUPED_MFMA_BLOCKS = (1024, 2048)
 
 # Runtime-channel MFMA dgrad (direct_dgrad_mfma_rtc). Its compute pass is the
 # forward kernel on the transposed problem, whose K reduction runs over the
@@ -263,6 +279,8 @@ def _grouped_channels() -> Tuple[int, ...]:
 _WGRAD_FILTERS = (3, 5, 7)
 # Input channel counts the dense small-C gather wgrad bakes (KW * C <= 96).
 _GATHER_CHANNELS = (3, 4, 8, 16, 24, 32)
+# Channels per group (= filters per group) the 4x4x4 MFMA grouped wgrad bakes.
+_GROUPED_MFMA_CHANNELS = (4, 8)
 # Dilated 3x3 rows of the non-grouped kernel: (PAD, stride, dil_h, dil_w). The
 # kernel bounds its rows and columns by the runtime output extents, so any
 # padding works; the rows bake the "same" padding of the symmetric dilations
@@ -457,6 +475,26 @@ def _col_pads(kh: int) -> Tuple[int, ...]:
             for pad in _col_pads(k)
         ),
     ),
+    # Depthwise wgrad on the batched 4x4x4 MFMA: every depthwise filter,
+    # stride 1, C runtime (a multiple of 8).
+    "direct_wgrad_dw_mfma": (
+        "wgrad",
+        tuple(
+            DirectCaps(KH=k, PAD=pad, stride=1, cpg=1, kpg=1)
+            for k in _DEPTHWISE_FILTERS
+            for pad in _col_pads(k)
+        ),
+    ),
+    # Grouped wgrad with 4 or 8 channels per group (cpg = kpg baked, groups
+    # runtime), 3x3 stride 1, on the batched 4x4x4 MFMA.
+    "direct_wgrad_grouped_mfma": (
+        "wgrad",
+        tuple(
+            DirectCaps(KH=3, PAD=pad, stride=1, cpg=c, kpg=c)
+            for c in _GROUPED_MFMA_CHANNELS
+            for pad in _col_pads(3)
+        ),
+    ),
     # Dense small-C wgrad, any stride: C is baked, K runtime.
     "direct_wgrad_gather": (
         "wgrad",
@@ -649,6 +687,34 @@ def _knob_grid(
             WGRAD_DW_COLS_BLOCKS,
         ):
             yield dict(channels=c, strip=st, threads=th, target_blocks=tb)
+    elif variant == "direct_wgrad_dw_mfma":
+        for (qt, un), rpc, pf, tb in itertools.product(
+            WGRAD_DW_MFMA_TILES,
+            WGRAD_DW_MFMA_RPC,
+            WGRAD_DW_MFMA_PREFETCH,
+            WGRAD_DW_MFMA_BLOCKS,
+        ):
+            yield dict(
+                q_tiles=qt,
+                rows_per_chunk=rpc,
+                prefetch=pf,
+                target_blocks=tb,
+                unfold_n=un,
+            )
+    elif variant == "direct_wgrad_grouped_mfma":
+        for (wq, q4), rpc, pf, tb in itertools.product(
+            WGRAD_GROUPED_MFMA_WAVES,
+            WGRAD_GROUPED_MFMA_RPC,
+            WGRAD_GROUPED_MFMA_PREFETCH,
+            WGRAD_GROUPED_MFMA_BLOCKS,
+        ):
+            yield dict(
+                waves_q=wq,
+                wave_q4=q4,
+                rows_per_chunk=rpc,
+                prefetch=pf,
+                target_blocks=tb,
+            )
     elif variant == "direct_wgrad_gather":
         for (wk16, wk), pf, rps, tb in itertools.product(
             WGRAD_GATHER_WAVES,
@@ -830,6 +896,20 @@ def _dw_cols_probe_width(caps: DirectCaps, channels: int, strip: int) -> int:
     return 64
 
 
+def _dw_mfma_probe_width(caps: DirectCaps, knobs: dict) -> int:
+    """The narrowest W whose ``unfold_n`` packed images fit the depthwise
+    MFMA wgrad's column tile (W is a runtime extent, so any such width probes
+    the same binary), or 64 when none does."""
+    for w in range(1, 65):
+        wo = w + 2 * caps.PAD - caps.KH + 1
+        if wo < 1:
+            continue
+        pitch = max(wo, wo + caps.KH - 1 - caps.PAD, caps.PAD + w)
+        if knobs["unfold_n"] * pitch <= 4 * knobs["q_tiles"]:
+            return w
+    return 64
+
+
 def probe_problem(
     caps: DirectCaps,
     dtype: str = DIRECT_DTYPE,
@@ -872,6 +952,21 @@ def probe_problem(
             stride=caps.stride,
             dtype=dtype,
         )
+    if variant == "direct_wgrad_dw_mfma" and knobs.get("unfold_n", 1) > 1:
+        # Packed images: the narrowest image ``unfold_n`` of fit the tile.
+        return DirectConvProblem(
+            N=knobs["unfold_n"],
+            H=64,
+            W=_dw_mfma_probe_width(caps, knobs),
+            groups=256,
+            cpg=1,
+            kpg=1,
+            KH=caps.KH,
+            KW=caps.KW,
+            PAD=caps.PAD,
+            stride=caps.stride,
+            dtype=dtype,
+        )
     if variant == "direct_wgrad_gather":
         # Dense, C baked, K runtime: a K wider than every block_k.
         return DirectConvProblem(
@@ -999,8 +1094,12 @@ def _caps_jobs(
     n_raw = 0
     rtc_dgrad = variant == "direct_dgrad_mfma_rtc"
     # The runtime-channel dgrad probe follows each knob's tile counts, the
-    # small-C depthwise wgrad's its baked channel count.
-    knob_probe = rtc_dgrad or variant == "direct_wgrad_dw_cols"
+    # small-C depthwise wgrad's its baked channel count, the depthwise MFMA
+    # wgrad's its packed images (a narrow enough image).
+    knob_probe = rtc_dgrad or variant in (
+        "direct_wgrad_dw_cols",
+        "direct_wgrad_dw_mfma",
+    )
     probe = None if knob_probe else probe_problem(caps, dtype, variant)
     for knobs in _knob_grid(variant, arch, dtype, caps):
         n_raw += 1
diff --git a/dnn-providers/hip-kernel-provider/rocke/library/kernels/conv/direct_variants.py b/dnn-providers/hip-kernel-provider/rocke/library/kernels/conv/direct_variants.py
index 3aca472670d3..a7047b8c789b 100644
--- a/dnn-providers/hip-kernel-provider/rocke/library/kernels/conv/direct_variants.py
+++ b/dnn-providers/hip-kernel-provider/rocke/library/kernels/conv/direct_variants.py
@@ -68,7 +68,9 @@
 DIRECT_WGRAD_FOLDED = (
     "direct_wgrad_dw",
     "direct_wgrad_dw_cols",
+    "direct_wgrad_dw_mfma",
     "direct_wgrad_gather",
+    "direct_wgrad_grouped_mfma",
 )
 _HELPERS = (
     DIRECT_DGRAD_TRANSPOSE,
@@ -179,6 +181,20 @@ def make_direct_spec(variant: str, problem, knobs: dict):
         return WgradGatherSpec(
             problem=problem, name="rocke_direct_wgrad_gather", **knobs
         )
+    if variant == "direct_wgrad_dw_mfma":
+        from kernels.common.conv_direct_dw_wgrad_mfma import DwWgradMfmaSpec
+
+        return DwWgradMfmaSpec(
+            problem=problem, name="rocke_direct_wgrad_dw_mfma", **knobs
+        )
+    if variant == "direct_wgrad_grouped_mfma":
+        from kernels.common.conv_direct_grouped_wgrad_mfma import (
+            GroupedWgradMfmaSpec,
+        )
+
+        return GroupedWgradMfmaSpec(
+            problem=problem, name="rocke_direct_wgrad_grouped_mfma", **knobs
+        )
     if variant == DIRECT_WGRAD_FOLD:
         from kernels.common.conv_direct_dw_wgrad import DwWgradFoldSpec
 
@@ -328,6 +344,18 @@ def validate_direct_spec(variant: str, spec, arch: str) -> tuple[bool, str]:
             )
 
             return is_valid_wgrad_gather_spec(spec, arch=arch)
+        if variant == "direct_wgrad_dw_mfma":
+            from kernels.common.conv_direct_dw_wgrad_mfma import (
+                is_valid_dw_wgrad_mfma_spec,
+            )
+
+            return is_valid_dw_wgrad_mfma_spec(spec, arch=arch)
+        if variant == "direct_wgrad_grouped_mfma":
+            from kernels.common.conv_direct_grouped_wgrad_mfma import (
+                is_valid_grouped_wgrad_mfma_spec,
+            )
+
+            return is_valid_grouped_wgrad_mfma_spec(spec, arch=arch)
         if variant in _HELPERS:
             return True, "ok"
         if variant == DIRECT_SPLIT_REDUCE:
@@ -348,7 +376,9 @@ def select_direct_kernel_builder(variant: str, spec, arch: str):
     an extra ``variant``.
     """
     from kernels.common import conv_direct_dw_wgrad as dw
+    from kernels.common import conv_direct_dw_wgrad_mfma as dwm
     from kernels.common import conv_direct_grouped as dc
+    from kernels.common import conv_direct_grouped_wgrad_mfma as gwm
     from kernels.common import conv_direct_nongrouped as dn
     from kernels.common import conv_direct_wgrad_gather as wg
     from kernels.common import conv_direct_wgrad_tiled as wt
@@ -358,6 +388,8 @@ def select_direct_kernel_builder(variant: str, spec, arch: str):
         "direct_wgrad_dw": dw.build_dw_wgrad_lds,
         "direct_wgrad_dw_cols": dw.build_dw_wgrad_cols,
         "direct_wgrad_gather": wg.build_wgrad_gather,
+        "direct_wgrad_dw_mfma": dwm.build_dw_wgrad_mfma,
+        "direct_wgrad_grouped_mfma": gwm.build_grouped_wgrad_mfma,
         DIRECT_WGRAD_FOLD: dw.build_dw_wgrad_fold,
         "direct_grouped": dc.build_direct_conv,
         "direct_nongrouped": dn.build_direct_conv_nongrouped,
@@ -439,9 +471,11 @@ def direct_wgrad_fold_knobs(variant: str, KH: int, KW: int, cpg: int) -> dict:
     """Knobs of the fold a partial-slot wgrad ``variant`` runs after its pass.
 
     ``taps`` is the dW extent per output channel: ``KH * KW`` for depthwise,
-    ``KH * KW * C`` for the dense gather kernel (cpg baked).
+    ``KH * KW * cpg`` for the dense gather and the grouped MFMA kernels (cpg
+    baked).
     """
     if variant not in DIRECT_WGRAD_FOLDED:
         raise ValueError(f"{variant!r} does not fold partial slots")
-    taps = KH * KW * (cpg if variant == "direct_wgrad_gather" else 1)
+    per_channel = variant in ("direct_wgrad_gather", "direct_wgrad_grouped_mfma")
+    taps = KH * KW * (cpg if per_channel else 1)
     return {"taps": taps}
diff --git a/dnn-providers/hip-kernel-provider/rocke/library/tests/parity/conv_direct_grouped_emit.c b/dnn-providers/hip-kernel-provider/rocke/library/tests/parity/conv_direct_grouped_emit.c
index a56a659fde82..2ce63d011d21 100644
--- a/dnn-providers/hip-kernel-provider/rocke/library/tests/parity/conv_direct_grouped_emit.c
+++ b/dnn-providers/hip-kernel-provider/rocke/library/tests/parity/conv_direct_grouped_emit.c
@@ -15,8 +15,10 @@
  * baked and runtime-channel; from index 95 the split-C non-grouped spec and its
  * rocke_direct_conv_split_reduce_spec_t second pass; from index 103 the grouped /
  * group-merge non-grouped spec; from index 114 the direct wgrad kernels:
- * rocke_direct_wgrad_tiled_spec_t, rocke_dw_wgrad_{fold,cols,lds}_spec_t and
- * rocke_direct_wgrad_gather_spec_t) identically to the Python emitter
+ * rocke_direct_wgrad_tiled_spec_t, rocke_dw_wgrad_{fold,cols,lds}_spec_t,
+ * rocke_direct_wgrad_gather_spec_t and, from index 148, the 4x4x4 MFMA
+ * rocke_dw_wgrad_mfma_spec_t and rocke_grouped_wgrad_mfma_spec_t) identically
+ * to the Python emitter
  * conv_direct_grouped_emit.py, builds the
  * kernel via the matching rocke_build_direct_conv_*_new function and lowers via
  * rocke_lower_kernel_to_llvm (per-config arch, flavor AUTO) and prints the .ll
@@ -27,7 +29,9 @@
 #include <string.h>
 
 #include "rocke/instance_conv_direct_dw_wgrad.h"
+#include "rocke/instance_conv_direct_dw_wgrad_mfma.h"
 #include "rocke/instance_conv_direct_grouped.h"
+#include "rocke/instance_conv_direct_grouped_wgrad_mfma.h"
 #include "rocke/instance_conv_direct_nongrouped.h"
 #include "rocke/instance_conv_direct_wgrad_gather.h"
 #include "rocke/instance_conv_direct_wgrad_tiled.h"
@@ -58,7 +62,9 @@ enum
     KIND_DW_WGRAD_FOLD = 17,
     KIND_DW_WGRAD_COLS = 18,
     KIND_DW_WGRAD_LDS = 19,
-    KIND_WGRAD_GATHER = 20
+    KIND_WGRAD_GATHER = 20,
+    KIND_DW_WGRAD_MFMA = 21,
+    KIND_GROUPED_WGRAD_MFMA = 22
 };
 
 #define NONGROUPED_CFG_BASE 42
@@ -1556,6 +1562,8 @@ typedef struct wgrad_direct_specs
     rocke_dw_wgrad_cols_spec_t cols;
     rocke_dw_wgrad_lds_spec_t lds;
     rocke_direct_wgrad_gather_spec_t gather;
+    rocke_dw_wgrad_mfma_spec_t dw_mfma;
+    rocke_grouped_wgrad_mfma_spec_t grouped_mfma;
 } wgrad_direct_specs_t;
 
 /* _wp(): N=2 H=16 W=32 groups=1 cpg=kpg=64 3x3 PAD=1 stride=1 bf16. */
@@ -1606,6 +1614,10 @@ static int make_wgrad_direct_cfg(int idx, int* kind, wgrad_direct_specs_t* w, co
     *cs = rocke_dw_wgrad_cols_spec_default();
     *ls = rocke_dw_wgrad_lds_spec_default();
     *gs = rocke_direct_wgrad_gather_spec_default();
+    rocke_dw_wgrad_mfma_spec_t* dm = &w->dw_mfma;
+    rocke_grouped_wgrad_mfma_spec_t* gm = &w->grouped_mfma;
+    *dm = rocke_dw_wgrad_mfma_spec_default();
+    *gm = rocke_grouped_wgrad_mfma_spec_default();
     *arch = "gfx950";
     if(idx <= 6 || idx == 20)
         *kind = KIND_WGRAD_TILED;
@@ -1615,10 +1627,15 @@ static int make_wgrad_direct_cfg(int idx, int* kind, wgrad_direct_specs_t* w, co
         *kind = KIND_DW_WGRAD_COLS;
     else if(idx <= 15)
         *kind = KIND_DW_WGRAD_LDS;
-    else
+    else if(idx <= 33)
         *kind = KIND_WGRAD_GATHER;
+    else if(idx <= 40)
+        *kind = KIND_DW_WGRAD_MFMA;
+    else
+        *kind = KIND_GROUPED_WGRAD_MFMA;
     ts->problem = mk_wp();
     gs->problem = mk_wp();
+    gm->problem = mk_wp();
     switch(idx)
     {
     case 0:
@@ -1851,6 +1868,87 @@ static int make_wgrad_direct_cfg(int idx, int* kind, wgrad_direct_specs_t* w, co
         gs->waves_k = 3;
         gs->prefetch = 1;
         return 0;
+    case 34: /* depthwise 4x4x4 MFMA: 7x7, preshifted window, both prologue bodies */
+        dm->problem = mk_dwp(64);
+        mk_filter(&dm->problem, 7, 7, 3, 1);
+        dm->problem.W = 20;
+        dm->q_tiles = 4;
+        return 0;
+    case 35: /* 3x3 fp16 (no preshift), a partial channel tile, prefetch 1 */
+        dm->problem = mk_dwp(72);
+        dm->problem.H = 9;
+        dm->problem.W = 30;
+        dm->problem.dtype = "fp16";
+        dm->rows_per_chunk = 4;
+        dm->prefetch = 1;
+        return 0;
+    case 36: /* three 7x7 images packed in one 32-wide tile */
+        dm->problem = mk_dwp(128);
+        mk_filter(&dm->problem, 7, 7, 3, 1);
+        dm->problem.N = 5;
+        dm->problem.H = 7;
+        dm->problem.W = 7;
+        dm->rows_per_chunk = 14;
+        dm->unfold_n = 3;
+        return 0;
+    case 37: /* 5x5 "valid" padding, two packed images, three staging buffers */
+        dm->problem = mk_dwp(64);
+        mk_filter(&dm->problem, 5, 5, 0, 1);
+        dm->problem.H = 10;
+        dm->problem.W = 6;
+        dm->q_tiles = 5;
+        dm->unfold_n = 2;
+        dm->prefetch = 3;
+        return 0;
+    case 38: /* 11x11: 3x3 tap groups */
+        dm->problem = mk_dwp(64);
+        mk_filter(&dm->problem, 11, 11, 5, 1);
+        dm->problem.H = 12;
+        dm->problem.W = 12;
+        dm->q_tiles = 2;
+        dm->rows_per_chunk = 8;
+        dm->target_blocks = 64;
+        return 0;
+    case 39: /* PAD = KH - 1: one prologue body only */
+        dm->problem = mk_dwp(64);
+        dm->problem.PAD = 2;
+        dm->q_tiles = 4;
+        return 0;
+    case 40: /* must reject: stride 2 */
+        dm->problem = mk_dwp(64);
+        dm->problem.stride = 2;
+        return 0;
+    case 41: /* grouped 4x4x4 MFMA: 4 channels per group */
+        gm->problem.groups = 16;
+        gm->problem.cpg = 4;
+        gm->problem.kpg = 4;
+        return 0;
+    case 42: /* 8 channels per group, fp16, a partial channel tile, 2x2 wave split */
+        gm->problem.groups = 9;
+        gm->problem.cpg = 8;
+        gm->problem.kpg = 8;
+        gm->problem.dtype = "fp16";
+        gm->problem.W = 24;
+        gm->waves_q = 2;
+        gm->wave_q4 = 2;
+        gm->rows_per_chunk = 8;
+        gm->prefetch = 3;
+        return 0;
+    case 43: /* 5x5, one wave over four column groups, four staging buffers */
+        gm->problem.groups = 16;
+        gm->problem.cpg = 4;
+        gm->problem.kpg = 4;
+        mk_filter(&gm->problem, 5, 5, 2, 1);
+        gm->waves_q = 1;
+        gm->wave_q4 = 4;
+        gm->prefetch = 4;
+        gm->target_blocks = 64;
+        return 0;
+    case 44: /* must reject: 16 channels per group */
+        gm->problem.groups = 4;
+        gm->problem.cpg = 16;
+        gm->problem.kpg = 16;
+        return 0;
     default:
         return -1;
     }
@@ -1980,6 +2078,10 @@ int main(int argc, char** argv)
         kernel = rocke_build_direct_wgrad_gather_new(&b, &swd.gather, arch);
     else if(kind == KIND_SPLIT_REDUCE)
         kernel = rocke_build_direct_conv_split_reduce_new(&b, &sred, arch);
+    else if(kind == KIND_DW_WGRAD_MFMA)
+        kernel = rocke_build_dw_wgrad_mfma_new(&b, &swd.dw_mfma, arch);
+    else if(kind == KIND_GROUPED_WGRAD_MFMA)
+        kernel = rocke_build_grouped_wgrad_mfma_new(&b, &swd.grouped_mfma, arch);
     else
         kernel = rocke_build_direct_depthwise_new(&b, &sdw, arch);
     if(kernel == NULL)
diff --git a/dnn-providers/hip-kernel-provider/rocke/library/tests/parity/conv_direct_grouped_emit.py b/dnn-providers/hip-kernel-provider/rocke/library/tests/parity/conv_direct_grouped_emit.py
index e5188d388efa..67696a345b70 100644
--- a/dnn-providers/hip-kernel-provider/rocke/library/tests/parity/conv_direct_grouped_emit.py
+++ b/dnn-providers/hip-kernel-provider/rocke/library/tests/parity/conv_direct_grouped_emit.py
@@ -17,7 +17,8 @@
 # DirectNongroupedConvSpec and its DirectSplitReduceSpec second pass, from
 # index 103 its grouped / group-merge mode, and from index 114 the direct
 # wgrad kernels: DirectWgradTiledSpec, DwWgradFoldSpec, DwWgradColsSpec,
-# DwWgradLdsSpec and WgradGatherSpec), builds the kernel
+# DwWgradLdsSpec, WgradGatherSpec, DwWgradMfmaSpec and GroupedWgradMfmaSpec),
+# builds the kernel
 # via the matching build_direct_conv_* function (arch=<cfg arch>) and prints
 # _native_lower(arch=<cfg arch>) to stdout so it can be byte-compared with
 # the C emitter conv_direct_grouped_emit.c.
@@ -68,6 +69,14 @@
     build_dw_wgrad_fold,
     build_dw_wgrad_lds,
 )
+from kernels.common.conv_direct_dw_wgrad_mfma import (
+    DwWgradMfmaSpec,
+    build_dw_wgrad_mfma,
+)
+from kernels.common.conv_direct_grouped_wgrad_mfma import (
+    GroupedWgradMfmaSpec,
+    build_grouped_wgrad_mfma,
+)
 from kernels.common.conv_direct_wgrad_gather import (
     WgradGatherSpec,
     build_wgrad_gather,
@@ -121,7 +130,7 @@
 
 
 # ---------------------------------------------------------------------------
-# The direct wgrad kernels, indices 114..147: DirectWgradTiledSpec (wave tiles,
+# The direct wgrad kernels, indices 114..158: DirectWgradTiledSpec (wave tiles,
 # unfold_n, row_unroll, q_groups with and without pingpong, 1x1 and 5x5), the
 # DwWgradFoldSpec partial-slot fold, DwWgradColsSpec (small C), DwWgradLdsSpec
 # (row rings, strided, cols_per_lane, a partial channel tile), WgradGatherSpec
@@ -129,8 +138,12 @@
 # both engines must reject; a small-C depthwise spec with more outputs than
 # threads; then eight more rejections (LDS footprints over budget, targets
 # without 16-byte LDS DMA, a negative row_unroll, no blocks, no dtype); then
-# WgradGatherSpec at valid padding over three k-waves. Follows the 11 grouped
-# non-grouped configs.
+# WgradGatherSpec at valid padding over three k-waves; from 148 DwWgradMfmaSpec
+# (preshifted and plain windows, packed images, 3x3 tap groups, the
+# single-prologue PAD = KH - 1 case, a partial channel tile) and
+# GroupedWgradMfmaSpec (4 and 8 channels per group, wave splits, a partial
+# channel tile), each with a reject. Follows the 11 grouped non-grouped
+# configs.
 # ---------------------------------------------------------------------------
 _WGRAD_DIRECT_CFG_BASE = 114
 
@@ -159,11 +172,13 @@ def _dwp(C, **over):
 
 def _wgrad_direct_spec(idx: int):
     """Return (kind, spec, arch) for config ``_WGRAD_DIRECT_CFG_BASE + idx``."""
-    ts, ls, cs, gs = (
+    ts, ls, cs, gs, dm, gm = (
         DirectWgradTiledSpec,
         DwWgradLdsSpec,
         DwWgradColsSpec,
         WgradGatherSpec,
+        DwWgradMfmaSpec,
+        GroupedWgradMfmaSpec,
     )
     table = {
         0: ("wgrad_tiled", ts(problem=_wp(cpg=128, kpg=128, H=30))),
@@ -349,6 +364,79 @@ def _wgrad_direct_spec(idx: int):
                 prefetch=1,
             ),
         ),
+        # Depthwise 4x4x4 MFMA: 7x7 (preshifted window, both prologue bodies).
+        34: (
+            "dw_wgrad_mfma",
+            dm(problem=_dwp(64, KH=7, KW=7, PAD=3, W=20), q_tiles=4),
+        ),
+        # 3x3 fp16 (no preshift), a partial channel tile, prefetch 1.
+        35: (
+            "dw_wgrad_mfma",
+            dm(
+                problem=_dwp(72, H=9, W=30, dtype="fp16"),
+                rows_per_chunk=4,
+                prefetch=1,
+            ),
+        ),
+        # Three 7x7 images packed in one 32-wide tile.
+        36: (
+            "dw_wgrad_mfma",
+            dm(
+                problem=_dwp(128, KH=7, KW=7, PAD=3, N=5, H=7, W=7),
+                rows_per_chunk=14,
+                unfold_n=3,
+            ),
+        ),
+        # 5x5 "valid" padding, two packed images, three staging buffers.
+        37: (
+            "dw_wgrad_mfma",
+            dm(
+                problem=_dwp(64, KH=5, KW=5, PAD=0, H=10, W=6),
+                q_tiles=5,
+                unfold_n=2,
+                prefetch=3,
+            ),
+        ),
+        # 11x11: 3x3 tap groups.
+        38: (
+            "dw_wgrad_mfma",
+            dm(
+                problem=_dwp(64, KH=11, KW=11, PAD=5, H=12, W=12),
+                q_tiles=2,
+                rows_per_chunk=8,
+                target_blocks=64,
+            ),
+        ),
+        # PAD = KH - 1: one prologue body only.
+        39: ("dw_wgrad_mfma", dm(problem=_dwp(64, PAD=2), q_tiles=4)),
+        # Must reject: stride 2.
+        40: ("dw_wgrad_mfma", dm(problem=_dwp(64, stride=2))),
+        # Grouped 4x4x4 MFMA: 4 channels per group.
+        41: ("grouped_wgrad_mfma", gm(problem=_wp(groups=16, cpg=4, kpg=4))),
+        # 8 channels per group, fp16, a partial channel tile, 2x2 wave split.
+        42: (
+            "grouped_wgrad_mfma",
+            gm(
+                problem=_wp(groups=9, cpg=8, kpg=8, dtype="fp16", W=24),
+                waves_q=2,
+                wave_q4=2,
+                rows_per_chunk=8,
+                prefetch=3,
+            ),
+        ),
+        # 5x5, one wave over four column groups, four staging buffers.
+        43: (
+            "grouped_wgrad_mfma",
+            gm(
+                problem=_wp(groups=16, cpg=4, kpg=4, KH=5, KW=5, PAD=2),
+                waves_q=1,
+                wave_q4=4,
+                prefetch=4,
+                target_blocks=64,
+            ),
+        ),
+        # Must reject: 16 channels per group.
+        44: ("grouped_wgrad_mfma", gm(problem=_wp(groups=4, cpg=16, kpg=16))),
     }
     sel = table.get(idx)
     if sel is None:
@@ -1811,6 +1899,10 @@ def main() -> int:
         kernel = build_wgrad_gather(spec, arch=arch)
     elif kind == "split_reduce":
         kernel = build_direct_conv_split_reduce(spec, arch=arch)
+    elif kind == "dw_wgrad_mfma":
+        kernel = build_dw_wgrad_mfma(spec, arch=arch)
+    elif kind == "grouped_wgrad_mfma":
+        kernel = build_grouped_wgrad_mfma(spec, arch=arch)
     else:
         kernel = build_direct_depthwise(spec, arch=arch)
     if mode == "ll":
diff --git a/dnn-providers/hip-kernel-provider/rocke/library/tests/test_builder_signature_contract.py b/dnn-providers/hip-kernel-provider/rocke/library/tests/test_builder_signature_contract.py
index 3876542c0651..5835baae5aad 100644
--- a/dnn-providers/hip-kernel-provider/rocke/library/tests/test_builder_signature_contract.py
+++ b/dnn-providers/hip-kernel-provider/rocke/library/tests/test_builder_signature_contract.py
@@ -342,11 +342,13 @@
     "kernels.common.conv_direct_nongrouped.DirectNongroupedConvSpec": ("problem",),
     "kernels.common.conv_direct_wgrad_tiled.DirectWgradTiledSpec": ("problem",),
     "kernels.common.conv_direct_dw_wgrad.DwWgradLdsSpec": ("problem",),
+    "kernels.common.conv_direct_dw_wgrad_mfma.DwWgradMfmaSpec": ("problem",),
     "kernels.common.conv_direct_dw_wgrad.DwWgradColsSpec": ("problem",),
     # The dW extent per output channel the fold sums partial slots over.
     "kernels.common.conv_direct_dw_wgrad.DwWgradFoldSpec": ("taps",),
     "kernels.common.conv_direct_wgrad_gather.WgradGatherSpec": ("problem",),
     "kernels.common.conv_direct_nongrouped.DirectSplitReduceSpec": (),
+    "kernels.common.conv_direct_grouped_wgrad_mfma.GroupedWgradMfmaSpec": ("problem",),
     "kernels.common.deep_fused_conv_pool.FusedConvPoolProblem": ("conv",),
     "kernels.common.deep_fused_conv_pool.DeepFusedConvPoolSpec": ("problem",),
     "kernels.common.img2col.Img2ColSpec": ("problem",),
diff --git a/dnn-providers/hip-kernel-provider/rocke/library/tests/test_direct_conv_correctness.py b/dnn-providers/hip-kernel-provider/rocke/library/tests/test_direct_conv_correctness.py
index 6cbc09936dcc..ff1f8096081b 100644
--- a/dnn-providers/hip-kernel-provider/rocke/library/tests/test_direct_conv_correctness.py
+++ b/dnn-providers/hip-kernel-provider/rocke/library/tests/test_direct_conv_correctness.py
@@ -3219,6 +3219,48 @@ def test_mfma_dgrad_runtime_channels(self):
         dict(N=2, H=16, W=32, groups=8, cpg=1, kpg=1, KH=9, KW=9, PAD=4),
         dict(channels=8, strip=8, threads=64, target_blocks=9),
     ),
+    # Depthwise 4x4x4 MFMA, 7x7 in row chunks, a partial channel tile.
+    (
+        "dw_mfma_7x7_chunks",
+        "direct_wgrad_dw_mfma",
+        dict(N=2, H=17, W=23, groups=72, cpg=1, kpg=1, KH=7, KW=7, PAD=3),
+        dict(q_tiles=4, rows_per_chunk=7, prefetch=2, target_blocks=1024, unfold_n=1),
+    ),
+    # Whole images per unit (the PAD-row prologue), one staging buffer.
+    (
+        "dw_mfma_7x7_whole",
+        "direct_wgrad_dw_mfma",
+        dict(N=3, H=9, W=30, groups=64, cpg=1, kpg=1, KH=7, KW=7, PAD=3),
+        dict(q_tiles=8, rows_per_chunk=14, prefetch=1, target_blocks=1024, unfold_n=1),
+    ),
+    # Three 7x7 images packed in one tile, a batch that leaves one short.
+    (
+        "dw_mfma_unfold3",
+        "direct_wgrad_dw_mfma",
+        dict(N=5, H=7, W=7, groups=128, cpg=1, kpg=1, KH=7, KW=7, PAD=3),
+        dict(q_tiles=8, rows_per_chunk=14, prefetch=2, target_blocks=1024, unfold_n=3),
+    ),
+    # 3x3 valid padding: no preshift, both prologue bodies differ.
+    (
+        "dw_mfma_3x3_pad0",
+        "direct_wgrad_dw_mfma",
+        dict(N=2, H=12, W=18, groups=64, cpg=1, kpg=1, PAD=0),
+        dict(q_tiles=4, rows_per_chunk=7, prefetch=2, target_blocks=2048, unfold_n=1),
+    ),
+    # Grouped 4x4x4 MFMA, 4 channels per group.
+    (
+        "grouped_mfma_c4",
+        "direct_wgrad_grouped_mfma",
+        dict(N=2, H=13, W=21, groups=32, cpg=4, kpg=4),
+        dict(waves_q=4, wave_q4=2, rows_per_chunk=28, prefetch=2, target_blocks=1024),
+    ),
+    # 8 channels per group, a partial channel tile.
+    (
+        "grouped_mfma_c8",
+        "direct_wgrad_grouped_mfma",
+        dict(N=2, H=13, W=21, groups=9, cpg=8, kpg=8),
+        dict(waves_q=2, wave_q4=1, rows_per_chunk=28, prefetch=3, target_blocks=1024),
+    ),
 ]
 
 
@@ -3252,6 +3294,8 @@ def _run_aot_wgrad_one(arch: str, variant: str, problem, knobs: dict) -> float:
         caps = DirectCaps(KH=p.KH, PAD=p.PAD, stride=p.stride, cpg=0, kpg=0)
     elif variant == "direct_wgrad_gather":
         caps = DirectCaps(KH=p.KH, PAD=p.PAD, stride=p.stride, cpg=p.cpg, kpg=0)
+    elif variant == "direct_wgrad_grouped_mfma":
+        caps = DirectCaps(KH=p.KH, PAD=p.PAD, stride=p.stride, cpg=p.cpg, kpg=p.kpg)
     else:
         caps = DirectCaps(KH=p.KH, PAD=p.PAD, stride=p.stride, cpg=1, kpg=1)
     probe = probe_problem(caps, p.dtype, variant, knobs)
@@ -3268,7 +3312,7 @@ def _run_aot_wgrad_one(arch: str, variant: str, problem, knobs: dict) -> float:
     torch.manual_seed(7)
     td = torch.bfloat16 if p.dtype == "bf16" else torch.float16
     total_c, total_k = p.groups * p.cpg, p.groups * p.kpg
-    if folded and variant != "direct_wgrad_gather":
+    if folded and p.cpg == 1 and p.kpg == 1:
         total_k = total_c = p.groups  # depthwise
     X = torch.empty(p.N, p.H, p.W, total_c, dtype=td).uniform_(-0.5, 0.5)
     dY = torch.empty(p.N, p.Ho, p.Wo, total_k, dtype=td).uniform_(-0.5, 0.5)
@@ -3326,7 +3370,7 @@ def _run_aot_wgrad_one(arch: str, variant: str, problem, knobs: dict) -> float:
                 arch=arch,
             )
             channels = (
-                spec.fold_channels() if variant == "direct_wgrad_gather" else p.groups
+                spec.fold_channels() if hasattr(spec, "fold_channels") else p.groups
             )
             fgrid, fblock = fold_spec.launch_geometry(channels)
             fold = KernelLauncher(
diff --git a/dnn-providers/hip-kernel-provider/rocke/library/tests/test_direct_kernel_cache.py b/dnn-providers/hip-kernel-provider/rocke/library/tests/test_direct_kernel_cache.py
index 0bbe1b4825f4..fe68dc01ed30 100644
--- a/dnn-providers/hip-kernel-provider/rocke/library/tests/test_direct_kernel_cache.py
+++ b/dnn-providers/hip-kernel-provider/rocke/library/tests/test_direct_kernel_cache.py
@@ -856,6 +856,17 @@ def test_dgrad_rtc_rejects_what_no_binary_fits(cache, groups, cpg, kpg, why):
             "direct_wgrad_gather",
             _problem(groups=1, cpg=24, kpg=96, stride=2, H=30, W=40),
         ),
+        (
+            "direct_wgrad_dw_mfma",
+            _problem(groups=192, cpg=1, kpg=1, KH=7, KW=7, PAD=3, H=56, W=56),
+        ),
+        # Small images: the packed-image binaries join in.
+        (
+            "direct_wgrad_dw_mfma",
+            _problem(N=128, groups=1024, cpg=1, kpg=1, KH=7, KW=7, PAD=3, H=7, W=7),
+        ),
+        ("direct_wgrad_grouped_mfma", _problem(groups=32, cpg=4, kpg=4, H=56, W=56)),
+        ("direct_wgrad_grouped_mfma", _problem(groups=32, cpg=8, kpg=8, H=56, W=56)),
     ],
 )
 def test_folded_wgrad_plans_run_the_pass_then_the_fold(cache, variant, problem):
@@ -875,7 +886,7 @@ def test_folded_wgrad_plans_run_the_pass_then_the_fold(cache, variant, problem):
         assert fold.buffers == ("ws", None, "dw")
         spec = dks.make_spec(variant, problem, json.loads(plan.identity.knobs))
         assert (main.grid, main.block) == spec.launch_geometry()
-        if variant == "direct_wgrad_gather":
+        if hasattr(spec, "fold_channels"):
             channels, taps = spec.fold_channels(), spec.fold_taps()
         else:
             channels, taps = problem.groups, problem.KH * problem.KW
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/arch/data.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/arch/data.cpp
index 3bde4b2b6104..0650f942170a 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/arch/data.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/arch/data.cpp
@@ -965,6 +965,7 @@ static const rocke_ati_mma_frag_row_t rocke_ati_mma_frag[] = {
     {"mfma_f32_32x32x8_f16", 16},
     {"mfma_f32_32x32x16_f16", 16},
     {"mfma_f32_4x4x4_f16", 4},
+    {"mfma_f32_4x4x4_bf16", 4},
     /* --- MFMA bf16 (wave64) --- */
     {"mfma_f32_16x16x16_bf16", 4},
     {"mfma_f32_16x16x32_bf16", 4},
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/ir_tile.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/ir_tile.cpp
index c32c1839f5f8..bebc6594cbca 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/ir_tile.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/ir_tile.cpp
@@ -696,6 +696,7 @@ ROCKE_MMA_WRAP(rocke_b_mfma_f32_32x32x16_bf16, "mfma_f32_32x32x16_bf16")
 ROCKE_MMA_WRAP(rocke_b_mfma_f32_32x32x16_fp8, "mfma_f32_32x32x16_fp8")
 ROCKE_MMA_WRAP(rocke_b_mfma_f32_32x32x16_bf8, "mfma_f32_32x32x16_bf8")
 ROCKE_MMA_WRAP(rocke_b_mfma_f32_4x4x4_f16, "mfma_f32_4x4x4_f16")
+ROCKE_MMA_WRAP(rocke_b_mfma_f32_4x4x4_bf16, "mfma_f32_4x4x4_bf16")
 ROCKE_MMA_WRAP(rocke_b_mfma_f32_16x16x128_fp4, "mfma_f32_16x16x128_fp4")
 ROCKE_MMA_WRAP(rocke_b_mfma_f32_16x16x96_fp6, "mfma_f32_16x16x96_fp6")
 ROCKE_MMA_WRAP(rocke_b_wmma_f32_16x16x16_f16, "wmma_f32_16x16x16_f16")
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/data.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/data.cpp
index b289ffdeac88..aee3cc1de1ce 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/data.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/data.cpp
@@ -262,6 +262,9 @@ const rocke_ll_decl_t ROCKE_LL_INTRINSIC_DECLS[] = {
     {"mfma.f32.4x4x4f16",
      "declare <4 x float> @llvm.amdgcn.mfma.f32.4x4x4f16(<4 x half>, <4 x half>, <4 x float>, i32 "
      "immarg, i32 immarg, i32 immarg)"},
+    {"mfma.f32.4x4x4bf16.1k",
+     "declare <4 x float> @llvm.amdgcn.mfma.f32.4x4x4bf16.1k(<4 x i16>, <4 x i16>, <4 x float>, "
+     "i32 immarg, i32 immarg, i32 immarg)"},
     {"mfma.f32.16x16x32.fp8.fp8",
      "declare <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.fp8.fp8(<2 x i32>, <2 x i32>, <4 x "
      "float>, i32 immarg, i32 immarg, i32 immarg)"},
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/mma.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/mma.cpp
index 9fad95493119..10197f585b43 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/mma.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/mma.cpp
@@ -1149,6 +1149,12 @@ static const _mfma_spec_t MFMA_SPECS[] = {
      "<4 x half>",
      "<4 x float>",
      NULL},
+    {"mfma_f32_4x4x4_bf16",
+     "mfma.f32.4x4x4bf16.1k",
+     "llvm.amdgcn.mfma.f32.4x4x4bf16.1k",
+     "<4 x bfloat>",
+     "<4 x float>",
+     "<4 x i16>"},
 };
 static const int MFMA_SPECS_N = (int)(sizeof(MFMA_SPECS) / sizeof(MFMA_SPECS[0]));
 
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_conv_direct_dw_wgrad_mfma.h b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_conv_direct_dw_wgrad_mfma.h
new file mode 100644
index 000000000000..30cec3fc1d06
--- /dev/null
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_conv_direct_dw_wgrad_mfma.h
@@ -0,0 +1,101 @@
+/* Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
+ * SPDX-License-Identifier: MIT
+ *
+ * rocke/instance_conv_direct_dw_wgrad_mfma.h -- C++ engine port of the
+ * depthwise wgrad on the batched 4x4x4 MFMA in
+ * library/kernels/common/conv_direct_dw_wgrad_mfma.py.
+ *
+ * mfma_f32_4x4x4 runs 16 independent 4x4x4 products per wave, one channel
+ * each: per input row and 4-column group, A[s][k] = X[y][q0 + k + s - PAD]
+ * (a Hankel window built from an aligned LDS window with v_perm_b32) times
+ * B[k][r] = dY[y + PAD - r][q0 + k] (a register ring of dY rows rotated by
+ * quad_perm). Rows reach LDS by DMA in NHWC order and are transposed
+ * channel-major with ds_read_tr16_b64. The block stores one partial slot
+ * ws[part][tap][c], folded into dW by the DwWgradFoldSpec kernel
+ * (rocke/instance_conv_direct_dw_wgrad.h) with taps = KH * KW, channels = C.
+ *
+ *   Python (conv_direct_dw_wgrad_mfma.py)      C (this header)
+ *   -------------------------------------   -----------------------------------
+ *   @dataclass(frozen) DwWgradMfmaSpec       rocke_dw_wgrad_mfma_spec_t
+ *     host-side geometry                     rocke_dw_wgrad_mfma_{lds_bytes,
+ *                                              image_pitch,n_partitions,
+ *                                              workspace_bytes,launch_geometry}
+ *     .kernel_name()                         rocke_dw_wgrad_mfma_kernel_name
+ *   is_valid_dw_wgrad_mfma_spec(spec, arch)  rocke_dw_wgrad_mfma_is_valid_spec
+ *   build_dw_wgrad_mfma(spec, arch)          rocke_build_dw_wgrad_mfma(_new)
+ */
+#ifndef ROCKE_INSTANCE_CONV_DIRECT_DW_WGRAD_MFMA_H
+#define ROCKE_INSTANCE_CONV_DIRECT_DW_WGRAD_MFMA_H
+
+#include <stdbool.h>
+#include <stddef.h>
+
+#include "rocke/instance_conv_direct_grouped.h" /* rocke_direct_conv_problem_t */
+#include "rocke/ir.h"
+#include "rocke/lower_llvm.h"
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+/* ===================================================================== *
+ *  DwWgradMfmaSpec
+ *
+ *  @dataclass(frozen=True)
+ *  class DwWgradMfmaSpec:
+ *      problem: DirectConvProblem          # depthwise, stride 1
+ *      name: str = "dw_wgrad_mfma"
+ *      q_tiles: int = 8
+ *      rows_per_chunk: int = 16
+ *      prefetch: int = 2
+ *      target_blocks: int = 2048
+ *      unfold_n: int = 1
+ * ===================================================================== */
+typedef struct rocke_dw_wgrad_mfma_spec
+{
+    rocke_direct_conv_problem_t problem;
+    const char* name; /* default "dw_wgrad_mfma" */
+    int q_tiles; /* default 8 */
+    int rows_per_chunk; /* default 16 */
+    int prefetch; /* default 2 */
+    int target_blocks; /* default 2048 */
+    int unfold_n; /* default 1 */
+} rocke_dw_wgrad_mfma_spec_t;
+
+rocke_dw_wgrad_mfma_spec_t rocke_dw_wgrad_mfma_spec_default(void);
+long rocke_dw_wgrad_mfma_lds_bytes(const rocke_dw_wgrad_mfma_spec_t* spec);
+/* image_pitch() = max(Wo, Wo + KW - 1 - PAD, PAD + W). */
+int rocke_dw_wgrad_mfma_image_pitch(const rocke_dw_wgrad_mfma_spec_t* spec);
+int rocke_dw_wgrad_mfma_n_partitions(const rocke_dw_wgrad_mfma_spec_t* spec);
+long rocke_dw_wgrad_mfma_workspace_bytes(const rocke_dw_wgrad_mfma_spec_t* spec);
+void rocke_dw_wgrad_mfma_launch_geometry(const rocke_dw_wgrad_mfma_spec_t* spec,
+                                         int grid[3],
+                                         int block[3]);
+/* kernel_name_join(name, f{KH}x{KW}p{PAD}s{stride}, qt{n}, rpc{n}, pf{n},
+ * un{n}, tb{n}, flags [bf16]). */
+rocke_status_t rocke_dw_wgrad_mfma_kernel_name(const rocke_dw_wgrad_mfma_spec_t* spec,
+                                               char* out,
+                                               size_t out_cap);
+/* is_valid_dw_wgrad_mfma_spec(spec, arch) -> (ok, reason). arch NULL => "gfx950". */
+bool rocke_dw_wgrad_mfma_is_valid_spec(const rocke_dw_wgrad_mfma_spec_t* spec,
+                                       const char* arch,
+                                       char* reason,
+                                       size_t reason_cap);
+rocke_kernel_def_t* rocke_build_dw_wgrad_mfma(rocke_ir_builder_t* b,
+                                              const rocke_dw_wgrad_mfma_spec_t* spec,
+                                              const char* arch);
+rocke_kernel_def_t* rocke_build_dw_wgrad_mfma_new(rocke_ir_builder_t* b,
+                                                  const rocke_dw_wgrad_mfma_spec_t* spec,
+                                                  const char* arch);
+rocke_status_t rocke_dw_wgrad_mfma_lower_to_llvm(const rocke_dw_wgrad_mfma_spec_t* spec,
+                                                 const char* arch,
+                                                 rocke_llvm_flavor_t flavor,
+                                                 char** out_ll,
+                                                 char* err,
+                                                 size_t err_cap);
+
+#ifdef __cplusplus
+} /* extern "C" */
+#endif
+
+#endif /* ROCKE_INSTANCE_CONV_DIRECT_DW_WGRAD_MFMA_H */
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_conv_direct_grouped_wgrad_mfma.h b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_conv_direct_grouped_wgrad_mfma.h
new file mode 100644
index 000000000000..4c89c47342f4
--- /dev/null
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_conv_direct_grouped_wgrad_mfma.h
@@ -0,0 +1,106 @@
+/* Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
+ * SPDX-License-Identifier: MIT
+ *
+ * rocke/instance_conv_direct_grouped_wgrad_mfma.h -- C++ engine port of the
+ * grouped wgrad for 4 and 8 channels per group on the batched 4x4x4 MFMA in
+ * library/kernels/common/conv_direct_grouped_wgrad_mfma.py.
+ *
+ * A wave runs 16 groups of 4 channels as the 16 batches of mfma_f32_4x4x4:
+ * per input row, tap (r, s) and 4-column group, A = X (lane = channel, four
+ * columns shifted by s) and B = dY (the row of tap r, from a register ring),
+ * both read straight from the NHWC rows staged in LDS by ds_read_tr16_b64. With
+ * 8 channels per group a second dY read with the group's two halves swapped
+ * gives the off-diagonal 4x4 blocks. The waves reduce through LDS and the block
+ * stores one partial slot ws[part][tap][o] (tap = (r * KW + s) * cpg + c),
+ * folded into dW by the DwWgradFoldSpec kernel
+ * (rocke/instance_conv_direct_dw_wgrad.h) with taps = KH * KW * cpg and
+ * channels = K.
+ *
+ *   Python (conv_direct_grouped_wgrad_mfma.py)       C (this header)
+ *   ---------------------------------------------   ---------------------------
+ *   @dataclass(frozen) GroupedWgradMfmaSpec          rocke_grouped_wgrad_mfma_spec_t
+ *     host-side geometry                             rocke_grouped_wgrad_mfma_{
+ *                                                      threads,lds_bytes,
+ *                                                      n_partitions,fold_taps,
+ *                                                      workspace_bytes,
+ *                                                      launch_geometry}
+ *     .kernel_name()                                 rocke_grouped_wgrad_mfma_kernel_name
+ *   is_valid_grouped_wgrad_mfma_spec(spec, arch)     rocke_grouped_wgrad_mfma_is_valid_spec
+ *   build_grouped_wgrad_mfma(spec, arch)             rocke_build_grouped_wgrad_mfma(_new)
+ */
+#ifndef ROCKE_INSTANCE_CONV_DIRECT_GROUPED_WGRAD_MFMA_H
+#define ROCKE_INSTANCE_CONV_DIRECT_GROUPED_WGRAD_MFMA_H
+
+#include <stdbool.h>
+#include <stddef.h>
+
+#include "rocke/instance_conv_direct_grouped.h" /* rocke_direct_conv_problem_t */
+#include "rocke/ir.h"
+#include "rocke/lower_llvm.h"
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+/* ===================================================================== *
+ *  GroupedWgradMfmaSpec
+ *
+ *  @dataclass(frozen=True)
+ *  class GroupedWgradMfmaSpec:
+ *      problem: DirectConvProblem          # cpg = kpg in (4, 8), stride 1
+ *      name: str = "grouped_wgrad_mfma"
+ *      waves_q: int = 4
+ *      wave_q4: int = 1
+ *      rows_per_chunk: int = 56
+ *      prefetch: int = 2
+ *      target_blocks: int = 1024
+ * ===================================================================== */
+typedef struct rocke_grouped_wgrad_mfma_spec
+{
+    rocke_direct_conv_problem_t problem;
+    const char* name; /* default "grouped_wgrad_mfma" */
+    int waves_q; /* default 4 */
+    int wave_q4; /* default 1 */
+    int rows_per_chunk; /* default 56 */
+    int prefetch; /* default 2 */
+    int target_blocks; /* default 1024 */
+} rocke_grouped_wgrad_mfma_spec_t;
+
+rocke_grouped_wgrad_mfma_spec_t rocke_grouped_wgrad_mfma_spec_default(void);
+int rocke_grouped_wgrad_mfma_threads(const rocke_grouped_wgrad_mfma_spec_t* spec);
+long rocke_grouped_wgrad_mfma_lds_bytes(const rocke_grouped_wgrad_mfma_spec_t* spec);
+int rocke_grouped_wgrad_mfma_n_partitions(const rocke_grouped_wgrad_mfma_spec_t* spec);
+/* taps() = KH * KW * cpg; fold_channels() = K. */
+int rocke_grouped_wgrad_mfma_fold_taps(const rocke_grouped_wgrad_mfma_spec_t* spec);
+long rocke_grouped_wgrad_mfma_workspace_bytes(const rocke_grouped_wgrad_mfma_spec_t* spec);
+void rocke_grouped_wgrad_mfma_launch_geometry(const rocke_grouped_wgrad_mfma_spec_t* spec,
+                                              int grid[3],
+                                              int block[3]);
+/* kernel_name_join(name, f{KH}x{KW}p{PAD}s{stride}c{cpg}, wq{n}, q4{n}, rpc{n},
+ * pf{n}, tb{n}, flags [bf16]). */
+rocke_status_t rocke_grouped_wgrad_mfma_kernel_name(const rocke_grouped_wgrad_mfma_spec_t* spec,
+                                                    char* out,
+                                                    size_t out_cap);
+/* is_valid_grouped_wgrad_mfma_spec(spec, arch) -> (ok, reason). arch NULL => "gfx950". */
+bool rocke_grouped_wgrad_mfma_is_valid_spec(const rocke_grouped_wgrad_mfma_spec_t* spec,
+                                            const char* arch,
+                                            char* reason,
+                                            size_t reason_cap);
+rocke_kernel_def_t* rocke_build_grouped_wgrad_mfma(rocke_ir_builder_t* b,
+                                                   const rocke_grouped_wgrad_mfma_spec_t* spec,
+                                                   const char* arch);
+rocke_kernel_def_t* rocke_build_grouped_wgrad_mfma_new(rocke_ir_builder_t* b,
+                                                       const rocke_grouped_wgrad_mfma_spec_t* spec,
+                                                       const char* arch);
+rocke_status_t rocke_grouped_wgrad_mfma_lower_to_llvm(const rocke_grouped_wgrad_mfma_spec_t* spec,
+                                                      const char* arch,
+                                                      rocke_llvm_flavor_t flavor,
+                                                      char** out_ll,
+                                                      char* err,
+                                                      size_t err_cap);
+
+#ifdef __cplusplus
+} /* extern "C" */
+#endif
+
+#endif /* ROCKE_INSTANCE_CONV_DIRECT_GROUPED_WGRAD_MFMA_H */
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/ir.h b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/ir.h
index 8a15803e49b0..633020a38771 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/ir.h
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/ir.h
@@ -1237,6 +1237,10 @@ rocke_value_t* rocke_b_mfma_f32_4x4x4_f16(rocke_ir_builder_t* b,
                                           rocke_value_t* a,
                                           rocke_value_t* bb,
                                           rocke_value_t* c);
+rocke_value_t* rocke_b_mfma_f32_4x4x4_bf16(rocke_ir_builder_t* b,
+                                           rocke_value_t* a,
+                                           rocke_value_t* bb,
+                                           rocke_value_t* c);
 rocke_value_t* rocke_b_mfma_f32_16x16x128_fp4(rocke_ir_builder_t* b,
                                               rocke_value_t* a,
                                               rocke_value_t* bb,
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/conv_direct_dw_wgrad_mfma.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/conv_direct_dw_wgrad_mfma.cpp
new file mode 100644
index 000000000000..e96331b01829
--- /dev/null
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/conv_direct_dw_wgrad_mfma.cpp
@@ -0,0 +1,939 @@
+// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
+/*
+ * conv_direct_dw_wgrad_mfma.cpp -- C++ engine port of
+ * library/kernels/common/conv_direct_dw_wgrad_mfma.py: the depthwise wgrad on
+ * the batched 4x4x4 MFMA (build_dw_wgrad_mfma).
+ *
+ * Every IR op (including each arith.constant) is emitted in the order the
+ * Python builder emits it; see the contract in
+ * rocke/instance_conv_direct_wgrad_internal.hpp.
+ */
+#include "rocke/instance_conv_direct_dw_wgrad_mfma.h"
+
+#include <stdint.h>
+#include <stdio.h>
+#include <string.h>
+
+#include <string>
+#include <vector>
+
+#include "rocke/helper_rocke.helpers.spec.h" /* rocke_kernel_name_join */
+#include "rocke/instance_conv_direct_grouped_internal.h" /* rocke_dconv_emit_params */
+#include "rocke/instance_conv_direct_wgrad_internal.hpp"
+
+using namespace rocke_dwg;
+
+namespace
+{
+
+constexpr int kNW = 4; /* waves per block */
+constexpr int kBC = 64; /* channels per block: 16 per wave */
+constexpr int kTaps = 4; /* M = N of the 4x4x4 atom */
+constexpr int kCH = kBC / 8; /* 16-byte chunks per staged pixel */
+
+const char* bf16_flag[1] = {"bf16"};
+
+int imax(int a, int b)
+{
+    return a > b ? a : b;
+}
+
+int imin(int a, int b)
+{
+    return a < b ? a : b;
+}
+
+int block_q(const rocke_dw_wgrad_mfma_spec_t* s)
+{
+    return 4 * s->q_tiles;
+}
+
+int rg(const rocke_dw_wgrad_mfma_spec_t* s)
+{
+    return cdiv(s->problem.KH, kTaps);
+}
+
+int sg(const rocke_dw_wgrad_mfma_spec_t* s)
+{
+    return cdiv(s->problem.KW, kTaps);
+}
+
+/* A lane's tap shift reaches past one v_perm (3 halves) from KW 4 on. */
+bool preshift(const rocke_dw_wgrad_mfma_spec_t* s)
+{
+    return imin(kTaps, s->problem.KW) - 1 > 2;
+}
+
+int win_words(const rocke_dw_wgrad_mfma_spec_t* s)
+{
+    return 2 * (s->q_tiles - 1) + 2 * (sg(s) - 1) + 3;
+}
+
+int win_reads(const rocke_dw_wgrad_mfma_spec_t* s)
+{
+    return cdiv(win_words(s) + (preshift(s) ? 1 : 0), 4);
+}
+
+/* main_in elements per channel: the lane's whole aligned window. */
+int in_pitch(const rocke_dw_wgrad_mfma_spec_t* s)
+{
+    return 8 * win_reads(s);
+}
+
+/* 4-column groups of X staged and transposed per row. */
+int in_groups(const rocke_dw_wgrad_mfma_spec_t* s)
+{
+    int cols = 2 * (win_words(s) + (preshift(s) ? 1 : 0));
+    return cdiv(cols, 4);
+}
+
+int d_pitch(const rocke_dw_wgrad_mfma_spec_t* s)
+{
+    return block_q(s) + 4;
+}
+
+int dnew_elems(const rocke_dw_wgrad_mfma_spec_t* s)
+{
+    return kBC * d_pitch(s);
+}
+
+/* DMA instructions per wave per staged X row. */
+int ix(const rocke_dw_wgrad_mfma_spec_t* s)
+{
+    return cdiv(cdiv(4 * in_groups(s) * kCH, kWave), kNW);
+}
+
+int iy(const rocke_dw_wgrad_mfma_spec_t* s)
+{
+    return cdiv(cdiv(block_q(s) * kCH, kWave), kNW);
+}
+
+int n_channel_tiles(const rocke_dw_wgrad_mfma_spec_t* s)
+{
+    return cdiv(s->problem.groups, kBC);
+}
+
+int n_q(const rocke_dw_wgrad_mfma_spec_t* s)
+{
+    return cdiv(problem_wo(&s->problem), block_q(s));
+}
+
+int n_units(const rocke_dw_wgrad_mfma_spec_t* s)
+{
+    const rocke_direct_conv_problem_t* p = &s->problem;
+    return cdiv(p->N, s->unfold_n) * cdiv(p->H, s->rows_per_chunk);
+}
+
+int n_z(const rocke_dw_wgrad_mfma_spec_t* s)
+{
+    int planes = n_channel_tiles(s) * n_q(s);
+    return imin(n_units(s), imax(1, cdiv(s->target_blocks, planes)));
+}
+
+} // namespace
+
+extern "C" rocke_dw_wgrad_mfma_spec_t rocke_dw_wgrad_mfma_spec_default(void)
+{
+    rocke_dw_wgrad_mfma_spec_t s;
+    memset(&s, 0, sizeof(s));
+    s.problem = rocke_direct_conv_problem_default();
+    s.name = "dw_wgrad_mfma";
+    s.q_tiles = 8;
+    s.rows_per_chunk = 16;
+    s.prefetch = 2;
+    s.target_blocks = 2048;
+    s.unfold_n = 1;
+    return s;
+}
+
+extern "C" long rocke_dw_wgrad_mfma_lds_bytes(const rocke_dw_wgrad_mfma_spec_t* s)
+{
+    long stage = (long)s->prefetch * kNW * (ix(s) + iy(s)) * kWave * 16;
+    return stage + ((long)kBC * in_pitch(s) + dnew_elems(s)) * 2;
+}
+
+extern "C" int rocke_dw_wgrad_mfma_image_pitch(const rocke_dw_wgrad_mfma_spec_t* s)
+{
+    const rocke_direct_conv_problem_t* p = &s->problem;
+    int wo = problem_wo(p);
+    return imax(imax(wo, wo + p->KW - 1 - p->PAD), p->PAD + p->W);
+}
+
+extern "C" int rocke_dw_wgrad_mfma_n_partitions(const rocke_dw_wgrad_mfma_spec_t* s)
+{
+    return n_q(s) * n_z(s);
+}
+
+extern "C" long rocke_dw_wgrad_mfma_workspace_bytes(const rocke_dw_wgrad_mfma_spec_t* s)
+{
+    const rocke_direct_conv_problem_t* p = &s->problem;
+    return (long)rocke_dw_wgrad_mfma_n_partitions(s) * p->KH * p->KW * p->groups * 4;
+}
+
+extern "C" void rocke_dw_wgrad_mfma_launch_geometry(const rocke_dw_wgrad_mfma_spec_t* s,
+                                                    int grid[3],
+                                                    int block[3])
+{
+    grid[0] = n_channel_tiles(s);
+    grid[1] = n_q(s);
+    grid[2] = n_z(s);
+    block[0] = kNW * kWave;
+    block[1] = 1;
+    block[2] = 1;
+}
+
+extern "C" rocke_status_t
+    rocke_dw_wgrad_mfma_kernel_name(const rocke_dw_wgrad_mfma_spec_t* s, char* out, size_t out_cap)
+{
+    if(s == NULL || out == NULL || out_cap == 0)
+    {
+        return ROCKE_ERR_VALUE;
+    }
+    const rocke_direct_conv_problem_t* p = &s->problem;
+    char f_buf[64], qt_buf[24], rpc_buf[24], pf_buf[24], un_buf[24], tb_buf[24];
+    snprintf(f_buf, sizeof f_buf, "f%dx%dp%ds%d", p->KH, p->KW, p->PAD, p->stride);
+    snprintf(qt_buf, sizeof qt_buf, "qt%d", s->q_tiles);
+    snprintf(rpc_buf, sizeof rpc_buf, "rpc%d", s->rows_per_chunk);
+    snprintf(pf_buf, sizeof pf_buf, "pf%d", s->prefetch);
+    snprintf(un_buf, sizeof un_buf, "un%d", s->unfold_n);
+    snprintf(tb_buf, sizeof tb_buf, "tb%d", s->target_blocks);
+    const char* parts[6] = {f_buf, qt_buf, rpc_buf, pf_buf, un_buf, tb_buf};
+    int on[1] = {problem_is_bf16(p) ? 1 : 0};
+    return rocke_kernel_name_join(s->name, parts, 6, bf16_flag, on, 1, out, out_cap, NULL);
+}
+
+extern "C" bool rocke_dw_wgrad_mfma_is_valid_spec(const rocke_dw_wgrad_mfma_spec_t* s,
+                                                  const char* arch,
+                                                  char* reason,
+                                                  size_t reason_cap)
+{
+    char msg[ROCKE_ERR_MSG_CAP];
+    if(s == NULL)
+    {
+        set_reason(reason, reason_cap, "null spec");
+        return false;
+    }
+    if(arch == NULL)
+    {
+        arch = "gfx950";
+    }
+    const rocke_direct_conv_problem_t* p = &s->problem;
+    if(dilation_reason(p, msg, sizeof msg))
+    {
+        set_reason(reason, reason_cap, msg);
+        return false;
+    }
+    const char* why = NULL;
+    if(p->cpg != 1 || p->kpg != 1)
+    {
+        why = "depthwise only";
+    }
+    else if(!problem_dtype_ok(p))
+    {
+        why = "fp16/bf16 only";
+    }
+    else if(p->stride != 1)
+    {
+        why = "stride 1 only";
+    }
+    else if(p->PAD > p->KH - 1 || p->PAD > p->KW - 1)
+    {
+        why = "PAD <= K - 1";
+    }
+    else if(p->groups % 8)
+    {
+        why = "C must be a multiple of 8";
+    }
+    else if(s->q_tiles < 1 || s->q_tiles > 8)
+    {
+        why = "q_tiles in 1..8";
+    }
+    else if(s->prefetch < 1 || s->prefetch > 4)
+    {
+        why = "prefetch in 1..4";
+    }
+    else if(s->unfold_n < 1 || s->unfold_n > 4)
+    {
+        why = "unfold_n in 1..4";
+    }
+    else if(s->unfold_n > 1 && s->unfold_n * rocke_dw_wgrad_mfma_image_pitch(s) > block_q(s))
+    {
+        ROCKE_ERR_SNPRINTF(msg,
+                           sizeof msg,
+                           "%d images at pitch %d exceed the %d-column tile",
+                           s->unfold_n,
+                           rocke_dw_wgrad_mfma_image_pitch(s),
+                           block_q(s));
+        why = msg;
+    }
+    else if(s->rows_per_chunk < 1)
+    {
+        why = "rows_per_chunk >= 1";
+    }
+    else if(rocke_dw_wgrad_mfma_lds_bytes(s) > 160 * 1024)
+    {
+        ROCKE_ERR_SNPRINTF(
+            msg, sizeof msg, "LDS %ld B > 160 KiB", rocke_dw_wgrad_mfma_lds_bytes(s));
+        why = msg;
+    }
+    else if(strcmp(arch, "gfx950") != 0)
+    {
+        why = "ds_read_tr16_b64 needs gfx950";
+    }
+    if(why != NULL)
+    {
+        set_reason(reason, reason_cap, why);
+        return false;
+    }
+    set_reason(reason, reason_cap, "ok");
+    return true;
+}
+
+namespace
+{
+
+struct Lane
+{
+    V col, ok, cc;
+};
+
+struct Off
+{
+    V off, ok, j;
+};
+
+/* ring[g][qt] = the two i32 words of a <4 x 16-bit> dY fragment. */
+using Ring = std::vector<std::vector<std::vector<V>>>;
+
+rocke_kernel_def_t*
+    build_mfma(rocke_ir_builder_t* b, const rocke_dw_wgrad_mfma_spec_t* spec, const char* arch)
+{
+    char reason[ROCKE_ERR_MSG_CAP];
+    if(!rocke_dw_wgrad_mfma_is_valid_spec(spec, arch, reason, sizeof reason))
+    {
+        fail_invalid(b, "invalid DwWgradMfmaSpec", reason);
+        return NULL;
+    }
+    Em e{b};
+    const rocke_direct_conv_problem_t& p = spec->problem;
+    const int KH = p.KH, KW = p.KW, PAD = p.PAD;
+    const int TAPS = KH * KW;
+    const int QT = spec->q_tiles, BQ = block_q(spec);
+    const int RG = rg(spec), SG = sg(spec);
+    const int D = spec->prefetch;
+    const int IX = ix(spec), IY = iy(spec);
+    const int GI = in_groups(spec);
+    const int XC = 4 * GI; /* staged X columns */
+    const int REAL_COLS = BQ + KW - 1;
+    const int HALF = in_pitch(spec);
+    const int DP = d_pitch(spec);
+    const int WW = win_words(spec);
+    const int T = kNW * kWave;
+    const int UN = spec->unfold_n;
+    const bool bf16 = problem_is_bf16(&p);
+    const rocke_type_t* io_type = bf16 ? rocke_bf16() : rocke_f16();
+
+    rocke_attr_set_int(b, &b->kernel->attrs, "max_workgroup_size", T);
+    rocke_attr_set_bool(b, &b->kernel->attrs, "lds_alias_scopes", true);
+    rocke_dconv_params_t P;
+    rocke_dconv_emit_params(b, &P, "wgrad", io_type);
+    if(!rocke_ir_builder_ok(b))
+    {
+        return NULL;
+    }
+
+    V c0 = e.ci(0);
+    V c1 = e.ci(1);
+    V c2 = e.ci(2);
+    V oob = e.ci(((int64_t)1 << 31) - 1);
+    V p_N = P.p_N;
+    V p_C = P.p_groups;
+    V p_H = P.p_Hi, p_W = P.p_Wi;
+    V p_Ho = P.p_Ho, p_Wo = P.p_Wo;
+    V sdy_n = P.p_A_stride_n, sdy_h = P.p_A_stride_hi, sdy_w = P.p_A_stride_wi;
+    V sx_n = P.p_B_stride_n, sx_h = P.p_B_stride_hi, sx_w = P.p_B_stride_wi;
+
+    V tid = rocke_b_thread_id_x(b);
+    V lane = e.modi(tid, kWave);
+    V wave = rocke_b_readfirstlane(b, e.divi(tid, kWave));
+    V ctile = rocke_b_block_id_x(b);
+    V qtile = rocke_b_block_id_y(b);
+    V zslot = rocke_b_block_id_z(b);
+    V c_blk0 = e.muli(ctile, kBC);
+    V q0 = e.muli(qtile, BQ);
+    V n_q_v = e.cdivi(p_Wo, BQ);
+    V n_chunks = e.cdivi(p_H, spec->rows_per_chunk);
+    V n_ngrp = e.cdivi(p_N, UN);
+    V n_units_v = e.mul(n_ngrp, n_chunks);
+    /* Same formula as DwWgradMfmaSpec.n_z, which sizes the grid. */
+    V n_ct = e.cdivi(p_C, kBC);
+    V planes = e.mul(n_ct, n_q_v);
+    V nz = e.ci(spec->target_blocks);
+    nz = e.add(nz, e.sub(planes, c1));
+    nz = e.div(nz, planes);
+    nz = e.smax(nz, c1);
+    V n_z_v = e.smin(n_units_v, nz);
+
+    V a_rsrc = rocke_b_buffer_rsrc(b, P.A, P.A_bytes); /* dY */
+    V b_rsrc = rocke_b_buffer_rsrc(b, P.Bp, P.B_bytes); /* X */
+
+    /* ---- LDS: one allocation per staging buffer (one alias scope each) ---- */
+    const int X_ELEMS = kNW * IX * kWave * 8;
+    const int Y_ELEMS = kNW * IY * kWave * 8;
+    std::vector<V> x_stage, y_stage, x_addr, y_addr;
+    for(int i = 0; i < D; ++i)
+    {
+        std::string h = "xs" + std::to_string(i);
+        x_stage.push_back(e.smem_alloc_1d(io_type, X_ELEMS, h.c_str()));
+    }
+    for(int i = 0; i < D; ++i)
+    {
+        std::string h = "ys" + std::to_string(i);
+        y_stage.push_back(e.smem_alloc_1d(io_type, Y_ELEMS, h.c_str()));
+    }
+    for(V s : x_stage)
+    {
+        x_addr.push_back(rocke_b_smem_addr_of(b, s));
+    }
+    for(V s : y_stage)
+    {
+        y_addr.push_back(rocke_b_smem_addr_of(b, s));
+    }
+    const int MIN = kBC * HALF;
+    const int DNEW = dnew_elems(spec);
+    V main_in = e.smem_alloc_1d(io_type, MIN, "min");
+    /* The newest dY row, channel-major; older rows live in registers. */
+    V dnew = e.smem_alloc_1d(io_type, DNEW, "dnew");
+
+    /* DMA lane geometry: slot L = (wave * I + i) * 64 + lane -> (col, chunk). */
+    auto dma_lanes = [&](int n_instr, int cols) {
+        std::vector<Lane> out;
+        for(int i = 0; i < n_instr; ++i)
+        {
+            V slot = e.muli(wave, n_instr);
+            slot = e.addi(slot, i);
+            slot = e.muli(slot, kWave);
+            slot = e.add(slot, lane);
+            Lane l;
+            l.col = e.divi(slot, kCH);
+            V ch = e.modi(slot, kCH);
+            ch = e.muli(ch, 8);
+            l.cc = e.add(c_blk0, ch);
+            V a = e.lti(l.col, cols);
+            l.ok = e.land(a, e.lt(l.cc, p_C));
+            out.push_back(l);
+        }
+        return out;
+    };
+
+    /* Lane offsets: (offset without the image and row, ok, packed image j). */
+    V pitch = NULL;
+    if(UN > 1)
+    {
+        /* Image j owns tile columns [j * P, (j + 1) * P). */
+        V t = e.addi(p_Wo, KW - 1 - PAD);
+        pitch = e.smax(p_Wo, t);
+        t = e.addi(p_W, PAD);
+        pitch = e.smax(pitch, t);
+    }
+    auto split = [&](V col, V* j, V* u) {
+        if(UN > 1)
+        {
+            *j = e.div(col, pitch);
+            V m = e.mul(*j, pitch);
+            *u = e.sub(col, m);
+        }
+        else
+        {
+            *j = c0;
+            *u = e.add(q0, col);
+        }
+    };
+
+    std::vector<Off> x_offs, y_offs;
+    {
+        std::vector<Lane> xl = dma_lanes(IX, imin(XC, REAL_COLS));
+        for(const Lane& l : xl)
+        {
+            V j, u;
+            split(l.col, &j, &u);
+            V wi = e.subi(u, PAD);
+            V a = e.ge(wi, c0);
+            a = e.land(a, e.lt(wi, p_W));
+            V okc = e.land(l.ok, a);
+            okc = e.land(okc, e.lti(j, UN));
+            V off = e.mul(wi, sx_w);
+            off = e.add(off, l.cc);
+            off = e.add(off, e.mul(j, sx_n));
+            x_offs.push_back({off, okc, j});
+        }
+        std::vector<Lane> yl = dma_lanes(IY, BQ);
+        for(const Lane& l : yl)
+        {
+            V j, wo;
+            split(l.col, &j, &wo);
+            V ok = e.land(l.ok, e.lt(wo, p_Wo));
+            ok = e.land(ok, e.lti(j, UN));
+            V off = e.mul(wo, sdy_w);
+            off = e.add(off, l.cc);
+            off = e.add(off, e.mul(j, sdy_n));
+            y_offs.push_back({off, ok, j});
+        }
+    }
+
+    auto dma = [&](V rsrc, V addr, int n_instr, int i, V voff) {
+        V o = e.muli(wave, n_instr);
+        o = e.addi(o, i);
+        o = e.muli(o, kWave * 16);
+        V lds = rocke_b_smem_ptr_add(b, addr, e.zext64(o));
+        rocke_b_async_buffer_load_lds(b, rsrc, lds, voff, c0, 4, 0);
+    };
+
+    /* ---- transpose geometry (lane L ends up with channel L) ---- */
+    V tr_row = e.modi(e.divi(lane, 4), 4);
+    V tr_cg = e.muli(e.divi(lane, 16), 4);
+    tr_cg = e.add(tr_cg, e.modi(lane, 4));
+    V tr_base = e.muli(tr_row, kBC);
+    tr_base = e.add(tr_base, e.muli(tr_cg, 4));
+
+    auto transpose_rows = [&](V stage_x, V stage_y, bool with_x) {
+        struct Read
+        {
+            int g;
+            V gg;
+            bool is_x;
+        };
+        std::vector<Read> reads;
+        if(with_x)
+        {
+            for(int g = 0; g < cdiv(GI, kNW); ++g)
+            {
+                reads.push_back({g, e.addi(wave, g * kNW), true});
+            }
+        }
+        for(int g = 0; g < cdiv(QT, kNW); ++g)
+        {
+            reads.push_back({g, e.addi(wave, g * kNW), false});
+        }
+        auto one = [&](V gg, bool is_x) {
+            V src = e.muli(gg, 4 * kBC);
+            src = e.add(tr_base, src);
+            V v = e.ds_read_tr(is_x ? stage_x : stage_y, c0, src, io_type);
+            V dst = e.muli(lane, is_x ? HALF : DP);
+            dst = e.add(dst, e.muli(gg, 4));
+            e.smem_store(is_x ? main_in : dnew, c0, dst, v, 4);
+        };
+        for(const Read& r : reads)
+        {
+            int n_groups = r.is_x ? GI : QT;
+            if((r.g + 1) * kNW > n_groups)
+            {
+                e.if_(e.lti(r.gg, n_groups), [&]() { one(r.gg, r.is_x); });
+            }
+            else
+            {
+                one(r.gg, r.is_x);
+            }
+        }
+    };
+
+    /* ---- per-lane MFMA operand geometry ---- */
+    V lane4 = e.modi(lane, 4);
+    V ch_w = e.muli(wave, 16);
+    ch_w = e.add(ch_w, e.divi(lane, 4)); /* channel in the tile */
+    V lane_s = e.smini(lane4, imin(kTaps, KW) - 1);
+    V lane_word = preshift(spec) ? e.div(lane_s, c2) : c0;
+    V lane_res = e.sub(lane_s, e.mul(lane_word, c2));
+    V sel0 = e.ci(0x03020100);
+    V sel1 = e.ci(0x01010101);
+    V pre_sel = e.mul(sel1, e.muli(lane_word, 4));
+    pre_sel = e.add(sel0, pre_sel);
+    V res_sel = e.mul(sel1, e.mul(lane_res, c2));
+    res_sel = e.add(sel0, res_sel);
+    const rocke_type_t* frag_t = rocke_vector_type(b, io_type, 4);
+    /* The four lanes of a channel each read a quarter of its main_in window. */
+    const int NWR = win_reads(spec);
+    std::vector<V> a_reads, d_reads;
+    for(int k = 0; k < cdiv(NWR, 4); ++k)
+    {
+        V a = e.muli(ch_w, HALF);
+        V c = e.addi(lane4, 4 * k);
+        c = e.smini(c, NWR - 1);
+        c = e.muli(c, 8);
+        a_reads.push_back(e.add(a, c));
+    }
+    /* Likewise the newest dY row: lane j of a channel reads q tiles j, j + 4. */
+    for(int k = 0; k < cdiv(QT, 4); ++k)
+    {
+        V a = e.muli(ch_w, DP);
+        V c = e.addi(lane4, 4 * k);
+        c = e.smini(c, QT - 1);
+        c = e.muli(c, 4);
+        d_reads.push_back(e.add(a, c));
+    }
+    V head = e.eq(lane4, c0);
+
+    auto bcast = [&](V v, int src) { return rocke_b_quad_perm(b, v, src, src, src, src); };
+    auto words = [&](V v, int n) {
+        V w = rocke_b_vec_bitcast(b, v, rocke_vector_type(b, rocke_i32(), n));
+        std::vector<V> out;
+        for(int el = 0; el < n; ++el)
+        {
+            out.push_back(rocke_b_vec_extract(b, w, el));
+        }
+        return out;
+    };
+
+    /* Advance the dY ring one row: lane (g, r) takes (g, r - 1)'s row, lane
+     * (g, 0) takes (g - 1, 3)'s, and lane (0, 0) the newest row. */
+    auto shift_ring = [&](const Ring& ring) {
+        std::vector<std::vector<V>> fresh;
+        for(V idx : d_reads)
+        {
+            fresh.push_back(words(e.smem_load(dnew, c0, idx, io_type, 4), 2));
+        }
+        Ring rot(RG, std::vector<std::vector<V>>(QT));
+        for(int g = 0; g < RG; ++g)
+        {
+            for(int qt = 0; qt < QT; ++qt)
+            {
+                for(V d : ring[g][qt])
+                {
+                    rot[g][qt].push_back(rocke_b_quad_perm(b, d, 3, 0, 1, 2));
+                }
+            }
+        }
+        Ring out(RG, std::vector<std::vector<V>>(QT));
+        for(int g = 0; g < RG; ++g)
+        {
+            for(int qt = 0; qt < QT; ++qt)
+            {
+                std::vector<V> src;
+                if(g == 0)
+                {
+                    for(V d : fresh[qt / 4])
+                    {
+                        src.push_back(bcast(d, qt % 4));
+                    }
+                }
+                else
+                {
+                    src = rot[g - 1][qt];
+                }
+                for(int el = 0; el < 2; ++el)
+                {
+                    out[g][qt].push_back(e.select(head, src[el], rot[g][qt][el]));
+                }
+            }
+        }
+        return out;
+    };
+
+    /* MFMA the X row in main_in against the dY rows in ``ring``. */
+    auto compute = [&](const Ring& ring, std::vector<V>& accs) {
+        std::vector<std::vector<V>> part;
+        for(V idx : a_reads)
+        {
+            part.push_back(words(e.smem_load(main_in, c0, idx, io_type, 8), 4));
+        }
+        std::vector<V> win;
+        for(int m = 0; m < NWR; ++m)
+        {
+            for(int el = 0; el < 4; ++el)
+            {
+                win.push_back(bcast(part[m / 4][el], m % 4));
+            }
+        }
+        if(preshift(spec))
+        {
+            std::vector<V> w2;
+            for(int j = 0; j < WW; ++j)
+            {
+                w2.push_back(rocke_b_perm_b32(b, win[j + 1], win[j], pre_sel));
+            }
+            win = w2;
+        }
+        std::vector<std::vector<V>> bfr(RG);
+        for(int g = 0; g < RG; ++g)
+        {
+            for(int qt = 0; qt < QT; ++qt)
+            {
+                V pk = rocke_b_vec_pack(b, ring[g][qt].data(), 2, rocke_i32());
+                bfr[g].push_back(rocke_b_vec_bitcast(b, pk, frag_t));
+            }
+        }
+        for(int sgi = 0; sgi < SG; ++sgi)
+        {
+            for(int qt = 0; qt < QT; ++qt)
+            {
+                int k = 2 * qt + 2 * sgi;
+                V a0 = rocke_b_perm_b32(b, win[k + 1], win[k], res_sel);
+                V a1 = rocke_b_perm_b32(b, win[k + 2], win[k + 1], res_sel);
+                V pr[2] = {a0, a1};
+                V af = rocke_b_vec_bitcast(b, rocke_b_vec_pack(b, pr, 2, rocke_i32()), frag_t);
+                for(int rgi = 0; rgi < RG; ++rgi)
+                {
+                    int i = rgi * SG + sgi;
+                    accs[i] = bf16 ? rocke_b_mfma_f32_4x4x4_bf16(b, af, bfr[rgi][qt], accs[i])
+                                   : rocke_b_mfma_f32_4x4x4_f16(b, af, bfr[rgi][qt], accs[i]);
+                }
+            }
+        }
+    };
+
+    V zero_acc = rocke_b_zero_vec_f32(b, 4);
+    const int n_acc = RG * SG;
+    const int PER_STEP = IX + IY;
+
+    auto pack = [&](const std::vector<V>& accs, const Ring& ring) {
+        std::vector<V> out(accs);
+        for(const auto& g : ring)
+        {
+            for(const auto& qt : g)
+            {
+                for(V d : qt)
+                {
+                    out.push_back(d);
+                }
+            }
+        }
+        return out;
+    };
+    auto unpack = [&](const std::vector<V>& vals, std::vector<V>* accs, Ring* ring) {
+        accs->assign(vals.begin(), vals.begin() + n_acc);
+        ring->assign(RG, std::vector<std::vector<V>>(QT));
+        for(int g = 0; g < RG; ++g)
+        {
+            for(int qt = 0; qt < QT; ++qt)
+            {
+                int o = n_acc + (g * QT + qt) * 2;
+                (*ring)[g][qt] = {vals[o], vals[o + 1]};
+            }
+        }
+    };
+
+    /* Unit loop and epilogue, PRE dY rows staged before a unit's first X row. */
+    auto emit_body = [&](int PRE) {
+        std::string acc_prefix = "acc" + std::to_string(PRE) + "_";
+        std::string unit_name = "unit" + std::to_string(PRE);
+        std::vector<V> accs = e.for_iter(
+            zslot,
+            n_units_v,
+            n_z_v,
+            iter_names(acc_prefix.c_str(), n_acc),
+            std::vector<V>(n_acc, zero_acc),
+            unit_name.c_str(),
+            [&](V unit, std::vector<V> carried) {
+                /* A fresh ring per unit: its zeros stand in for the dY rows
+                 * above the image, which the whole-image prologue skips. */
+                std::vector<V> la_accs = carried;
+                Ring ring(RG, std::vector<std::vector<V>>(QT, std::vector<V>{c0, c0}));
+                V n_img = e.muli(e.mod(unit, n_ngrp), UN);
+                V chunk = e.div(unit, n_ngrp);
+                V y0 = e.muli(chunk, spec->rows_per_chunk);
+                V k_rpc = e.ci(spec->rows_per_chunk);
+                V n_rows = e.smin(k_rpc, e.sub(p_H, y0));
+                V n_steps = e.addi(n_rows, PRE);
+                V x_img = e.mul(n_img, sx_n);
+                V y_img = e.mul(n_img, sdy_n);
+
+                /* DMA step t: X row y0 - PRE + t (steps >= PRE), dY row y + PAD. */
+                auto issue = [&](int buf, V t) {
+                    V y = e.add(e.subi(y0, PRE), t);
+                    V x_ok = e.gei(t, PRE);
+                    x_ok = e.land(x_ok, e.lt(t, n_steps));
+                    V x_row = e.add(x_img, e.mul(y, sx_h));
+                    for(int i = 0; i < (int)x_offs.size(); ++i)
+                    {
+                        V ok = e.land(x_offs[i].ok, x_ok);
+                        if(UN > 1)
+                        {
+                            V a = e.add(n_img, x_offs[i].j);
+                            ok = e.land(ok, e.lt(a, p_N));
+                        }
+                        V voff = e.mul(e.add(x_offs[i].off, x_row), c2);
+                        voff = e.select(ok, voff, oob);
+                        dma(b_rsrc, x_addr[buf], IX, i, voff);
+                    }
+                    V pr = e.addi(y, PAD);
+                    V y_ok = e.ge(pr, c0);
+                    y_ok = e.land(y_ok, e.lt(pr, p_Ho));
+                    y_ok = e.land(y_ok, e.lt(t, n_steps));
+                    V y_row = e.add(y_img, e.mul(pr, sdy_h));
+                    for(int i = 0; i < (int)y_offs.size(); ++i)
+                    {
+                        V ok = e.land(y_offs[i].ok, y_ok);
+                        if(UN > 1)
+                        {
+                            V a = e.add(n_img, y_offs[i].j);
+                            ok = e.land(ok, e.lt(a, p_N));
+                        }
+                        V voff = e.mul(e.add(y_offs[i].off, y_row), c2);
+                        voff = e.select(ok, voff, oob);
+                        dma(a_rsrc, y_addr[buf], IY, i, voff);
+                    }
+                };
+
+                auto step
+                    = [&](V t, int buf, std::vector<V>& accs_s, Ring& ring_s, bool do_compute) {
+                          /* Issue the row D - 1 ahead into the buffer the last step drained. */
+                          issue((buf + D - 1) % D, e.addi(t, D - 1));
+                          rocke_b_s_waitcnt(b, (D - 1) * PER_STEP, 0, -1);
+                          rocke_b_s_barrier_bare(b);
+                          transpose_rows(x_stage[buf], y_stage[buf], do_compute);
+                          rocke_b_s_waitcnt(b, -1, 0, -1);
+                          rocke_b_s_barrier_bare(b);
+                          ring_s = shift_ring(ring_s);
+                          if(do_compute)
+                          {
+                              compute(ring_s, accs_s);
+                          }
+                      };
+
+                for(int t = 0; t < D - 1; ++t)
+                {
+                    issue(t % D, e.ci(t));
+                }
+                /* Prologue: the PRE dY rows the first input row pairs with. */
+                for(int t = 0; t < PRE; ++t)
+                {
+                    step(e.ci(t), t % D, la_accs, ring, false);
+                }
+                /* Main: one input row per step, unrolled by D. */
+                V trips = e.cdivi(n_rows, D);
+                trips = e.smax(trips, c1);
+                std::string a_prefix = "a" + std::to_string(PRE) + "_";
+                std::string rb_name = "rb" + std::to_string(PRE);
+                std::vector<V> init = pack(la_accs, ring);
+                std::vector<V> res = e.for_iter(c0,
+                                                trips,
+                                                c1,
+                                                iter_names(a_prefix.c_str(), init.size()),
+                                                init,
+                                                rb_name.c_str(),
+                                                [&](V rb, std::vector<V> la) {
+                                                    std::vector<V> l_acc;
+                                                    Ring l_ring;
+                                                    unpack(la, &l_acc, &l_ring);
+                                                    for(int u = 0; u < D; ++u)
+                                                    {
+                                                        V t = e.muli(rb, D);
+                                                        t = e.addi(t, PRE);
+                                                        t = e.addi(t, u);
+                                                        step(t, (PRE + u) % D, l_acc, l_ring, true);
+                                                    }
+                                                    return pack(l_acc, l_ring);
+                                                });
+                /* Drain the tail DMAs before the next unit reuses the buffers. */
+                rocke_b_s_waitcnt(b, 0, 0, -1);
+                rocke_b_s_barrier_bare(b);
+                if((int)res.size() < n_acc)
+                {
+                    return std::vector<V>();
+                }
+                return std::vector<V>(res.begin(), res.begin() + n_acc);
+            });
+        if((int)accs.size() != n_acc)
+        {
+            return;
+        }
+
+        /* ---- epilogue: ws[part][tap][c], part = zslot * n_q + qtile ---- */
+        V part = e.add(e.mul(zslot, n_q_v), qtile);
+        V c_glb = e.add(c_blk0, ch_w);
+        V c_ok = e.lt(c_glb, p_C);
+        for(int rgi = 0; rgi < RG; ++rgi)
+        {
+            V r = e.addi(lane4, rgi * kTaps);
+            V row0 = e.muli(part, TAPS);
+            row0 = e.add(row0, e.muli(r, KW));
+            row0 = e.mul(row0, p_C);
+            row0 = e.add(row0, c_glb);
+            e.if_(e.land(c_ok, e.lti(r, KH)), [&]() {
+                for(int sgi = 0; sgi < SG; ++sgi)
+                {
+                    V acc = accs[rgi * SG + sgi];
+                    for(int i = 0; i < kTaps; ++i)
+                    {
+                        int s = sgi * kTaps + i;
+                        if(s < KW)
+                        {
+                            V ks = e.ci(s);
+                            V off = e.add(row0, e.mul(ks, p_C));
+                            rocke_b_global_store(b, P.D, off, rocke_b_vec_extract(b, acc, i), 4);
+                        }
+                    }
+                }
+            });
+        }
+    };
+
+    if(PAD == KH - 1)
+    {
+        emit_body(PAD);
+    }
+    else
+    {
+        /* Whole images per unit (one row chunk) skip the KH - 1 - PAD zero rows. */
+        rocke_if_else_t ife = rocke_b_scf_if_else(b, e.eq(n_chunks, c1));
+        if(ife.then_region == NULL || ife.else_region == NULL)
+        {
+            return NULL;
+        }
+        rocke_b_region_enter(b, ife.then_region);
+        emit_body(PAD);
+        rocke_b_region_leave(b);
+        rocke_b_region_enter(b, ife.else_region);
+        emit_body(KH - 1);
+        rocke_b_region_leave(b);
+    }
+    return rocke_ir_builder_ok(b) ? b->kernel : NULL;
+}
+
+} // namespace
+
+extern "C" rocke_kernel_def_t* rocke_build_dw_wgrad_mfma(rocke_ir_builder_t* b,
+                                                         const rocke_dw_wgrad_mfma_spec_t* spec,
+                                                         const char* arch)
+{
+    if(b == NULL || spec == NULL)
+    {
+        return NULL;
+    }
+    return build_mfma(b, spec, arch ? arch : "gfx950");
+}
+
+extern "C" rocke_kernel_def_t* rocke_build_dw_wgrad_mfma_new(rocke_ir_builder_t* b,
+                                                             const rocke_dw_wgrad_mfma_spec_t* spec,
+                                                             const char* arch)
+{
+    return build_new(
+        b,
+        spec != NULL,
+        "build_dw_wgrad_mfma",
+        [&](char* out, size_t cap) { return rocke_dw_wgrad_mfma_kernel_name(spec, out, cap); },
+        [&]() { return rocke_build_dw_wgrad_mfma(b, spec, arch); });
+}
+
+extern "C" rocke_status_t rocke_dw_wgrad_mfma_lower_to_llvm(const rocke_dw_wgrad_mfma_spec_t* spec,
+                                                            const char* arch,
+                                                            rocke_llvm_flavor_t flavor,
+                                                            char** out_ll,
+                                                            char* err,
+                                                            size_t err_cap)
+{
+    if(arch == NULL)
+    {
+        arch = "gfx950";
+    }
+    return lower_to_llvm(
+        spec != NULL,
+        arch,
+        flavor,
+        out_ll,
+        err,
+        err_cap,
+        "build_dw_wgrad_mfma failed",
+        [&](rocke_ir_builder_t* b) { return rocke_build_dw_wgrad_mfma_new(b, spec, arch); });
+}
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/conv_direct_grouped_wgrad_mfma.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/conv_direct_grouped_wgrad_mfma.cpp
new file mode 100644
index 000000000000..7d382c4b48a1
--- /dev/null
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/conv_direct_grouped_wgrad_mfma.cpp
@@ -0,0 +1,699 @@
+// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
+/*
+ * conv_direct_grouped_wgrad_mfma.cpp -- C++ engine port of
+ * library/kernels/common/conv_direct_grouped_wgrad_mfma.py: the grouped wgrad
+ * for 4 and 8 channels per group on the batched 4x4x4 MFMA
+ * (build_grouped_wgrad_mfma).
+ *
+ * Every IR op (including each arith.constant) is emitted in the order the
+ * Python builder emits it; see the contract in
+ * rocke/instance_conv_direct_wgrad_internal.hpp.
+ */
+#include "rocke/instance_conv_direct_grouped_wgrad_mfma.h"
+
+#include <stdint.h>
+#include <stdio.h>
+#include <string.h>
+
+#include <string>
+#include <vector>
+
+#include "rocke/helper_rocke.helpers.spec.h" /* rocke_kernel_name_join */
+#include "rocke/instance_conv_direct_grouped_internal.h" /* rocke_dconv_emit_params */
+#include "rocke/instance_conv_direct_wgrad_internal.hpp"
+
+using namespace rocke_dwg;
+
+namespace
+{
+
+constexpr int kBC = 64; /* channels per block */
+constexpr int kCH = kBC / 8; /* 16-byte chunks per staged pixel */
+
+const char* bf16_flag[1] = {"bf16"};
+
+int total_c(const rocke_direct_conv_problem_t* p)
+{
+    return p->groups * p->cpg;
+}
+
+int total_k(const rocke_direct_conv_problem_t* p)
+{
+    return p->groups * p->kpg;
+}
+
+int block_q(const rocke_grouped_wgrad_mfma_spec_t* s)
+{
+    return 4 * s->wave_q4 * s->waves_q;
+}
+
+/* 4x4 blocks along k per group: 1 for cpg 4, 2 for cpg 8. */
+int halves(const rocke_grouped_wgrad_mfma_spec_t* s)
+{
+    return s->problem.cpg / 4;
+}
+
+int x_cols(const rocke_grouped_wgrad_mfma_spec_t* s)
+{
+    return block_q(s) + s->problem.KW - 1;
+}
+
+int ix(const rocke_grouped_wgrad_mfma_spec_t* s)
+{
+    return cdiv(cdiv(x_cols(s) * kCH, kWave), s->waves_q);
+}
+
+int iy(const rocke_grouped_wgrad_mfma_spec_t* s)
+{
+    return cdiv(cdiv(block_q(s) * kCH, kWave), s->waves_q);
+}
+
+int n_channel_tiles(const rocke_grouped_wgrad_mfma_spec_t* s)
+{
+    return cdiv(total_c(&s->problem), kBC);
+}
+
+int n_q(const rocke_grouped_wgrad_mfma_spec_t* s)
+{
+    return cdiv(problem_wo(&s->problem), block_q(s));
+}
+
+int n_units(const rocke_grouped_wgrad_mfma_spec_t* s)
+{
+    return s->problem.N * cdiv(s->problem.H, s->rows_per_chunk);
+}
+
+int n_z(const rocke_grouped_wgrad_mfma_spec_t* s)
+{
+    int planes = n_channel_tiles(s) * n_q(s);
+    int c = cdiv(s->target_blocks, planes);
+    if(c < 1)
+    {
+        c = 1;
+    }
+    int u = n_units(s);
+    return u < c ? u : c;
+}
+
+int n_acc(const rocke_grouped_wgrad_mfma_spec_t* s)
+{
+    return s->problem.KH * s->problem.KW * halves(s);
+}
+
+} // namespace
+
+extern "C" rocke_grouped_wgrad_mfma_spec_t rocke_grouped_wgrad_mfma_spec_default(void)
+{
+    rocke_grouped_wgrad_mfma_spec_t s;
+    memset(&s, 0, sizeof(s));
+    s.problem = rocke_direct_conv_problem_default();
+    s.name = "grouped_wgrad_mfma";
+    s.waves_q = 4;
+    s.wave_q4 = 1;
+    s.rows_per_chunk = 56;
+    s.prefetch = 2;
+    s.target_blocks = 1024;
+    return s;
+}
+
+extern "C" int rocke_grouped_wgrad_mfma_threads(const rocke_grouped_wgrad_mfma_spec_t* s)
+{
+    return s->waves_q * kWave;
+}
+
+extern "C" long rocke_grouped_wgrad_mfma_lds_bytes(const rocke_grouped_wgrad_mfma_spec_t* s)
+{
+    long stage = (long)s->prefetch * s->waves_q * (ix(s) + iy(s)) * kWave * 16;
+    long red = (long)s->waves_q * kWave * 4 * 4;
+    return stage + red;
+}
+
+extern "C" int rocke_grouped_wgrad_mfma_n_partitions(const rocke_grouped_wgrad_mfma_spec_t* s)
+{
+    return n_q(s) * n_z(s);
+}
+
+extern "C" int rocke_grouped_wgrad_mfma_fold_taps(const rocke_grouped_wgrad_mfma_spec_t* s)
+{
+    return s->problem.KH * s->problem.KW * s->problem.cpg;
+}
+
+extern "C" long rocke_grouped_wgrad_mfma_workspace_bytes(const rocke_grouped_wgrad_mfma_spec_t* s)
+{
+    return (long)rocke_grouped_wgrad_mfma_n_partitions(s) * rocke_grouped_wgrad_mfma_fold_taps(s)
+           * total_k(&s->problem) * 4;
+}
+
+extern "C" void rocke_grouped_wgrad_mfma_launch_geometry(const rocke_grouped_wgrad_mfma_spec_t* s,
+                                                         int grid[3],
+                                                         int block[3])
+{
+    grid[0] = n_channel_tiles(s);
+    grid[1] = n_q(s);
+    grid[2] = n_z(s);
+    block[0] = rocke_grouped_wgrad_mfma_threads(s);
+    block[1] = 1;
+    block[2] = 1;
+}
+
+extern "C" rocke_status_t rocke_grouped_wgrad_mfma_kernel_name(
+    const rocke_grouped_wgrad_mfma_spec_t* s, char* out, size_t out_cap)
+{
+    if(s == NULL || out == NULL || out_cap == 0)
+    {
+        return ROCKE_ERR_VALUE;
+    }
+    const rocke_direct_conv_problem_t* p = &s->problem;
+    char f_buf[64], wq_buf[24], q4_buf[24], rpc_buf[24], pf_buf[24], tb_buf[24];
+    snprintf(f_buf, sizeof f_buf, "f%dx%dp%ds%dc%d", p->KH, p->KW, p->PAD, p->stride, p->cpg);
+    snprintf(wq_buf, sizeof wq_buf, "wq%d", s->waves_q);
+    snprintf(q4_buf, sizeof q4_buf, "q4%d", s->wave_q4);
+    snprintf(rpc_buf, sizeof rpc_buf, "rpc%d", s->rows_per_chunk);
+    snprintf(pf_buf, sizeof pf_buf, "pf%d", s->prefetch);
+    snprintf(tb_buf, sizeof tb_buf, "tb%d", s->target_blocks);
+    const char* parts[6] = {f_buf, wq_buf, q4_buf, rpc_buf, pf_buf, tb_buf};
+    int on[1] = {problem_is_bf16(p) ? 1 : 0};
+    return rocke_kernel_name_join(s->name, parts, 6, bf16_flag, on, 1, out, out_cap, NULL);
+}
+
+extern "C" bool rocke_grouped_wgrad_mfma_is_valid_spec(const rocke_grouped_wgrad_mfma_spec_t* s,
+                                                       const char* arch,
+                                                       char* reason,
+                                                       size_t reason_cap)
+{
+    char msg[ROCKE_ERR_MSG_CAP];
+    if(s == NULL)
+    {
+        set_reason(reason, reason_cap, "null spec");
+        return false;
+    }
+    if(arch == NULL)
+    {
+        arch = "gfx950";
+    }
+    const rocke_direct_conv_problem_t* p = &s->problem;
+    if(dilation_reason(p, msg, sizeof msg))
+    {
+        set_reason(reason, reason_cap, msg);
+        return false;
+    }
+    const char* why = NULL;
+    if(p->cpg != p->kpg || (p->cpg != 4 && p->cpg != 8))
+    {
+        why = "cpg = kpg in (4, 8)";
+    }
+    else if(!problem_dtype_ok(p))
+    {
+        why = "fp16/bf16 only";
+    }
+    else if(p->stride != 1)
+    {
+        why = "stride 1 only";
+    }
+    else if(p->PAD > p->KH - 1 || p->PAD > p->KW - 1)
+    {
+        why = "PAD <= K - 1";
+    }
+    else if(total_c(p) % 8)
+    {
+        why = "C must be a multiple of 8";
+    }
+    else if(s->waves_q != 1 && s->waves_q != 2 && s->waves_q != 4 && s->waves_q != 8)
+    {
+        why = "waves_q in (1, 2, 4, 8)";
+    }
+    else if(s->wave_q4 != 1 && s->wave_q4 != 2 && s->wave_q4 != 4)
+    {
+        why = "wave_q4 in (1, 2, 4)";
+    }
+    else if(s->prefetch < 2 || s->prefetch > 4)
+    {
+        why = "prefetch in 2..4";
+    }
+    else if(s->rows_per_chunk < 1)
+    {
+        why = "rows_per_chunk >= 1";
+    }
+    else if(n_acc(s) * 4 > 128)
+    {
+        why = "too many accumulators";
+    }
+    else if(rocke_grouped_wgrad_mfma_lds_bytes(s) > 160 * 1024)
+    {
+        ROCKE_ERR_SNPRINTF(
+            msg, sizeof msg, "LDS %ld B > 160 KiB", rocke_grouped_wgrad_mfma_lds_bytes(s));
+        why = msg;
+    }
+    else if(strcmp(arch, "gfx950") != 0)
+    {
+        why = "ds_read_tr16_b64 needs gfx950";
+    }
+    if(why != NULL)
+    {
+        set_reason(reason, reason_cap, why);
+        return false;
+    }
+    set_reason(reason, reason_cap, "ok");
+    return true;
+}
+
+namespace
+{
+
+struct Lane
+{
+    V col, ok, cc;
+};
+
+struct Off
+{
+    V off, ok;
+};
+
+/* ring[row][q4][half]: dY fragments, newest row first. */
+using Ring = std::vector<std::vector<std::vector<V>>>;
+
+rocke_kernel_def_t* build_grouped(rocke_ir_builder_t* b,
+                                  const rocke_grouped_wgrad_mfma_spec_t* spec,
+                                  const char* arch)
+{
+    char reason[ROCKE_ERR_MSG_CAP];
+    if(!rocke_grouped_wgrad_mfma_is_valid_spec(spec, arch, reason, sizeof reason))
+    {
+        fail_invalid(b, "invalid GroupedWgradMfmaSpec", reason);
+        return NULL;
+    }
+    Em e{b};
+    const rocke_direct_conv_problem_t& p = spec->problem;
+    const int KH = p.KH, KW = p.KW, PAD = p.PAD, CPG = p.cpg;
+    const int HV = halves(spec);
+    const int TAPS = rocke_grouped_wgrad_mfma_fold_taps(spec);
+    const int WQ = spec->waves_q, Q4 = spec->wave_q4;
+    const int BQ = block_q(spec);
+    const int XC = x_cols(spec);
+    const int D = spec->prefetch;
+    const int IX = ix(spec), IY = iy(spec);
+    const int T = rocke_grouped_wgrad_mfma_threads(spec);
+    const int PRE = KH - 1; /* dY rows a unit needs before its first input row */
+    const bool bf16 = problem_is_bf16(&p);
+    const rocke_type_t* io_type = bf16 ? rocke_bf16() : rocke_f16();
+
+    rocke_attr_set_int(b, &b->kernel->attrs, "max_workgroup_size", T);
+    rocke_attr_set_bool(b, &b->kernel->attrs, "lds_alias_scopes", true);
+    rocke_dconv_params_t P;
+    rocke_dconv_emit_params(b, &P, "wgrad", io_type);
+    if(!rocke_ir_builder_ok(b))
+    {
+        return NULL;
+    }
+
+    V c0 = e.ci(0);
+    V c1 = e.ci(1);
+    V c2 = e.ci(2);
+    V oob = e.ci(((int64_t)1 << 31) - 1);
+    V p_N = P.p_N;
+    V p_C = P.p_total_c;
+    V p_K = P.p_total_k;
+    V p_H = P.p_Hi, p_W = P.p_Wi;
+    V p_Ho = P.p_Ho, p_Wo = P.p_Wo;
+    V sdy_n = P.p_A_stride_n, sdy_h = P.p_A_stride_hi, sdy_w = P.p_A_stride_wi;
+    V sx_n = P.p_B_stride_n, sx_h = P.p_B_stride_hi, sx_w = P.p_B_stride_wi;
+
+    V tid = rocke_b_thread_id_x(b);
+    V lane = e.modi(tid, kWave);
+    V wave = rocke_b_readfirstlane(b, e.divi(tid, kWave));
+    V ctile = rocke_b_block_id_x(b);
+    V qtile = rocke_b_block_id_y(b);
+    V zslot = rocke_b_block_id_z(b);
+    V c_blk0 = e.muli(ctile, kBC);
+    V q0 = e.muli(qtile, BQ);
+    V n_q_v = e.cdivi(p_Wo, BQ);
+    V n_chunks = e.cdivi(p_H, spec->rows_per_chunk);
+    V n_units_v = e.mul(p_N, n_chunks);
+    /* Same formula as GroupedWgradMfmaSpec.n_z, which sizes the grid. */
+    V n_ct = e.cdivi(p_C, kBC);
+    V planes = e.mul(n_ct, n_q_v);
+    V nz = e.ci(spec->target_blocks);
+    nz = e.add(nz, e.sub(planes, c1));
+    nz = e.div(nz, planes);
+    nz = e.smax(nz, c1);
+    V n_z_v = e.smin(n_units_v, nz);
+
+    V a_rsrc = rocke_b_buffer_rsrc(b, P.A, P.A_bytes); /* dY */
+    V b_rsrc = rocke_b_buffer_rsrc(b, P.Bp, P.B_bytes); /* X */
+
+    /* ---- LDS: one allocation per staging buffer (one alias scope each) ---- */
+    const int X_ELEMS = WQ * IX * kWave * 8;
+    const int Y_ELEMS = WQ * IY * kWave * 8;
+    std::vector<V> x_stage, y_stage, x_addr, y_addr;
+    for(int i = 0; i < D; ++i)
+    {
+        std::string h = "xs" + std::to_string(i);
+        x_stage.push_back(e.smem_alloc_1d(io_type, X_ELEMS, h.c_str()));
+    }
+    for(int i = 0; i < D; ++i)
+    {
+        std::string h = "ys" + std::to_string(i);
+        y_stage.push_back(e.smem_alloc_1d(io_type, Y_ELEMS, h.c_str()));
+    }
+    for(V s : x_stage)
+    {
+        x_addr.push_back(rocke_b_smem_addr_of(b, s));
+    }
+    for(V s : y_stage)
+    {
+        y_addr.push_back(rocke_b_smem_addr_of(b, s));
+    }
+
+    /* DMA lane geometry: slot L = (wave * I + i) * 64 + lane -> (col, chunk). */
+    auto dma_lanes = [&](int n_instr, int cols) {
+        std::vector<Lane> out;
+        for(int i = 0; i < n_instr; ++i)
+        {
+            V slot = e.muli(wave, n_instr);
+            slot = e.addi(slot, i);
+            slot = e.muli(slot, kWave);
+            slot = e.add(slot, lane);
+            Lane l;
+            l.col = e.divi(slot, kCH);
+            V ch = e.modi(slot, kCH);
+            ch = e.muli(ch, 8);
+            l.cc = e.add(c_blk0, ch);
+            V a = e.lti(l.col, cols);
+            l.ok = e.land(a, e.lt(l.cc, p_C));
+            out.push_back(l);
+        }
+        return out;
+    };
+
+    std::vector<Off> x_offs, y_offs;
+    {
+        std::vector<Lane> xl = dma_lanes(IX, XC);
+        for(const Lane& l : xl)
+        {
+            V wi = e.add(e.subi(q0, PAD), l.col);
+            V a = e.ge(wi, c0);
+            a = e.land(a, e.lt(wi, p_W));
+            V okc = e.land(l.ok, a);
+            V off = e.add(e.mul(wi, sx_w), l.cc);
+            x_offs.push_back({off, okc});
+        }
+        std::vector<Lane> yl = dma_lanes(IY, BQ);
+        for(const Lane& l : yl)
+        {
+            V wo = e.add(q0, l.col);
+            V off = e.add(e.mul(wo, sdy_w), l.cc);
+            V ok = e.land(l.ok, e.lt(wo, p_Wo));
+            y_offs.push_back({off, ok});
+        }
+    }
+
+    auto dma = [&](V rsrc, V addr, int n_instr, int i, V voff) {
+        V o = e.muli(wave, n_instr);
+        o = e.addi(o, i);
+        o = e.muli(o, kWave * 16);
+        V lds = rocke_b_smem_ptr_add(b, addr, e.zext64(o));
+        rocke_b_async_buffer_load_lds(b, rsrc, lds, voff, c0, 4, 0);
+    };
+
+    /* ---- operand reads: out lane L gets channel L, four columns ---- */
+    V tr_row = e.modi(e.divi(lane, 4), 4);
+    V tr_cg = e.muli(e.divi(lane, 16), 4);
+    tr_cg = e.add(tr_cg, e.modi(lane, 4));
+    V tr_base = e.muli(tr_row, kBC);
+    tr_base = e.add(tr_base, e.muli(tr_cg, 4));
+    /* The partner half of an 8-channel group (channel ^ 4). */
+    V tr_swap = e.muli(tr_row, kBC);
+    tr_swap = e.add(tr_swap, e.muli(e.bxor(tr_cg, c1), 4));
+    V col_w = e.muli(wave, 4 * Q4); /* this wave's first column in the tile */
+
+    auto read = [&](V stage, V base, int col) {
+        V idx = e.addi(col_w, col);
+        idx = e.muli(idx, kBC);
+        idx = e.add(base, idx);
+        return e.ds_read_tr(stage, c0, idx, io_type);
+    };
+
+    V zero_acc = rocke_b_zero_vec_f32(b, 4);
+    V zero_frag = rocke_b_zero_vec(b, io_type, 4);
+    const int NACC = KH * KW * HV;
+    /* The carried ring holds the dY fragments of the KH - 1 rows before the
+     * current one, newest first: [row][q4][half]. */
+    const int n_ring = (KH - 1) * Q4 * HV;
+    auto acc_i = [&](int r, int s, int h) { return (r * KW + s) * HV + h; };
+
+    auto pack = [&](const std::vector<V>& accs, const Ring& ring) {
+        std::vector<V> out(accs);
+        for(const auto& row : ring)
+        {
+            for(const auto& q : row)
+            {
+                for(V v : q)
+                {
+                    out.push_back(v);
+                }
+            }
+        }
+        return out;
+    };
+    auto unpack = [&](const std::vector<V>& vals, std::vector<V>* accs, Ring* ring) {
+        accs->assign(vals.begin(), vals.begin() + NACC);
+        ring->assign(KH - 1, std::vector<std::vector<V>>(Q4));
+        for(int j = 0; j < KH - 1; ++j)
+        {
+            for(int q = 0; q < Q4; ++q)
+            {
+                int o = NACC + (j * Q4 + q) * HV;
+                (*ring)[j][q].assign(vals.begin() + o, vals.begin() + o + HV);
+            }
+        }
+    };
+
+    std::vector<std::string> unit_names = iter_names("acc", NACC);
+    std::vector<std::string> ring_names = iter_names("ring", n_ring);
+    unit_names.insert(unit_names.end(), ring_names.begin(), ring_names.end());
+    std::vector<V> unit_init(NACC, zero_acc);
+    unit_init.insert(unit_init.end(), n_ring, zero_frag);
+
+    std::vector<V> unit_out = e.for_iter(
+        zslot,
+        n_units_v,
+        n_z_v,
+        unit_names,
+        unit_init,
+        "unit",
+        [&](V unit, std::vector<V> carried) {
+            std::vector<V> accs;
+            Ring ring;
+            unpack(carried, &accs, &ring);
+            V n_img = e.mod(unit, p_N);
+            V chunk = e.div(unit, p_N);
+            V y0 = e.muli(chunk, spec->rows_per_chunk);
+            V k_rpc = e.ci(spec->rows_per_chunk);
+            V n_rows = e.smin(k_rpc, e.sub(p_H, y0));
+            V n_steps = e.addi(n_rows, PRE);
+            V x_img = e.mul(n_img, sx_n);
+            V y_img = e.mul(n_img, sdy_n);
+            const int PER_STEP = IX + IY;
+
+            /* DMA step t: X row y0 - PRE + t (steps >= PRE), dY row y + PAD. */
+            auto issue = [&](int buf, V t) {
+                V y = e.add(e.subi(y0, PRE), t);
+                V x_ok = e.gei(t, PRE);
+                x_ok = e.land(x_ok, e.lt(t, n_steps));
+                V x_row = e.add(x_img, e.mul(y, sx_h));
+                for(int i = 0; i < (int)x_offs.size(); ++i)
+                {
+                    V ok = e.land(x_offs[i].ok, x_ok);
+                    V voff = e.mul(e.add(x_offs[i].off, x_row), c2);
+                    voff = e.select(ok, voff, oob);
+                    dma(b_rsrc, x_addr[buf], IX, i, voff);
+                }
+                V pr = e.addi(y, PAD);
+                V y_ok = e.ge(pr, c0);
+                y_ok = e.land(y_ok, e.lt(pr, p_Ho));
+                y_ok = e.land(y_ok, e.lt(t, n_steps));
+                V y_row = e.add(y_img, e.mul(pr, sdy_h));
+                for(int i = 0; i < (int)y_offs.size(); ++i)
+                {
+                    V ok = e.land(y_offs[i].ok, y_ok);
+                    V voff = e.mul(e.add(y_offs[i].off, y_row), c2);
+                    voff = e.select(ok, voff, oob);
+                    dma(a_rsrc, y_addr[buf], IY, i, voff);
+                }
+            };
+
+            /* Barrier, DMA step t + D - 1, read and MFMA step t. */
+            auto step = [&](V t, int buf, std::vector<V>& accs_s, Ring& ring_s, bool do_compute) {
+                rocke_b_s_waitcnt(b, (D - 2) * PER_STEP, 0, -1);
+                rocke_b_s_barrier_bare(b);
+                issue((buf + D - 1) % D, e.addi(t, D - 1));
+                std::vector<std::vector<V>> newest;
+                for(int q = 0; q < Q4; ++q)
+                {
+                    std::vector<V> h;
+                    h.push_back(read(y_stage[buf], tr_base, 4 * q));
+                    if(HV == 2)
+                    {
+                        h.push_back(read(y_stage[buf], tr_swap, 4 * q));
+                    }
+                    newest.push_back(h);
+                }
+                Ring full;
+                full.push_back(newest);
+                full.insert(full.end(), ring_s.begin(), ring_s.end());
+                if(do_compute)
+                {
+                    for(int s = 0; s < KW; ++s)
+                    {
+                        for(int q = 0; q < Q4; ++q)
+                        {
+                            V a = read(x_stage[buf], tr_base, 4 * q + s);
+                            for(int r = 0; r < KH; ++r)
+                            {
+                                for(int h = 0; h < HV; ++h)
+                                {
+                                    int i = acc_i(r, s, h);
+                                    accs_s[i] = bf16 ? rocke_b_mfma_f32_4x4x4_bf16(
+                                                           b, a, full[r][q][h], accs_s[i])
+                                                     : rocke_b_mfma_f32_4x4x4_f16(
+                                                           b, a, full[r][q][h], accs_s[i]);
+                                }
+                            }
+                        }
+                    }
+                }
+                ring_s.assign(full.begin(), full.begin() + (KH - 1));
+            };
+
+            for(int t = 0; t < D - 1; ++t)
+            {
+                issue(t, e.ci(t));
+            }
+            /* Prologue: the PRE dY rows the first input row pairs with. */
+            for(int t = 0; t < PRE; ++t)
+            {
+                step(e.ci(t), t % D, accs, ring, false);
+            }
+            /* Main: one input row per step, unrolled by D. */
+            V trips = e.cdivi(n_rows, D);
+            trips = e.smax(trips, c1);
+            std::vector<V> init = pack(accs, ring);
+            std::vector<V> res = e.for_iter(c0,
+                                            trips,
+                                            c1,
+                                            iter_names("a", init.size()),
+                                            init,
+                                            "rb",
+                                            [&](V rb, std::vector<V> la) {
+                                                std::vector<V> l_acc;
+                                                Ring l_ring;
+                                                unpack(la, &l_acc, &l_ring);
+                                                for(int u = 0; u < D; ++u)
+                                                {
+                                                    V t = e.muli(rb, D);
+                                                    t = e.addi(t, PRE);
+                                                    t = e.addi(t, u);
+                                                    step(t, (PRE + u) % D, l_acc, l_ring, true);
+                                                }
+                                                return pack(l_acc, l_ring);
+                                            });
+            /* Drain the tail DMAs before the next unit reuses the buffers. */
+            rocke_b_s_waitcnt(b, 0, 0, -1);
+            rocke_b_s_barrier_bare(b);
+            return res;
+        });
+    if((int)unit_out.size() != NACC + n_ring)
+    {
+        return NULL;
+    }
+    std::vector<V> accs(unit_out.begin(), unit_out.begin() + NACC);
+
+    /* ---- epilogue: sum the waves in LDS, store ws[part][tap][o] ---- */
+    V red = e.smem_alloc_1d(rocke_f32(), WQ * kWave * 4, "red");
+    V part = e.add(e.mul(zslot, n_q_v), qtile);
+    V lane4 = e.modi(lane, 4);
+    V quad = e.muli(e.divi(lane, 4), 4);
+    for(int i_acc = 0; i_acc < NACC; ++i_acc)
+    {
+        int r = i_acc / (KW * HV), s = (i_acc / HV) % KW, h = i_acc % HV;
+        e.smem_store_f32(red, c0, e.muli(tid, 4), accs[i_acc], 4);
+        rocke_b_sync(b);
+        e.if_(e.lti(tid, kWave), [&]() {
+            V tot = e.smem_load_f32(red, c0, e.muli(lane, 4), 4);
+            for(int w = 1; w < WQ; ++w)
+            {
+                V idx = e.addi(lane, w * kWave);
+                idx = e.muli(idx, 4);
+                V v = e.smem_load_f32(red, c0, idx, 4);
+                tot = rocke_b_vector_add(b, tot, v);
+            }
+            V o_loc = e.bxor(quad, e.ci(4 * h));
+            o_loc = e.add(o_loc, lane4);
+            V o_glb = e.add(c_blk0, o_loc);
+            /* Channel c of the group: (4 * (L // 4) + i) % cpg. */
+            e.if_(e.lt(o_glb, p_K), [&]() {
+                for(int el = 0; el < 4; ++el)
+                {
+                    V c_in = e.modi(e.addi(quad, el), CPG);
+                    V tap = e.ci((r * KW + s) * CPG);
+                    tap = e.add(tap, c_in);
+                    V off = e.muli(part, TAPS);
+                    off = e.add(off, tap);
+                    off = e.mul(off, p_K);
+                    off = e.add(off, o_glb);
+                    rocke_b_global_store(b, P.D, off, rocke_b_vec_extract(b, tot, el), 4);
+                }
+            });
+        });
+        rocke_b_sync(b);
+    }
+    return rocke_ir_builder_ok(b) ? b->kernel : NULL;
+}
+
+} // namespace
+
+extern "C" rocke_kernel_def_t* rocke_build_grouped_wgrad_mfma(
+    rocke_ir_builder_t* b, const rocke_grouped_wgrad_mfma_spec_t* spec, const char* arch)
+{
+    if(b == NULL || spec == NULL)
+    {
+        return NULL;
+    }
+    return build_grouped(b, spec, arch ? arch : "gfx950");
+}
+
+extern "C" rocke_kernel_def_t* rocke_build_grouped_wgrad_mfma_new(
+    rocke_ir_builder_t* b, const rocke_grouped_wgrad_mfma_spec_t* spec, const char* arch)
+{
+    return build_new(
+        b,
+        spec != NULL,
+        "build_grouped_wgrad_mfma",
+        [&](char* out, size_t cap) { return rocke_grouped_wgrad_mfma_kernel_name(spec, out, cap); },
+        [&]() { return rocke_build_grouped_wgrad_mfma(b, spec, arch); });
+}
+
+extern "C" rocke_status_t
+    rocke_grouped_wgrad_mfma_lower_to_llvm(const rocke_grouped_wgrad_mfma_spec_t* spec,
+                                           const char* arch,
+                                           rocke_llvm_flavor_t flavor,
+                                           char** out_ll,
+                                           char* err,
+                                           size_t err_cap)
+{
+    if(arch == NULL)
+    {
+        arch = "gfx950";
+    }
+    return lower_to_llvm(
+        spec != NULL,
+        arch,
+        flavor,
+        out_ll,
+        err,
+        err_cap,
+        "build_grouped_wgrad_mfma failed",
+        [&](rocke_ir_builder_t* b) { return rocke_build_grouped_wgrad_mfma_new(b, spec, arch); });
+}
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/architecture/kernel_taxonomy.md b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/architecture/kernel_taxonomy.md
index feebce4ae0f8..d8ef89b34128 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/architecture/kernel_taxonomy.md
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/architecture/kernel_taxonomy.md
@@ -91,6 +91,8 @@ and exact gfx targets.
 | `build_direct_wgrad_tiled` | [L/common/conv_direct_wgrad_tiled.py](../../../library/kernels/common/conv_direct_wgrad_tiled.py) | direct wgrad, LDS-DMA row pipeline, fp32 atomics | `mfma_f32_16x16x32_{f16,bf16}` | owning validator |
 | `build_dw_wgrad_lds`, `build_dw_wgrad_cols`, `build_dw_wgrad_fold` | [L/common/conv_direct_dw_wgrad.py](../../../library/kernels/common/conv_direct_dw_wgrad.py) | depthwise wgrad into partial slots, and their fold | VALU FMA | owning validator |
 | `build_wgrad_gather` | [L/common/conv_direct_wgrad_gather.py](../../../library/kernels/common/conv_direct_wgrad_gather.py) | dense small-C wgrad into partial slots | `mfma_f32_16x16x32_{f16,bf16}` | owning validator |
+| `build_dw_wgrad_mfma` | [L/common/conv_direct_dw_wgrad_mfma.py](../../../library/kernels/common/conv_direct_dw_wgrad_mfma.py) | depthwise wgrad into partial slots, one channel per MFMA batch | direct `mfma_f32_4x4x4_{f16,bf16}` intrinsic outside `MmaCatalog` | owning validator |
+| `build_grouped_wgrad_mfma` | [L/common/conv_direct_grouped_wgrad_mfma.py](../../../library/kernels/common/conv_direct_grouped_wgrad_mfma.py) | 4/8-channel grouped wgrad into partial slots, one group per MFMA batch | direct `mfma_f32_4x4x4_{f16,bf16}` intrinsic outside `MmaCatalog` | owning validator |
 | `build_deep_fused_conv_pool` | [L/common/deep_fused_conv_pool.py](../../../library/kernels/common/deep_fused_conv_pool.py), [L/gfx1151/deep_fused_conv_pool.py](../../../library/kernels/gfx1151/deep_fused_conv_pool.py), [L/gfx1201/deep_fused_conv_pool.py](../../../library/kernels/gfx1201/deep_fused_conv_pool.py), [L/gfx950/deep_fused_conv_pool.py](../../../library/kernels/gfx950/deep_fused_conv_pool.py) | convolution plus pooling | target-specific matrix MMA plus VALU/window reduction | exact implementation selected by target |
 | `build_moe_gate_up_silu_gemm`, `build_moe_interleaved_gate_up_silu_gemm`, `build_moe_down_reduce_gemm`, `build_moe_down_silu_reduce_gemm` | [P/common/moe_gemm_fused.py](../../python/rocke/instances/common/moe_gemm_fused.py) | expert GEMM plus activation or reduction | catalog-selected matrix MMA plus fused VALU/reduction | owning validator |
 | `build_moe_fused_mega_gemm`, `build_moe_fused_mega_gemm_fp8`, `build_moe_fused_mega_wmma` | [P/common/moe_fused_mega.py](../../python/rocke/instances/common/moe_fused_mega.py), [P/common/moe_fused_mega_fp8.py](../../python/rocke/instances/common/moe_fused_mega_fp8.py), [P/gfx1250/fused_moe_mega_wmma.py](../../python/rocke/instances/gfx1250/fused_moe_mega_wmma.py) | persistent expert GEMM pipeline | target-selected matrix MMA plus fused routing/activation | owning validator or exact-target module |
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/instances/convolution.md b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/instances/convolution.md
index 15e8df4b1c3f..cc5f78730f11 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/instances/convolution.md
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/instances/convolution.md
@@ -7,7 +7,9 @@ This page covers:
 - `library/kernels/common/conv_direct_nongrouped.py`
 - `library/kernels/common/conv_direct_wgrad_tiled.py`,
   `library/kernels/common/conv_direct_dw_wgrad.py`,
-  `library/kernels/common/conv_direct_wgrad_gather.py` (direct wgrad)
+  `library/kernels/common/conv_direct_wgrad_gather.py`,
+  `library/kernels/common/conv_direct_dw_wgrad_mfma.py`,
+  `library/kernels/common/conv_direct_grouped_wgrad_mfma.py` (direct wgrad)
 - `library/kernels/common/img2col.py`
 - `instances/common/pooling.py`
 
@@ -774,7 +776,8 @@ C++ engine twin: `rocke_build_direct_conv_nongrouped` in
 ## Direct Wgrad Kernels
 
 Sources: `library/kernels/common/conv_direct_wgrad_tiled.py`,
-`conv_direct_dw_wgrad.py`, `conv_direct_wgrad_gather.py`. Weight-gradient
+`conv_direct_dw_wgrad.py`, `conv_direct_wgrad_gather.py`,
+`conv_direct_dw_wgrad_mfma.py`, `conv_direct_grouped_wgrad_mfma.py`. Weight-gradient
 kernels of the direct-conv AOT sweep (`--directions wgrad`), gfx950; all take
 the `conv_direct_arg_names("wgrad")` kernargs (dY on A, X on B) and accumulate
 in fp32.
@@ -785,6 +788,8 @@ in fp32.
 | `DwWgradLdsSpec` | depthwise, `C % 8 == 0`, 3x3, stride 1/2 | partial slots `ws[part][tap][c]` |
 | `DwWgradColsSpec` | depthwise, `C <= 8` (stem layers), any filter | partial slots `ws[part][tap][c]` |
 | `WgradGatherSpec` | dense, small `C` (`KW * C <= 96`), any stride | partial slots `ws[part][(r, s, c)][k]` |
+| `DwWgradMfmaSpec` | depthwise, `C % 8 == 0`, any filter, stride 1 | partial slots `ws[part][tap][c]` |
+| `GroupedWgradMfmaSpec` | grouped, `cpg = kpg` in (4, 8), stride 1 | partial slots `ws[part][(r, s, c)][k]` |
 | `DwWgradFoldSpec` | the fold after a partial-slot pass | `dW` written whole (no zeroing) |
 
 - **Tiled.** A wave owns `wave_k16*16 x wave_c16*16` channels for all taps;
@@ -799,9 +804,26 @@ in fp32.
   `DwWgradFoldSpec` sums the partitions. In the sweep they are two-step plans:
   the pass writes the scratch `ws`, the fold reads it (kernargs `A, D, A_bytes,
   D_bytes, parts, channels`).
-- All four row-streaming kernels set the kernel attr `lds_alias_scopes`: every
-  smem allocation gets its own alias scope, so the waitcnt pass does not drain
-  the in-flight LDS DMA before each `ds_read`.
+- **Batched 4x4x4 MFMA.** `mfma_f32_4x4x4_{f16,bf16}` runs 16 independent
+  4x4x4 products per wave. `DwWgradMfmaSpec` gives each batch one channel: per
+  input row and 4-column group, `A[s][k] = X[y][q0 + k + s - PAD]` (a Hankel
+  window: the four lanes of a channel read a quarter each of its channel-major
+  LDS row, share it with `quad_perm`, and build every A fragment with two
+  `v_perm_b32`, the lane's tap shift being its byte selector) times
+  `B[k][r] = dY[y + PAD - r][q0 + k]` from a register ring rotated one row per
+  step. Rows arrive by LDS DMA in NHWC order and are transposed channel-major
+  with `ds_read_b64_tr_b16`. `unfold_n` packs small images side by side in one
+  column tile (pitch `max(Wo, Wo + KW - 1 - PAD, PAD + W)`, shared zero
+  padding); when a unit is a whole image the prologue stages only PAD dY rows.
+  `GroupedWgradMfmaSpec` gives each batch one 4-channel group and reads both
+  operands straight from the staged NHWC rows (the tap shift is a row offset of
+  the transpose-read addresses); `cpg = 8` adds a dY read with the group's
+  halves swapped for the off-diagonal blocks.
+- All the row-streaming kernels set the kernel attr `lds_alias_scopes`: every
+  smem allocation gets its own alias scope, and the LDS DMA, the transpose
+  reads and the vector LDS loads/stores (`smem_load_vN` / `smem_store_vN`) are
+  tagged with it, so the waitcnt pass does not drain the in-flight LDS DMA
+  before each LDS access.
 - The validators gate on `lds_bytes()`, which counts what the builder
   allocates: ring slots rounded up to whole DMA instructions of every wave
   (plus the gather's 16-element transpose tail) and the reduction buffers.
@@ -811,8 +833,10 @@ in fp32.
 C++ engine twins: `rocke_build_direct_wgrad_tiled` (`cpp/instances/common/conv_direct_wgrad_tiled.cpp`,
 also `rocke_direct_wgrad_pick_swizzle`), `rocke_build_dw_wgrad_{fold,cols,lds}`
 (`conv_direct_dw_wgrad.cpp`), `rocke_build_direct_wgrad_gather`
-(`conv_direct_wgrad_gather.cpp`); byte-identity gated by
-`library/tests/parity/conv_direct_grouped_emit.{c,py}` (configs 95+), Python IR
+(`conv_direct_wgrad_gather.cpp`), `rocke_build_dw_wgrad_mfma`
+(`conv_direct_dw_wgrad_mfma.cpp`), `rocke_build_grouped_wgrad_mfma`
+(`conv_direct_grouped_wgrad_mfma.cpp`); byte-identity gated by
+`library/tests/parity/conv_direct_grouped_emit.{c,py}` (configs 114+), Python IR
 golden `conv_direct_wgrad/*` in `tests/instances/rocke_ir_parity_harness.py`.
 On-GPU numeric coverage: `TestDirectWgradAotKernelsCorrectness` in
 `library/tests/test_direct_conv_correctness.py` (built on the cache probe,
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/instances/index.md b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/instances/index.md
index a79e0aa975ad..f90409d0c483 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/instances/index.md
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/instances/index.md
@@ -46,6 +46,8 @@ implementation.
 | `conv_direct_wgrad_tiled.py` | `DirectWgradTiledSpec` (direct wgrad, atomics) | `instances/convolution.md` |
 | `conv_direct_dw_wgrad.py` | `DwWgradLdsSpec`, `DwWgradColsSpec`, `DwWgradFoldSpec` (depthwise wgrad) | `instances/convolution.md` |
 | `conv_direct_wgrad_gather.py` | `WgradGatherSpec` (dense small-C wgrad) | `instances/convolution.md` |
+| `conv_direct_dw_wgrad_mfma.py` | `DwWgradMfmaSpec` (depthwise wgrad, 4x4x4 MFMA) | `instances/convolution.md` |
+| `conv_direct_grouped_wgrad_mfma.py` | `GroupedWgradMfmaSpec` (4/8-channel grouped wgrad, 4x4x4 MFMA) | `instances/convolution.md` |
 | `img2col.py` | `Img2ColSpec` | `instances/convolution.md` |
 | `pooling.py` | `PoolingProblem`, `Pooling2DSpec`, `PoolOp` | `instances/convolution.md` |
 
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/ir_lowering/ir_model.md b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/ir_lowering/ir_model.md
index e7ee4a13e269..4d612534b07e 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/ir_lowering/ir_model.md
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/ir_lowering/ir_model.md
@@ -237,6 +237,7 @@ mfma_f32_16x16x32_bf16(a, b, c)    # <8 x bfloat>, <8 x bfloat>, <4 x float>
 mfma_f32_32x32x8_f16(a, b, c)      # <4 x half>, <4 x half>, <16 x float>
 mfma_f32_32x32x16_f16(a, b, c)     # <8 x half>, <8 x half>, <16 x float>  (gfx950 only)
 mfma_f32_4x4x4_f16(a, b, c)        # 16 independent 4x4 matmuls per wave
+mfma_f32_4x4x4_bf16(a, b, c)       # same, via the *_1k variant with <4 x i16> operands
 ```
 
 The shipped `MFMA_F16_ATOMS` catalog in `helpers/atoms.py` exposes the five f16 variants with full lane mapping and dispatch. See `reference/mfma_atom_catalog.md`. These MFMA atoms target the CDNA (wave64) backends. For RDNA (wave32) targets the builder also exposes WMMA matrix ops (`wmma_f32_16x16x16_f16` / `wmma_f32_16x16x16_bf16`, and the gfx12 variants) backed by `WmmaAtom` in `helpers/atoms.py`; all matrix ops route through the generic `mma()` builder method.
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/reference/op_vocabulary.md b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/reference/op_vocabulary.md
index 167a2401739b..f7b1952ae960 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/reference/op_vocabulary.md
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/reference/op_vocabulary.md
@@ -142,6 +142,7 @@ Complete reference of operations recognized by `core/ir.py` and lowered to AMDGP
 | `tile.mfma_f32_32x32x8_f16`      | `<4xhalf>`, `<4xhalf>`, `<16xfloat>` | `<16xfloat>`   |
 | `tile.mfma_f32_32x32x16_f16`     | `<8xhalf>`, `<8xhalf>`, `<16xfloat>` | `<16xfloat>`   |
 | `tile.mfma_f32_4x4x4_f16`        | `<4xhalf>`, `<4xhalf>`, `<4xfloat>` | `<4xfloat>`     |
+| `tile.mfma_f32_4x4x4_bf16`       | `<4xbfloat>`, same, `<4xfloat>`  | `<4xfloat>`        |
 
 All MFMAs take three literal immarg constants (`cbsz=0, abid=0, blgp=0`) at the LLVM level.
 
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/arch/target.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/arch/target.py
index eaeb255e1f4c..6cd949166ae8 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/arch/target.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/arch/target.py
@@ -925,6 +925,7 @@ class _FragInfo:
         8, 8, 16, 64, _mfma_a_32x32x16, _mfma_b_32x32x16, _mfma_acc_32x32
     ),
     "mfma_f32_4x4x4_f16": _FragInfo(4, 4, 4, 64),
+    "mfma_f32_4x4x4_bf16": _FragInfo(4, 4, 4, 64),
     # --- MFMA bf16 (wave64) ----------------------------------------------
     "mfma_f32_16x16x16_bf16": _FragInfo(
         4, 4, 4, 64, _mfma_a_16x16, _mfma_b_16x16, _mfma_acc_16x16
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/ir.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/ir.py
index cc2e3dd377ff..63dba600af11 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/ir.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/ir.py
@@ -3302,6 +3302,16 @@ def mfma_f32_4x4x4_f16(self, a: Value, b: Value, c: Value) -> Value:
         """
         return self.mma("mfma_f32_4x4x4_f16", a, b, c)
 
+    def mfma_f32_4x4x4_bf16(self, a: Value, b: Value, c: Value) -> Value:
+        """The bf16 twin of :meth:`mfma_f32_4x4x4_f16` (CDNA2+).
+
+        Per lane: A `<4 x bfloat>`, B `<4 x bfloat>`, acc `<4 x float>`; 16
+        independent 4x4x4 products per wave, batch ``lane / 4``. Lowers to
+        `@llvm.amdgcn.mfma.f32.4x4x4bf16.1k` with the operands bitcast to
+        `<4 x i16>`, like the other bf16 `_1k` atoms.
+        """
+        return self.mma("mfma_f32_4x4x4_bf16", a, b, c)
+
     # ----- vector type casts (for packed buffer-load + LDS reads) -----
 
     def vec_bitcast(self, v: Value, target: Type) -> Value:
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py
index bd0279deff7a..bf6b23948247 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py
@@ -831,6 +831,13 @@ def _op_tile_mfma_f32_4x4x4_f16(self, op: Op) -> None:
             f"{_name(a)}, {_name(b)}, {_name(c)}, 0, 0, 0);"
         )
 
+    def _op_tile_mfma_f32_4x4x4_bf16(self, op: Op) -> None:
+        a, b, c = op.operands
+        self._emit(
+            f"f32x4 {_name(op.result)} = __builtin_amdgcn_mfma_f32_4x4x4bf16_1k("
+            f"{_name(a)}, {_name(b)}, {_name(c)}, 0, 0, 0);"
+        )
+
     # ---- WMMA f16 (RDNA3/3.5, gfx11, wave32) ----
     #
     # The RDNA half of the neutral-MMA contract. A WMMA ``op_id`` reaches this
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py
index 6d9deb4c9b42..f57fd24e4198 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py
@@ -692,6 +692,11 @@ def _resolve_llvm_flavor() -> str:
         "<4 x half>, <4 x half>, <4 x float>, "
         "i32 immarg, i32 immarg, i32 immarg)"
     ),
+    "mfma.f32.4x4x4bf16.1k": (
+        "declare <4 x float> @llvm.amdgcn.mfma.f32.4x4x4bf16.1k("
+        "<4 x i16>, <4 x i16>, <4 x float>, "
+        "i32 immarg, i32 immarg, i32 immarg)"
+    ),
     # FP8 / BF8 MFMA (gfx940+). Operands are packed as <2 x i32>
     # at the intrinsic boundary (= 8 fp8 / bf8 bytes per lane); the
     # lowering helper bitcasts the <8 x fp8e4m3> / <8 x bf8e5m2> IR
@@ -4594,6 +4599,25 @@ def _op_tile_mfma_f32_4x4x4_f16(self, op: Op) -> None:
             f"i32 0, i32 0, i32 0)"
         )
 
+    def _op_tile_mfma_f32_4x4x4_bf16(self, op: Op) -> None:
+        a, b, c = op.operands
+        self._need("mfma.f32.4x4x4bf16.1k")
+        a_cast = self._fresh("mfma_a_i16")
+        b_cast = self._fresh("mfma_b_i16")
+        self._current().emit(
+            f"  {a_cast} = bitcast <4 x bfloat> {self._operand(a)} to <4 x i16>"
+        )
+        self._current().emit(
+            f"  {b_cast} = bitcast <4 x bfloat> {self._operand(b)} to <4 x i16>"
+        )
+        self._current().emit(
+            f"  {op.result.name} = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x4bf16.1k("
+            f"<4 x i16> {a_cast}, "
+            f"<4 x i16> {b_cast}, "
+            f"<4 x float> {self._operand(c)}, "
+            f"i32 0, i32 0, i32 0)"
+        )
+
     def _op_vector_bitcast(self, op: Op) -> None:
         (v,) = op.operands
         target = op.result.type
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/SUPPORT_MATRIX.md b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/SUPPORT_MATRIX.md
index 32acf749a114..1d8fd950285e 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/SUPPORT_MATRIX.md
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/SUPPORT_MATRIX.md
@@ -101,6 +101,7 @@ and [packed FP6 input contract](../examples/gfx1250/gemm/FP6.md).
 | `conv_direct_nongrouped` | ✅ | ✅ | ❌ | `groups=1`; gfx942 uses 32x32x8 / 16x16x16 atoms, gfx950 also 32x32x16 / 16x16x32; MFMA-only; the only direct kernel with dilation. Built from `library/kernels`; not yet wired into dispatch |
 | `direct_wgrad_tiled` | ❌ | ✅ | ❌ | dense and grouped wgrad, stride 1; needs `ds_read_b64_tr_b16` and `mfma_f32_16x16x32`. Built from `library/kernels`; AOT sweep only, not yet wired into dispatch |
 | `dw_wgrad_lds` / `dw_wgrad_cols` / `direct_wgrad_gather` | ❌ | ✅ | ❌ | depthwise (`C % 8 == 0` / `C <= 8`) and dense small-C wgrad, partial slots + `dw_wgrad_fold`; validated on gfx950 only (the LDS kernels DMA 16 bytes per lane). AOT sweep only, not yet wired into dispatch |
+| `dw_wgrad_mfma` / `grouped_wgrad_mfma` | ❌ | ✅ | ❌ | depthwise (`C % 8 == 0`, any filter) and 4/8-channel grouped wgrad, stride 1, on `mfma_f32_4x4x4`; partial slots + `dw_wgrad_fold`; need `ds_read_b64_tr_b16`. AOT sweep only, not yet wired into dispatch |
 | `direct_depthwise_tiled` | ✅ | ✅ | ❌ | `cpg=kpg=1`, filters up to 31x31, any padding; scalar FMA, wave64. Built from `library/kernels`; not yet wired into dispatch |
 
 ---
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256.json b/dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256.json
index cd2dfb08b99a..b392e4efbf90 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256.json
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256.json
@@ -1178,6 +1178,22 @@
           "sha256": "3ef756d5ce85d6a9761e59b04681a1a01229ae78a5a9b53ef8e77cb838895611",
           "status": "ok"
         },
+        "conv_direct_wgrad/gfx950/dw_mfma_7x7_c192": {
+          "arch": "gfx950",
+          "bytes": 306323,
+          "family": "conv_direct_wgrad",
+          "kernel_name": "irhash_direct_wgrad_dw_mfma_7x7_c192_f7x7p3s1_qt8_rpc28_pf2_un1_tb2048_bf16",
+          "sha256": "90e4b1bd507127e491804489a5cf401f0f60242ffe592e4a8dcc3f31808b14f8",
+          "status": "ok"
+        },
+        "conv_direct_wgrad/gfx950/dw_mfma_7x7_unfold3": {
+          "arch": "gfx950",
+          "bytes": 312421,
+          "family": "conv_direct_wgrad",
+          "kernel_name": "irhash_direct_wgrad_dw_mfma_7x7_unfold3_f7x7p3s1_qt8_rpc14_pf2_un3_tb2048_bf16",
+          "sha256": "356d3b1384970b2b52311e4953c93e33cda1635b7f50d4d5a39de671481e77ed",
+          "status": "ok"
+        },
         "conv_direct_wgrad/gfx950/fold_t9": {
           "arch": "gfx950",
           "bytes": 11967,
@@ -1202,6 +1218,22 @@
           "sha256": "b6d37178d5e31f5c397aaed864dd61d12c9931872fe9b69f4034d7c442f23d61",
           "status": "ok"
         },
+        "conv_direct_wgrad/gfx950/grouped_mfma_c4": {
+          "arch": "gfx950",
+          "bytes": 84396,
+          "family": "conv_direct_wgrad",
+          "kernel_name": "irhash_direct_wgrad_grouped_mfma_c4_f3x3p1s1c4_wq4_q42_rpc56_pf2_tb1024_bf16",
+          "sha256": "37f6d5ba2f1f42075e925c944dcdd896effc6d7413de3121acab7e02efbcfa77",
+          "status": "ok"
+        },
+        "conv_direct_wgrad/gfx950/grouped_mfma_c8_fp16": {
+          "arch": "gfx950",
+          "bytes": 147462,
+          "family": "conv_direct_wgrad",
+          "kernel_name": "irhash_direct_wgrad_grouped_mfma_c8_fp16_f3x3p1s1c8_wq4_q42_rpc56_pf3_tb1024",
+          "sha256": "0ed2fe34f4ab088705940ca12bc4ae4b5030e6912bfafdc4b9baecdb01a7bf8d",
+          "status": "ok"
+        },
         "conv_direct_wgrad/gfx950/tiled_bf16_c128": {
           "arch": "gfx950",
           "bytes": 193857,
@@ -1965,12 +1997,12 @@
           "gfx1201": 8,
           "gfx1250": 4,
           "gfx90a": 5,
-          "gfx942": 42,
-          "gfx950": 159
+          "gfx942": 47,
+          "gfx950": 170
         },
         "expected_failures": 4,
         "families": {
-          "attention": 33,
+          "attention": 35,
           "attention_d128_swa": 1,
           "attention_d256": 2,
           "attention_dense": 44,
@@ -1980,7 +2012,7 @@
           "conv_direct_depthwise_tiled": 3,
           "conv_direct_dgrad": 5,
           "conv_direct_nongrouped": 7,
-          "conv_direct_wgrad": 8,
+          "conv_direct_wgrad": 12,
           "conv_wgrad": 21,
           "conv_wgrad_reduce": 3,
           "conv_wgrad_two_stage": 2,
@@ -1989,10 +2021,11 @@
           "gemm": 13,
           "grouped_gemm": 1,
           "kda_chunkwise": 13,
+          "lightning_indexer": 10,
           "moe": 9,
           "nontemporal": 2
         },
-        "ok_cases": 230
+        "ok_cases": 246
       }
     },
     "llvm22": {
@@ -3173,6 +3206,22 @@
           "sha256": "f420503d056420ab4ae36118f2fb032ffd725a6faa5985f4ac92b5bc9f46d385",
           "status": "ok"
         },
+        "conv_direct_wgrad/gfx950/dw_mfma_7x7_c192": {
+          "arch": "gfx950",
+          "bytes": 306419,
+          "family": "conv_direct_wgrad",
+          "kernel_name": "irhash_direct_wgrad_dw_mfma_7x7_c192_f7x7p3s1_qt8_rpc28_pf2_un1_tb2048_bf16",
+          "sha256": "49f46fc2d9dfea5d84e1043dbaede93adcf474419479284f346abb678e852356",
+          "status": "ok"
+        },
+        "conv_direct_wgrad/gfx950/dw_mfma_7x7_unfold3": {
+          "arch": "gfx950",
+          "bytes": 312517,
+          "family": "conv_direct_wgrad",
+          "kernel_name": "irhash_direct_wgrad_dw_mfma_7x7_unfold3_f7x7p3s1_qt8_rpc14_pf2_un3_tb2048_bf16",
+          "sha256": "8c189d730b4e243d2eb546516f94d49cac19e172dde997da1779ff43d2f7caaa",
+          "status": "ok"
+        },
         "conv_direct_wgrad/gfx950/fold_t9": {
           "arch": "gfx950",
           "bytes": 11974,
@@ -3197,6 +3246,22 @@
           "sha256": "0649e52910a0cae5a02cd8c9257f033d712aec4d9889067e9ab1f2c2dab08583",
           "status": "ok"
         },
+        "conv_direct_wgrad/gfx950/grouped_mfma_c4": {
+          "arch": "gfx950",
+          "bytes": 84492,
+          "family": "conv_direct_wgrad",
+          "kernel_name": "irhash_direct_wgrad_grouped_mfma_c4_f3x3p1s1c4_wq4_q42_rpc56_pf2_tb1024_bf16",
+          "sha256": "c38ff3d806f7efc1135d23ab5ff2279dfcec40ef21747b5a0dc6f06e14abb76d",
+          "status": "ok"
+        },
+        "conv_direct_wgrad/gfx950/grouped_mfma_c8_fp16": {
+          "arch": "gfx950",
+          "bytes": 147563,
+          "family": "conv_direct_wgrad",
+          "kernel_name": "irhash_direct_wgrad_grouped_mfma_c8_fp16_f3x3p1s1c8_wq4_q42_rpc56_pf3_tb1024",
+          "sha256": "3a831e710e5914181aa2cf7b571c948aedd79485e3ce7130c8562bb0cd26b801",
+          "status": "ok"
+        },
         "conv_direct_wgrad/gfx950/tiled_bf16_c128": {
           "arch": "gfx950",
           "bytes": 193953,
@@ -3960,12 +4025,12 @@
           "gfx1201": 8,
           "gfx1250": 4,
           "gfx90a": 5,
-          "gfx942": 42,
-          "gfx950": 159
+          "gfx942": 47,
+          "gfx950": 170
         },
         "expected_failures": 4,
         "families": {
-          "attention": 33,
+          "attention": 35,
           "attention_d128_swa": 1,
           "attention_d256": 2,
           "attention_dense": 44,
@@ -3975,7 +4040,7 @@
           "conv_direct_depthwise_tiled": 3,
           "conv_direct_dgrad": 5,
           "conv_direct_nongrouped": 7,
-          "conv_direct_wgrad": 8,
+          "conv_direct_wgrad": 12,
           "conv_wgrad": 21,
           "conv_wgrad_reduce": 3,
           "conv_wgrad_two_stage": 2,
@@ -3984,10 +4049,11 @@
           "gemm": 13,
           "grouped_gemm": 1,
           "kda_chunkwise": 13,
+          "lightning_indexer": 10,
           "moe": 9,
           "nontemporal": 2
         },
-        "ok_cases": 230
+        "ok_cases": 246
       }
     },
     "llvm23": {
@@ -5168,6 +5234,22 @@
           "sha256": "805e6b9db3d4aca2c6a1656159dcb50a35a35ffe09bbc32bb1d04474047355f7",
           "status": "ok"
         },
+        "conv_direct_wgrad/gfx950/dw_mfma_7x7_c192": {
+          "arch": "gfx950",
+          "bytes": 306483,
+          "family": "conv_direct_wgrad",
+          "kernel_name": "irhash_direct_wgrad_dw_mfma_7x7_c192_f7x7p3s1_qt8_rpc28_pf2_un1_tb2048_bf16",
+          "sha256": "f726f945fb7b84e6413e6819528db954ca41fd9e25e24dd6b9ffce9c871e77c7",
+          "status": "ok"
+        },
+        "conv_direct_wgrad/gfx950/dw_mfma_7x7_unfold3": {
+          "arch": "gfx950",
+          "bytes": 312581,
+          "family": "conv_direct_wgrad",
+          "kernel_name": "irhash_direct_wgrad_dw_mfma_7x7_unfold3_f7x7p3s1_qt8_rpc14_pf2_un3_tb2048_bf16",
+          "sha256": "1ab085a89838c6fb4b43c2d484aad358309bf8a903b18c755e4b1dde25003d94",
+          "status": "ok"
+        },
         "conv_direct_wgrad/gfx950/fold_t9": {
           "arch": "gfx950",
           "bytes": 12038,
@@ -5192,6 +5274,22 @@
           "sha256": "44512065b1f099343444bc836b155b65df48f5cf1f7a25c529afe8faaf64ee39",
           "status": "ok"
         },
+        "conv_direct_wgrad/gfx950/grouped_mfma_c4": {
+          "arch": "gfx950",
+          "bytes": 84556,
+          "family": "conv_direct_wgrad",
+          "kernel_name": "irhash_direct_wgrad_grouped_mfma_c4_f3x3p1s1c4_wq4_q42_rpc56_pf2_tb1024_bf16",
+          "sha256": "66c40c8f3515e0c768951e7b1bbc82daf47ff2d932c16ec4cf23115717f05029",
+          "status": "ok"
+        },
+        "conv_direct_wgrad/gfx950/grouped_mfma_c8_fp16": {
+          "arch": "gfx950",
+          "bytes": 147627,
+          "family": "conv_direct_wgrad",
+          "kernel_name": "irhash_direct_wgrad_grouped_mfma_c8_fp16_f3x3p1s1c8_wq4_q42_rpc56_pf3_tb1024",
+          "sha256": "b7f365f4cfe09352e15c8a614a4448cfb890fec9bfe2e859547de90d2c97bf4a",
+          "status": "ok"
+        },
         "conv_direct_wgrad/gfx950/tiled_bf16_c128": {
           "arch": "gfx950",
           "bytes": 194017,
@@ -5970,12 +6068,12 @@
           "gfx1201": 8,
           "gfx1250": 8,
           "gfx90a": 5,
-          "gfx942": 42,
-          "gfx950": 159
+          "gfx942": 47,
+          "gfx950": 170
         },
         "expected_failures": 0,
         "families": {
-          "attention": 33,
+          "attention": 35,
           "attention_d128_swa": 1,
           "attention_d256": 2,
           "attention_dense": 44,
@@ -5985,7 +6083,7 @@
           "conv_direct_depthwise_tiled": 3,
           "conv_direct_dgrad": 5,
           "conv_direct_nongrouped": 7,
-          "conv_direct_wgrad": 8,
+          "conv_direct_wgrad": 12,
           "conv_wgrad": 21,
           "conv_wgrad_reduce": 3,
           "conv_wgrad_two_stage": 2,
@@ -5994,10 +6092,11 @@
           "gemm": 17,
           "grouped_gemm": 1,
           "kda_chunkwise": 13,
+          "lightning_indexer": 10,
           "moe": 9,
           "nontemporal": 2
         },
-        "ok_cases": 234
+        "ok_cases": 250
       }
     }
   },
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/rocke_ir_parity_harness.py b/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/rocke_ir_parity_harness.py
index c6c5ac94bbfc..a6b83d0895f0 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/rocke_ir_parity_harness.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/rocke_ir_parity_harness.py
@@ -1209,12 +1209,14 @@ def _build():
 
 def build_direct_wgrad(name, arch, kind, problem, **spec_kw):
     """A direct wgrad kernel: ``kind`` is tiled / dw_lds / dw_cols / gather /
-    fold; ``problem`` overrides the DirectConvProblem fields (unused by the
-    fold). Configs mirror parity/conv_direct_grouped_emit.* (indices 95+),
+    dw_mfma / grouped_mfma / fold; ``problem`` overrides the DirectConvProblem
+    fields (unused by the fold). Configs mirror parity/conv_direct_grouped_emit.* (indices 95+),
     where the C++ engine is gated byte-identical by check_byte_identity.py."""
 
     def _build():
         from kernels.common import conv_direct_dw_wgrad as dw
+        from kernels.common import conv_direct_dw_wgrad_mfma as dwm
+        from kernels.common import conv_direct_grouped_wgrad_mfma as gwm
         from kernels.common import conv_direct_wgrad_gather as wg
         from kernels.common import conv_direct_wgrad_tiled as wt
         from kernels.common.conv_direct_grouped import DirectConvProblem
@@ -1243,6 +1245,11 @@ def _build():
             "dw_lds": (dw.DwWgradLdsSpec, dw.build_dw_wgrad_lds),
             "dw_cols": (dw.DwWgradColsSpec, dw.build_dw_wgrad_cols),
             "gather": (wg.WgradGatherSpec, wg.build_wgrad_gather),
+            "dw_mfma": (dwm.DwWgradMfmaSpec, dwm.build_dw_wgrad_mfma),
+            "grouped_mfma": (
+                gwm.GroupedWgradMfmaSpec,
+                gwm.build_grouped_wgrad_mfma,
+            ),
         }[kind]
         return build(spec_cls(problem=p, name=name, **spec_kw), arch=arch)
 
@@ -4118,7 +4125,8 @@ def add(family, case_id, arch, build):
 
     # --- conv_direct_wgrad: the direct wgrad kernels on gfx950 ---
     # Tiled (LDS-DMA row pipeline, swizzled transpose reads, atomics), the
-    # depthwise partial-slot kernels and their fold, and the small-C gather.
+    # depthwise partial-slot kernels and their fold, the small-C gather, and
+    # the 4x4x4 MFMA depthwise and 4/8-channel grouped kernels.
     # The full branch matrix (and the C++ twins) lives in
     # library/tests/parity/conv_direct_grouped_emit.* (indices 95+).
     for _case_id, _kind, _prob, _over in (
@@ -4162,6 +4170,30 @@ def add(family, case_id, arch, build):
             dict(cpg=24, kpg=96, stride=2, H=30, W=40),
             dict(wave_k16=1, waves_k=6, rows_per_seg=8, target_blocks=256),
         ),
+        (
+            "dw_mfma_7x7_c192",
+            "dw_mfma",
+            dict(groups=192, cpg=1, kpg=1, KH=7, KW=7, PAD=3, H=56, W=56),
+            dict(q_tiles=8, rows_per_chunk=28),
+        ),
+        (
+            "dw_mfma_7x7_unfold3",
+            "dw_mfma",
+            dict(groups=1024, cpg=1, kpg=1, KH=7, KW=7, PAD=3, N=8, H=7, W=7),
+            dict(rows_per_chunk=14, unfold_n=3),
+        ),
+        (
+            "grouped_mfma_c4",
+            "grouped_mfma",
+            dict(groups=32, cpg=4, kpg=4, H=56, W=56),
+            dict(wave_q4=2),
+        ),
+        (
+            "grouped_mfma_c8_fp16",
+            "grouped_mfma",
+            dict(groups=32, cpg=8, kpg=8, H=56, W=56, dtype="fp16"),
+            dict(wave_q4=2, prefetch=3),
+        ),
     ):
         add(
             "conv_direct_wgrad",

From d3d9ee71f2350aa1383a11861647548009afb199 Mon Sep 17 00:00:00 2001
From: Bartlomiej Kocot <barkocot@amd.com>
Date: Thu, 22 Oct 2026 17:57:16 -0500
Subject: [PATCH 2/6] fix(rocke): reject nonpositive target_blocks in 4x4x4
 wgrad; C++ HIP 4x4x4 atoms

The depthwise and grouped batched-4x4x4 wgrad validators accepted a
target_blocks of 0 or -1, which n_z() and the emitted smax then turned into
one z partition. Both engines now reject it as the other direct-wgrad
validators do; parity configs 159..162 pin the rejection in both engines.

The C++ HIP lowerer had no handler for mfma_f32_4x4x4_{f16,bf16}: tile.mma
fell through to the legacy dispatch and failed with "no HIP lowering". Emit
the same builtin calls as the Python handlers, with a both-engines HIP test.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
---
 .../common/conv_direct_dw_wgrad_mfma.py       |  2 ++
 .../common/conv_direct_grouped_wgrad_mfma.py  |  2 ++
 .../tests/parity/conv_direct_grouped_emit.c   | 22 +++++++++++-
 .../tests/parity/conv_direct_grouped_emit.py  | 17 +++++++--
 .../library/tests/test_direct_wgrad_specs.py  | 24 +++++++++++++
 .../cpp/core/lower_hip/lower_hip_mma.cpp      | 14 ++++++++
 .../common/conv_direct_dw_wgrad_mfma.cpp      |  4 +++
 .../common/conv_direct_grouped_wgrad_mfma.cpp |  4 +++
 .../core/test_mma_operand_register_type.py    | 36 +++++++++++++++++--
 9 files changed, 118 insertions(+), 7 deletions(-)

diff --git a/dnn-providers/hip-kernel-provider/rocke/library/kernels/common/conv_direct_dw_wgrad_mfma.py b/dnn-providers/hip-kernel-provider/rocke/library/kernels/common/conv_direct_dw_wgrad_mfma.py
index 1daf8e5aa2a8..3615530abe61 100644
--- a/dnn-providers/hip-kernel-provider/rocke/library/kernels/common/conv_direct_dw_wgrad_mfma.py
+++ b/dnn-providers/hip-kernel-provider/rocke/library/kernels/common/conv_direct_dw_wgrad_mfma.py
@@ -198,6 +198,8 @@ def is_valid_dw_wgrad_mfma_spec(spec: DwWgradMfmaSpec, arch: str = "gfx950"):
         )
     if spec.rows_per_chunk < 1:
         return False, "rows_per_chunk >= 1"
+    if spec.target_blocks < 1:
+        return False, "target_blocks >= 1"
     if spec.lds_bytes() > 160 * 1024:
         return False, f"LDS {spec.lds_bytes()} B > 160 KiB"
     if arch not in ("gfx950",):
diff --git a/dnn-providers/hip-kernel-provider/rocke/library/kernels/common/conv_direct_grouped_wgrad_mfma.py b/dnn-providers/hip-kernel-provider/rocke/library/kernels/common/conv_direct_grouped_wgrad_mfma.py
index ff382f08e5ce..0e054aa1a506 100644
--- a/dnn-providers/hip-kernel-provider/rocke/library/kernels/common/conv_direct_grouped_wgrad_mfma.py
+++ b/dnn-providers/hip-kernel-provider/rocke/library/kernels/common/conv_direct_grouped_wgrad_mfma.py
@@ -159,6 +159,8 @@ def is_valid_grouped_wgrad_mfma_spec(spec: GroupedWgradMfmaSpec, arch: str = "gf
         return False, "prefetch in 2..4"
     if spec.rows_per_chunk < 1:
         return False, "rows_per_chunk >= 1"
+    if spec.target_blocks < 1:
+        return False, "target_blocks >= 1"
     if spec.n_acc() * 4 > 128:
         return False, "too many accumulators"
     if spec.lds_bytes() > 160 * 1024:
diff --git a/dnn-providers/hip-kernel-provider/rocke/library/tests/parity/conv_direct_grouped_emit.c b/dnn-providers/hip-kernel-provider/rocke/library/tests/parity/conv_direct_grouped_emit.c
index 2ce63d011d21..2435c21dd7e8 100644
--- a/dnn-providers/hip-kernel-provider/rocke/library/tests/parity/conv_direct_grouped_emit.c
+++ b/dnn-providers/hip-kernel-provider/rocke/library/tests/parity/conv_direct_grouped_emit.c
@@ -1629,7 +1629,7 @@ static int make_wgrad_direct_cfg(int idx, int* kind, wgrad_direct_specs_t* w, co
         *kind = KIND_DW_WGRAD_LDS;
     else if(idx <= 33)
         *kind = KIND_WGRAD_GATHER;
-    else if(idx <= 40)
+    else if(idx <= 40 || idx == 45 || idx == 46)
         *kind = KIND_DW_WGRAD_MFMA;
     else
         *kind = KIND_GROUPED_WGRAD_MFMA;
@@ -1949,6 +1949,26 @@ static int make_wgrad_direct_cfg(int idx, int* kind, wgrad_direct_specs_t* w, co
         gm->problem.cpg = 16;
         gm->problem.kpg = 16;
         return 0;
+    case 45: /* must reject: no blocks */
+        dm->problem = mk_dwp(64);
+        dm->target_blocks = 0;
+        return 0;
+    case 46: /* must reject: a negative block budget */
+        dm->problem = mk_dwp(64);
+        dm->target_blocks = -1;
+        return 0;
+    case 47: /* must reject: no blocks */
+        gm->problem.groups = 16;
+        gm->problem.cpg = 4;
+        gm->problem.kpg = 4;
+        gm->target_blocks = 0;
+        return 0;
+    case 48: /* must reject: a negative block budget */
+        gm->problem.groups = 16;
+        gm->problem.cpg = 4;
+        gm->problem.kpg = 4;
+        gm->target_blocks = -1;
+        return 0;
     default:
         return -1;
     }
diff --git a/dnn-providers/hip-kernel-provider/rocke/library/tests/parity/conv_direct_grouped_emit.py b/dnn-providers/hip-kernel-provider/rocke/library/tests/parity/conv_direct_grouped_emit.py
index 67696a345b70..9669ff01c242 100644
--- a/dnn-providers/hip-kernel-provider/rocke/library/tests/parity/conv_direct_grouped_emit.py
+++ b/dnn-providers/hip-kernel-provider/rocke/library/tests/parity/conv_direct_grouped_emit.py
@@ -130,7 +130,7 @@
 
 
 # ---------------------------------------------------------------------------
-# The direct wgrad kernels, indices 114..158: DirectWgradTiledSpec (wave tiles,
+# The direct wgrad kernels, indices 114..162: DirectWgradTiledSpec (wave tiles,
 # unfold_n, row_unroll, q_groups with and without pingpong, 1x1 and 5x5), the
 # DwWgradFoldSpec partial-slot fold, DwWgradColsSpec (small C), DwWgradLdsSpec
 # (row rings, strided, cols_per_lane, a partial channel tile), WgradGatherSpec
@@ -142,8 +142,8 @@
 # (preshifted and plain windows, packed images, 3x3 tap groups, the
 # single-prologue PAD = KH - 1 case, a partial channel tile) and
 # GroupedWgradMfmaSpec (4 and 8 channels per group, wave splits, a partial
-# channel tile), each with a reject. Follows the 11 grouped non-grouped
-# configs.
+# channel tile), each with a reject; then both with no blocks and with a
+# negative block budget. Follows the 11 grouped non-grouped configs.
 # ---------------------------------------------------------------------------
 _WGRAD_DIRECT_CFG_BASE = 114
 
@@ -437,6 +437,17 @@ def _wgrad_direct_spec(idx: int):
         ),
         # Must reject: 16 channels per group.
         44: ("grouped_wgrad_mfma", gm(problem=_wp(groups=4, cpg=16, kpg=16))),
+        # Must reject: no blocks, and a negative block budget.
+        45: ("dw_wgrad_mfma", dm(problem=_dwp(64), target_blocks=0)),
+        46: ("dw_wgrad_mfma", dm(problem=_dwp(64), target_blocks=-1)),
+        47: (
+            "grouped_wgrad_mfma",
+            gm(problem=_wp(groups=16, cpg=4, kpg=4), target_blocks=0),
+        ),
+        48: (
+            "grouped_wgrad_mfma",
+            gm(problem=_wp(groups=16, cpg=4, kpg=4), target_blocks=-1),
+        ),
     }
     sel = table.get(idx)
     if sel is None:
diff --git a/dnn-providers/hip-kernel-provider/rocke/library/tests/test_direct_wgrad_specs.py b/dnn-providers/hip-kernel-provider/rocke/library/tests/test_direct_wgrad_specs.py
index 70c5e46468dd..26a37f68f525 100644
--- a/dnn-providers/hip-kernel-provider/rocke/library/tests/test_direct_wgrad_specs.py
+++ b/dnn-providers/hip-kernel-provider/rocke/library/tests/test_direct_wgrad_specs.py
@@ -17,7 +17,15 @@
     is_valid_dw_wgrad_cols_spec,
     is_valid_dw_wgrad_lds_spec,
 )
+from kernels.common.conv_direct_dw_wgrad_mfma import (
+    DwWgradMfmaSpec,
+    is_valid_dw_wgrad_mfma_spec,
+)
 from kernels.common.conv_direct_grouped import DirectConvProblem
+from kernels.common.conv_direct_grouped_wgrad_mfma import (
+    GroupedWgradMfmaSpec,
+    is_valid_grouped_wgrad_mfma_spec,
+)
 from kernels.common.conv_direct_wgrad_gather import (
     WgradGatherSpec,
     build_wgrad_gather,
@@ -204,6 +212,22 @@ def test_unknown_target_is_rejected(spec, valid):
             DirectWgradTiledSpec(problem=_p(), target_blocks=0),
             is_valid_wgrad_tiled_spec,
         ),
+        (
+            DwWgradMfmaSpec(problem=_dw(64), target_blocks=0),
+            is_valid_dw_wgrad_mfma_spec,
+        ),
+        (
+            DwWgradMfmaSpec(problem=_dw(64), target_blocks=-1),
+            is_valid_dw_wgrad_mfma_spec,
+        ),
+        (
+            GroupedWgradMfmaSpec(problem=_p(groups=16, cpg=4, kpg=4), target_blocks=0),
+            is_valid_grouped_wgrad_mfma_spec,
+        ),
+        (
+            GroupedWgradMfmaSpec(problem=_p(groups=16, cpg=4, kpg=4), target_blocks=-1),
+            is_valid_grouped_wgrad_mfma_spec,
+        ),
         # -3 % 3 == 0 and 30 % -3 == 0: without the sign check the row loop
         # runs no iterations and dW stays zero.
         (DirectWgradTiledSpec(problem=_p(), row_unroll=-3), is_valid_wgrad_tiled_spec),
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mma.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mma.cpp
index 7978ce07af49..75fa99e8fec4 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mma.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mma.cpp
@@ -255,6 +255,20 @@ static rocke_status_t rocke_h_op_tile_mma(rocke_h_lowerer_t* lw, const rocke_op_
                       rocke_h_name(lw, op->operands[2]));
         return ROCKE_OK;
     }
+    /* Batched 4x4x4 atoms. No legacy opcode either; Python reaches them as the
+     * tile.mfma_f32_4x4x4_{f16,bf16} handlers, which pass the operands as they are. */
+    if(strcmp(op_id, "mfma_f32_4x4x4_f16") == 0 || strcmp(op_id, "mfma_f32_4x4x4_bf16") == 0)
+    {
+        bool bf16 = strcmp(op_id, "mfma_f32_4x4x4_bf16") == 0;
+        rocke_h_emitf(lw,
+                      "f32x4 %s = __builtin_amdgcn_%s(%s, %s, %s, 0, 0, 0);",
+                      rocke_h_name(lw, op->results[0]),
+                      bf16 ? "mfma_f32_4x4x4bf16_1k" : "mfma_f32_4x4x4f16",
+                      rocke_h_name(lw, op->operands[0]),
+                      rocke_h_name(lw, op->operands[1]),
+                      rocke_h_name(lw, op->operands[2]));
+        return ROCKE_OK;
+    }
     if(rocke_gfx1250_scaled_wmma(op_id))
     {
         return h_emit_gfx1250_scaled_wmma(lw, op);
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/conv_direct_dw_wgrad_mfma.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/conv_direct_dw_wgrad_mfma.cpp
index e96331b01829..9a7735142b88 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/conv_direct_dw_wgrad_mfma.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/conv_direct_dw_wgrad_mfma.cpp
@@ -271,6 +271,10 @@ extern "C" bool rocke_dw_wgrad_mfma_is_valid_spec(const rocke_dw_wgrad_mfma_spec
     {
         why = "rows_per_chunk >= 1";
     }
+    else if(s->target_blocks < 1)
+    {
+        why = "target_blocks >= 1";
+    }
     else if(rocke_dw_wgrad_mfma_lds_bytes(s) > 160 * 1024)
     {
         ROCKE_ERR_SNPRINTF(
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/conv_direct_grouped_wgrad_mfma.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/conv_direct_grouped_wgrad_mfma.cpp
index 7d382c4b48a1..4c039f4b1dfd 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/conv_direct_grouped_wgrad_mfma.cpp
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/conv_direct_grouped_wgrad_mfma.cpp
@@ -235,6 +235,10 @@ extern "C" bool rocke_grouped_wgrad_mfma_is_valid_spec(const rocke_grouped_wgrad
     {
         why = "rows_per_chunk >= 1";
     }
+    else if(s->target_blocks < 1)
+    {
+        why = "target_blocks >= 1";
+    }
     else if(n_acc(s) * 4 > 128)
     {
         why = "too many accumulators";
diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_mma_operand_register_type.py b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_mma_operand_register_type.py
index cc6d1a4582c9..e971a0fdaeef 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_mma_operand_register_type.py
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_mma_operand_register_type.py
@@ -11,9 +11,10 @@
 in the LLVM and HIP lowerings.
 
 Every LLVM case lowers with both engines and requires the same IR, or the same
-error. The HIP cases do the same for the f32 MFMAs; the C++ HIP lowerer has no
-integer or gfx1250 fp8 WMMA lowering, so those cases run the Python HIP lowering
-and require the LLVM lowering's error text.
+error. The HIP cases do the same for the f32 MFMAs and the batched 4x4x4
+f16/bf16 MFMAs; the C++ HIP lowerer has no integer or gfx1250 fp8 WMMA lowering,
+so those cases run the Python HIP lowering and require the LLVM lowering's error
+text.
 """
 
 import re
@@ -24,6 +25,7 @@
 from rocke.core.arch import ArchTarget
 from rocke.core.ir import (
     BF8E5M2,
+    BF16,
     F16,
     F32,
     FP8E4M3,
@@ -386,6 +388,34 @@ def test_hip_f32_mfma_reinterprets_only_the_one_element_vector(
     _assert_hipcc_accepts(source, arch, tmp_path)
 
 
+@pytest.mark.parametrize(
+    ("op_id", "elem", "builtin"),
+    [
+        ("mfma_f32_4x4x4_f16", F16, "mfma_f32_4x4x4f16"),
+        ("mfma_f32_4x4x4_bf16", BF16, "mfma_f32_4x4x4bf16_1k"),
+    ],
+)
+@pytest.mark.parametrize("arch", ["gfx942", "gfx950"])
+def test_hip_batched_4x4x4_mfma_lowers_in_both_engines(
+    arch, op_id, elem, builtin, tmp_path
+):
+    # Not in the arch MMA catalog (the kernels emit them by op_id), so the op_id
+    # goes to b.mma as it is.
+    b = IRBuilder(f"mma_operand_{op_id}")
+    out = b.param("D", PtrType(F32, "global"))
+    a_value, b_value, c_value = (
+        make(b) for make in (_loaded(elem, 4), _loaded(elem, 4), _loaded(F32, 4))
+    )
+    result = b.mma(op_id, a_value, b_value, c_value)
+    b.global_store(out, b.thread_id_x(), b.vec_extract(result, 0))
+    source = lower_hip_with_both_engines(b.kernel, arch)
+    assert re.search(
+        rf"= __builtin_amdgcn_{builtin}\(\w+, \w+, \w+, 0, 0, 0\);",
+        _mma_builtin_call(source),
+    )
+    _assert_hipcc_accepts(source, arch, tmp_path)
+
+
 def test_hip_integer_wmma_reinterprets_only_the_packed_byte_operand(tmp_path):
     kernel = _mma_kernel(
         "gfx1151",

From 2760d87fdc9ac9d269288901b6793f2270881b43 Mon Sep 17 00:00:00 2001
From: Bartlomiej Kocot <barkocot@amd.com>
Date: Fri, 23 Oct 2026 07:26:54 -0500
Subject: [PATCH 3/6] fix(rocke): measure mfma.f32.4x4x4bf16.1k in the llvm20
 arch domain

The bf16 4x4x4 atom added the decl key mfma.f32.4x4x4bf16.1k but the
llvm20 arch-domain column had no row for it (nor for the
amdgcn.fdot2.{f16,bf16} keys from develop), so
test_every_decl_key_is_covered_on_every_wired_arch failed and the
arch-domain check stayed silent when the atom was used on RDNA.

Regenerated with tools/gen_arch_domain.py against ROCm 7.1.1's LLVM 20,
the toolchain the column was measured with. Additions only: the bf16 atom
is ok on gfx90a/gfx942/gfx950 and arch_absent on RDNA, matching its f16
twin. The llvm22 column already has these rows from develop.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
---
 .../data/intrinsic_arch_domain.llvm20.json    | 114 ++++++++++++++++++
 1 file changed, 114 insertions(+)

diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/arch/data/intrinsic_arch_domain.llvm20.json b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/arch/data/intrinsic_arch_domain.llvm20.json
index bd1e0425972b..80d6418f7b26 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/arch/data/intrinsic_arch_domain.llvm20.json
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/arch/data/intrinsic_arch_domain.llvm20.json
@@ -12,6 +12,8 @@
     "amdgcn.cvt.scalef32.pk.f32.fp8": "llvm.amdgcn.cvt.scalef32.pk.f32.fp8",
     "amdgcn.ds.swizzle": "llvm.amdgcn.ds.swizzle",
     "amdgcn.exp2.f32": "llvm.amdgcn.exp2.f32",
+    "amdgcn.fdot2.bf16": "llvm.amdgcn.fdot2.f32.bf16",
+    "amdgcn.fdot2.f16": "llvm.amdgcn.fdot2",
     "amdgcn.perm": "llvm.amdgcn.perm",
     "amdgcn.permlane16": "llvm.amdgcn.permlane16.i32",
     "amdgcn.permlane32.swap": "llvm.amdgcn.permlane32.swap",
@@ -71,6 +73,7 @@
     "mfma.f32.32x32x4.xf32": "llvm.amdgcn.mfma.f32.32x32x4.xf32",
     "mfma.f32.32x32x8bf16.1k": "llvm.amdgcn.mfma.f32.32x32x8bf16.1k",
     "mfma.f32.32x32x8f16": "llvm.amdgcn.mfma.f32.32x32x8f16",
+    "mfma.f32.4x4x4bf16.1k": "llvm.amdgcn.mfma.f32.4x4x4bf16.1k",
     "mfma.f32.4x4x4f16": "llvm.amdgcn.mfma.f32.4x4x4f16",
     "mfma.scale.f32.16x16x128.f8f6f4": "llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4",
     "mfma.scale.f32.32x32x64.f8f6f4.v8i32.v8i32": "llvm.amdgcn.mfma.scale.f32.32x32x64.f8f6f4.v8i32.v8i32",
@@ -556,6 +559,78 @@
         "verified_on": "llvm20"
       }
     },
+    "amdgcn.fdot2.bf16": {
+      "gfx11-generic": {
+        "status": "ok",
+        "verified_on": "llvm20"
+      },
+      "gfx1151": {
+        "status": "ok",
+        "verified_on": "llvm20"
+      },
+      "gfx12-generic": {
+        "status": "ok",
+        "verified_on": "llvm20"
+      },
+      "gfx1201": {
+        "status": "ok",
+        "verified_on": "llvm20"
+      },
+      "gfx1250": {
+        "evidence": "invalid target ID",
+        "status": "target_unsupported",
+        "verified_on": "llvm20"
+      },
+      "gfx90a": {
+        "evidence": "fatal error: error in backend: Cannot select: intrinsic %llvm.amdgcn.fdot2.f32.bf16",
+        "status": "arch_absent",
+        "verified_on": "llvm20"
+      },
+      "gfx942": {
+        "evidence": "fatal error: error in backend: Cannot select: intrinsic %llvm.amdgcn.fdot2.f32.bf16",
+        "status": "arch_absent",
+        "verified_on": "llvm20"
+      },
+      "gfx950": {
+        "status": "ok",
+        "verified_on": "llvm20"
+      }
+    },
+    "amdgcn.fdot2.f16": {
+      "gfx11-generic": {
+        "status": "ok",
+        "verified_on": "llvm20"
+      },
+      "gfx1151": {
+        "status": "ok",
+        "verified_on": "llvm20"
+      },
+      "gfx12-generic": {
+        "status": "ok",
+        "verified_on": "llvm20"
+      },
+      "gfx1201": {
+        "status": "ok",
+        "verified_on": "llvm20"
+      },
+      "gfx1250": {
+        "evidence": "invalid target ID",
+        "status": "target_unsupported",
+        "verified_on": "llvm20"
+      },
+      "gfx90a": {
+        "status": "ok",
+        "verified_on": "llvm20"
+      },
+      "gfx942": {
+        "status": "ok",
+        "verified_on": "llvm20"
+      },
+      "gfx950": {
+        "status": "ok",
+        "verified_on": "llvm20"
+      }
+    },
     "amdgcn.perm": {
       "gfx11-generic": {
         "status": "ok",
@@ -4346,6 +4421,45 @@
         "verified_on": "llvm20"
       }
     },
+    "mfma.f32.4x4x4bf16.1k": {
+      "gfx11-generic": {
+        "evidence": "fatal error: error in backend: Cannot select: intrinsic %llvm.amdgcn.mfma.f32.4x4x4bf16.1k",
+        "status": "arch_absent",
+        "verified_on": "llvm20"
+      },
+      "gfx1151": {
+        "evidence": "fatal error: error in backend: Cannot select: intrinsic %llvm.amdgcn.mfma.f32.4x4x4bf16.1k",
+        "status": "arch_absent",
+        "verified_on": "llvm20"
+      },
+      "gfx12-generic": {
+        "evidence": "fatal error: error in backend: Cannot select: intrinsic %llvm.amdgcn.mfma.f32.4x4x4bf16.1k",
+        "status": "arch_absent",
+        "verified_on": "llvm20"
+      },
+      "gfx1201": {
+        "evidence": "fatal error: error in backend: Cannot select: intrinsic %llvm.amdgcn.mfma.f32.4x4x4bf16.1k",
+        "status": "arch_absent",
+        "verified_on": "llvm20"
+      },
+      "gfx1250": {
+        "evidence": "invalid target ID",
+        "status": "target_unsupported",
+        "verified_on": "llvm20"
+      },
+      "gfx90a": {
+        "status": "ok",
+        "verified_on": "llvm20"
+      },
+      "gfx942": {
+        "status": "ok",
+        "verified_on": "llvm20"
+      },
+      "gfx950": {
+        "status": "ok",
+        "verified_on": "llvm20"
+      }
+    },
     "mfma.f32.4x4x4f16": {
       "gfx11-generic": {
         "evidence": "fatal error: error in backend: Cannot select: intrinsic %llvm.amdgcn.mfma.f32.4x4x4f16",

From becfb94a65fb509e4b9b308b73fb3e7698514ec9 Mon Sep 17 00:00:00 2001
From: Bartlomiej Kocot <barkocot@amd.com>
Date: Fri, 23 Oct 2026 07:33:53 -0500
Subject: [PATCH 4/6] fix(rocke): measure mfma.f32.4x4x4bf16.1k in the llvm23
 arch domain

Same gap as the llvm20 column: no rows for mfma.f32.4x4x4bf16.1k or
amdgcn.fdot2.{f16,bf16}. Regenerated with tools/gen_arch_domain.py
against clang 23 at 8f497e0992fb (the toolchain the column was measured
with, from the rocm/pytorch rocm10.0 image, CPU only). The bf16 atom is
ok on gfx90a/gfx942/gfx950 and arch_absent on RDNA and gfx1250, like its
f16 twin. Four existing rows change evidence wording only (this llvm-mc
names the target in its message); their status is unchanged.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
---
 .../data/intrinsic_arch_domain.llvm23.json    | 122 +++++++++++++++++-
 1 file changed, 118 insertions(+), 4 deletions(-)

diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/arch/data/intrinsic_arch_domain.llvm23.json b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/arch/data/intrinsic_arch_domain.llvm23.json
index b55f9265e396..c885416e2046 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/arch/data/intrinsic_arch_domain.llvm23.json
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/arch/data/intrinsic_arch_domain.llvm23.json
@@ -12,6 +12,8 @@
     "amdgcn.cvt.scalef32.pk.f32.fp8": "llvm.amdgcn.cvt.scalef32.pk.f32.fp8",
     "amdgcn.ds.swizzle": "llvm.amdgcn.ds.swizzle",
     "amdgcn.exp2.f32": "llvm.amdgcn.exp2.f32",
+    "amdgcn.fdot2.bf16": "llvm.amdgcn.fdot2.f32.bf16",
+    "amdgcn.fdot2.f16": "llvm.amdgcn.fdot2",
     "amdgcn.perm": "llvm.amdgcn.perm",
     "amdgcn.permlane16": "llvm.amdgcn.permlane16.i32",
     "amdgcn.permlane32.swap": "llvm.amdgcn.permlane32.swap",
@@ -106,6 +108,7 @@
     "mfma.f32.32x32x4.xf32": "llvm.amdgcn.mfma.f32.32x32x4.xf32",
     "mfma.f32.32x32x8bf16.1k": "llvm.amdgcn.mfma.f32.32x32x8bf16.1k",
     "mfma.f32.32x32x8f16": "llvm.amdgcn.mfma.f32.32x32x8f16",
+    "mfma.f32.4x4x4bf16.1k": "llvm.amdgcn.mfma.f32.4x4x4bf16.1k",
     "mfma.f32.4x4x4f16": "llvm.amdgcn.mfma.f32.4x4x4f16",
     "mfma.scale.f32.16x16x128.f8f6f4": "llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32",
     "mfma.scale.f32.32x32x64.f8f6f4.v8i32.v8i32": "llvm.amdgcn.mfma.scale.f32.32x32x64.f8f6f4.v8i32.v8i32",
@@ -484,12 +487,12 @@
         "verified_on": "llvm23"
       },
       "gfx90a": {
-        "evidence": "assembler: instruction not supported on this GPU (v_cvt_scalef32_pk_f32_bf8)",
+        "evidence": "assembler: instruction not supported on this GPU (gfx90a): v_cvt_scalef32_pk_f32_bf8 (v_cvt_scalef32_pk_f32_bf8)",
         "status": "arch_absent",
         "verified_on": "llvm23"
       },
       "gfx942": {
-        "evidence": "assembler: instruction not supported on this GPU (v_cvt_scalef32_pk_f32_bf8)",
+        "evidence": "assembler: instruction not supported on this GPU (gfx942): v_cvt_scalef32_pk_f32_bf8 (v_cvt_scalef32_pk_f32_bf8)",
         "status": "arch_absent",
         "verified_on": "llvm23"
       },
@@ -525,12 +528,12 @@
         "verified_on": "llvm23"
       },
       "gfx90a": {
-        "evidence": "assembler: instruction not supported on this GPU (v_cvt_scalef32_pk_f32_fp8)",
+        "evidence": "assembler: instruction not supported on this GPU (gfx90a): v_cvt_scalef32_pk_f32_fp8 (v_cvt_scalef32_pk_f32_fp8)",
         "status": "arch_absent",
         "verified_on": "llvm23"
       },
       "gfx942": {
-        "evidence": "assembler: instruction not supported on this GPU (v_cvt_scalef32_pk_f32_fp8)",
+        "evidence": "assembler: instruction not supported on this GPU (gfx942): v_cvt_scalef32_pk_f32_fp8 (v_cvt_scalef32_pk_f32_fp8)",
         "status": "arch_absent",
         "verified_on": "llvm23"
       },
@@ -607,6 +610,78 @@
         "verified_on": "llvm23"
       }
     },
+    "amdgcn.fdot2.bf16": {
+      "gfx11-generic": {
+        "status": "ok",
+        "verified_on": "llvm23"
+      },
+      "gfx1151": {
+        "status": "ok",
+        "verified_on": "llvm23"
+      },
+      "gfx12-generic": {
+        "status": "ok",
+        "verified_on": "llvm23"
+      },
+      "gfx1201": {
+        "status": "ok",
+        "verified_on": "llvm23"
+      },
+      "gfx1250": {
+        "evidence": "fatal error: error in backend: Cannot select: intrinsic %llvm.amdgcn.fdot2.f32.bf16",
+        "status": "arch_absent",
+        "verified_on": "llvm23"
+      },
+      "gfx90a": {
+        "evidence": "fatal error: error in backend: Cannot select: intrinsic %llvm.amdgcn.fdot2.f32.bf16",
+        "status": "arch_absent",
+        "verified_on": "llvm23"
+      },
+      "gfx942": {
+        "evidence": "fatal error: error in backend: Cannot select: intrinsic %llvm.amdgcn.fdot2.f32.bf16",
+        "status": "arch_absent",
+        "verified_on": "llvm23"
+      },
+      "gfx950": {
+        "status": "ok",
+        "verified_on": "llvm23"
+      }
+    },
+    "amdgcn.fdot2.f16": {
+      "gfx11-generic": {
+        "status": "ok",
+        "verified_on": "llvm23"
+      },
+      "gfx1151": {
+        "status": "ok",
+        "verified_on": "llvm23"
+      },
+      "gfx12-generic": {
+        "status": "ok",
+        "verified_on": "llvm23"
+      },
+      "gfx1201": {
+        "status": "ok",
+        "verified_on": "llvm23"
+      },
+      "gfx1250": {
+        "evidence": "fatal error: error in backend: Cannot select: 0xADDR: f32 = AMDGPUISD::FDOT2 0xADDR, 0xADDR, 0xADDR, TargetConstant:i1<0>",
+        "status": "arch_absent",
+        "verified_on": "llvm23"
+      },
+      "gfx90a": {
+        "status": "ok",
+        "verified_on": "llvm23"
+      },
+      "gfx942": {
+        "status": "ok",
+        "verified_on": "llvm23"
+      },
+      "gfx950": {
+        "status": "ok",
+        "verified_on": "llvm23"
+      }
+    },
     "amdgcn.perm": {
       "gfx11-generic": {
         "status": "ok",
@@ -4219,6 +4294,45 @@
         "verified_on": "llvm23"
       }
     },
+    "mfma.f32.4x4x4bf16.1k": {
+      "gfx11-generic": {
+        "evidence": "fatal error: error in backend: Cannot select: intrinsic %llvm.amdgcn.mfma.f32.4x4x4bf16.1k",
+        "status": "arch_absent",
+        "verified_on": "llvm23"
+      },
+      "gfx1151": {
+        "evidence": "fatal error: error in backend: Cannot select: intrinsic %llvm.amdgcn.mfma.f32.4x4x4bf16.1k",
+        "status": "arch_absent",
+        "verified_on": "llvm23"
+      },
+      "gfx12-generic": {
+        "evidence": "fatal error: error in backend: Cannot select: intrinsic %llvm.amdgcn.mfma.f32.4x4x4bf16.1k",
+        "status": "arch_absent",
+        "verified_on": "llvm23"
+      },
+      "gfx1201": {
+        "evidence": "fatal error: error in backend: Cannot select: intrinsic %llvm.amdgcn.mfma.f32.4x4x4bf16.1k",
+        "status": "arch_absent",
+        "verified_on": "llvm23"
+      },
+      "gfx1250": {
+        "evidence": "fatal error: error in backend: Cannot select: intrinsic %llvm.amdgcn.mfma.f32.4x4x4bf16.1k",
+        "status": "arch_absent",
+        "verified_on": "llvm23"
+      },
+      "gfx90a": {
+        "status": "ok",
+        "verified_on": "llvm23"
+      },
+      "gfx942": {
+        "status": "ok",
+        "verified_on": "llvm23"
+      },
+      "gfx950": {
+        "status": "ok",
+        "verified_on": "llvm23"
+      }
+    },
     "mfma.f32.4x4x4f16": {
       "gfx11-generic": {
         "evidence": "fatal error: error in backend: Cannot select: intrinsic %llvm.amdgcn.mfma.f32.4x4x4f16",

From ac9fcf2bcd1769e71aead15a583b287ca8dc39a6 Mon Sep 17 00:00:00 2001
From: Bartlomiej Kocot <barkocot@amd.com>
Date: Sun, 25 Oct 2026 11:37:29 +0000
Subject: [PATCH 5/6] ir

---
 .../golden/rocke_representative_ir_sha256.json | 18 +++++++++---------
 1 file changed, 9 insertions(+), 9 deletions(-)

diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256.json b/dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256.json
index c02b0506c601..77d8639d4500 100644
--- a/dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256.json
+++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256.json
@@ -2205,8 +2205,8 @@
           "gfx1201": 8,
           "gfx1250": 12,
           "gfx90a": 5,
-          "gfx942": 50,
-          "gfx950": 181
+          "gfx942": 52,
+          "gfx950": 182
         },
         "expected_failures": 6,
         "families": {
@@ -2233,7 +2233,7 @@
           "moe": 9,
           "nontemporal": 2
         },
-        "ok_cases": 268
+        "ok_cases": 271
       }
     },
     "llvm22": {
@@ -4441,8 +4441,8 @@
           "gfx1201": 8,
           "gfx1250": 12,
           "gfx90a": 5,
-          "gfx942": 50,
-          "gfx950": 181
+          "gfx942": 52,
+          "gfx950": 182
         },
         "expected_failures": 6,
         "families": {
@@ -4469,7 +4469,7 @@
           "moe": 9,
           "nontemporal": 2
         },
-        "ok_cases": 268
+        "ok_cases": 271
       }
     },
     "llvm23": {
@@ -6700,8 +6700,8 @@
           "gfx1201": 8,
           "gfx1250": 18,
           "gfx90a": 5,
-          "gfx942": 50,
-          "gfx950": 181
+          "gfx942": 52,
+          "gfx950": 182
         },
         "expected_failures": 0,
         "families": {
@@ -6729,7 +6729,7 @@
           "moe": 9,
           "nontemporal": 2
         },
-        "ok_cases": 272
+        "ok_cases": 277
       }
     }
   },

From 862c9f4735a343c104e13dd268dfbb0cbac372f3 Mon Sep 17 00:00:00 2001
From: Bartlomiej Kocot <barkocot@amd.com>
Date: Sun, 25 Oct 2026 18:20:58 -0500
Subject: [PATCH 6/6] fix(rocke): split the dwm alias shared by WMMA dgrad and
 MFMA dw wgrad

The develop merge left conv_direct_dgrad_wmma and conv_direct_dw_wgrad_mfma
both imported as dwm in select_direct_kernel_builder, so the second shadowed
the first and every direct conv build raised AttributeError on
build_direct_dgrad_wmma. Import the MFMA dw wgrad module as dwmf.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
---
 .../rocke/library/kernels/conv/direct_variants.py             | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/dnn-providers/hip-kernel-provider/rocke/library/kernels/conv/direct_variants.py b/dnn-providers/hip-kernel-provider/rocke/library/kernels/conv/direct_variants.py
index 5f04b1399f47..7965351f4255 100644
--- a/dnn-providers/hip-kernel-provider/rocke/library/kernels/conv/direct_variants.py
+++ b/dnn-providers/hip-kernel-provider/rocke/library/kernels/conv/direct_variants.py
@@ -491,7 +491,7 @@ def select_direct_kernel_builder(variant: str, spec, arch: str):
     """
     from kernels.common import conv_direct_dgrad_wmma as dwm
     from kernels.common import conv_direct_dw_wgrad as dw
-    from kernels.common import conv_direct_dw_wgrad_mfma as dwm
+    from kernels.common import conv_direct_dw_wgrad_mfma as dwmf
     from kernels.common import conv_direct_grouped as dc
     from kernels.common import conv_direct_grouped_wgrad_mfma as gwm
     from kernels.common import conv_direct_nongrouped as dn
@@ -504,7 +504,7 @@ def select_direct_kernel_builder(variant: str, spec, arch: str):
         "direct_wgrad_dw": dw.build_dw_wgrad_lds,
         "direct_wgrad_dw_cols": dw.build_dw_wgrad_cols,
         "direct_wgrad_gather": wg.build_wgrad_gather,
-        "direct_wgrad_dw_mfma": dwm.build_dw_wgrad_mfma,
+        "direct_wgrad_dw_mfma": dwmf.build_dw_wgrad_mfma,
         "direct_wgrad_grouped_mfma": gwm.build_grouped_wgrad_mfma,
         DIRECT_WGRAD_FOLD: dw.build_dw_wgrad_fold,
         "direct_grouped": dc.build_direct_conv,