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

ISSUE ID : AICK-1513

## What this is

PR 1 of a 4-PR stack that splits the GDN gfx950 work (originally one
51-file PR) into reviewable layers. This is the foundation layer: it
adds a new crosslane hardware primitive and a launch-path speedup, with
**no GDN code** and no dependency on the other three PRs.

Stack (bottom → top): **quad_perm (this)** → dispatch.core hoist → GDN
decode → GDN prefill.

## What changed

**1. `quad_perm` — an intra-quad DPP permute intrinsic, added to both
engines.**
`quad_perm(data, [p0,p1,p2,p3])` lets lane `4q+i` read `data` from lane
`4q+perm[i]` on the VALU (via `v_mov_b32_dpp` /
`llvm.amdgcn.update.dpp.i32`), with no LDS crossbar and no `lgkmcnt`
wait. Also adds `warp_shuffle_xor_quad`, a fast path for the two
xor-masks that stay inside a quad: masks 1 and 2 lower to `quad_perm`,
and **any other mask is rejected** — callers needing a wider mask use
`warp_shuffle_xor`, which goes through `ds_swizzle`.

- Python engine: `core/ir.py` (builder + validation),
`core/lower_hip.py`, `core/lower_llvm.py`.
- C++ engine (live): `core/ir/ir_flow.cpp`, `core/ir/core_types.cpp`
(opcode + purity), `core/lower_hip/lower_hip_mma.cpp`,
`core/lower_llvm/crosslane.cpp`, `include/rocke/ir.h`.
- Coverage: Python `tests/test_rocke.py`, C++
`tests/core/future_intrinsic_lowering.cpp`, and cross-engine parity
`tests/instances/parity/target_intrinsics_emit.{c,py}` (identical IR
from both lowerers).

**2. Kernarg-packer precompile — a launch-path speedup for every
kernel.**
`runtime/packing.py` gains `compile_packer(signature)`;
`runtime/launcher.py` compiles the packer once at module-load and calls
it per launch instead of re-deriving the layout each time.
Byte-identical to `pack_args`.

What it removes is **kernarg layout reconstruction** — the
offset/alignment walk, the per-argument type dispatch, and the
format-string assembly. It is *not* a saving on format compilation:
CPython's `struct` module already caches recently used format strings,
so re-packing the same format is a cache lookup, not a recompile.

## Review round 1 — what changed since the first push

Four items were raised by @yraparti and @tenpercent. All four are
addressed.

**1. Lowerers accepted out-of-range `ctrl` (`bffb5fd911f`).** All four
`quad_perm` lowering sites masked the control word with `0xFF` instead
of validating it. `ctrl` packs four two-bit lane selectors (`p0 | p1<<2
| p2<<4 | p3<<6`), so `0..255` is the whole legal range — and masking
silently rewrote malformed IR into a *different, valid* permute: `256`
became `0` (`[0,0,0,0]`, a lane-0 broadcast) and `-1` became `255`
(`[3,3,3,3]`). A wrong reduction then computed wrong numbers instead of
failing.

The builders already validate selectors, but IR reaching a lowerer by
another route (deserialized, rewritten by a pass, hand-built) skips
them. Now rejected in the Python lowerers (`ValueError`) and the C++
ones (`ROCKE_ERR_VALUE`), with the mask dropped so the check cannot be
bypassed. Tests on both sides are **mutation-verified**: with the mask
restored, the Python subtests and the four C++ assertions fail.

**2. Wave-size semantics undocumented (`cfeec0a1315`, `8cea213942f`).**
`quad_perm` makes no wave-size assumption, and both engines now say so:
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 (8 in wave32, 16 in wave64).

Deliberately scoped — wave-size-independent is **not**
architecture-independent. The op still needs DPP-capable hardware; the
useful point is that base-DPP `quad_perm` is available on CDNA where
`dpp_xor`'s RDNA-only `row_xmask` is not. Also recorded: the op has no
lane targeting (the control is broadcast to every quad, row/bank masks
fixed at `15, 15`), so selecting a subset of quads is the caller's job.

A first draft of this docstring cited `dpp_xor` as an op whose partner
lane can leave the wave; `8cea213942f` corrects that. `dpp_xor` caps
`xor_mask` at `1..15` and its partner stays inside a 16-lane row, which
also divides both wave sizes. The accurate counterexample is
`warp_shuffle_xor` at `lane_xor = 32`.

**3. "The dominant Python cost" was unbacked (`786f6c827b0`).** Correct:
a packing-only microbenchmark cannot establish a share of launcher
overhead. Both sites now describe the mechanism instead of a magnitude
-- the packer precomputes the fixed argument layout once per launcher,
so a launch does not rebuild the offset table, re-dispatch on argument
types, or re-assemble the format string. `struct` already caches
recently used format strings, so the saving is that surrounding work
rather than the format compile itself.

The profiler written to measure the share now lives in its own PR
(#12168, draft), and performance claims are deferred until its gate is
fixed: its chunk-size-sensitivity check is a ratio, so a chunk size
large enough to saturate both comparison windows accepts a
back-pressured run.

**4. PR description misdescribed `warp_shuffle_xor_quad`.** It said the
helper "leaves wider masks on `ds_swizzle`", implying a dispatcher. It
is a specialist: masks other than 1 and 2 raise. Corrected in the
What-changed section above.

## Scope note (disclosed)

Item 2 is functionally independent of item 1. It is kept here rather
than split for one concrete reason: its byte-identity test
(`test_compile_packer_matches_pack_args_byte_for_byte`) lives in
`tests/test_rocke.py`, the same file that holds the `quad_perm` tests —
splitting would put two PRs on one test file. Reviewers should weigh the
packer's **global blast radius** (it changes the launch path for every
kernel), not only its decode benefit.

## Why it's safe

- No kernel IR changes — this only adds an opcode and a runtime fast
path.
- `quad_perm` is validated at the builder (selectors ∈ 0..3, i32-only)
**and now at every lowerer** (`ctrl` ∈ 0..255), and is marked pure in
both engines.
- `llvm.amdgcn.update.dpp.i32` is an already-shipped overload (used by
`mov_dpp8`/`ds_swizzle`), not a new intrinsic surface.
- The packer is asserted byte-identical to the existing `pack_args`.

## Why it's first in the stack

GDN decode's in-quad xor-butterfly reduction lowers to `quad_perm` (3
call sites), so this must land before the decode PR. It touches only
`platform/`, so it reviews without any GDN context.

## Verification (run, not planned)

Both engines built and exercised locally on this diff:

- Python, with the C++ extension importable: `pytest tests/test_rocke.py
-k "quad_perm or warp_shuffle_xor_quad or compile_packer"` → **8 passed,
0 skipped** (the cross-engine assertions no longer skip). Full file:
**291 passed, 14 skipped, 43 subtests**.
- C++ engine: `rocke_future_intrinsic_lowering` → **31 case(s) OK**.
- Cross-engine byte-identity: `tools/check_byte_identity.py --only
target_intrinsics` → **GREEN, configs=8, bad=0** — Python and C++ emit
identical `.ll`.
- `test_rocke_ci_static.py` **5 passed**.
- Formatters: `black` clean; `clang-format` applied (local binary is
v20; repo pre-commit pins v18.1.4, so CI may adjust whitespace).

GPU numerical execution of `quad_perm` itself and the full all-family
byte-identity gate were **not** run here.

## Stacking / base

#11807 has merged, and this PR is now based directly on `develop`.

Author: AviralGoelAMD

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

@@ -327,6 +327,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)
 {
@@ -1346,6 +1369,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,
@@ -1137,6 +1138,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

@@ -2845,6 +2845,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+).
@@ -4477,6 +4535,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

@@ -1971,6 +1971,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

@@ -3800,6 +3800,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]: