feat(rocke): quad_perm crosslane intrinsic + kernarg-packer precompile

dnn-providers/hip-kernel-provider/rocke/library/builders/gfx950/kda/CASE_STUDY.md

@@ -116,11 +116,12 @@ the source.
 
 3. **Blocked triangular solve.** The solve's arithmetic splits into per-block
    forward substitution (irreducibly scalar VALU) and a rank update against
-   already-solved blocks (a matmul, so MFMA). The substitution shrinks as the
-   *square* of the block size, so smaller blocks move more of the `O(C^3)` work
-   onto the MFMA pipe at the cost of one more block step. `solve_block` exposes
-   the trade; the degenerate `solve_block == chunk` is the original unblocked
-   scalar solve.
+   already-solved blocks (a matmul, so MFMA). One block of size `b` costs
+   `O(b²)` scalar work, but there are `C / b` blocks, so the total is `≈ C·b/2`
+   — **linear** in the block size, not quadratic. Smaller blocks therefore move
+   more of the `O(C^3)` work onto the MFMA pipe roughly in proportion, at the
+   cost of one more block step. `solve_block` exposes the trade; the degenerate
+   `solve_block == chunk` is the original unblocked scalar solve.
 
    This step introduced a race: `ds_write_b128` followed by a `ds_read_b128`
    from a *different* lane needs an explicit `s_waitcnt lgkmcnt(0)` between

dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/core_types.cpp

@@ -581,6 +581,7 @@ static const char* const rocke_opcode_names[ROCKE_OP__COUNT] = {
     "tile.ds_swizzle_xor",
     "tile.ds_swizzle",
     "tile.mov_dpp8",
+    "tile.quad_perm",
     "tile.wave_reduce",
     "tile.readlane",
     "tile.writelane",
@@ -826,6 +827,7 @@ static const bool rocke_opcode_pure[ROCKE_OP__COUNT] = {
     /* ds_swizzle_xor    */ true,
     /* ds_swizzle        */ true,
     /* mov_dpp8          */ true,
+    /* quad_perm         */ true,
     /* wave_reduce       */ true,
     /* readlane          */ true,
     /* writelane         */ true,

dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/ir_flow.cpp

@@ -182,6 +182,32 @@ rocke_value_t* rocke_b_mov_dpp8(rocke_ir_builder_t* b, rocke_value_t* data, int
     return rocke_i_op1(b, ROCKE_OP_TILE_MOV_DPP8, &data, 1, data->type, &attrs, "dpp8");
 }
 
+rocke_value_t*
+    rocke_b_quad_perm(rocke_ir_builder_t* b, rocke_value_t* data, int p0, int p1, int p2, int p3)
+{
+    rocke_attr_map_t attrs;
+    int ctrl;
+    if(!rocke_i_live(b))
+        return NULL;
+    if(!data)
+        return (rocke_value_t*)rocke_i_set_err(b, ROCKE_ERR_VALUE, "quad_perm: NULL data");
+    if(!rocke_flow_is_i32(data->type))
+        return (rocke_value_t*)rocke_i_set_err(b, ROCKE_ERR_VALUE, "quad_perm requires i32 data");
+    if(p0 < 0 || p0 > 3 || p1 < 0 || p1 > 3 || p2 < 0 || p2 > 3 || p3 < 0 || p3 > 3)
+        return (rocke_value_t*)rocke_i_set_err(
+            b,
+            ROCKE_ERR_VALUE,
+            "quad_perm lane selectors must be in 0..3, got [%d,%d,%d,%d]",
+            p0,
+            p1,
+            p2,
+            p3);
+    ctrl = p0 | (p1 << 2) | (p2 << 4) | (p3 << 6);
+    attrs = rocke_i_attrs(b);
+    rocke_attr_set_int(b, &attrs, "ctrl", (int64_t)ctrl);
+    return rocke_i_op1(b, ROCKE_OP_TILE_QUAD_PERM, &data, 1, rocke_i32(), &attrs, "qperm");
+}
+
 static bool flow_wave_reduce_allowed(const char* reduce_op, const char* ty)
 {
     if(!reduce_op || !ty)

dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mma.cpp

@@ -287,6 +287,29 @@ static rocke_status_t rocke_h_op_tile_ds_swizzle_xor(rocke_h_lowerer_t* lw, cons
     return lw->status;
 }
 
+static rocke_status_t rocke_h_op_tile_quad_perm(rocke_h_lowerer_t* lw, const rocke_op_t* op)
+{
+    const rocke_value_t* data = op->operands[0];
+    const rocke_value_t* r = h_res(op);
+    int64_t ctrl = 0;
+    if(!rocke_attr_get_int(&op->attrs, "ctrl", &ctrl))
+        return rocke_h_fail(lw, ROCKE_ERR_KEY, "tile.quad_perm: missing 'ctrl'");
+    /* 0..255 is the whole legal range: ctrl packs four two-bit lane
+     * selectors. Reject instead of masking (see lower_llvm/crosslane.cpp). */
+    if(ctrl < 0 || ctrl > 255)
+        return rocke_h_fail(lw,
+                            ROCKE_ERR_VALUE,
+                            "tile.quad_perm: ctrl must be in 0..255, got %lld",
+                            (long long)ctrl);
+    rocke_h_emitf(lw,
+                  "int %s = __builtin_amdgcn_update_dpp(%s, %s, %lld, 15, 15, 1);",
+                  rocke_h_name(lw, r),
+                  rocke_h_name(lw, data),
+                  rocke_h_name(lw, data),
+                  (long long)ctrl);
+    return lw->status;
+}
+
 /* def _op_tile_mov_dpp(self, op): row_shr/row_shl -> dpp_ctrl, update_dpp. */
 static rocke_status_t rocke_h_op_tile_mov_dpp(rocke_h_lowerer_t* lw, const rocke_op_t* op)
 {
@@ -1306,6 +1329,7 @@ const rocke_h_handler_entry_t* rocke_h_handlers_mma(void)
         {ROCKE_OP_TILE_DS_BPERMUTE, rocke_h_op_tile_ds_bpermute},
         {ROCKE_OP_TILE_DS_BPERMUTE_B64, rocke_h_op_tile_ds_bpermute_b64},
         {ROCKE_OP_TILE_DS_SWIZZLE_XOR, rocke_h_op_tile_ds_swizzle_xor},
+        {ROCKE_OP_TILE_QUAD_PERM, rocke_h_op_tile_quad_perm},
         {ROCKE_OP_TILE_MOV_DPP, rocke_h_op_tile_mov_dpp},
         {ROCKE_OP_TILE_PERMLANE32_SWAP, rocke_h_op_tile_permlane32_swap},
         {ROCKE_OP_TILE_PERM_B32, rocke_h_op_tile_perm_b32},

dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/crosslane.cpp

@@ -325,6 +325,30 @@ static void _op_tile_mov_dpp8(rocke_lower_t* L, const rocke_op_t* op)
                    (long long)(sel & 0xFFFFFF));
 }
 
+static void _op_tile_quad_perm(rocke_lower_t* L, const rocke_op_t* op)
+{
+    const rocke_value_t* data = op->operands[0];
+    int64_t ctrl = 0;
+    if(!rocke_attr_get_int(&op->attrs, "ctrl", &ctrl))
+        rocke_ll_fail(L, ROCKE_ERR_KEY, "tile.quad_perm: missing 'ctrl'");
+    /* ctrl packs four two-bit lane selectors (p0 | p1<<2 | p2<<4 | p3<<6),
+     * so 0..255 is the whole legal range. Reject instead of masking: a
+     * truncated control is a different, silently valid permutation. */
+    if(ctrl < 0 || ctrl > 255)
+        rocke_ll_fail(L,
+                      ROCKE_ERR_VALUE,
+                      "tile.quad_perm: ctrl must be in 0..255, got %lld",
+                      (long long)ctrl);
+    rocke_ll_need(L, "update.dpp.i32");
+    rocke_ll_emitf(L,
+                   "  %s = call i32 @llvm.amdgcn.update.dpp.i32("
+                   "i32 %s, i32 %s, i32 %lld, i32 15, i32 15, i1 true)",
+                   ll_result_name(op),
+                   rocke_ll_operand(L, data),
+                   rocke_ll_operand(L, data),
+                   (long long)ctrl);
+}
+
 static void _op_tile_wave_reduce(rocke_lower_t* L, const rocke_op_t* op)
 {
     const rocke_value_t* v = op->operands[0];
@@ -1050,6 +1074,7 @@ void rocke_ll_register_crosslane(void)
     rocke_ll_set_handler(ROCKE_OP_TILE_DS_SWIZZLE_XOR, _op_tile_ds_swizzle_xor);
     rocke_ll_set_handler(ROCKE_OP_TILE_DS_SWIZZLE, _op_tile_ds_swizzle);
     rocke_ll_set_handler(ROCKE_OP_TILE_MOV_DPP8, _op_tile_mov_dpp8);
+    rocke_ll_set_handler(ROCKE_OP_TILE_QUAD_PERM, _op_tile_quad_perm);
     rocke_ll_set_handler(ROCKE_OP_TILE_WAVE_REDUCE, _op_tile_wave_reduce);
     rocke_ll_set_handler(ROCKE_OP_TILE_READLANE, _op_tile_readlane);
     rocke_ll_set_handler(ROCKE_OP_TILE_WRITELANE, _op_tile_writelane);

dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/ir.h

@@ -360,6 +360,7 @@ typedef enum rocke_opcode
     ROCKE_OP_TILE_DS_SWIZZLE_XOR,
     ROCKE_OP_TILE_DS_SWIZZLE,
     ROCKE_OP_TILE_MOV_DPP8,
+    ROCKE_OP_TILE_QUAD_PERM,
     ROCKE_OP_TILE_WAVE_REDUCE,
     ROCKE_OP_TILE_READLANE,
     ROCKE_OP_TILE_WRITELANE,
@@ -1125,6 +1126,15 @@ rocke_value_t* rocke_b_permlane16(rocke_ir_builder_t* b,
                                   rocke_value_t* src2,
                                   bool fi,
                                   bool bound_ctrl);
+/* quad_perm: each pN selects source lane 0..3 for destination lane N.
+ * Wave-size-independent: the control word applies within every four-lane
+ * group and four divides both 32 and 64, so a lane never addresses outside
+ * its own quad; wave size changes only the number of quads. Requires
+ * DPP-capable hardware (base DPP, so CDNA as well as RDNA). No lane
+ * targeting -- the control is broadcast to every quad, row/bank masks
+ * fixed at 15, 15 by the lowerers. */
+rocke_value_t*
+    rocke_b_quad_perm(rocke_ir_builder_t* b, rocke_value_t* data, int p0, int p1, int p2, int p3);
 rocke_value_t* rocke_b_permlane64(rocke_ir_builder_t* b, rocke_value_t* src);
 rocke_value_t* rocke_b_alignbyte(rocke_ir_builder_t* b,
                                  rocke_value_t* a,

dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/ir.py

@@ -2812,6 +2812,64 @@ def dpp_xor(self, data: Value, xor_mask: int) -> Value:
             result_name_hint="dppx",
         ).result
 
+    def quad_perm(self, data: Value, perm) -> Value:
+        """Intra-quad ``v_mov_b32_dpp`` permutation on the VALU.
+
+        Lane ``4q + i`` reads ``data`` from lane ``4q + perm[i]``.
+        ``perm`` is encoded in the low eight bits of the DPP control word.
+
+        **Wave size.** The mapping is wave-size-independent: the same
+        control word applies within every four-lane group, and four
+        divides both 32 and 64, so a lane never addresses outside its own
+        quad. Wave size changes only the *number* of quads (8 in wave32,
+        16 in wave64), never the permutation a quad performs. Contrast
+        :meth:`warp_shuffle_xor`, whose ``lane_xor = 32`` partner is a
+        real lane in wave64 and does not exist in wave32.
+
+        That is a property of the quad, not a claim about every target:
+        the op still requires DPP-capable hardware. Base-DPP
+        ``quad_perm`` is available on CDNA, where the RDNA-only
+        ``row_xmask`` of :meth:`dpp_xor` is not.
+
+        The op carries no lane targeting -- the control word is broadcast
+        to every quad in the wave, with row and bank masks fixed at
+        ``15, 15`` (all enabled) by the lowerers. Selecting a subset of
+        quads is the caller's job.
+        """
+        perm = list(perm)
+        if len(perm) != 4 or any(not (0 <= p <= 3) for p in perm):
+            raise ValueError(f"quad_perm perm must be 4 values in 0..3, got {perm}")
+        if data.type.name != "i32":
+            raise ValueError("quad_perm requires i32 data")
+        ctrl = perm[0] | (perm[1] << 2) | (perm[2] << 4) | (perm[3] << 6)
+        return self._op(
+            "tile.quad_perm",
+            [data],
+            [I32],
+            attrs={"ctrl": int(ctrl)},
+            result_name_hint="qperm",
+        ).result
+
+    def warp_shuffle_xor_quad(self, v: Value, xor_mask: int) -> Value:
+        """XOR shuffle within a four-lane quad.
+
+        Masks 1 and 2 stay inside the quad and use :meth:`quad_perm`. Larger
+        masks require :meth:`warp_shuffle_xor`, which uses ``ds_swizzle``.
+        """
+        if xor_mask == 1:
+            perm = [1, 0, 3, 2]
+        elif xor_mask == 2:
+            perm = [2, 3, 0, 1]
+        else:
+            raise ValueError(
+                f"warp_shuffle_xor_quad supports xor_mask 1 or 2, got {xor_mask}"
+            )
+        if v.type.name == "f32":
+            return self.bitcast(self.quad_perm(self.bitcast(v, I32), perm), F32)
+        if v.type.name == "i32":
+            return self.quad_perm(v, perm)
+        raise ValueError(f"warp_shuffle_xor_quad: unsupported type {v.type.name}")
+
     def ds_bpermute_b64(self, addr: Value, data: Value) -> Value:
         """Packed 64-bit ``ds_bpermute`` — single LDS op for paired
         ``(val, idx)`` cross-lane shuffles (gfx9+).
@@ -4444,6 +4502,7 @@ def scf_if_else(self, cond: Value):
     "tile.ds_swizzle_xor",
     "tile.ds_swizzle",
     "tile.mov_dpp8",
+    "tile.quad_perm",
     "tile.wave_reduce",
     "tile.readlane",
     "tile.writelane",

dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py

@@ -1930,6 +1930,22 @@ def _op_tile_dpp_xor(self, op: Op) -> None:
             f"{_name(data)}, {_name(data)}, {dpp_ctrl}, 15, 15, 1);"
         )
 
+    def _op_tile_quad_perm(self, op: Op) -> None:
+        """Lower an eight-bit DPP quad-permute control word.
+
+        See :meth:`_op_tile_quad_perm` in ``lower_llvm.py``: ``ctrl``
+        packs four two-bit lane selectors, so ``0..255`` is the whole
+        legal range and out-of-range values are malformed IR.
+        """
+        (data,) = op.operands
+        ctrl = int(op.attrs["ctrl"])
+        if not 0 <= ctrl <= 255:
+            raise ValueError(f"tile.quad_perm: ctrl must be in 0..255, got {ctrl}")
+        self._emit(
+            f"int {_name(op.result)} = __builtin_amdgcn_update_dpp("
+            f"{_name(data)}, {_name(data)}, {ctrl}, 15, 15, 1);"
+        )
+
     def _op_tile_ds_swizzle_xor(self, op: Op) -> None:
         """``ds_swizzle_b32`` XOR butterfly via SWAP-mode encoding.
 

dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py

@@ -3788,6 +3788,26 @@ def _op_tile_dpp_xor(self, op: Op) -> None:
             f"i32 {dpp_ctrl}, i32 15, i32 15, i1 true)"
         )
 
+    def _op_tile_quad_perm(self, op: Op) -> None:
+        """Lower an eight-bit DPP quad-permute control word.
+
+        ``ctrl`` packs four two-bit lane selectors
+        (``p0 | p1 << 2 | p2 << 4 | p3 << 6``), so every value in
+        ``0..255`` is legal and anything outside it is malformed IR.
+        Reject rather than mask: truncation would turn an out-of-range
+        control into a different, silently valid permutation.
+        """
+        (data,) = op.operands
+        self._need("update.dpp.i32")
+        ctrl = int(op.attrs["ctrl"])
+        if not 0 <= ctrl <= 255:
+            raise ValueError(f"tile.quad_perm: ctrl must be in 0..255, got {ctrl}")
+        self._current().emit(
+            f"  {op.result.name} = call i32 @llvm.amdgcn.update.dpp.i32("
+            f"i32 {self._operand(data)}, i32 {self._operand(data)}, "
+            f"i32 {ctrl}, i32 15, i32 15, i1 true)"
+        )
+
     def _op_tile_ds_swizzle_xor(self, op: Op) -> None:
         """``ds_swizzle_b32`` with XOR butterfly via SWAP-mode encoding.
 

dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/runtime/launcher.py

@@ -95,7 +95,7 @@
 from typing import Any, Callable, Dict, Iterator, Mapping, Optional, Sequence, Tuple
 
 from .hip_module import Runtime
-from .packing import pack_args
+from .packing import compile_packer
 from .torch_interop import resolve_stream
 
 __all__ = [
@@ -420,6 +420,14 @@ def __init__(
         rt = _runtime()
         self._module = rt.load_module(hsaco)
         self._fn = self._module.get_function(kernel_name)
+        # Precompiled hot-path kernarg packer (signature is immutable for
+        # the launcher's lifetime). Byte-identical to ``pack_args``: it
+        # precomputes the fixed argument layout once here so a launch does
+        # not rebuild the offset table, re-dispatch on argument types, or
+        # re-assemble the format string. Note ``struct`` already caches
+        # recently used formats, so the saving is that surrounding work,
+        # not the format compile itself.
+        self._packer = compile_packer(self._signature)
 
     @property
     def kernel_name(self) -> str:
@@ -436,7 +444,7 @@ def __call__(
         config: LaunchConfig,
     ) -> LaunchSummary:
         rt = _runtime()
-        args = pack_args(self._signature, values)
+        args = self._packer(values)
         stream = resolve_stream(config.stream)
         fence = _resolved_fence(config.fence)
 

dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/runtime/packing.py

@@ -81,6 +81,67 @@ def pack_args(
     return struct.pack("".join(fmt_parts), *packed)
 
 
+def compile_packer(signature: Sequence[Mapping[str, Any]]):
+    """Precompile a hot-path kernarg packer for a fixed ``signature``.
+
+    Returns ``packer(values) -> bytes`` that is **byte-identical** to
+    :func:`pack_args` for the same signature, but hoists the invariant
+    work (type dispatch, alignment padding, format-string build, and
+    ``struct`` format compile) out of the per-call path, leaving only
+    the per-call work that genuinely varies: reading each value and
+    coercing it by kind. ``struct`` already caches recently used
+    formats, so the saving is the layout reconstruction around that
+    compile rather than the compile itself. The signature is immutable
+    per kernel, so a launcher can build this once at construction and
+    reuse it; the effect is proportionally largest where per-launch
+    host work is smallest.
+    """
+    _TY_FMT: Mapping[str, Tuple[str, int, int]] = {
+        "i32": ("i", 4, 4),
+        "i64": ("q", 8, 8),
+        "f32": ("f", 4, 4),
+    }
+    fmt_parts: List[str] = ["<"]
+    # plan: (name, kind) with kind in {"ptr", "f32", "int"} -- the only
+    # per-call work is reading values[name] and coercing by kind.
+    plan: List[Tuple[str, str]] = []
+    offset = 0
+    for arg in signature:
+        name = str(arg["name"])
+        ty = str(arg["type"])
+        if ty.startswith("ptr<"):
+            fmt_char, size, align, kind = "Q", 8, 8, "ptr"
+        elif ty in _TY_FMT:
+            fmt_char, size, align = _TY_FMT[ty]
+            kind = "f32" if ty == "f32" else "int"
+        else:
+            raise ValueError(f"unsupported kernel arg type {ty!r} for {name}")
+        pad = (-offset) % align
+        if pad:
+            fmt_parts.append(f"{pad}x")
+            offset += pad
+        fmt_parts.append(fmt_char)
+        plan.append((name, kind))
+        offset += size
+    packer_struct = struct.Struct("".join(fmt_parts))
+
+    def packer(values: Mapping[str, Any]) -> bytes:
+        out: List[Any] = []
+        for name, kind in plan:
+            if name not in values:
+                raise KeyError(f"missing kernel arg {name!r}")
+            v = values[name]
+            if kind == "ptr":
+                out.append(_as_ptr(v))
+            elif kind == "f32":
+                out.append(float(v))
+            else:
+                out.append(int(v))
+        return packer_struct.pack(*out)
+
+    return packer
+
+
 def pack_args_kernelparams(
     signature: Sequence[Mapping[str, Any]], values: Mapping[str, Any]
 ) -> List[Any]: