diff --git a/projects/hipblaslt/tensilelite/Tensile/Common/GlobalParameters.py b/projects/hipblaslt/tensilelite/Tensile/Common/GlobalParameters.py
index fe1a9b00089d..76c8110f3c3a 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Common/GlobalParameters.py
+++ b/projects/hipblaslt/tensilelite/Tensile/Common/GlobalParameters.py
@@ -572,6 +572,7 @@
     {"StreamK": [0]},
     {"StreamKForceDPOnly": [0]},
     {"StreamKAtomic": [0]},
+    {"StreamKWorkStealing": [0]},
     {"StreamKXCCMapping": [0]},
     {"StreamKFixupTreeReduction": [0]},
     {"DebugStreamK": [0]},
diff --git a/projects/hipblaslt/tensilelite/Tensile/Common/RequiredParameters.py b/projects/hipblaslt/tensilelite/Tensile/Common/RequiredParameters.py
index 1082950e4d79..732a2a55e342 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Common/RequiredParameters.py
+++ b/projects/hipblaslt/tensilelite/Tensile/Common/RequiredParameters.py
@@ -122,6 +122,7 @@ def getRequiredParametersMin() -> set:
         'StoreVectorWidth',
         'StreamK',
         'StreamKForceDPOnly',
+        'StreamKWorkStealing',
         'StreamKXCCMapping',
         'StreamKFixupTreeReduction',
         'SuppressNoLoadLoop',
diff --git a/projects/hipblaslt/tensilelite/Tensile/Common/ValidParameters.py b/projects/hipblaslt/tensilelite/Tensile/Common/ValidParameters.py
index 988b172950ec..67893009170f 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Common/ValidParameters.py
+++ b/projects/hipblaslt/tensilelite/Tensile/Common/ValidParameters.py
@@ -828,6 +828,14 @@ def makeValidMatrixInstructions():
     # 0: uses workspace to store partial tiles, accumulate in deterministic fix-up step
     # 1: uses atomics to accumulate partial tiles
     "StreamKAtomic": [0, 1],
+    # Codegen-time toggle for single-hop next-neighbor work stealing in the
+    # dynamic-queue StreamK fetch (SK4 / SK5-dynamic). Queue count =
+    # archCaps['NumXCD'] (8 on gfx942/gfx950). When a workgroup's home queue
+    # empties, it makes one atomic attempt on its next-neighbor per-XCD queue.
+    # Valid only for StreamK in (4, 5).
+    #  0: off
+    #  1: on
+    "StreamKWorkStealing": [0, 1],
     # Enables XCC-based remapping of workgroups, set the value to the number of XCCs
     # for the device/configuration being used
     #  0: uses default workgroup assignment
diff --git a/projects/hipblaslt/tensilelite/Tensile/Components/StreamK.py b/projects/hipblaslt/tensilelite/Tensile/Components/StreamK.py
index 7b5c93d7a346..cbb6d6156fac 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Components/StreamK.py
+++ b/projects/hipblaslt/tensilelite/Tensile/Components/StreamK.py
@@ -26,7 +26,7 @@
     VOP3PModifiers, ContinuousRegister, DSModifiers
 from rocisa.instruction import GlobalInv, GlobalWb, SAddCU32, SAddU32, SAndB32, SBarrier, \
     SBranch, SCBranchSCC0, SCBranchSCC1, SCMovB32, SCSelectB32, SCmpEQU32, SCmpEQU64, \
-    SCmpGtU32, SCmpLeU32, SCmpLtU32, SLShiftLeftB32, SLShiftLeftB64, SLShiftRightB32, VLShiftLeftB32, SLoadB32, \
+    SCmpGeU32, SCmpGtU32, SCmpLeU32, SCmpLtU32, SLShiftLeftB32, SLShiftLeftB64, SLShiftRightB32, VLShiftLeftB32, SLoadB32, \
     SMaxI32, SMinU32, SMovB32, SMovB64, SMulI32, SNop, SOrB32, SSleep, SStoreB32, SSubU32, \
     SWaitCnt, SWaitXCnt, VAddF32, VAddF64, VAddPKF16, VAddU32, VSubU32, VLShiftRightB32, VMovB32, \
     VReadfirstlaneB32, VCmpXEqU32, VCvtBF16toFP32, GlobalAtomicIncU32Saddr, BufferLoadB32, BufferStoreB32, \
@@ -382,6 +382,145 @@ def _skv(self, writer, name):
         """Return the VGPR index holding a StreamK constant."""
         return writer.states.skConstVgprs[name]
 
+    # ------------------------------------------------------------------
+    # Single-hop next-neighbor work stealing (codegen-time, off by default)
+    #
+    # Queues are one per XCD: numQueues = archCaps["NumXCD"], a power of two so
+    # queue mapping uses shift/AND fast masking (queueIdx = StreamKIdx & mask).
+    # A queue whose home fetch is empty steals once from its next neighbor
+    # s = (q+1) & mask; each queue has exactly one predecessor p = (q-1) & mask.
+    #
+    # AddressFlags buffer layout (per problem):
+    #   [0, numQueues*stride)     per-queue counters, one per cache line
+    #   [numQueues*stride, ...)   partials/fixup ready flags (one word per tile)
+    # Counter stride == archCaps["CacheLineBytes"] (128B on gfx942/gfx950), so
+    # each per-XCD counter sits on its own line. Each counter's atomic_inc uses a
+    # static predecessor-inclusive auto-reset bound, so it self-zeroes every
+    # launch -- there is no explicit end-of-kernel reset.
+
+    def _wsQueueConstants(self, writer, kernel):
+        """Return (numQueues, mask, log2Queues, cacheLineLog2) for this arch.
+
+        ``numQueues`` = archCaps["NumXCD"] and the counter stride =
+        archCaps["CacheLineBytes"]; both must be powers of two for the shift/AND
+        queue masking and queue-address shift to be valid (asserted below).
+        """
+        numQueues = writer.states.archCaps["NumXCD"]
+        assert numQueues > 0 and (numQueues & (numQueues - 1)) == 0, (
+            "StreamK dynamic-queue fast masking requires a power-of-two queue count "
+            "(got %d for ISA %s)" % (numQueues, tuple(kernel["ISA"][:2])))
+        strideBytes = writer.states.archCaps["CacheLineBytes"]
+        assert strideBytes > 0 and (strideBytes & (strideBytes - 1)) == 0, (
+            "StreamK per-queue counter stride must be a power-of-two cache-line "
+            "size (got %d for ISA %s)" % (strideBytes, tuple(kernel["ISA"][:2])))
+        return numQueues, numQueues - 1, log2(numQueues), log2(strideBytes)
+
+    def _wsFlagsBaseOffset(self, writer, kernel):
+        """Byte offset where the partials/fixup ready flags begin.
+
+        The flags region starts right after the per-queue counters, i.e. after
+        ``numQueues * strideBytes`` bytes (8 * 128 = 1024 on gfx942/gfx950).
+        """
+        numQueues, _, _, cacheLineLog2 = self._wsQueueConstants(writer, kernel)
+        return numQueues << cacheLineLog2
+
+    def _wsStructuralCount(self, mod, mask, log2Queues, sDst, sTotal, sQueue, sTmp, comment):
+        """Emit sDst = (sTotal >> log2Queues) + [sQueue < (sTotal & mask)].
+
+        Reuses the shift/and(mask)/cmp/cselect idiom already used for
+        tilesInQueue / workgroupsInQueue in graWorkGroup so the per-queue
+        structural share (tiles or workgroups) can be recomputed for an
+        arbitrary queue index. ``sTmp`` is a caller-owned scratch SGPR.
+        """
+        mod.add(SLShiftRightB32(dst=sgpr(sDst), src=sgpr(sTotal), shiftHex=log2Queues, comment=comment))
+        mod.add(SAndB32(dst=sgpr(sTmp), src0=sgpr(sTotal), src1=mask, comment="Remainder"))
+        mod.add(SCmpLtU32(src0=sgpr(sQueue), src1=sgpr(sTmp), comment="Queue gets a structural extra?"))
+        mod.add(SCSelectB32(dst=sgpr(sTmp), src0=1, src1=0))
+        mod.add(SAddU32(dst=sgpr(sDst), src0=sgpr(sDst), src1=sgpr(sTmp)))
+
+    def streamKWorkStealingHomeBound(self, writer, mod, kernel, sBound, sQueueIdx, sGrid):
+        """Fold the predecessor's workgroup count into the home auto-reset bound.
+
+        Adds the predecessor term W_p to the ``tiles_q + W_q - 1`` already in
+        ``sBound``, giving the stealing bound ``tiles_q + W_q + W_p - 1`` (queue q
+        also absorbs W_p increments from its one predecessor p = (q-1) & mask).
+        Caller gates on kernel["StreamKWorkStealing"] and passes the grid SGPR
+        name ("skGrid" for SK4, "SKGrid" for SK5-dynamic); ``sQueueIdx`` is
+        preserved. Exact only when W_q >= 1 whenever tiles_q >= 1 (skGrid >=
+        numQueues); the Solution layer rejects debug overrides that break this.
+        """
+        _, mask, log2Queues, _ = self._wsQueueConstants(writer, kernel)
+        sPred = writer.sgprPool.checkOut(1, "wsPredQueue")
+        sWp = writer.sgprPool.checkOut(1, "wsPredWorkgroups")
+        sTmp = writer.sgprPool.checkOut(1, "wsPredTmp")
+        # p = (q - 1) & mask  (wraps 0 -> numQueues-1 for unsigned subtract)
+        mod.add(SSubU32(dst=sgpr(sPred), src0=sgpr(sQueueIdx), src1=1, comment="Predecessor queue (q-1)"))
+        mod.add(SAndB32(dst=sgpr(sPred), src0=sgpr(sPred), src1=mask, comment="Wrap predecessor index"))
+        # W_p = (skGrid >> log2) + [p < (skGrid & mask)]
+        self._wsStructuralCount(mod, mask, log2Queues, sWp, sGrid, sPred, sTmp,
+                                comment="Predecessor workgroups W_(q-1)")
+        mod.add(SAddU32(dst=sgpr(sBound), src0=sgpr(sBound), src1=sgpr(sWp),
+                        comment="Home auto-reset bound += predecessor workgroups (next-neighbor steal)"))
+        writer.sgprPool.checkIn(sTmp)
+        writer.sgprPool.checkIn(sWp)
+        writer.sgprPool.checkIn(sPred)
+
+    def streamKWorkStealingSteal(self, writer, mod, kernel, sQueueIdx, sWorkItemIdx, sGrid, mkLabel):
+        """Single-hop next-neighbor steal on the per-XCD queue topology.
+
+        On entry sQueueIdx holds home queue q and sWorkItemIdx holds the home
+        fetch result (both live). If the home fetch was valid (index <
+        TotalItems) this is a no-op; otherwise one s_atomic_inc steals from the
+        next neighbor s = (q+1) & mask and the global tile index is recomputed
+        from s. A lost race leaves sWorkItemIdx >= TotalItems, so the downstream
+        valid-index check turns this WG into a no-op. sQueueIdx is clobbered
+        (advanced to s). Caller gates on kernel["StreamKWorkStealing"] and passes
+        the grid SGPR name ("skGrid" for SK4, "SKGrid" for SK5-dynamic). The
+        steal atomic uses the stolen queue's bound ``tiles_s + W_s + W_q - 1``.
+        """
+        _, mask, log2Queues, cacheLineLog2 = self._wsQueueConstants(writer, kernel)
+        skFetchDone = mkLabel("SK_FetchDone")
+        mod.add(SCmpLtU32(src0=sgpr(sWorkItemIdx), src1=sgpr("TotalItems"), comment="Home fetch valid?"))
+        mod.add(SCBranchSCC1(labelName=skFetchDone.getLabelName(), comment="Valid work fetched; no steal"))
+
+        # Build the steal auto-reset bound tiles_s + W_s + W_q - 1 into
+        # sWorkItemIdx (dead here). W_q is the stealer's own workgroup count and
+        # must be computed while sQueueIdx still holds q, before advancing to s.
+        sTmp = writer.sgprPool.checkOut(1, "wsStealTmp")
+        sWq = writer.sgprPool.checkOut(1, "wsStealerWorkgroups")
+        self._wsStructuralCount(mod, mask, log2Queues, sWq, sGrid, sQueueIdx, sTmp,
+                                comment="Stealer workgroups W_q")
+
+        # Walk to the immediate next queue (wrap within the per-XCD queues, single-hop next-neighbor).
+        mod.add(SAddU32(dst=sgpr(sQueueIdx), src0=sgpr(sQueueIdx), src1=1, comment="Next queue"))
+        mod.add(SAndB32(dst=sgpr(sQueueIdx), src0=sgpr(sQueueIdx), src1=mask, comment="Wrap queue index"))
+
+        # tiles_s into sWorkItemIdx, then += W_s and += W_q, then -1.
+        self._wsStructuralCount(mod, mask, log2Queues, sWorkItemIdx, "TotalItems", sQueueIdx, sTmp,
+                                comment="Stolen-queue tiles tiles_s")
+        sWs = writer.sgprPool.checkOut(1, "wsStolenWorkgroups")
+        self._wsStructuralCount(mod, mask, log2Queues, sWs, sGrid, sQueueIdx, sTmp,
+                                comment="Stolen-queue workgroups W_s")
+        mod.add(SAddU32(dst=sgpr(sWorkItemIdx), src0=sgpr(sWorkItemIdx), src1=sgpr(sWs), comment="tiles_s + W_s"))
+        writer.sgprPool.checkIn(sWs)
+        mod.add(SAddU32(dst=sgpr(sWorkItemIdx), src0=sgpr(sWorkItemIdx), src1=sgpr(sWq), comment="+ W_q (stealer)"))
+        mod.add(SSubU32(dst=sgpr(sWorkItemIdx), src0=sgpr(sWorkItemIdx), src1=1, comment="Steal auto-reset bound"))
+        writer.sgprPool.checkIn(sWq)
+        writer.sgprPool.checkIn(sTmp)
+
+        # One atomic on the neighbor's counter with the static self-reset bound.
+        sAddress = writer.sgprPool.checkOutAligned(2, 2, "wsStealAddress")
+        mod.add(SLShiftLeftB32(dst=sgpr(sAddress), src=sgpr(sQueueIdx), shiftHex=cacheLineLog2, comment="Stride queues to cache lines (stolen queue)"))
+        mod.add(SAddU32(dst=sgpr(sAddress+0), src0=sgpr(sAddress+0), src1=sgpr("AddressFlags+0")))
+        mod.add(SAddCU32(dst=sgpr(sAddress+1), src0=0, src1=sgpr("AddressFlags+1")))
+        mod.add(SAtomicInc(dst=sgpr(sWorkItemIdx), base=sgpr(sAddress, 2), soffset=0, smem=SMEMModifiers(glc=True), comment="Fetch stolen work item index"))
+        mod.add(SWaitCnt(kmcnt=0, comment="Wait for scalar memory op"))
+        writer.sgprPool.checkIn(sAddress)
+        # Recompute global tile index from the neighbor's queue.
+        mod.add(SLShiftLeftB32(dst=sgpr(sWorkItemIdx), src=sgpr(sWorkItemIdx), shiftHex=log2Queues))
+        mod.add(SAddU32(dst=sgpr(sWorkItemIdx), src0=sgpr(sWorkItemIdx), src1=sgpr(sQueueIdx)))
+        mod.add(skFetchDone)
+
     @abc.abstractmethod
     def preLoop(self, writer, kernel):
         pass
@@ -1354,7 +1493,7 @@ def partialsWriteProcedure(self, writer, kernel, vectorWidths, elements, alpha,
                 # TODO modularize this section into abstract function
                 module.add(self.calculatePartialIdx(tmpSgpr))
                 module.add(SLShiftLeftB32(dst=sgpr(tmpSgpr), src=sgpr(tmpSgpr), shiftHex=log2(4), comment="flag offset based on partial index"))
-                module.add(SAddU32(dst=sgpr(tmpSgpr), src0=sgpr(tmpSgpr), src1=(256*8), comment="Offset flags to come after the work queues"))
+                module.add(SAddU32(dst=sgpr(tmpSgpr), src0=sgpr(tmpSgpr), src1=self._wsFlagsBaseOffset(writer, kernel), comment="Offset flags to come after the work queues"))
             elif kernel["StreamK"] == 5:
                 # SK5 hybrid: dispatch on StreamKHybridMode bit
                 # (0 = static SK3 -> use StreamKIdx, 1 = dynamic SK4 -> use calculatePartialIdx).
@@ -1368,7 +1507,7 @@ def partialsWriteProcedure(self, writer, kernel, vectorWidths, elements, alpha,
                 module.add(self.calculatePartialIdx(tmpSgpr))
                 module.add(SLShiftLeftB32(dst=sgpr(tmpSgpr), src=sgpr(tmpSgpr), shiftHex=log2(4),
                                           comment="SK5/SK4: flag offset based on partial index"))
-                module.add(SAddU32(dst=sgpr(tmpSgpr), src0=sgpr(tmpSgpr), src1=(256*8),
+                module.add(SAddU32(dst=sgpr(tmpSgpr), src0=sgpr(tmpSgpr), src1=self._wsFlagsBaseOffset(writer, kernel),
                                    comment="SK5/SK4: offset flags to come after the work queues"))
                 module.add(SBranch(labelName=sk5FlagDone.getLabelName(),
                                    comment="SK5: skip static flag offset"))
@@ -2916,6 +3055,9 @@ def preLoop(self, writer, kernel):
             module.add(SLShiftRightB32(dst=sgpr("WorkGroup2"), shiftHex=hex(0x10), src="ttmp7", comment="workaround"))
 
         module.add(SMovB32(dst=sgpr("StreamKIdx"), src=sgpr("WorkGroup0"), comment="Save original StreamK index"))
+        # Work stealing: this WG has not yet seen its home queue empty.
+        if kernel["StreamKWorkStealing"]:
+            module.add(SMovB32(dst=sgpr("StreamKStickyEmpty"), src=0, comment="WS: home not yet empty"))
         # Two-tile SK (DP first)
         # Do DP tiles before SK
         skInitDone = Label("SK_InitDone", "")
@@ -2952,23 +3094,27 @@ def graWorkGroup(self, writer, kernel, tPA, tPB):
         module.add(SCBranchSCC0(labelName=skSkipWorkItem.getLabelName(), comment="Skip work item"))
         writer.sgprPool.checkIn(sWave)
 
+        # Per-arch dynamic-queue fast-mask constants: log2(numQueues) for the
+        # StreamKIdx/queue divisions, log2(cache-line size) for the counter stride.
+        _, _, wsLog2Queues, wsCacheLineLog2 = self._wsQueueConstants(writer, kernel)
+
         # Default queue index
         sQueueIdx = writer.sgprPool.checkOut(1, "QueueIdx")
-        module.add(SLShiftRightB32(dst=sgpr(sQueueIdx), src=sgpr("StreamKIdx"), shiftHex=log2(8)))
-        module.add(SLShiftLeftB32(dst=sgpr(sQueueIdx), src=sgpr(sQueueIdx), shiftHex=log2(8)))
+        module.add(SLShiftRightB32(dst=sgpr(sQueueIdx), src=sgpr("StreamKIdx"), shiftHex=wsLog2Queues))
+        module.add(SLShiftLeftB32(dst=sgpr(sQueueIdx), src=sgpr(sQueueIdx), shiftHex=wsLog2Queues))
         module.add(SSubU32(dst=sgpr(sQueueIdx), src0=sgpr("StreamKIdx"), src1=sgpr(sQueueIdx), comment="Default queue index"))
 
         # Queue address
         sAddress = writer.sgprPool.checkOutAligned(2, 2, "Address")
-        module.add(SLShiftLeftB32(dst=sgpr(sAddress), src=sgpr(sQueueIdx), shiftHex=log2(256), comment="Stride queues to different cache lines"))
+        module.add(SLShiftLeftB32(dst=sgpr(sAddress), src=sgpr(sQueueIdx), shiftHex=wsCacheLineLog2, comment="Stride queues to different cache lines"))
         module.add(SAddU32(dst=sgpr(sAddress+0), src0=sgpr(sAddress+0), src1=sgpr("AddressFlags+0")))
         module.add(SAddCU32(dst=sgpr(sAddress+1), src0=0, src1=sgpr("AddressFlags+1")))
 
         # Tiles in queue
         sTilesInQueue = writer.sgprPool.checkOut(1, "tilesInQueue")
-        module.add(SLShiftRightB32(dst=sgpr(sTilesInQueue), src=sgpr("TotalItems"), shiftHex=log2(8)))
+        module.add(SLShiftRightB32(dst=sgpr(sTilesInQueue), src=sgpr("TotalItems"), shiftHex=wsLog2Queues))
         sRemainder = writer.sgprPool.checkOut(1, "remainder tiles")
-        module.add(SLShiftLeftB32(dst=sgpr(sRemainder), src=sgpr(sTilesInQueue), shiftHex=log2(8)))
+        module.add(SLShiftLeftB32(dst=sgpr(sRemainder), src=sgpr(sTilesInQueue), shiftHex=wsLog2Queues))
         module.add(SSubU32(dst=sgpr(sRemainder), src0=sgpr("TotalItems"), src1=sgpr(sRemainder), comment="Remainder tiles"))
         module.add(SCmpLtU32(src0=sgpr(sQueueIdx), src1=sgpr(sRemainder), comment="Check if queue gets an extra tile"))
         module.add(SCSelectB32(dst=sgpr(sRemainder), src0=1, src1=0))
@@ -2977,9 +3123,9 @@ def graWorkGroup(self, writer, kernel, tPA, tPB):
 
         # Workgroups in queue
         sWorkgroupsInQueue = writer.sgprPool.checkOut(1, "workgroupsInQueue")
-        module.add(SLShiftRightB32(dst=sgpr(sWorkgroupsInQueue), src=sgpr("skGrid"), shiftHex=log2(8)))
+        module.add(SLShiftRightB32(dst=sgpr(sWorkgroupsInQueue), src=sgpr("skGrid"), shiftHex=wsLog2Queues))
         sRemainder = writer.sgprPool.checkOut(1, "remainder workgroups")
-        module.add(SLShiftLeftB32(dst=sgpr(sRemainder), src=sgpr(sWorkgroupsInQueue), shiftHex=log2(8)))
+        module.add(SLShiftLeftB32(dst=sgpr(sRemainder), src=sgpr(sWorkgroupsInQueue), shiftHex=wsLog2Queues))
         module.add(SSubU32(dst=sgpr(sRemainder), src0=sgpr("skGrid"), src1=sgpr(sRemainder), comment="Remainder workgroups"))
         module.add(SCmpLtU32(src0=sgpr(sQueueIdx), src1=sgpr(sRemainder), comment="Check if queue gets an extra tile"))
         module.add(SCSelectB32(dst=sgpr(sRemainder), src0=1, src1=0))
@@ -2993,13 +3139,39 @@ def graWorkGroup(self, writer, kernel, tPA, tPB):
         writer.sgprPool.checkIn(sTilesInQueue)
         writer.sgprPool.checkIn(sWorkgroupsInQueue)
 
+        # Work stealing: fold the predecessor's workgroup count into the home
+        # auto-reset bound so the counter still self-resets under next-neighbor stealing.
+        if kernel["StreamKWorkStealing"]:
+            self.streamKWorkStealingHomeBound(writer, module, kernel, sWorkItemIdx, sQueueIdx, "skGrid")
+
+        # Work stealing: once this WG has seen its home queue empty (sticky), it
+        # never touches the home counter again -- skip the home fetch and force
+        # the steal path with an invalid sentinel index (>= TotalItems).
+        if kernel["StreamKWorkStealing"]:
+            skStealOnly = Label(writer.labels.getNameInc("SK_StealOnly"), "")
+            skHomeFetched = Label(writer.labels.getNameInc("SK_HomeFetched"), "")
+            module.add(SCmpEQU32(src0=sgpr("StreamKStickyEmpty"), src1=0, comment="Home not yet empty?"))
+            module.add(SCBranchSCC0(labelName=skStealOnly.getLabelName(), comment="Sticky: skip home fetch, steal only"))
+
         # Fetch next work item
         module.add(self._fetchNextWorkItem(writer, kernel, sWorkItemIdx, sAddress))
         writer.sgprPool.checkIn(sAddress)
 
         # Convert to global work item index
-        module.add(SLShiftLeftB32(dst=sgpr(sWorkItemIdx), src=sgpr(sWorkItemIdx), shiftHex=log2(8)))
+        module.add(SLShiftLeftB32(dst=sgpr(sWorkItemIdx), src=sgpr(sWorkItemIdx), shiftHex=wsLog2Queues))
         module.add(SAddU32(dst=sgpr(sWorkItemIdx), src0=sgpr(sWorkItemIdx), src1=sgpr(sQueueIdx)))
+
+        # Work stealing: latch the sticky-empty flag on the first empty home
+        # fetch, then fall through to the steal (or, when already sticky, jump
+        # straight to the steal with the sentinel index).
+        if kernel["StreamKWorkStealing"]:
+            module.add(SCmpGeU32(src0=sgpr(sWorkItemIdx), src1=sgpr("TotalItems"), comment="Home fetch empty?"))
+            module.add(SCSelectB32(dst=sgpr("StreamKStickyEmpty"), src0=1, src1=0, comment="Latch sticky-empty on empty home"))
+            module.add(SBranch(labelName=skHomeFetched.getLabelName(), comment="Home fetched; try one steal"))
+            module.add(skStealOnly)
+            module.add(SMovB32(dst=sgpr(sWorkItemIdx), src=sgpr("TotalItems"), comment="Sentinel index (>= TotalItems) forces steal"))
+            module.add(skHomeFetched)
+            self.streamKWorkStealingSteal(writer, module, kernel, sQueueIdx, sWorkItemIdx, "skGrid", lambda base: Label(writer.labels.getNameInc(base), ""))
         writer.sgprPool.checkIn(sQueueIdx)
 
         # Share work item index with all waves
@@ -3194,7 +3366,7 @@ def storeBranches(self, writer, kernel, skPartialsLabel, vectorWidths, elements,
 
             # Check flag
             module.add(SLShiftLeftB32(dst=sgpr(tmpSgpr), src=sgpr(sPartialIdx), shiftHex=log2(4), comment="flag offset based on partial index"))
-            module.add(SAddU32(dst=sgpr(tmpSgpr), src0=sgpr(tmpSgpr), src1=(256*8), comment="Offset flags to come after the work queues"))
+            module.add(SAddU32(dst=sgpr(tmpSgpr), src0=sgpr(tmpSgpr), src1=self._wsFlagsBaseOffset(writer, kernel), comment="Offset flags to come after the work queues"))
             module.add(SLoadB32(dst=sgpr(tmpSgpr+2), base=sgpr("AddressFlags", 2), soffset=sgpr(tmpSgpr), smem=SMEMModifiers(glc=True, dlc=True, scope=CacheScope.SCOPE_DEV), comment="get flag"))
 
             module.add(SWaitCnt(kmcnt=0, comment="wait for flag load"))
@@ -3278,10 +3450,7 @@ def routeToGeneralBatchedOrStridedBatched(self, stridedBatchedGemmLoad, generalB
     def kernelEnd(self, writer, kernel):
         module = Module("StreamK Dynamic kernelEnd")
 
-        # We don't need to track completed kernels if we know total tiles and grid size
-        # Reset is baked into the atomic_inc at the top of the loop
-        # TODO will need to reset the rest of the synchronizer if tiles were split
-        # Remaining reset can be done if workitem = grid + total - 1
+        # Per-queue atomic_inc auto-resets; no kernelEnd reset needed.
 
         return module
 
@@ -3387,6 +3556,10 @@ def preLoop(self, writer, kernel):
         # so save directly to the StreamKIdx SGPR (no VGPR-cache path).
         module.add(SMovB32(dst=sgpr("StreamKIdx"), src=sgpr("WorkGroup0"),
                            comment="SK5: save original StreamK index"))
+        # Work stealing: this WG has not yet seen its home queue empty.
+        if kernel["StreamKWorkStealing"]:
+            module.add(SMovB32(dst=sgpr("StreamKStickyEmpty"), src=0,
+                               comment="WS: home not yet empty"))
 
         # ----- Extract the mode bit once for the whole kernel -----
         module.add(self._emitModeExtraction(writer, kernel))
@@ -3576,19 +3749,23 @@ def emitDynamicGRA(mod):
                                     comment="Skip work item"))
             writer.sgprPool.checkIn(sWave)
 
+            # Per-arch dynamic-queue fast-mask constants: log2(numQueues) for the
+            # StreamKIdx/queue divisions, log2(cache-line size) for the counter stride.
+            _, _, wsLog2Queues, wsCacheLineLog2 = self._wsQueueConstants(writer, kernel)
+
             # Default queue index
             sQueueIdx = writer.sgprPool.checkOut(1, "QueueIdx")
             mod.add(SLShiftRightB32(dst=sgpr(sQueueIdx), src=sgpr("StreamKIdx"),
-                                       shiftHex=log2(8)))
+                                       shiftHex=wsLog2Queues))
             mod.add(SLShiftLeftB32(dst=sgpr(sQueueIdx), src=sgpr(sQueueIdx),
-                                      shiftHex=log2(8)))
+                                      shiftHex=wsLog2Queues))
             mod.add(SSubU32(dst=sgpr(sQueueIdx), src0=sgpr("StreamKIdx"),
                                src1=sgpr(sQueueIdx), comment="Default queue index"))
 
             # Queue address
             sAddress = writer.sgprPool.checkOutAligned(2, 2, "Address")
             mod.add(SLShiftLeftB32(dst=sgpr(sAddress), src=sgpr(sQueueIdx),
-                                      shiftHex=log2(256),
+                                      shiftHex=wsCacheLineLog2,
                                       comment="Stride queues to different cache lines"))
             mod.add(SAddU32(dst=sgpr(sAddress+0), src0=sgpr(sAddress+0),
                                src1=sgpr("AddressFlags+0")))
@@ -3597,10 +3774,10 @@ def emitDynamicGRA(mod):
             # Tiles in queue
             sTilesInQueue = writer.sgprPool.checkOut(1, "tilesInQueue")
             mod.add(SLShiftRightB32(dst=sgpr(sTilesInQueue), src=sgpr("TotalItems"),
-                                       shiftHex=log2(8)))
+                                       shiftHex=wsLog2Queues))
             sRemainder = writer.sgprPool.checkOut(1, "remainder tiles")
             mod.add(SLShiftLeftB32(dst=sgpr(sRemainder), src=sgpr(sTilesInQueue),
-                                      shiftHex=log2(8)))
+                                      shiftHex=wsLog2Queues))
             mod.add(SSubU32(dst=sgpr(sRemainder), src0=sgpr("TotalItems"),
                                src1=sgpr(sRemainder), comment="Remainder tiles"))
             mod.add(SCmpLtU32(src0=sgpr(sQueueIdx), src1=sgpr(sRemainder),
@@ -3614,10 +3791,10 @@ def emitDynamicGRA(mod):
             sWorkgroupsInQueue = writer.sgprPool.checkOut(1, "workgroupsInQueue")
             # SK5: SKGrid is the SK4-dedicated grid SGPR (uppercase).
             mod.add(SLShiftRightB32(dst=sgpr(sWorkgroupsInQueue), src=sgpr("SKGrid"),
-                                       shiftHex=log2(8)))
+                                       shiftHex=wsLog2Queues))
             sRemainder = writer.sgprPool.checkOut(1, "remainder workgroups")
             mod.add(SLShiftLeftB32(dst=sgpr(sRemainder), src=sgpr(sWorkgroupsInQueue),
-                                      shiftHex=log2(8)))
+                                      shiftHex=wsLog2Queues))
             mod.add(SSubU32(dst=sgpr(sRemainder), src0=sgpr("SKGrid"),
                                src1=sgpr(sRemainder), comment="Remainder workgroups"))
             mod.add(SCmpLtU32(src0=sgpr(sQueueIdx), src1=sgpr(sRemainder),
@@ -3635,14 +3812,50 @@ def emitDynamicGRA(mod):
             writer.sgprPool.checkIn(sTilesInQueue)
             writer.sgprPool.checkIn(sWorkgroupsInQueue)
 
+            # Work stealing: fold the predecessor's workgroup count into the
+            # home auto-reset bound so the counter still self-resets under
+            # next-neighbor stealing.
+            if kernel["StreamKWorkStealing"]:
+                self.streamKWorkStealingHomeBound(writer, mod, kernel, sWorkItemIdx,
+                                                  sQueueIdx, "SKGrid")
+
+            # Work stealing: once this WG has seen its home queue empty (sticky),
+            # never touch the home counter again -- skip the home fetch and force
+            # the steal path with an invalid sentinel index (>= TotalItems).
+            if kernel["StreamKWorkStealing"]:
+                skStealOnly = Label(writer.labels.getNameInc("SK_StealOnly"), "")
+                skHomeFetched = Label(writer.labels.getNameInc("SK_HomeFetched"), "")
+                mod.add(SCmpEQU32(src0=sgpr("StreamKStickyEmpty"), src1=0,
+                                  comment="Home not yet empty?"))
+                mod.add(SCBranchSCC0(labelName=skStealOnly.getLabelName(),
+                                     comment="Sticky: skip home fetch, steal only"))
+
             mod.add(self._fetchNextWorkItem(writer, kernel, sWorkItemIdx, sAddress))
             writer.sgprPool.checkIn(sAddress)
 
             # Convert to global work item index
             mod.add(SLShiftLeftB32(dst=sgpr(sWorkItemIdx), src=sgpr(sWorkItemIdx),
-                                      shiftHex=log2(8)))
+                                      shiftHex=wsLog2Queues))
             mod.add(SAddU32(dst=sgpr(sWorkItemIdx), src0=sgpr(sWorkItemIdx),
                                src1=sgpr(sQueueIdx)))
+
+            # Work stealing: latch the sticky-empty flag on the first empty home
+            # fetch, then fall through to one steal (or, when already sticky,
+            # jump straight to the steal with the sentinel index).
+            if kernel["StreamKWorkStealing"]:
+                mod.add(SCmpGeU32(src0=sgpr(sWorkItemIdx), src1=sgpr("TotalItems"),
+                                  comment="Home fetch empty?"))
+                mod.add(SCSelectB32(dst=sgpr("StreamKStickyEmpty"), src0=1, src1=0,
+                                    comment="Latch sticky-empty on empty home"))
+                mod.add(SBranch(labelName=skHomeFetched.getLabelName(),
+                                comment="Home fetched; try one steal"))
+                mod.add(skStealOnly)
+                mod.add(SMovB32(dst=sgpr(sWorkItemIdx), src=sgpr("TotalItems"),
+                                comment="Sentinel index (>= TotalItems) forces steal"))
+                mod.add(skHomeFetched)
+                self.streamKWorkStealingSteal(writer, mod, kernel, sQueueIdx, sWorkItemIdx,
+                                              "SKGrid",
+                                              lambda base: Label(writer.labels.getNameInc(base), ""))
             writer.sgprPool.checkIn(sQueueIdx)
 
             # Share work item index with all waves
@@ -3974,7 +4187,7 @@ def emitDynamicStore(mod):
                 mod.add(SLShiftLeftB32(dst=sgpr(tmpSgpr), src=sgpr(sPartialIdx),
                                        shiftHex=log2(4),
                                        comment="flag offset based on partial index"))
-                mod.add(SAddU32(dst=sgpr(tmpSgpr), src0=sgpr(tmpSgpr), src1=(256*8),
+                mod.add(SAddU32(dst=sgpr(tmpSgpr), src0=sgpr(tmpSgpr), src1=self._wsFlagsBaseOffset(writer, kernel),
                                 comment="Offset flags to come after the work queues"))
                 mod.add(SLoadB32(dst=sgpr(tmpSgpr+2), base=sgpr("AddressFlags", 2),
                                  soffset=sgpr(tmpSgpr),
@@ -4083,6 +4296,9 @@ def routeToGeneralBatchedOrStridedBatched(self, stridedBatchedGemmLoad, generalB
 
     def kernelEnd(self, writer, kernel):
         module = Module("StreamK Hybrid kernelEnd")
+
+        # Per-queue atomic_inc auto-resets; no kernelEnd reset needed.
+
         return module
 
 
diff --git a/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py b/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py
index 83587a108a43..8e72cf02c502 100644
--- a/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py
+++ b/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py
@@ -5256,6 +5256,11 @@ def kernelBodySubtile(self, kernel, tensorParametersA, tensorParametersB):
       module.add(kernelEndLabel)
       if kernel["ProblemType"]["OutputAmaxD"]:
         module.add(self.insertAmaxD(kernel))
+      # Mirror functionEnd's StreamK kernelEnd on the deferred-blocks epilogue.
+      # Single-hop next-neighbor work stealing self-resets its per-queue counters via the
+      # atomic_inc auto-reset bounds, so kernelEnd emits no explicit reset.
+      skComponent = Component.StreamK.find(self)
+      module.add(skComponent.kernelEnd(self, kernel))
       module.add(SEndpgm(comment="Kernel End"))
     else:
       # If activation was deferred but no other deferred blocks exist, emit it before functionEnd
@@ -9474,6 +9479,11 @@ def vgprAllocationImplSubtile():
         "StreamKLocalStart",
         "StreamKLocalEnd",
       ]
+      # Work stealing: per-WG sticky-empty flag. Persists across the persistent
+      # loop back-edge (added to nonPostLoopSgpr below) so each WG only touches
+      # its home counter for its valid dispenses + exactly one empty fetch.
+      if kernel["StreamKWorkStealing"]:
+        requiredUnalignedSgprVar.append("StreamKStickyEmpty")
       if kernel["StreamKAtomic"] == 0:
         requiredAligned4SgprVar.append("SrdWS")
     elif kernel["StreamK"] == 5:
@@ -9495,6 +9505,10 @@ def vgprAllocationImplSubtile():
         "StreamKLocalEnd",
         "StreamKHybridMode",
       ]
+      # Work stealing: per-WG sticky-empty flag (see SK4 note above). Only the
+      # SK4 sub-path uses it, but it must persist for the whole persistent loop.
+      if kernel["StreamKWorkStealing"]:
+        requiredUnalignedSgprVar.append("StreamKStickyEmpty")
       if len(kernel["SpaceFillingAlgo"]):
         requiredUnalignedSgprVar.append("StreamKTileID")
       if kernel["StreamKAtomic"] == 0:
diff --git a/projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py b/projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py
index f1074dee9726..2a8149471c05 100644
--- a/projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py
+++ b/projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py
@@ -1783,6 +1783,41 @@ def assignDerivedParameters(
         reject(state, printRejectionReason,
                "DebugPersistentKernelLoopForever requires StreamK=3 (got %d)"
                % state["StreamK"])
+      if state["StreamKWorkStealing"]:
+        # Codegen-time rejections only; there is no hardware context here
+        # (MI300A/MI300X both compile as gfx942). The kernel bakes a power-of-two
+        # per-XCD queue count from origami; the host (ContractionSolution.cpp)
+        # enforces the device's runtime NUM_XCD against that baked count and
+        # otherwise serves a non-work-stealing solution. Work stealing only exists
+        # in the dynamic-queue fetch (auto-mode SK4 and the SK4 sub-path of SK5).
+        if state["StreamK"] not in (4, 5):
+          reject(state, printRejectionReason,
+                 "StreamKWorkStealing requires StreamK in {4,5} (got %d)"
+                 % state["StreamK"])
+        # Stealing is only defined for the non-atomic partials+fixup path.
+        # Atomic SK4/SK5 is already rejected above; keep this explicit guard so
+        # the combination can never slip through.
+        if state["StreamKAtomic"]:
+          reject(state, printRejectionReason,
+                 "StreamKWorkStealing is not supported with StreamKAtomic")
+        # Reject DebugStreamK with work stealing: it can leave a tile-owning
+        # queue with no home workgroup, breaking the W_q>=1 auto-reset precondition.
+        if state["DebugStreamK"]:
+          reject(state, printRejectionReason,
+                 "StreamKWorkStealing requires DebugStreamK=0 (the per-queue "
+                 "auto-reset relies on W_q>=1 whenever tiles_q>=1); got %d"
+                 % state["DebugStreamK"])
+        # The steal path (streamKWorkStealingSteal) emits s_atomic_inc
+        # unconditionally, but the home fetch (_fetchNextWorkItem) only falls
+        # back to a returning vector atomic when scalar atomics are absent. On
+        # arches without scalar atomics (HasSAtomic=false, e.g. gfx1250) the
+        # steal would emit an unsupported s_atomic_inc with no vector fallback,
+        # so reject work stealing there.
+        if not isaInfoMap[isa].asmCaps["HasSAtomic"]:
+          reject(state, printRejectionReason,
+                 "StreamKWorkStealing requires scalar atomics (HasSAtomic); the "
+                 "work-stealing steal path emits s_atomic_inc with no vector "
+                 "fallback (e.g. gfx1250 has HasSAtomic=false)")
       if not state["Valid"]:
         print2("in assignDerivedParameters, state['Valid'] = False")
         return
@@ -1790,6 +1825,7 @@ def assignDerivedParameters(
       # If not using StreamK, clear other stream-k settings to avoid duplicate kernels
       state["StreamKForceDPOnly"] = 0
       state["StreamKAtomic"] = 0
+      state["StreamKWorkStealing"] = 0
       state["StreamKXCCMapping"] = 0
       state["StreamKFixupTreeReduction"] = 0
       state["DebugStreamK"] = 0
@@ -2814,6 +2850,7 @@ def evaluateExpertSchedulingMode() -> Literal[0, 1, 2]:
         "GroupLoadStore": not state["GroupLoadStore"],
         "StreamK": not state["StreamK"],
         "StreamKAtomic": not state["StreamKAtomic"],
+        "StreamKWorkStealing": not state["StreamKWorkStealing"],
         "StreamKXCCMapping": not state["StreamKXCCMapping"],
         "StreamKFixupTreeReduction": not state["StreamKFixupTreeReduction"],
         "DebugStreamK": not state["DebugStreamK"],
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/sk_dynamic_work_stealing.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/sk_dynamic_work_stealing.yaml
new file mode 100644
index 000000000000..107677cc23fd
--- /dev/null
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/sk_dynamic_work_stealing.yaml
@@ -0,0 +1,103 @@
+# Copyright © Advanced Micro Devices, Inc., or its affiliates.
+# SPDX-License-Identifier: MIT
+TestParameters:
+  marks: [skip-gfx900, skip-gfx906, skip-gfx908, skip-gfx90a, skip-gfx1010, skip-gfx1011, skip-gfx1012, skip-gfx1030, skip-gfx1100, skip-gfx1101, skip-gfx1102, skip-gfx1200, skip-gfx1201, skip-gfx1250] # not supported by arch
+
+GlobalParameters:
+  SyncsPerBenchmark: 0
+  NumElementsToValidate: -1
+  BoundsCheck: 0
+  KernelTime: False
+  DataInitTypeAlpha: 1
+  DataInitTypeBeta: 1
+  DataInitTypeA: 12
+  DataInitTypeB: 13
+  DataInitTypeC: 12
+  MaxWorkspaceSize: 134217728
+  # Back-to-back enqueues reuse the same Synchronizer workspace, exercising the
+  # per-queue auto-reset across launches.
+  NumWarmups: 2
+  EnqueuesPerSync: 4
+  # Variable per-WG delay drains queues unevenly so the next-neighbor steal fires under contention.
+  SleepPercent: 50
+
+# Shared SK4-dynamic fork entries; StreamKWorkStealing forked over [0, 1] to
+# validate both the stealing-off and stealing-on paths.
+.SkDynamicCommonFork: &sk_dynamic_common_fork
+  - KernelLanguage: ["Assembly"]
+  - PrefetchLocalRead: [1]
+  - 1LDSBuffer: [1]
+  - ExpandPointerSwap: [False]
+  - GlobalSplitU: [0]
+  - MIArchVgpr: [False]
+  - PrefetchGlobalRead: [2]
+  - ScheduleIterAlg: [3]
+  - SourceSwap: [True]
+  - StoreRemapVectorWidth: [0]
+  - StreamK: [4]
+  - StreamKWorkStealing: [0, 1]
+  - TransposeLDS: [0]
+  - WorkGroupMapping: [1]
+
+.SkDynamicBpsg: &sk_dynamic_bpsg
+  InitialSolutionParameters:
+  BenchmarkCommonParameters: *sk_dynamic_common_fork
+  BenchmarkForkParameters:
+  JoinParameters:
+  BenchmarkJoinParameters:
+  BenchmarkFinalParameters:
+    # Aligned and unaligned sizes plus a batched case exercise the neighbor
+    # steal and per-queue auto-reset across repeated launches.
+    - ProblemSizes:
+      - Exact: [512, 512, 1, 512]
+      - Exact: [640, 640, 1, 512]
+      - Exact: [896, 896, 1, 512]
+      - Exact: [384, 384, 3, 256]
+
+BenchmarkProblems:
+
+  - # SGEMM NT - StreamK=4 dynamic with single-hop work stealing
+    - # ProblemType
+      OperationType: GEMM
+      DataType: s
+      TransposeA: False
+      TransposeB: True
+      UseBeta: True
+      Batched: True
+
+    - # SK4 dynamic - small MI matrix to keep kernel-object size bounded
+      <<: *sk_dynamic_bpsg
+      ForkParameters:
+        - DepthU: [16]
+        - GlobalReadVectorWidthA: [1]
+        - GlobalReadVectorWidthB: [1]
+        - LocalReadVectorWidth: [1]
+        - MatrixInstruction:
+          - [32, 32, 2, 1, 1, 2,2, 2,2]
+          - [16, 16, 4, 1, 1, 2,2, 2,2]
+        - PrefetchLocalRead: [1]
+        - VectorWidthA: [1]
+        - VectorWidthB: [1]
+
+  - # HGEMM NT - StreamK=4 dynamic with single-hop work stealing
+    - # ProblemType
+      OperationType: GEMM
+      DataType: b
+      DestDataType: b
+      ComputeDataType: s
+      HighPrecisionAccumulate: True
+      TransposeA: True
+      TransposeB: False
+      UseBeta: True
+      Batched: True
+
+    - # SK4 dynamic - small MI matrix
+      <<: *sk_dynamic_bpsg
+      ForkParameters:
+        - DepthU: [64]
+        - GlobalReadVectorWidthA: [8]
+        - GlobalReadVectorWidthB: [8]
+        - MatrixInstruction:
+          - [32, 32, 8, 1, 1, 2,2, 2,2]
+          - [16, 16, 16, 1, 1, 2,2, 2,2]
+        - PrefetchLocalRead: [1]
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/sk_hybrid_work_stealing.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/sk_hybrid_work_stealing.yaml
new file mode 100644
index 000000000000..ad69afca1fe1
--- /dev/null
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/sk_hybrid_work_stealing.yaml
@@ -0,0 +1,106 @@
+# Copyright © Advanced Micro Devices, Inc., or its affiliates.
+# SPDX-License-Identifier: MIT
+TestParameters:
+  marks: [skip-gfx900, skip-gfx906, skip-gfx908, skip-gfx90a, skip-gfx1010, skip-gfx1011, skip-gfx1012, skip-gfx1030, skip-gfx1100, skip-gfx1101, skip-gfx1102, skip-gfx1200, skip-gfx1201, skip-gfx1250] # SK5 requires MFMA + gfx94x/950
+
+GlobalParameters:
+  SyncsPerBenchmark: 0
+  NumElementsToValidate: -1
+  BoundsCheck: 0
+  KernelTime: False
+  DataInitTypeAlpha: 1
+  DataInitTypeBeta: 1
+  DataInitTypeA: 12
+  DataInitTypeB: 13
+  DataInitTypeC: 12
+  MaxWorkspaceSize: 134217728
+  # Back-to-back enqueues reuse the same Synchronizer workspace, exercising the
+  # per-queue auto-reset across launches.
+  NumWarmups: 2
+  EnqueuesPerSync: 4
+  # Variable per-WG delay drains queues unevenly so the next-neighbor steal fires under contention.
+  SleepPercent: 50
+  # SK5 hybrid-mode paths: 0 = static (SK3 sub-path), 1 = dynamic per-queue
+  # (SK4 sub-path, where work stealing applies). Crossed with StreamKWorkStealing below.
+  StreamKHybridMode: [0, 1]
+
+# Shared SK5-hybrid fork entries; StreamKWorkStealing forked over [0, 1] to
+# validate both stealing-off and stealing-on across both hybrid sub-paths.
+.SkHybridCommonFork: &sk_hybrid_common_fork
+  - KernelLanguage: ["Assembly"]
+  - PrefetchLocalRead: [1]
+  - 1LDSBuffer: [1]
+  - ExpandPointerSwap: [False]
+  - GlobalSplitU: [0]
+  - MIArchVgpr: [False]
+  - PrefetchGlobalRead: [2]
+  - ScheduleIterAlg: [3]
+  - SourceSwap: [True]
+  - StoreRemapVectorWidth: [0]
+  - StreamK: [5]
+  - StreamKWorkStealing: [0, 1]
+  - TransposeLDS: [0]
+  - WorkGroupMapping: [1]
+
+.SkHybridBpsg: &sk_hybrid_bpsg
+  InitialSolutionParameters:
+  BenchmarkCommonParameters: *sk_hybrid_common_fork
+  BenchmarkForkParameters:
+  JoinParameters:
+  BenchmarkJoinParameters:
+  BenchmarkFinalParameters:
+    # Aligned and unaligned sizes plus a batched case exercise the neighbor
+    # steal and per-queue auto-reset across repeated launches.
+    - ProblemSizes:
+      - Exact: [512, 512, 1, 512]
+      - Exact: [640, 640, 1, 512]
+      - Exact: [896, 896, 1, 512]
+      - Exact: [384, 384, 3, 256]
+
+BenchmarkProblems:
+
+  - # SGEMM NT - StreamK=5 hybrid with single-hop work stealing (dynamic sub-path)
+    - # ProblemType
+      OperationType: GEMM
+      DataType: s
+      TransposeA: False
+      TransposeB: True
+      UseBeta: True
+      Batched: True
+
+    - # SK5 hybrid - small MI matrix to keep kernel-object size bounded
+      <<: *sk_hybrid_bpsg
+      ForkParameters:
+        - DepthU: [16]
+        - GlobalReadVectorWidthA: [1]
+        - GlobalReadVectorWidthB: [1]
+        - LocalReadVectorWidth: [1]
+        - MatrixInstruction:
+          - [32, 32, 2, 1, 1, 2,2, 2,2]
+          - [16, 16, 4, 1, 1, 2,2, 2,2]
+        - PrefetchLocalRead: [1]
+        - VectorWidthA: [1]
+        - VectorWidthB: [1]
+
+  - # HGEMM NT - StreamK=5 hybrid with single-hop work stealing
+    - # ProblemType
+      OperationType: GEMM
+      DataType: b
+      DestDataType: b
+      ComputeDataType: s
+      HighPrecisionAccumulate: True
+      TransposeA: True
+      TransposeB: False
+      UseBeta: True
+      Batched: True
+
+    - # SK5 hybrid - small MI matrix
+      <<: *sk_hybrid_bpsg
+      ForkParameters:
+        - DepthU: [64]
+        - GlobalReadVectorWidthA: [8]
+        - GlobalReadVectorWidthB: [8]
+        - MatrixInstruction:
+          - [32, 32, 8, 1, 1, 2,2, 2,2]
+          - [16, 16, 16, 1, 1, 2,2, 2,2]
+        - PrefetchLocalRead: [1]
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/DECISIONS.md b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/DECISIONS.md
index ff2727a15fd6..1d688e0b3f71 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/DECISIONS.md
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/DECISIONS.md
@@ -379,3 +379,12 @@ paths): DataType → `Tensile/Common/DataType.py`; CommonTypes →
 
 ## D16 — BufferLoad/BufferStore promoted to Required Parameters
 **Context** kernel basename hash changes across all archs; assembly verified unchanged/correct; no err or kernel-count changes."
+
+## D17 — StreamKWorkStealing added to the required (min-naming) parameter set
+**Decision:** Promote `StreamKWorkStealing` to the required (min-naming) parameter set in
+`Common/RequiredParameters.py` and accept the regenerated `_codegen` / SolutionClass /
+ValidParameters goldens.
+**Why:** without it, two solutions differing only in `StreamKWorkStealing` would collide on the
+same kernel identity name/hash.
+**Verification:** only `basename` hashes + the `SKWS0` name token + the roster/valid-values entry
+change (`num_keys` 334→335); no `err`, instruction-count, or emitted-assembly changes.
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/SolutionClass/__snapshots__/test_solution_class_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/SolutionClass/__snapshots__/test_solution_class_char.ambr
index 1cd33e82b59c..0e124a014f1b 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/SolutionClass/__snapshots__/test_solution_class_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/SolutionClass/__snapshots__/test_solution_class_char.ambr
@@ -3,7 +3,7 @@
   dict({
     'count': 1,
     'names': list([
-      'Cijk_Alik_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT32x32x32_MI16x16x1_SN_LDSB0_AA0_AFC1_AF1_AG0_AGGSUA0_AGNTAB0_AAIGTEn1_AAILTEn1_AFEM1_AFEM1_ASEM1_BL1_BS1_CD1_1_CLR1_CLS0_CADS0_DPKLF0_DSK0_DU32_DTL0_DTLM0_DTVA0_DTVB0_DTVMXSA0_DTVMXSB0_DTVSM0_DPLB0_EPS0_ELFLR0_EMLLn1_FDSI0_GRPM1_GRVWA1_GRVWB1_GSU1_GSUAMB_GSUC0_GSUWGMRR0_GLS0_HPLR0_ISA942_ICIW0_IU1_IA0_KLA_LDSSI0_LDSTI0_LBSPPA128_LBSPPB256_LBSPPMXSA0_LBSPPMXSB0_LBSPPM0_LPA4_LPB16_LPMXSA0_LPMXSB0_LPM0_LRVW4_LRVWA4_LRVWB4_LWPMn1_MIAV0_MIWT1_1_MXLIBL_MXSFNS_MDA2_MI16_16_16_1_MLDS65536_MO40_MPM0_MGRIPM1_NR0_NTn1_NTA0_NTB0_NTC0_NTD4_NTE0_NTG0_NTMXSA0_NTMXSB0_NTM0_NTWS0_NVn1_NVA0_NVB0_NVC0_NVD0_NVE0_NVMXSA0_NVMXSB0_NVM0_NVWS0_NEPBS2_NLCA1_NLCB1_ONLL1_PAP0_PGL0_PGR2_PLR1_PKA1_SFCWGM1_1_1_1_SGROB0_SGR1_SIA3_SLW1_SS1_SU32_SUM1_SUS256_SPO1_SRVW0_SSO0_SVW1_SK0_SKA0_SKFTR0_SKFDPO0_SKXCCM0_SNLL0_SIP1_SGRO0_TDMI0_TDMIM0_TDMS0_TIN0_THn1_THA0_THB0_THC0_THD0_THE0_THMXSA0_THMXSB0_THM0_THWS0_TT1_1_TLDS1_TLDSMn1_ULSGRO0_USL1_USLMX0_UCMLS0_UDFMAC0_UIOFGRO0_UPLRP0_USFGROn1_USI0_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WSK0_WS64_WG32_8_1_WGM8_WGMXCC1_WGMXCCGn1_WGR0',
+      'Cijk_Alik_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT32x32x32_MI16x16x1_SN_LDSB0_AA0_AFC1_AF1_AG0_AGGSUA0_AGNTAB0_AAIGTEn1_AAILTEn1_AFEM1_AFEM1_ASEM1_BL1_BS1_CD1_1_CLR1_CLS0_CADS0_DPKLF0_DSK0_DU32_DTL0_DTLM0_DTVA0_DTVB0_DTVMXSA0_DTVMXSB0_DTVSM0_DPLB0_EPS0_ELFLR0_EMLLn1_FDSI0_GRPM1_GRVWA1_GRVWB1_GSU1_GSUAMB_GSUC0_GSUWGMRR0_GLS0_HPLR0_ISA942_ICIW0_IU1_IA0_KLA_LDSSI0_LDSTI0_LBSPPA128_LBSPPB256_LBSPPMXSA0_LBSPPMXSB0_LBSPPM0_LPA4_LPB16_LPMXSA0_LPMXSB0_LPM0_LRVW4_LRVWA4_LRVWB4_LWPMn1_MIAV0_MIWT1_1_MXLIBL_MXSFNS_MDA2_MI16_16_16_1_MLDS65536_MO40_MPM0_MGRIPM1_NR0_NTn1_NTA0_NTB0_NTC0_NTD4_NTE0_NTG0_NTMXSA0_NTMXSB0_NTM0_NTWS0_NVn1_NVA0_NVB0_NVC0_NVD0_NVE0_NVMXSA0_NVMXSB0_NVM0_NVWS0_NEPBS2_NLCA1_NLCB1_ONLL1_PAP0_PGL0_PGR2_PLR1_PKA1_SFCWGM1_1_1_1_SGROB0_SGR1_SIA3_SLW1_SS1_SU32_SUM1_SUS256_SPO1_SRVW0_SSO0_SVW1_SK0_SKA0_SKFTR0_SKFDPO0_SKWS0_SKXCCM0_SNLL0_SIP1_SGRO0_TDMI0_TDMIM0_TDMS0_TIN0_THn1_THA0_THB0_THC0_THD0_THE0_THMXSA0_THMXSB0_THM0_THWS0_TT1_1_TLDS1_TLDSMn1_ULSGRO0_USL1_USLMX0_UCMLS0_UDFMAC0_UIOFGRO0_UPLRP0_USFGROn1_USI0_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WSK0_WS64_WG32_8_1_WGM8_WGMXCC1_WGMXCCGn1_WGR0',
     ]),
   })
 # ---
@@ -55,7 +55,7 @@
     'getitem_kernel_language': 'Assembly',
     'iter_matches_keys': True,
     'keys_is_list': True,
-    'len': 338,
+    'len': 339,
   })
 # ---
 # name: test_solution_construction
@@ -296,6 +296,7 @@
       'StreamKAtomic',
       'StreamKFixupTreeReduction',
       'StreamKForceDPOnly',
+      'StreamKWorkStealing',
       'StreamKXCCMapping',
       'SubGroup0',
       'SubGroup1',
@@ -400,8 +401,8 @@
       'tailLoopOptA',
       'tailLoopOptB',
     ]),
-    'name': 'Cijk_Alik_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT32x32x32_MI16x16x1_SN_LDSB0_AA0_AFC1_AF1_AG0_AGGSUA0_AGNTAB0_AAIGTEn1_AAILTEn1_AFEM1_AFEM1_ASEM1_BL1_BS1_CD1_1_CLR1_CLS0_CADS0_DPKLF0_DSK0_DU32_DTL0_DTLM0_DTVA0_DTVB0_DTVMXSA0_DTVMXSB0_DTVSM0_DPLB0_EPS0_ELFLR0_EMLLn1_FDSI0_GRPM1_GRVWA1_GRVWB1_GSU1_GSUAMB_GSUC0_GSUWGMRR0_GLS0_HPLR0_ISA942_ICIW0_IU1_IA0_KLA_LDSSI0_LDSTI0_LBSPPA128_LBSPPB256_LBSPPMXSA0_LBSPPMXSB0_LBSPPM0_LPA4_LPB16_LPMXSA0_LPMXSB0_LPM0_LRVW4_LRVWA4_LRVWB4_LWPMn1_MIAV0_MIWT1_1_MXLIBL_MXSFNS_MDA2_MI16_16_16_1_MLDS65536_MO40_MPM0_MGRIPM1_NR0_NTn1_NTA0_NTB0_NTC0_NTD4_NTE0_NTG0_NTMXSA0_NTMXSB0_NTM0_NTWS0_NVn1_NVA0_NVB0_NVC0_NVD0_NVE0_NVMXSA0_NVMXSB0_NVM0_NVWS0_NEPBS2_NLCA1_NLCB1_ONLL1_PAP0_PGL0_PGR2_PLR1_PKA1_SFCWGM1_1_1_1_SGROB0_SGR1_SIA3_SLW1_SS1_SU32_SUM1_SUS256_SPO1_SRVW0_SSO0_SVW1_SK0_SKA0_SKFTR0_SKFDPO0_SKXCCM0_SNLL0_SIP1_SGRO0_TDMI0_TDMIM0_TDMS0_TIN0_THn1_THA0_THB0_THC0_THD0_THE0_THMXSA0_THMXSB0_THM0_THWS0_TT1_1_TLDS1_TLDSMn1_ULSGRO0_USL1_USLMX0_UCMLS0_UDFMAC0_UIOFGRO0_UPLRP0_USFGROn1_USI0_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WSK0_WS64_WG32_8_1_WGM8_WGMXCC1_WGMXCCGn1_WGR0',
-    'num_keys': 338,
+    'name': 'Cijk_Alik_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT32x32x32_MI16x16x1_SN_LDSB0_AA0_AFC1_AF1_AG0_AGGSUA0_AGNTAB0_AAIGTEn1_AAILTEn1_AFEM1_AFEM1_ASEM1_BL1_BS1_CD1_1_CLR1_CLS0_CADS0_DPKLF0_DSK0_DU32_DTL0_DTLM0_DTVA0_DTVB0_DTVMXSA0_DTVMXSB0_DTVSM0_DPLB0_EPS0_ELFLR0_EMLLn1_FDSI0_GRPM1_GRVWA1_GRVWB1_GSU1_GSUAMB_GSUC0_GSUWGMRR0_GLS0_HPLR0_ISA942_ICIW0_IU1_IA0_KLA_LDSSI0_LDSTI0_LBSPPA128_LBSPPB256_LBSPPMXSA0_LBSPPMXSB0_LBSPPM0_LPA4_LPB16_LPMXSA0_LPMXSB0_LPM0_LRVW4_LRVWA4_LRVWB4_LWPMn1_MIAV0_MIWT1_1_MXLIBL_MXSFNS_MDA2_MI16_16_16_1_MLDS65536_MO40_MPM0_MGRIPM1_NR0_NTn1_NTA0_NTB0_NTC0_NTD4_NTE0_NTG0_NTMXSA0_NTMXSB0_NTM0_NTWS0_NVn1_NVA0_NVB0_NVC0_NVD0_NVE0_NVMXSA0_NVMXSB0_NVM0_NVWS0_NEPBS2_NLCA1_NLCB1_ONLL1_PAP0_PGL0_PGR2_PLR1_PKA1_SFCWGM1_1_1_1_SGROB0_SGR1_SIA3_SLW1_SS1_SU32_SUM1_SUS256_SPO1_SRVW0_SSO0_SVW1_SK0_SKA0_SKFTR0_SKFDPO0_SKWS0_SKXCCM0_SNLL0_SIP1_SGRO0_TDMI0_TDMIM0_TDMS0_TIN0_THn1_THA0_THB0_THC0_THD0_THE0_THMXSA0_THMXSB0_THM0_THWS0_TT1_1_TLDS1_TLDSMn1_ULSGRO0_USL1_USLMX0_UCMLS0_UDFMAC0_UIOFGRO0_UPLRP0_USFGROn1_USI0_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WSK0_WS64_WG32_8_1_WGM8_WGMXCC1_WGMXCCGn1_WGR0',
+    'num_keys': 339,
     'problem_type': 'Cijk_Alik_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs',
     'stable': dict({
       'DepthU': 32,
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/ValidParameters/__snapshots__/test_builders_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/ValidParameters/__snapshots__/test_builders_char.ambr
index 871b9859983b..e3c9c86d5546 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/ValidParameters/__snapshots__/test_builders_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/ValidParameters/__snapshots__/test_builders_char.ambr
@@ -1482,6 +1482,7 @@
     'StreamKAtomic',
     'StreamKFixupTreeReduction',
     'StreamKForceDPOnly',
+    'StreamKWorkStealing',
     'StreamKXCCMapping',
     'SuppressNoLoadLoop',
     'SwInstructionPrefetch',
@@ -2596,6 +2597,14 @@
       'len': 2,
       'type': 'list',
     }),
+    'StreamKWorkStealing': dict({
+      'head': list([
+        0,
+        1,
+      ]),
+      'len': 2,
+      'type': 'list',
+    }),
     'StreamKXCCMapping': dict({
       'head': list([
         0,
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_emit_bigfiles_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_emit_bigfiles_char.ambr
index c22038cb4017..0c90fb5cf336 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_emit_bigfiles_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_emit_bigfiles_char.ambr
@@ -2,27 +2,27 @@
 # name: test_bigfile_capped_emit[aqua_FreeSize_GSU9]
   list([
     dict({
-      'basename': 'Cijk_Ailk_Bljk_F8NH_HHS_BH_MT128x16x128_MI16x16xS-Gljw7Tndlcz4DyWBcnsgvq3a2XnFHzrLYjk74Pzdc=',
+      'basename': 'Cijk_Ailk_Bljk_F8NH_HHS_BH_MT128x16x128_MI16x16x0PEaqgf_3RNkFiTAYvEWBPz1xQTaRvKD9CCT0wpWzO8=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bljk_F8NH_HHS_BH_MT128x16x32_MI16x16x1qoMYuCs_HiKfJoyDVjJF73iEbywxwpKqyYZG-pAZZxM=',
+      'basename': 'Cijk_Ailk_Bljk_F8NH_HHS_BH_MT128x16x32_MI16x16x1YtOY4yg6GqEedQwTv-n9heqwEGq3xaaRM1XPu__B898=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bljk_F8NH_HHS_BH_MT128x16x64_MI16x16x1QuEx_LLnYKhBzZS3anZFPLWuXBLYopd4QyxslRfomFg=',
+      'basename': 'Cijk_Ailk_Bljk_F8NH_HHS_BH_MT128x16x64_MI16x16x1hyUZ5k8XjqFqTa9F4lvmH3KVeG7MolFIfZBX4e-vAQg=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bljk_F8NH_HHS_BH_MT128x32x128_MI16x16x7Xuu72GxZdrwjLqko6M-grE0eUmaxIq67pliAXtFLsA=',
+      'basename': 'Cijk_Ailk_Bljk_F8NH_HHS_BH_MT128x32x128_MI16x16xgqZBDoCQlSnUDsyhEHDS_IywPtOFXyVyS9MrmTd2V_o=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bljk_F8NH_HHS_BH_MT128x32x32_MI16x16x1TB88pJVwEkDVdlgbF-lGwdUOsLEmhdOy147cN_NNt9k=',
+      'basename': 'Cijk_Ailk_Bljk_F8NH_HHS_BH_MT128x32x32_MI16x16x1XemHpkLDsG2jRjriI-QsPNhqOkIi1ChlH5RUqLCJOo0=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bljk_F8NH_HHS_BH_MT128x32x64_MI16x16x10DhfuxvbEVPlwpdPMyKHw58eOGsXDWtplJfAOlIe-Pc=',
+      'basename': 'Cijk_Ailk_Bljk_F8NH_HHS_BH_MT128x32x64_MI16x16x1CVv_Fn6a_cErQuelTtW-gz9OaFSRPryMGtZjoKyZJCA=',
       'err': 0,
     }),
   ])
@@ -30,27 +30,27 @@
 # name: test_bigfile_capped_emit[equality_gfx950_HSS_big]
   list([
     dict({
-      'basename': 'Cijk_Ailk_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT11XTh2yIemLgZzvT8Y55Aa0IvVgrAKL6JPGW9AK8iv_c=',
+      'basename': 'Cijk_Ailk_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT1AXp0F5JGbnx4dhbPfcfdiKLiZpzsdi2Qh-QdYKqJVds=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT2-GpJOvig6JRtx42eDC7p2ObeS9MhqXSKDjaLJiBHbdQ=',
+      'basename': 'Cijk_Ailk_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT2BG_Xr0IlyABFCKVD885PVZTijzjpldjXKvwYqpS2u5Y=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT24_bNGgNsfgaVnmsBLerK44Nb1f2KXpxWKkiJcXUA5_o=',
+      'basename': 'Cijk_Ailk_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT2Ixc5wNAiu0KxZpoJd5m1Njv5sCb21hdhWqFoUx3oFJw=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT2C_oVG0KdOGf-0Rg_3ZUYmJFVbUFw1dZ7yuUfJJ1HTBY=',
+      'basename': 'Cijk_Ailk_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT2LutL_WMujk_KDSlqnd6HHg28p89ARRyKdIMKnlqRsiw=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT2FFwlilLGG3fT5kv36Q8b2zsAmwmdK38pv44gtbsQurY=',
+      'basename': 'Cijk_Ailk_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT2_jfECdyYw388q4Z5NVWRhFa3Lzm4ezEjc4kBolTGXkE=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT2RZkVObdYiuO0nuXeONKEMF8qzIG4_hcwXc8cjFW2f3A=',
+      'basename': 'Cijk_Ailk_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT2aMtCgdbT7HrTb6g8kVjovM_zsuZwC23HONnAy8K-vUI=',
       'err': 0,
     }),
   ])
@@ -58,27 +58,27 @@
 # name: test_bigfile_capped_emit[freesize_gfx942_F8NH_GSU]
   list([
     dict({
-      'basename': 'Cijk_Ailk_Bljk_F8NH_HHS_BH_Bias_HA_S_SAB_SAV_MT133qRnATgZ6_SATxqR8eys5SQD0q41vstAoLMtxv1QWY=',
+      'basename': 'Cijk_Ailk_Bljk_F8NH_HHS_BH_Bias_HA_S_SAB_SAV_MT1E5aZ9_fehYw6fSt_tScq6myj_-d6eNVdGx2aC6pfeso=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bljk_F8NH_HHS_BH_Bias_HA_S_SAB_SAV_MT13Vej6vrkkNYN7Lr0t5xwlfy-q5SHfIaHsPXPEeK5JmE=',
+      'basename': 'Cijk_Ailk_Bljk_F8NH_HHS_BH_Bias_HA_S_SAB_SAV_MT1NO2HVcfYsfG84BvFWBHzQWBnOmF1w9dhZMfX5eVL4dw=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bljk_F8NH_HHS_BH_Bias_HA_S_SAB_SAV_MT17oFB9E_hglLQMLAOQeW-IP6AKk4eiEGeGPAnOP6P6Gc=',
+      'basename': 'Cijk_Ailk_Bljk_F8NH_HHS_BH_Bias_HA_S_SAB_SAV_MT1_JW277HsdkoUBIwc-fs5Rlh_ZvhPmDPhrJ_XiPaNLjk=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bljk_F8NH_HHS_BH_Bias_HA_S_SAB_SAV_MT1JpAcoj4hLJ615psm1yxpaeutLmA65Ez8u4DjMp6CXrk=',
+      'basename': 'Cijk_Ailk_Bljk_F8NH_HHS_BH_Bias_HA_S_SAB_SAV_MT1lVzYoGr91c8XWvfeEIZbdpQ5jvMhUIy5l1pPl_UmgQ8=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bljk_F8NH_HHS_BH_Bias_HA_S_SAB_SAV_MT1u1FRQ9WG5lIu1XxfDDKmSwL1EGW5eIYgO4JHv2A5kow=',
+      'basename': 'Cijk_Ailk_Bljk_F8NH_HHS_BH_Bias_HA_S_SAB_SAV_MT1mzGzNC--azxkizUIzkjlJR8msadgWFQ10afBUvVxuBM=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bljk_F8NH_HHS_BH_Bias_HA_S_SAB_SAV_MT1xaablpTT8n2AOqnnJG24vsPuioA6wNxLZqme7JYYEWQ=',
+      'basename': 'Cijk_Ailk_Bljk_F8NH_HHS_BH_Bias_HA_S_SAB_SAV_MT1rRExUpAg67uP8HjM0uHpxdoHGw5Y-g6VV_PAj0d7MF4=',
       'err': 0,
     }),
   ])
@@ -86,27 +86,27 @@
 # name: test_bigfile_capped_emit[gfx1201_I8II]
   list([
     dict({
-      'basename': 'Cijk_Ailk_Bljk_I8II_BH_HA_I_SAV_UserArgs_MT128x1t-w3mKQN4gPbAHvx7dEmtx5k2vVFzdhJvVVoZMNZlkU=',
+      'basename': 'Cijk_Ailk_Bljk_I8II_BH_HA_I_SAV_UserArgs_MT128x1VWV9dbHQKcIAhtN1me2X8z1f9JHeVFF-CYoSQrD1dTg=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bljk_I8II_BH_HA_I_SAV_UserArgs_MT16x64UqI7LKTAo8E27Z3PNZybELRT8OQMQHIuWDkWQPcC8As=',
+      'basename': 'Cijk_Ailk_Bljk_I8II_BH_HA_I_SAV_UserArgs_MT16x64Owh9Npi9kG0OFr53vXnmChtzzU6vTV1Pz3uil-bfsck=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bljk_I8II_BH_HA_I_SAV_UserArgs_MT16x64r-vXE83Tgqxz1GCSrKmwqRNS7NJ29DsjhklzuhU0JUE=',
+      'basename': 'Cijk_Ailk_Bljk_I8II_BH_HA_I_SAV_UserArgs_MT16x64eA9elviyPpe4aDf2sJGUfdSv5bIL1WC9c5GakgJjTx0=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bljk_I8II_BH_HA_I_SAV_UserArgs_MT256x11hOwt_IvHXQI_4oxZXmGBo0VrgdMZdx29D_13160H0Y=',
+      'basename': 'Cijk_Ailk_Bljk_I8II_BH_HA_I_SAV_UserArgs_MT256x11FGzDdqaXJHCGnOB7WIDkYA-p7NxZmdzO9DIWaXE730=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bljk_I8II_BH_HA_I_SAV_UserArgs_MT256x13VFHx09b6USykTyUU9-dvtP3ZFIrk7z424bSSKla7zw=',
+      'basename': 'Cijk_Ailk_Bljk_I8II_BH_HA_I_SAV_UserArgs_MT256x120Quts8nquDZEpPMPSV9CvZA35ufBLTCuIVjO61pZYw=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bljk_I8II_BH_HA_I_SAV_UserArgs_MT256x13va2QNOinbNB_T7QxVhpsn0KhKXiJB-BgEEGRgku7hU=',
+      'basename': 'Cijk_Ailk_Bljk_I8II_BH_HA_I_SAV_UserArgs_MT256x15ESidfh4mCdT4xoEf0fiFaOuH5sQAdlaQ40J1I9DwbA=',
       'err': 0,
     }),
   ])
@@ -114,15 +114,15 @@
 # name: test_bigfile_capped_emit[gfx1250_GG]
   list([
     dict({
-      'basename': 'Cijk_Alik_Bljk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT2I8Z4MnNhn0H6OkrgkCJB2REmBpHKVBnmyvU3B0pEk5o=',
+      'basename': 'Cijk_Alik_Bljk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT27R-Me2IAXS91gvesLlSE9It6N5WAoCUnDCgktX1SbT0=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT2On8jZq2r35NyCPrHV_DUmxEc59zqvQbrqb6HnMIdEFA=',
+      'basename': 'Cijk_Alik_Bljk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT2ox_DJro5lCW_-A7qhSVaKKLmiWJZRtkDZsjP1TDj0Bk=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT3qS1-XgGxY_ZW8LCkQ2h9qcTZZNFDrv93bs9HZOExcPM=',
+      'basename': 'Cijk_Alik_Bljk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT3iQRATBChzKmj1EIiwQhdonRchLYdnqrTnOjyxl1v-Nw=',
       'err': 0,
     }),
   ])
@@ -130,27 +130,27 @@
 # name: test_bigfile_capped_emit[gfx90a_HSS_big]
   list([
     dict({
-      'basename': 'Cijk_Ailk_Bljk_HSS_BH_UserArgs_MT128x128x32_MI32URRlBbZ_0D3f_0tbTrOj16A5kCb4eY_Qj7e89uwrRU8=',
+      'basename': 'Cijk_Ailk_Bljk_HSS_BH_UserArgs_MT128x128x32_MI32DLt8rifdTTwEiorPWkIA0f1r3Yn0EgSh2PCav5RnSWc=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bljk_HSS_BH_UserArgs_MT128x128x32_MI32wJQwzk7gLrdBXa4gk8QLXH-e6H-VV9NkcshCy7j3U1Q=',
+      'basename': 'Cijk_Ailk_Bljk_HSS_BH_UserArgs_MT128x128x32_MI32VeX7cJBv7G9NceBAXrpx-mKacNH0sm-NHzjdSpTD5rA=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bljk_HSS_BH_UserArgs_MT128x128x32_MI32xz67w3tBRw2b2OoNKoaOeNFnXAPua-muhWZCehLg7XY=',
+      'basename': 'Cijk_Ailk_Bljk_HSS_BH_UserArgs_MT128x128x32_MI32tgx4OkjbymJApkgI7_Sj2zmeA7YJSQDXivmS-G7Lby8=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bljk_HSS_BH_UserArgs_MT128x192x64_MI32X-_c4pYsuuvRP7JlvpuuONFQfNpg4b2YkHLzwDz01wA=',
+      'basename': 'Cijk_Ailk_Bljk_HSS_BH_UserArgs_MT128x192x64_MI32M9aFe5VVPZ_teiAeUhk_8QoWuEDm1s03_qMcDZ0h0sQ=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bljk_HSS_BH_UserArgs_MT128x32x64_MI32x7chYdE4VjFBaxocWCppI_6yoqGomi78hnzjCr0B8Nzo=',
+      'basename': 'Cijk_Ailk_Bljk_HSS_BH_UserArgs_MT128x32x64_MI32x1mDz2Fydou1101lZBUD6Ar8ijHp2c0lvRebadryheYg=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bljk_HSS_BH_UserArgs_MT128x32x64_MI32xgVjp1hBfMr_RfHk8OSzMRzghqb_eoxLgxDBaQL8mGLI=',
+      'basename': 'Cijk_Ailk_Bljk_HSS_BH_UserArgs_MT128x32x64_MI32x5b5fakZXtv75dqRUjtgqlptKIH32NWU16Osh5Sx8hTc=',
       'err': 0,
     }),
   ])
@@ -158,27 +158,27 @@
 # name: test_bigfile_capped_emit[gfx950_origami_MX]
   list([
     dict({
-      'basename': 'Cijk_Alik_Bjlk_S_MX_B_Bias_HA_S_SAV_UserArgs_MT15ELi53QSQPpRNXPhhP4pGjQ-VQOHGAaSQyjynFkCceg=',
+      'basename': 'Cijk_Alik_Bjlk_S_MX_B_Bias_HA_S_SAV_UserArgs_MT19C35JeeY-qDrrHdhgd7Va0C2swFtPIZpMjtAp3E6lAk=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bjlk_S_MX_B_Bias_HA_S_SAV_UserArgs_MT1FaIXyeyDH-rnatBNFsi1bGRpE9j9Y0bzfkxTgNoem-U=',
+      'basename': 'Cijk_Alik_Bjlk_S_MX_B_Bias_HA_S_SAV_UserArgs_MT19X7kPJSy59Wef9Z3WMgTsvLose5z_pkGDcDJa1OlHF0=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bjlk_S_MX_B_Bias_HA_S_SAV_UserArgs_MT1HSd8OaIgyIc0kO9VvvWo-rv8QC4ILmbLWUc9n3yNCDI=',
+      'basename': 'Cijk_Alik_Bjlk_S_MX_B_Bias_HA_S_SAV_UserArgs_MT1Ca23TPEModSClyjnn4g1E5qDzzSFI181SJ-eOQyZelo=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bjlk_S_MX_B_Bias_HA_S_SAV_UserArgs_MT1HgH3mKb2OiSx5w7r-7ekQRdnfTDS2-ju_YvgY0f7qnI=',
+      'basename': 'Cijk_Alik_Bjlk_S_MX_B_Bias_HA_S_SAV_UserArgs_MT1D3mWEpWvmyQnH2p33U0xAMtDpFY461JOXxtqEXYMn2A=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bjlk_S_MX_B_Bias_HA_S_SAV_UserArgs_MT1I5XdJ_aanho06clTD_WHmzIMINCtQGDIhEeyXsBYJqI=',
+      'basename': 'Cijk_Alik_Bjlk_S_MX_B_Bias_HA_S_SAV_UserArgs_MT1FzpkkGJIpIBVxsEPi161T7OS-SYgMalTiUeTAgdDqog=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bjlk_S_MX_B_Bias_HA_S_SAV_UserArgs_MT1NRoU4nLErdpJ9zNVlNGliHdrcWGeTEh3_2zCK0ga1CM=',
+      'basename': 'Cijk_Alik_Bjlk_S_MX_B_Bias_HA_S_SAV_UserArgs_MT1H3_nM0eqLE3jQSKKUJawz3C42KEGQmZccpPWRI6nglM=',
       'err': 0,
     }),
   ])
@@ -186,27 +186,27 @@
 # name: test_bigfile_capped_emit[navi31_HSS]
   list([
     dict({
-      'basename': 'Cijk_Alik_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT1B1L59GQFcbjjcFBhkHOu89GRxNBmkZ2MrQikGCLSt1c=',
+      'basename': 'Cijk_Alik_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT1EStc3x092hvUP7fmro5DgeLTNZ54Wui8d_jbUSHalWk=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT1lk047TNGyOQQsMfnkLNs9meWp0eyoFemwd9F4DW8JqQ=',
+      'basename': 'Cijk_Alik_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT1I4vZ6WDpwHg31N1-egbRfLG3F8vVFQGlJpbwiK2PU2U=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT3pSpXxwGvq311WN_Et9H1eYMTHWiETfIlWDUQqXB0k6U=',
+      'basename': 'Cijk_Alik_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT3kFPCvpHiPqMSGQfs66PnL_lZiJxAy9qVi9pw_1CdHvc=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT3uff6pT_n6SUbHos7kXoAyeaMGTt-vxtwniT31ekPN1o=',
+      'basename': 'Cijk_Alik_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT3sPwilggvpVIZG3DYqhMqnwY4OVJSl0E8_RavcQOI2-c=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT4-leGXNR2Cg8vmPrQnZu9cvWZFbOpzKwD0XjNwcNIQ4E=',
+      'basename': 'Cijk_Alik_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT4iHinGmpo6U9XZLki6M7rkVqcxQ6ZaeP3otCR8URNjNA=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT44hPvPwd2wfINL2uDgCL9jVjKRZWck1F-dPudxi7w1Ws=',
+      'basename': 'Cijk_Alik_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT4rzujJqyV9D4EA5Bal0CwBhJxQ9gEe-w10uIKXsTkNCs=',
       'err': 0,
     }),
   ])
@@ -214,19 +214,19 @@
 # name: test_bigfile_capped_emit[streamk_gfx942_Ailk]
   list([
     dict({
-      'basename': 'Cijk_Ailk_Bjlk_S_B_UserArgs_MT112x256x16_MI16x164488HlTT3XRdQ1iMiUWtbU-7BGJW-ZPExv7bEtaz3v4=',
+      'basename': 'Cijk_Ailk_Bjlk_S_B_UserArgs_MT112x256x16_MI16x16iamdKh77SiXB_tsBpkoYVJkqXz7yJXFcFX7zLAN-4e0=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bjlk_S_B_UserArgs_MT112x256x32_MI16x16nmlCwPFNaaWTdIJARwvCUATraAv01Bp_49325nQjMLw=',
+      'basename': 'Cijk_Ailk_Bjlk_S_B_UserArgs_MT112x256x32_MI16x16RtgaOEkRV8gTmNr339CCNhv41S6CH7M0IP134rKubs8=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bjlk_S_B_UserArgs_MT128x128x16_MI16x16F0GTsPMl_srsNgviNuetd6ghGCvbN5BAkQVNHgnroUE=',
+      'basename': 'Cijk_Ailk_Bjlk_S_B_UserArgs_MT128x128x16_MI16x16pQbK8lxueBbBe5oFnbPGmVGfvw-J71kHkCknwFIc_Lc=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bjlk_S_B_UserArgs_MT128x128x32_MI16x16ttjQ6-75OzKbfYcLAtGh0VSFgrSFnC7ceMtxxR66_ng=',
+      'basename': 'Cijk_Ailk_Bjlk_S_B_UserArgs_MT128x128x32_MI16x16HzSAmeO3iAdidAF6iz7WibCBqHVFeKoTR9uYOAPPEx8=',
       'err': 0,
     }),
   ])
@@ -234,27 +234,27 @@
 # name: test_bigfile_capped_emit[streamk_gfx942_S]
   list([
     dict({
-      'basename': 'Cijk_Alik_Bjlk_S_B_UserArgs_MT112x256x16_MI16x16LIPzO5RGOH-JAgrnAmNmDNstuT7ffdgJnNQwA_C5wEw=',
+      'basename': 'Cijk_Alik_Bjlk_S_B_UserArgs_MT112x256x16_MI16x16sUXwBn3XkfnISoWST8-w6dcxuX4c9gb1b-ACCXCmITQ=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bjlk_S_B_UserArgs_MT112x256x32_MI16x16shsJpGZCbngCiXB-1j-vOIKjWsvtroIwqzvnzhcpHsI=',
+      'basename': 'Cijk_Alik_Bjlk_S_B_UserArgs_MT112x256x32_MI16x16mICYU5z8Zeu9GVGXH6Q23-00wl4qQS3dpU1OEZeLH-c=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bjlk_S_B_UserArgs_MT128x128x16_MI16x16GxunDHjBJr9zSOOvIHwDzhlWONbdgnmhRbwhX7SpETA=',
+      'basename': 'Cijk_Alik_Bjlk_S_B_UserArgs_MT128x128x16_MI16x16iuBlVFr2UmLymPSjN3Tg2wGxkWAakfqmGQbhFQiSz3U=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bjlk_S_B_UserArgs_MT128x128x32_MI16x16bKJ1I-I9myKTFgC8TCcJ4cKXinqQexXtxx8Oeq6pIxw=',
+      'basename': 'Cijk_Alik_Bjlk_S_B_UserArgs_MT128x128x32_MI16x1626fyQeIWhR0VKyMZB16tLOIV6cOLZSsLIXN5XLBPg-k=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bjlk_S_B_UserArgs_MT128x160x16_MI16x16wh8AYgwInIbd9m_6RLeSGkJeByUk_8-ZVK25Uo7ijoc=',
+      'basename': 'Cijk_Alik_Bjlk_S_B_UserArgs_MT128x160x16_MI16x16j82k1qF_CMrsFJQJIwRAEjMkjf2SllvK-9iw23wVwaE=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bjlk_S_B_UserArgs_MT128x160x32_MI16x164lvqXbTOKTPbMjr0Nfo9cEi_2H-CV8-aI2Ed1z_dyJU=',
+      'basename': 'Cijk_Alik_Bjlk_S_B_UserArgs_MT128x160x32_MI16x16jvqTs3KuIGVtneIVcxMdIsn2Lyf6RfbrWJU5P3-v_9Q=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_emit_gfx1100_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_emit_gfx1100_char.ambr
index 780c1c67e8a6..165e16e85423 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_emit_gfx1100_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_emit_gfx1100_char.ambr
@@ -5,19 +5,19 @@
       'file': 'DB.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Ailk_Bjlk_D_B_UserArgs_MT16x16x16_SN_LDSB0_q3fwpwWSBfBQn_1RCsNBafSu9mOXjWxG4pp5iWAY2V8=',
+          'basename': 'Cijk_Ailk_Bjlk_D_B_UserArgs_MT16x16x16_SN_LDSB0_jVyggMAKLNhIJ3GHu3O_JlBcBt0LNRTeGa8ngP9kQEg=',
           'err': 0,
         }),
         dict({
-          'basename': 'Cijk_Ailk_Bjlk_D_B_UserArgs_MT16x16x8_SN_LDSB0_AcRA29MpDMmZWfX4IGyS8qXL0Tr7B_ghVvOwjR2KtzXo=',
+          'basename': 'Cijk_Ailk_Bjlk_D_B_UserArgs_MT16x16x8_SN_LDSB0_AinYY9xHMHLh5hRFYnka8n_XjzWehqlcR7aNIzE_hmIg=',
           'err': 0,
         }),
         dict({
-          'basename': 'Cijk_Ailk_Bjlk_D_B_UserArgs_MT8x8x16_SN_LDSB0_AFnM3UfCzzOoZEL_RmughTxk_1IJAok0YEIU1hI_I78kg=',
+          'basename': 'Cijk_Ailk_Bjlk_D_B_UserArgs_MT8x8x16_SN_LDSB0_AFmErPQciCYz_gK3UDWPl5jHZiRBW3m5xHjYh6o_q_4E8=',
           'err': 0,
         }),
         dict({
-          'basename': 'Cijk_Ailk_Bjlk_D_B_UserArgs_MT8x8x8_SN_LDSB0_AFCeUpNjM9_nfC6wiHAW8-6fxAB47PbMBn4os4NaZh9Ayw=',
+          'basename': 'Cijk_Ailk_Bjlk_D_B_UserArgs_MT8x8x8_SN_LDSB0_AFCWDKDLb0K8Knx0lgbTMUE-4Xdw575MOLSZefeZVoEu6Q=',
           'err': 0,
         }),
       ]),
@@ -26,15 +26,15 @@
       'file': 'HSS.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Alik_Bljk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT3HRNIYcxKt_cwtY72mIYtMNDoZ8q0Dh_pXB1h4pRZ1qg=',
+          'basename': 'Cijk_Alik_Bljk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT3jxpcIuMc7rLB3mgkcYyHssSNCtTBkb1dsJV1t-Kq9is=',
           'err': 0,
         }),
         dict({
-          'basename': 'Cijk_Alik_Bljk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT9_rZ_dYmRhJzyB4EBuMq_NO_dQIjB_hiSjvpSA5qEXQw=',
+          'basename': 'Cijk_Alik_Bljk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT9Jtq8NgnMA5Y7r_i5y-bHWmZcDdRk4Vbh9WuO7WMh324=',
           'err': 0,
         }),
         dict({
-          'basename': 'Cijk_Alik_Bljk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT9nkqe6FoUam5KUTUclYtJsnPv89_NmGynV5Dc4IYJ16Q=',
+          'basename': 'Cijk_Alik_Bljk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT9unSjMo48jqo78zR-NZUiKh3chqgHlOGt3RoDb_8h71c=',
           'err': 0,
         }),
       ]),
@@ -43,19 +43,19 @@
       'file': 'SB.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Ailk_Bjlk_S_B_Bias_HA_S_SAV_UserArgs_MT16x18T3p6KQcDOSn1fhQ1AYHmI7pdeFol5O5Hsb9xc27YNs=',
+          'basename': 'Cijk_Ailk_Bjlk_S_B_Bias_HA_S_SAV_UserArgs_MT16x1BdvJJ4dZfRxpJjAxL7HmxqjFW9aiyJ26xa9fScBtRXc=',
           'err': 0,
         }),
         dict({
-          'basename': 'Cijk_Ailk_Bjlk_S_B_Bias_HA_S_SAV_UserArgs_MT16x1kOgL_kb3pXsbrH1Nl1BRcKz7RAlgeRG7QNP3gkuWy1k=',
+          'basename': 'Cijk_Ailk_Bjlk_S_B_Bias_HA_S_SAV_UserArgs_MT16x1mxudZ-dlE6gRvjC5QHH-ePO6wBwMTHWtoW5S-V5nHGo=',
           'err': 0,
         }),
         dict({
-          'basename': 'Cijk_Ailk_Bjlk_S_B_Bias_HA_S_SAV_UserArgs_MT8x8xFG1IqJfQeqxw5DdWqUPA3iIubaUzocm0zjVkKfQEVgU=',
+          'basename': 'Cijk_Ailk_Bjlk_S_B_Bias_HA_S_SAV_UserArgs_MT8x8xYhoqEPI7oOLNJsJoEfx5ozaASk7Tgr-2_SI7US-EWdw=',
           'err': 0,
         }),
         dict({
-          'basename': 'Cijk_Ailk_Bjlk_S_B_Bias_HA_S_SAV_UserArgs_MT8x8xZFgHJZL1oBnT7hPJX4pEC3BnnqBP-Jctt_5KRXaWyL8=',
+          'basename': 'Cijk_Ailk_Bjlk_S_B_Bias_HA_S_SAV_UserArgs_MT8x8xj43yMzvq0C3LqB7jBnBC654wOqcuMcgrIz4x8ljNmCc=',
           'err': 0,
         }),
       ]),
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_emit_gfx1201_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_emit_gfx1201_char.ambr
index 6e3880b03d3b..361e5a34174e 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_emit_gfx1201_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_emit_gfx1201_char.ambr
@@ -5,7 +5,7 @@
       'file': 'B8B8S.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Alik_Bljk_B8B8S_BH_Bias_HA_S_SAV_UserArgs_MVSCCAfpf2R8QdTDYmUTUkRcTFd1IpfRAjLet-L5oqDI=',
+          'basename': 'Cijk_Alik_Bljk_B8B8S_BH_Bias_HA_S_SAV_UserArgs_M8vs6yMo_feMQNIFcBkD7w1e1VwWRutv8gTYybHubYXg=',
           'err': 0,
         }),
       ]),
@@ -14,7 +14,7 @@
       'file': 'BBS.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_SH_HA_S_SAV_UserArgs_8sV6H8NFpzuKk4RQtQUDFBuGXHOvwcVJZTQHj_x-Dn0=',
+          'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_SH_HA_S_SAV_UserArgs_DhUnfjRsWvXjpGzOWoG8mEanfHwaY18BXtpK9okDTYE=',
           'err': 0,
         }),
       ]),
@@ -23,7 +23,7 @@
       'file': 'HHS.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Alik_Bljk_HHS_BH_Bias_HA_S_SAV_UserArgs_MT1oPNOd09f6NDbRS-xCh2JnzflRtBmGH51JAsZHCRiX_g=',
+          'basename': 'Cijk_Alik_Bljk_HHS_BH_Bias_HA_S_SAV_UserArgs_MT1koFFGpVL9yYVJpUPO97fk4uqxWKENDykyPK2O0CQhXE=',
           'err': 0,
         }),
       ]),
@@ -32,7 +32,7 @@
       'file': 'HSS.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Alik_Bljk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT1oPNOd09f6NDbRS-xCh2JnzflRtBmGH51JAsZHCRiX_g=',
+          'basename': 'Cijk_Alik_Bljk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT1koFFGpVL9yYVJpUPO97fk4uqxWKENDykyPK2O0CQhXE=',
           'err': 0,
         }),
       ]),
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_emit_gfx1250_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_emit_gfx1250_char.ambr
index 87ca0975fe12..2e19eee1e285 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_emit_gfx1250_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_emit_gfx1250_char.ambr
@@ -5,7 +5,7 @@
       'file': 'BBS.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Ailk_Bjlk_BBS_BH_Bias_A_UserArgs_MT16x16x322B4FnC0zasYp4cXH3fT-GKmdbjeLh6UCDBMSbnvr9H8=',
+          'basename': 'Cijk_Ailk_Bjlk_BBS_BH_Bias_A_UserArgs_MT16x16x32fskvF-LdB9pCDb5Lk-YWNWT7yA1ZpgoZCFaV4S_LkaU=',
           'err': 0,
         }),
       ]),
@@ -14,7 +14,7 @@
       'file': 'F4_MX.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Alik_Bljk_F4BS_MXAE8B32_MXBE8B32_BH_Bias_HAz_J2wAOUGT1m_1qie4t6SOBIJJoMHRuEepJqyROT6IM=',
+          'basename': 'Cijk_Alik_Bljk_F4BS_MXAE8B32_MXBE8B32_BH_Bias_HAuVSexCu7NbiS_hMe_6fkxQ6QFaiFgTbZA-MnUSCg0NM=',
           'err': 0,
         }),
       ]),
@@ -23,7 +23,7 @@
       'file': 'HHS.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Ailk_Bjlk_HHS_BH_Bias_A_UserArgs_MT16x16x322B4FnC0zasYp4cXH3fT-GKmdbjeLh6UCDBMSbnvr9H8=',
+          'basename': 'Cijk_Ailk_Bjlk_HHS_BH_Bias_A_UserArgs_MT16x16x32fskvF-LdB9pCDb5Lk-YWNWT7yA1ZpgoZCFaV4S_LkaU=',
           'err': 0,
         }),
       ]),
@@ -32,7 +32,7 @@
       'file': 'SB.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Ailk_Bjlk_S_MX_B_Bias_SB_HA_S_SAB_SAV_UserA8iQvvl7Y4IR1E4EWOYJTgkvkJTtcL1ly4QaNKprHPhM=',
+          'basename': 'Cijk_Ailk_Bjlk_S_MX_B_Bias_SB_HA_S_SAB_SAV_UserAkNup8pM1Q7eqjciJzU8wtg4ohn_FHaMlYaCGFoZKsWE=',
           'err': 0,
         }),
       ]),
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_emit_gfx908_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_emit_gfx908_char.ambr
index b93f14eaab11..1a31a3e50ef0 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_emit_gfx908_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_emit_gfx908_char.ambr
@@ -5,7 +5,7 @@
       'file': 'BBS.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Ailk_Bljk_BBS_BH_Bias_HA_S_SAV_MT64x16x32_M9c3YtAWlA5MiQeenOldKlM36sy-9_qv1fabX_TyQwOU=',
+          'basename': 'Cijk_Ailk_Bljk_BBS_BH_Bias_HA_S_SAV_MT64x16x32_M-cY0TH6MR54ICDR3LjNMhONrN9-kuILW7N0mfGd6HKA=',
           'err': 0,
         }),
       ]),
@@ -14,7 +14,7 @@
       'file': 'HHS.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Ailk_Bljk_HHS_BH_Bias_HA_S_SAV_UserArgs_MT6J7G6Ae32LRMBxlwYJf1VrzBkP0AfJP_9i0mFOwPsoFg=',
+          'basename': 'Cijk_Ailk_Bljk_HHS_BH_Bias_HA_S_SAV_UserArgs_MT6TaZrDb6IMzNaPv813oPHMi-S5ExgkFTiVmVkPMh7J3Q=',
           'err': 0,
         }),
       ]),
@@ -23,7 +23,7 @@
       'file': 'HSS.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Ailk_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT19vruHpB_uR9ya8M6XdN43eODwaM1OI4IteY-_N_YgJo=',
+          'basename': 'Cijk_Ailk_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT1RrQWWzWtVFQrEfoAsedTyR4Z8TbHwF0iu7RfJqPKRoU=',
           'err': 0,
         }),
       ]),
@@ -32,7 +32,7 @@
       'file': 'SB.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Alik_Bjlk_S_B_Bias_HA_S_SAV_UserArgs_MT32x3EE94MjL7eMD990KTCHmdEl4BSA2y12dus4vSQpORhic=',
+          'basename': 'Cijk_Alik_Bjlk_S_B_Bias_HA_S_SAV_UserArgs_MT32x34rB42tf-gCPngpuwIzc7qmqlhwncmKEWdcwLzDbWKqM=',
           'err': 0,
         }),
       ]),
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_emit_gfx90a_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_emit_gfx90a_char.ambr
index 2ed4cf9acbfd..0bc872cfb3ba 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_emit_gfx90a_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_emit_gfx90a_char.ambr
@@ -5,7 +5,7 @@
       'file': 'BBS.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Alik_Bjlk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT3IY1ObBC6pEHccS919kMkRB6Kj6lv6C-F9uFZC6qUf-k=',
+          'basename': 'Cijk_Alik_Bjlk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT3wNNnJAHmSv8uZwexdB2wfx-J0uimzyTgdBhsDY2cuzo=',
           'err': 0,
         }),
       ]),
@@ -14,11 +14,11 @@
       'file': 'DB.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Ailk_Bjlk_D_B_UserArgs_MT16x16x16_MI16x16x1c1IYNt5lIMYw-G5u5TCiykF2xCPzGVxCn0W392imIC8=',
+          'basename': 'Cijk_Ailk_Bjlk_D_B_UserArgs_MT16x16x16_MI16x16x11IATzv_02jhIRP8HqFbgNdS58mJaKGFwV7OYDD70k58=',
           'err': 0,
         }),
         dict({
-          'basename': 'Cijk_Ailk_Bjlk_D_B_UserArgs_MT16x16x8_MI16x16x1_NbBe2NB4bgsvTqp-xs8yI70FoZYmBtwHlgYBybaOpro=',
+          'basename': 'Cijk_Ailk_Bjlk_D_B_UserArgs_MT16x16x8_MI16x16x1_TFeGaMnCQABy4V4zcPnXUp_Jwsu37bHaKwgO0jJzsuE=',
           'err': 0,
         }),
       ]),
@@ -27,7 +27,7 @@
       'file': 'HHS.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Alik_Bjlk_HHS_BH_Bias_HA_S_SAV_UserArgs_MT3IY1ObBC6pEHccS919kMkRB6Kj6lv6C-F9uFZC6qUf-k=',
+          'basename': 'Cijk_Alik_Bjlk_HHS_BH_Bias_HA_S_SAV_UserArgs_MT3wNNnJAHmSv8uZwexdB2wfx-J0uimzyTgdBhsDY2cuzo=',
           'err': 0,
         }),
       ]),
@@ -36,7 +36,7 @@
       'file': 'HSS.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Alik_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT3IY1ObBC6pEHccS919kMkRB6Kj6lv6C-F9uFZC6qUf-k=',
+          'basename': 'Cijk_Alik_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT3wNNnJAHmSv8uZwexdB2wfx-J0uimzyTgdBhsDY2cuzo=',
           'err': 0,
         }),
       ]),
@@ -45,7 +45,7 @@
       'file': 'SB.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Alik_Bjlk_S_B_Bias_HA_S_SAV_UserArgs_MT32x36gtYEFrtri4R9NbBAbxWIAPqCugrZj_OLUVY7DDsZLQ=',
+          'basename': 'Cijk_Alik_Bjlk_S_B_Bias_HA_S_SAV_UserArgs_MT32x3EQur3k3NcvZt0BJeK8GONSTxrqpK1kXVHAnHhtmA7Xc=',
           'err': 0,
         }),
       ]),
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_emit_gfx942_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_emit_gfx942_char.ambr
index 9cf50eb93366..de3fc2c62d6a 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_emit_gfx942_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_emit_gfx942_char.ambr
@@ -5,7 +5,7 @@
       'file': 'BBS_BH_Bias_Act.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Ailk_Bljk_BBS_BH_Bias_HA_S_UserArgs_MT256x9eMdVNAj9P08WWnGI3_kWn-IMCOVYTUnYMsPHGlvUmNE=',
+          'basename': 'Cijk_Ailk_Bljk_BBS_BH_Bias_HA_S_UserArgs_MT256x97Fe4FFRBn48fCcaOjC92WqNINj_mts7OTZ3qUqmunbc=',
           'err': 0,
         }),
       ]),
@@ -14,11 +14,11 @@
       'file': 'DB.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Ailk_Bljk_D_B_UserArgs_MT16x16x16_MI16x16x1XLF799yLEiNP2b6iJnAg_azQJ6BYpi3EWPoU4mUeZsI=',
+          'basename': 'Cijk_Ailk_Bljk_D_B_UserArgs_MT16x16x16_MI16x16x1dL3qX5RxCOlCt0IavzKj9ZFv_4QAdB7-eUtmtNfgsEU=',
           'err': 0,
         }),
         dict({
-          'basename': 'Cijk_Ailk_Bljk_D_B_UserArgs_MT16x16x8_MI16x16x1_qrEz6_f4QP_c_lns5bPZdTmKkwFVrLiBRsSgrl52JCA=',
+          'basename': 'Cijk_Ailk_Bljk_D_B_UserArgs_MT16x16x8_MI16x16x1_BLePoc7d_ClpkGqPFPy4akeEbHytPY2KZcxTSfd_vpg=',
           'err': 0,
         }),
       ]),
@@ -27,7 +27,7 @@
       'file': 'DTV.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Alik_Bljk_BBS_STA_BH_UserArgs_MT512x128x64_8cZ71mTFtrY8X6P-cUPliOl3J3mf0WSz5HBp8JkLUUQ=',
+          'basename': 'Cijk_Alik_Bljk_BBS_STA_BH_UserArgs_MT512x128x64_xwsrXfSE9CbMLV4XUyaBQVfXzq84ZGrJgxpdjEtdQlU=',
           'err': 0,
         }),
       ]),
@@ -36,19 +36,19 @@
       'file': 'F8N_multi.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Alik_Bljk_F8NF8NS_BH_Bias_H_AuxH_AmaxD_HA_S1Qc4vbvTw9FqBjQ-5S7ZIRN3TwuzLjExL27OAlUhio0=',
+          'basename': 'Cijk_Alik_Bljk_F8NF8NS_BH_Bias_H_AuxH_AmaxD_HA_S1r2eiXTGu-BeINB9TozaMzH2W5Ofh7G93y2G3jgFQQk=',
           'err': 0,
         }),
         dict({
-          'basename': 'Cijk_Alik_Bljk_F8NF8NS_BH_Bias_H_AuxH_AmaxD_HA_S1VcRHxkmupbG60ZB7x1pWveXxr4YXNWxXbEnktB_5YM=',
+          'basename': 'Cijk_Alik_Bljk_F8NF8NS_BH_Bias_H_AuxH_AmaxD_HA_S4CGElzg0fgamtrNK9ACCu504Xjl-zrBg9bW9rYnhdig=',
           'err': 0,
         }),
         dict({
-          'basename': 'Cijk_Alik_Bljk_F8NF8NS_BH_Bias_H_AuxH_AmaxD_HA_SqeVU7Haj7uWlbF5IqgUqPrjZ8cKT9_Wev9Gl4jF9TTY=',
+          'basename': 'Cijk_Alik_Bljk_F8NF8NS_BH_Bias_H_AuxH_AmaxD_HA_SOH4GivBu9PtOwaM5lu_itobACf5toizYVwWFnbbBM9A=',
           'err': 0,
         }),
         dict({
-          'basename': 'Cijk_Alik_Bljk_F8NF8NS_BH_Bias_H_AuxH_AmaxD_HA_SsbDUW0ffXSblcWbPIoIiEUMokkN_gCQIgS94OThhqP8=',
+          'basename': 'Cijk_Alik_Bljk_F8NF8NS_BH_Bias_H_AuxH_AmaxD_HA_Sdrx5sNOzh7RUMI3eSeVY4KOYOf62pulh-zyIw6S9R8I=',
           'err': 0,
         }),
       ]),
@@ -57,7 +57,7 @@
       'file': 'GG.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Ailk_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT16SeX6gx7xWuYbi3PUyxkZJIDKzWC4gacqbJeqrrNqY0=',
+          'basename': 'Cijk_Ailk_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT1Z8pMSIN_BBaLVpEHnA89glWRwxvLDg2N-7_hVf5dSyU=',
           'err': 0,
         }),
       ]),
@@ -66,11 +66,11 @@
       'file': 'GSU.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Ailk_Bljk_D_B_UserArgs_MT16x16x16_MI16x16x1XLF799yLEiNP2b6iJnAg_azQJ6BYpi3EWPoU4mUeZsI=',
+          'basename': 'Cijk_Ailk_Bljk_D_B_UserArgs_MT16x16x16_MI16x16x1dL3qX5RxCOlCt0IavzKj9ZFv_4QAdB7-eUtmtNfgsEU=',
           'err': 0,
         }),
         dict({
-          'basename': 'Cijk_Ailk_Bljk_D_B_UserArgs_MT16x16x8_MI16x16x1_qrEz6_f4QP_c_lns5bPZdTmKkwFVrLiBRsSgrl52JCA=',
+          'basename': 'Cijk_Ailk_Bljk_D_B_UserArgs_MT16x16x8_MI16x16x1_BLePoc7d_ClpkGqPFPy4akeEbHytPY2KZcxTSfd_vpg=',
           'err': 0,
         }),
       ]),
@@ -79,7 +79,7 @@
       'file': 'Grad.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Ailk_Bjlk_S_B_Bias_D_GradS_HA_S_SAV_UserArgQM2TnuRRC9Bc-J5AiR8DYNkP01YWewx-agM7_j4Lsl0=',
+          'basename': 'Cijk_Ailk_Bjlk_S_B_Bias_D_GradS_HA_S_SAV_UserArgSmQRW4Lzn_Mi_4Tyw3ZLm2bE2JXW7cyXBlRzcKFNdVs=',
           'err': 0,
         }),
       ]),
@@ -88,7 +88,7 @@
       'file': 'HHS_BH_Bias_GG.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Ailk_Bjlk_HHS_BH_Bias_HA_S_SAV_UserArgs_MT1gHOstBxLWntStkEmleJ-9nV3FdW5A_kY0iVeJWHdPDo=',
+          'basename': 'Cijk_Ailk_Bjlk_HHS_BH_Bias_HA_S_SAV_UserArgs_MT1i58JIcu4-V1ZHb6PTI_eUaSij6iZuAUcPyM4FPYFrRI=',
           'err': 0,
         }),
       ]),
@@ -97,7 +97,7 @@
       'file': 'HSS_BH_Bias.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Alik_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT3sU6CrhlAM13fvxNXoxtKcDzvMD3vel7lJiIzSBq5sHo=',
+          'basename': 'Cijk_Alik_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT3CcVTdHcpjVTzuY2D4vYNI3Nab1gIvp-I-f6e6TUY8yU=',
           'err': 0,
         }),
       ]),
@@ -106,11 +106,11 @@
       'file': 'LSU_MX.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Alik_Bjlk_S_MX_B_Bias_HA_S_SAV_UserArgs_MT1OLGfojG71ygY2zhDLj7J5awOtQCzCO93hHM9QBWfsh8=',
+          'basename': 'Cijk_Alik_Bjlk_S_MX_B_Bias_HA_S_SAV_UserArgs_MT1nuam_o97DX9CseOid8zDB7AGF7T5O4GVLql2mxWsmnc=',
           'err': 0,
         }),
         dict({
-          'basename': 'Cijk_Alik_Bjlk_S_MX_B_Bias_HA_S_SAV_UserArgs_MT1kRGa4UiRq9yR7bwHP63rdrMgYh9FQBcq-xjeuC4t63A=',
+          'basename': 'Cijk_Alik_Bjlk_S_MX_B_Bias_HA_S_SAV_UserArgs_MT1vxil4EIHlKWX4NEuo0YTupFj2w0dfnxKfgTN4Uw0d0w=',
           'err': 0,
         }),
       ]),
@@ -119,7 +119,7 @@
       'file': 'MX.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Ailk_Bljk_S_MX_B_Bias_HA_S_SAV_UserArgs_MT36W7iF9uqrs1KzkWctke-5VBnobNWqNh1TByECxNIl74=',
+          'basename': 'Cijk_Ailk_Bljk_S_MX_B_Bias_HA_S_SAV_UserArgs_MT3v1reSN_kr8DlzHOIXX_3k_IUNnFLZGCVsl4rnUYLc3s=',
           'err': 0,
         }),
       ]),
@@ -128,7 +128,7 @@
       'file': 'SB_Bias_Aux.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Ailk_Bjlk_S_B_Bias_AuxS_HA_S_SAV_UserArgs_MqG4Wwuz-wK6y3jvfKad0eovyp4Wa62uTnQHqH2JLWI4=',
+          'basename': 'Cijk_Ailk_Bjlk_S_B_Bias_AuxS_HA_S_SAV_UserArgs_Mk8r91wBeP0oLSvxEoF-I6KgkiY1hzt00X-flcDkIno8=',
           'err': 0,
         }),
       ]),
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_emit_gfx950_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_emit_gfx950_char.ambr
index 3daf6e1cabf2..7bf268303b92 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_emit_gfx950_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_emit_gfx950_char.ambr
@@ -5,7 +5,7 @@
       'file': 'BBS.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Alik_Bjlk_BBS_BH_Bias_AuxB_HA_S_SAV_UserArgx1JvbOyZcZnNiVD3Ka-e-uAVxfTykKnFazbIJ3FpU-o=',
+          'basename': 'Cijk_Alik_Bjlk_BBS_BH_Bias_AuxB_HA_S_SAV_UserArgBnKhTSS-cz2M_xqF7kjOyGKcaPXpV_rBZrIvmCput5I=',
           'err': 0,
         }),
       ]),
@@ -40,7 +40,7 @@
       'file': 'HHS.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Alik_Bljk_HHS_BH_Bias_HA_S_SAV_UserArgs_MT1UllvxMnRAK73ZgLdvBdjzP-HdaX-Kcjj_bw2GrCbsvI=',
+          'basename': 'Cijk_Alik_Bljk_HHS_BH_Bias_HA_S_SAV_UserArgs_MT1N43DK6MRSsCt4446DY1wfS2mLdCUO4voP5t_XKfJ4Cg=',
           'err': 0,
         }),
       ]),
@@ -49,7 +49,7 @@
       'file': 'HSS.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Alik_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT3bhCijMK38v052ojO8R1DlE9piSq0NkGI1jO7yGqdbgU=',
+          'basename': 'Cijk_Alik_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT3cnukg-tC8fL-fztmFiLNjUbk1RR_r1J5J8Joyp8NNdE=',
           'err': 0,
         }),
       ]),
@@ -58,7 +58,7 @@
       'file': 'I8_GSU.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Ailk_Bjlk_I8I8S_BH_Bias_HA_S_SAV_UserArgs_M_iQ0G1nh_AQJ6oxftoIDFID-p5c2KnAGsRUlvT3ZOEs=',
+          'basename': 'Cijk_Ailk_Bjlk_I8I8S_BH_Bias_HA_S_SAV_UserArgs_M8fmu5XvA9tLiYSDtPHKku-6DWDhWYiPiiN_D-EuUBYs=',
           'err': 0,
         }),
       ]),
@@ -67,7 +67,7 @@
       'file': 'MX.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Alik_Bjlk_S_MX_B_Bias_HA_S_SAV_UserArgs_MT4ePLfW5bNgk6m3ISxvEyYFuUo9MLq9x2ZAWZRn5TMf1c=',
+          'basename': 'Cijk_Alik_Bjlk_S_MX_B_Bias_HA_S_SAV_UserArgs_MT4HcjyU_fLuCDLXYlv9kgJ_5bwQK9sp72NVFu-okBy0co=',
           'err': 0,
         }),
       ]),
@@ -76,7 +76,7 @@
       'file': 'MX_StreamK.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Alik_Bjlk_S_MX_B_Bias_HA_S_SAV_UserArgs_MT4ePLfW5bNgk6m3ISxvEyYFuUo9MLq9x2ZAWZRn5TMf1c=',
+          'basename': 'Cijk_Alik_Bjlk_S_MX_B_Bias_HA_S_SAV_UserArgs_MT4HcjyU_fLuCDLXYlv9kgJ_5bwQK9sp72NVFu-okBy0co=',
           'err': 0,
         }),
       ]),
@@ -85,7 +85,7 @@
       'file': 'SB.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Ailk_Bjlk_S_B_Bias_AuxS_HA_S_SAV_UserArgs_MOVcSe7XYUy6V9YWUVZZg0aNNR2qHfpUEni1eulSUxA4=',
+          'basename': 'Cijk_Ailk_Bjlk_S_B_Bias_AuxS_HA_S_SAV_UserArgs_M3MShAdXhfloYgJE0x5JCpWqTusz1A0acbzftWW5Gr00=',
           'err': 0,
         }),
       ]),
@@ -94,23 +94,23 @@
       'file': 'StreamK_B8F8.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Alik_Bljk_B8F8_BH_Bias_HA_S_SAB_SAV_UserArg0usEK3cJaLP7R94jK3zxoYQWB8ojB0yVQLYeZVgQBaw=',
+          'basename': 'Cijk_Alik_Bljk_B8F8_BH_Bias_HA_S_SAB_SAV_UserArgE1TpHzRga8zRXz3tAwMnur5lG1Ttma6UHG4ANk8r3R0=',
           'err': 0,
         }),
         dict({
-          'basename': 'Cijk_Alik_Bljk_B8F8_BH_Bias_HA_S_SAB_SAV_UserArgLRFA3dZxQ3dkCHGO_f05TOCr6LPsBJX9sB2ePdI9a4A=',
+          'basename': 'Cijk_Alik_Bljk_B8F8_BH_Bias_HA_S_SAB_SAV_UserArgKcP26jeLWnuljyqds9Ac8C4zuGWo1EWh9cymkgS4JlE=',
           'err': 0,
         }),
         dict({
-          'basename': 'Cijk_Alik_Bljk_B8F8_BH_Bias_HA_S_SAB_SAV_UserArghQRxbukFSa_B_ypAqdQ-DAIuIlqKy5um6hoeNGAM4oE=',
+          'basename': 'Cijk_Alik_Bljk_B8F8_BH_Bias_HA_S_SAB_SAV_UserArgiWYhLE83jsaws22nCbBIOm4kZB0y1UrP-ugTcdJ6348=',
           'err': 0,
         }),
         dict({
-          'basename': 'Cijk_Alik_Bljk_B8F8_BH_Bias_HA_S_SAB_SAV_UserArgiaTGrYKcPaD56dxXizDCmQp6kIn9ZQCk0wk2J8g4IFg=',
+          'basename': 'Cijk_Alik_Bljk_B8F8_BH_Bias_HA_S_SAB_SAV_UserArgjERlLcdiqYq-m_IWyGNKOQ_YBA0bAGDw9PvFDjfvJ50=',
           'err': 0,
         }),
         dict({
-          'basename': 'Cijk_Alik_Bljk_B8F8_BH_Bias_HA_S_SAB_SAV_UserArgrSl-_6niMFLLouHSFcveQhT_0DOR0D_uoYgJUZjEbEc=',
+          'basename': 'Cijk_Alik_Bljk_B8F8_BH_Bias_HA_S_SAB_SAV_UserArgoFTmEwRO3zFvDNH90j_XJRIv-lvX_UXggWGgUbZ1zYk=',
           'err': 0,
         }),
       ]),
@@ -119,7 +119,7 @@
       'file': 'StreamK_F8F8S.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Alik_Bljk_F8F8S_BH_Bias_HA_S_SAV_UserArgs_MYZ1s6H47_UTltUp6fVHLX-Ojj9PHHiqOkOY5rniGgp0=',
+          'basename': 'Cijk_Alik_Bljk_F8F8S_BH_Bias_HA_S_SAV_UserArgs_M71K1smI3XgNpg9XQGoVWdOYoT5zCAL51p-uQ3kfJZCM=',
           'err': 0,
         }),
       ]),
@@ -128,7 +128,7 @@
       'file': 'WaveSplitK.yaml',
       'kernels': list([
         dict({
-          'basename': 'Cijk_Alik_Bjlk_HHS_BH_Bias_HA_S_SAV_UserArgs_MT3bhCijMK38v052ojO8R1DlE9piSq0NkGI1jO7yGqdbgU=',
+          'basename': 'Cijk_Alik_Bjlk_HHS_BH_Bias_HA_S_SAV_UserArgs_MT3cnukg-tC8fL-fztmFiLNjUbk1RR_r1J5J8Joyp8NNdE=',
           'err': 0,
         }),
       ]),
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_harness_smoke.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_harness_smoke.ambr
index c8eaa07617a9..a44ef12b7dc8 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_harness_smoke.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_harness_smoke.ambr
@@ -2,7 +2,7 @@
 # name: test_emit_golden_digest
   list([
     dict({
-      'basename': 'Cijk_Alik_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT3sU6CrhlAM13fvxNXoxtKcDzvMD3vel7lJiIzSBq5sHo=',
+      'basename': 'Cijk_Alik_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs_MT3CcVTdHcpjVTzuY2D4vYNI3Nab1gIvp-I-f6e6TUY8yU=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_activation_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_activation_char.ambr
index c7c98b6996c5..8e202acab202 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_activation_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_activation_char.ambr
@@ -2,7 +2,7 @@
 # name: test_r2_activation_gfx942_golden
   list([
     dict({
-      'basename': 'Cijk_Ailk_Bljk_HHS_BH_Bias_D_GradH_A_S_UserArgs_bwosqD0LND-p9G64ReX-Gvgt5qvaNDR-u1fqRh2tE_c=',
+      'basename': 'Cijk_Ailk_Bljk_HHS_BH_Bias_D_GradH_A_S_UserArgs_1Z81RGJJmFfLOUfeY-jI6155OYYMuthkC_hoP_WubK8=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_gsu_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_gsu_char.ambr
index 6f58076e2254..208e6c6cf608 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_gsu_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_gsu_char.ambr
@@ -2,11 +2,11 @@
 # name: test_r2_gsu_golden
   list([
     dict({
-      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_A_S_UserArgs_MT64x16x7Zj6Hz8AaBAwVyUqP6m_gm_pSBDp6bCD99NBVr-_eUQ=',
+      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_A_S_UserArgs_MT64x16x2MLCeKH-GsFV_p4XATG8MNz8Dj1fok2bc84PKj53Lo8=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_A_S_UserArgs_MT64x16xv21MMjZ0uaPycUD4xS5fWzxW5mFOoq_EWwFIAW6woL4=',
+      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_A_S_UserArgs_MT64x16xky_-dI0IsI6RVmAT5yZgno7FQduwNOfzNBuMDUThzIs=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_kwconv_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_kwconv_char.ambr
index 90e0a2d15453..e58afff881ed 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_kwconv_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_kwconv_char.ambr
@@ -2,7 +2,7 @@
 # name: test_kwconv_gfx942_golden
   list([
     dict({
-      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_B_A_S_UserArgs_MT64x12B3PIH_7a13CYp9BoxCSWQ49Ib4hw7R4Hvq9xnXW2Ns=',
+      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_B_A_S_UserArgs_MT64x1NY74UBIiLGloKBWKClGkhlFVWVo_o5QZ_ctWSeoq3Ls=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_lra_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_lra_char.ambr
index c74288b01501..b5574a986960 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_lra_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_lra_char.ambr
@@ -2,11 +2,11 @@
 # name: test_r2_lra_gfx942_golden
   list([
     dict({
-      'basename': 'Cijk_Alik_Bjlk_HHS_BH_UserArgs_MT16x8x64_SN_LDSBdXxZd3VDUWBq22zE0GV6SZVngZyzkxOTsJdNSUzy0mc=',
+      'basename': 'Cijk_Alik_Bjlk_HHS_BH_UserArgs_MT16x8x64_SN_LDSBANG4BDzwuRe_B49k-5O60hh9FIcd-f24j4PrmtQETas=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bjlk_HHS_BH_UserArgs_MT64x4x64_SN_LDSBRL0iX_IIchurOPm-7F9lCHABNxnurQeloU3o1yuKsAQ=',
+      'basename': 'Cijk_Alik_Bjlk_HHS_BH_UserArgs_MT64x4x64_SN_LDSBfPJQAoxs4jVpvCcytF79A5yxaXI1j5TdNH8orAxcGlc=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_mac_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_mac_char.ambr
index 64a8c1807c00..d673aeb07ad2 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_mac_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_mac_char.ambr
@@ -2,27 +2,27 @@
 # name: test_r2_mac_hhs_dot2_golden
   list([
     dict({
-      'basename': 'Cijk_Alik_Bljk_HHS_BH_UserArgs_MT16x32x16_SN_LDSL1ZOBtp_x4vcEhSTyncwXyQYn2PnYD6AAVAVJy6FHdk=',
+      'basename': 'Cijk_Alik_Bljk_HHS_BH_UserArgs_MT16x32x16_SN_LDSktHYQlLIe_io9uzFJYV5hvztqGY6uoDqheisc-Z-CQc=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_HHS_BH_UserArgs_MT16x64x16_SN_LDSgpJCbKZaU6WWwBt0Ur39jdC9y2qvHyZxSORp8c_CC4I=',
+      'basename': 'Cijk_Alik_Bljk_HHS_BH_UserArgs_MT16x64x16_SN_LDSngHJqUa-7Z9AtnaQkCfdoAYGNMtseB-eVQG5blWkPe0=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_HHS_BH_UserArgs_MT32x32x16_SN_LDSrWlu0oy0k0M4pAoSZ09XxCUHwXgqf1ucxA6AAlZGDJQ=',
+      'basename': 'Cijk_Alik_Bljk_HHS_BH_UserArgs_MT32x32x16_SN_LDS9v3mi9oBmz5I7BS8swfYqqQN5v04f3qJbCax_Cq9MvQ=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_HHS_BH_UserArgs_MT32x64x16_SN_LDSUtD1DZJ7fsyDCxYC3KfT6c8bCVkm07sZWYMX2-WCfOA=',
+      'basename': 'Cijk_Alik_Bljk_HHS_BH_UserArgs_MT32x64x16_SN_LDSr7jvmAcaWnA9e6OrSm0QZBUwkk8eJQb8ceP8H17GvA0=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_HHS_BH_UserArgs_MT64x32x16_SN_LDSkirfjiycGGwnUaNLZuxJEhQ6tJkqfjCuRllJOUobVfs=',
+      'basename': 'Cijk_Alik_Bljk_HHS_BH_UserArgs_MT64x32x16_SN_LDShzhObaCvRLolZA5g0VjAxYqMwEf0J7DZKCl1VkNcvvU=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_HHS_BH_UserArgs_MT8x64x16_SN_LDSBGP5TacojLNqPNy7YBE4GVNOzdhBqYgBAjbh84FzJ7UY=',
+      'basename': 'Cijk_Alik_Bljk_HHS_BH_UserArgs_MT8x64x16_SN_LDSB9DJ4DFqrDO9ObZ0LMn76bTK_m8DgxuHknz_SSwfKTPs=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_shiftvector_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_shiftvector_char.ambr
index 05dee520db8d..140b8944171b 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_shiftvector_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_shiftvector_char.ambr
@@ -2,15 +2,15 @@
 # name: test_r2_shiftvector_golden
   list([
     dict({
-      'basename': 'Cijk_Ailk_Bjlk_S_B_UserArgs_MT128x64x16_SN_LDSB0sH-JXCLh3dYIQKh2S9-3sK8e-YJXfJBxRDhEAq7cjvc=',
+      'basename': 'Cijk_Ailk_Bjlk_S_B_UserArgs_MT128x64x16_SN_LDSB0vmX4-WcG49AcrGdkyqtN5NzBoDQB3HS_BXqz7679mk4=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bjlk_S_B_UserArgs_MT64x128x16_SN_LDSB0vtvY3ZdWTUU7dpKQ-nw3Krj8AAFPb4XeJtvF1WzVWpY=',
+      'basename': 'Cijk_Ailk_Bjlk_S_B_UserArgs_MT64x128x16_SN_LDSB0wmz4t1jzHaj12PhITIMmOIAr2Y-MlSz8yxIQubvnN1w=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bjlk_S_B_UserArgs_MT64x64x16_SN_LDSB0_qhiqbvIc5-BiBghoWhTGWKQsdwkmiO0tSUz2mEkygYA=',
+      'basename': 'Cijk_Ailk_Bjlk_S_B_UserArgs_MT64x64x16_SN_LDSB0__TRUZxDp29c3oYMWfMeaUa1SGH7bkyS-eepzr8Y0LyA=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_solution_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_solution_char.ambr
index 83c4e4f16788..2a304a95845e 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_solution_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_solution_char.ambr
@@ -2,19 +2,19 @@
 # name: test_r2_solution_gfx942_golden
   list([
     dict({
-      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_A_S_UserArgs_MT64x16x447EqQAKhHgjmbCxtPk4kjwg_DwVa577jAQYMgNtNec=',
+      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_A_S_UserArgs_MT64x16x2MLCeKH-GsFV_p4XATG8MNz8Dj1fok2bc84PKj53Lo8=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_A_S_UserArgs_MT64x16xKmeIoDdY2LsndhIjTAbGI-UmoheOlhHCRltZBDNBRVQ=',
+      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_A_S_UserArgs_MT64x16x7Rbs63plMSaLHG8GOkxNKAfI8-jJiTYD0xzaQ8J1CDU=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_A_S_UserArgs_MT64x16xc_DWRLoALWf4g8aVDtbuxSs520zE4LbzQuxDSv-r7bM=',
+      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_A_S_UserArgs_MT64x16xBUfuCnk8_g7R3s3MGZuwljHRgZtOuwsltL4Gemck2Ow=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_A_S_UserArgs_MT64x16xv21MMjZ0uaPycUD4xS5fWzxW5mFOoq_EWwFIAW6woL4=',
+      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_A_S_UserArgs_MT64x16xZVrOmvLYMgPZmiPjgdOvvq1MxjTmNZU8CokLYpgZsLA=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_store_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_store_char.ambr
index add0a931b8d6..9777c2fc6916 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_store_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_store_char.ambr
@@ -2,19 +2,19 @@
 # name: test_r2_store_gfx942_golden
   list([
     dict({
-      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_SAV_UserArgs_MT64x16x2S-p4xE2gXGRZnYPUe0-kAwmTKNXSX7L2m2xMRdR8RA=',
+      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_SAV_UserArgs_MT64x16xHrLUf9oJusha_OLf_0Hf-hi5tOHN-xS5fUA_hK03gfE=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_SAV_UserArgs_MT64x16xkMZkjCAdpTnDPRU8DW-IAvVf2JaFdfJpl55eJ91c3WA=',
+      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_SAV_UserArgs_MT64x16xLGx7-Y2M1u2LKFYxqsZ-lWK0AJ422Dxb9dpHq8go4Yg=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_SAV_UserArgs_MT64x16xoKxwgcyHKTV17t9Fgvoiwt8mm98_ryOnIrLM-klewK4=',
+      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_SAV_UserArgs_MT64x16xjjpoTSMzs8AIpzYI-FPZ9M047InyaYG0FzxGGDBLRAE=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_SAV_UserArgs_MT64x16xuPGEtY_j4wICoDCMh3dJLAQ_f7kx_gscW0Xmdqkv0To=',
+      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_SAV_UserArgs_MT64x16xuYhuan24ICcO2xBgSRs45QT4UBBMSlAL-kTZpzBBylE=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_streamk_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_streamk_char.ambr
index f520a8cd58ff..e39c57a72769 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_streamk_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_streamk_char.ambr
@@ -2,19 +2,19 @@
 # name: test_r2_streamk_gfx942_golden
   list([
     dict({
-      'basename': 'Cijk_Ailk_Bjlk_S_B_UserArgs_MT128x128x16_MI16x16075MU-vd-b9VZWmcuz767MA0fw_ewfC8TIhe9ItqpG8=',
+      'basename': 'Cijk_Ailk_Bjlk_S_B_UserArgs_MT128x128x16_MI16x16G6FLshMSdDLhWRyC_S7RdbT6eMbZBDzNFTujs72RmkU=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bjlk_S_B_UserArgs_MT128x128x16_MI16x16EwuyvtgASua8VutE-bwv6T2KpyywjQb8RMH_j-hPlj4=',
+      'basename': 'Cijk_Ailk_Bjlk_S_B_UserArgs_MT128x128x16_MI16x16RCEEdxstCSPAGoGLQMl-9UKePT2VjQoh0IYwux58djE=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bjlk_S_B_UserArgs_MT128x128x16_MI16x16EwuyvtgASua8VutE-bwv6T2KpyywjQb8RMH_j-hPlj4=',
+      'basename': 'Cijk_Ailk_Bjlk_S_B_UserArgs_MT128x128x16_MI16x16p9JmY_wJcTLq8j_gkgl8hcHOKZ2IVjzIUtGBbyxIL24=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bjlk_S_B_UserArgs_MT128x128x16_MI16x16qS4SCfeWGXXQf4dRvX6IaBpri1HXUlrT7Ve41ghnDcw=',
+      'basename': 'Cijk_Ailk_Bjlk_S_B_UserArgs_MT128x128x16_MI16x16p9JmY_wJcTLq8j_gkgl8hcHOKZ2IVjzIUtGBbyxIL24=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_subtile_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_subtile_char.ambr
index 6b2695658b0c..d1f9a230bec5 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_subtile_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_subtile_char.ambr
@@ -2,35 +2,35 @@
 # name: test_r2_subtile_gfx950_golden
   list([
     dict({
-      'basename': 'Cijk_Alik_Bljk_BBS_BH_UserArgs_MT128x128x64_MI161trGTYH7BzR3j2Q2qNvJGaze8nFCRnHVJi2Ttqqaamo=',
+      'basename': 'Cijk_Alik_Bljk_BBS_BH_UserArgs_MT128x128x64_MI16EL7h3ZbQJsvMtJ0UnR4kfJ09QV8Iq1dzZRUhTnKxc9I=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_BBS_BH_UserArgs_MT128x128x64_MI16JgfylG0zuP56erxhFoAJb68f0p9I7DAxGwhknrLD6T4=',
+      'basename': 'Cijk_Alik_Bljk_BBS_BH_UserArgs_MT128x128x64_MI16sNPUjdOpu5msiHNeF6zXd5n7QjOzXx5bjkWd2FyO_OM=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_BBS_BH_UserArgs_MT128x64x64_MI16xZZ01gRfovJxi1MgOBgp8RNhjBz_NdOzJskAA9XazYhc=',
+      'basename': 'Cijk_Alik_Bljk_BBS_BH_UserArgs_MT128x64x64_MI16x4BxeC4SlPVoVxJaBEeA7t-4nZGSpy9KyJKXx5fnVOtE=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_BBS_BH_UserArgs_MT128x64x64_MI16xZt6rX3Mm_zeGOwIcaqUKIQ38ZhYN55LLmdgLGkAuP-8=',
+      'basename': 'Cijk_Alik_Bljk_BBS_BH_UserArgs_MT128x64x64_MI16xX0ABYkTRz48sKGiVop3wSe_f-1EtquhUW5RpKX7jxFE=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_BBS_BH_UserArgs_MT32x32x64_MI16x1FGsM3cSq1eaJoTRuKDLIkWvtVvOFT4yyGsgdExciSF8=',
+      'basename': 'Cijk_Alik_Bljk_BBS_BH_UserArgs_MT32x32x64_MI16x140uaxVoV8xR4s4UHLo57wXcMLE9XCcwqt8KaSssd6V0=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_BBS_BH_UserArgs_MT32x32x64_MI16x1bjFJZWMzSPzXyhs4xMtGmOyh3_TJsftel1ylI4lD72g=',
+      'basename': 'Cijk_Alik_Bljk_BBS_BH_UserArgs_MT32x32x64_MI16x1C2lw1AGjZ2J7gi08lGdfD0RgXRz8COwta8dLtiUdsYM=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_BBS_BH_UserArgs_MT64x128x64_MI16xEh7X_nhh2uaiZVbNxY0Y68N4BgqxYj9tdSHf_tucxAQ=',
+      'basename': 'Cijk_Alik_Bljk_BBS_BH_UserArgs_MT64x128x64_MI16xAsDlfivpqBCW5Xh_RFRvc6Z5Pflo4_cvvrbMkStR900=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_BBS_BH_UserArgs_MT64x128x64_MI16xx_GSxdHWhYFsVrOc16nbBjLhu8WH1hONQmf8h_w8jQ8=',
+      'basename': 'Cijk_Alik_Bljk_BBS_BH_UserArgs_MT64x128x64_MI16xp0DJQ5vxrKny2TOyLjrkyb_7EFKKsTAC_8V2x2U_4LQ=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_wgm_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_wgm_char.ambr
index 670d07f8974a..5f11b0ab7948 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_wgm_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r2_wgm_char.ambr
@@ -2,23 +2,23 @@
 # name: test_r2_wgm_gfx942_golden
   list([
     dict({
-      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_A_S_UserArgs_MT64x16x0VZcisUBFP3pbeOHUMG9XMO7c4tQoegdPPcz45zbzOU=',
+      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_A_S_UserArgs_MT64x16xHAaaMUTXFTpi-iJ9e35OTqPOiyzpaXmGTh8h2ZMteew=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_A_S_UserArgs_MT64x16x1cWk0bO4UQgkCoquRD5WV0Jb4XNXXpbk9KqiSvOV3yY=',
+      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_A_S_UserArgs_MT64x16xLlRF5Gygf_0qnJ0w8rZmtDWV2LoOu1ga_CHNwfe2FKU=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_A_S_UserArgs_MT64x16xClmgD-w1MpK1S24jgD_ecF_zi2TliEInCaBTuRrabDc=',
+      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_A_S_UserArgs_MT64x16xWWGTuS1YaVEA0g6pr3ipsSd5y_0Sn9SpkifCb-SDwwo=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_A_S_UserArgs_MT64x16xDqqhCShtbf5OwkSCYhiuvoOPyBujSvofMbF1oq2gTjk=',
+      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_A_S_UserArgs_MT64x16xlkYOoKwZ0YF34OD6IJe9_7CVmZG11dc4AkaJTv_rk9k=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_A_S_UserArgs_MT64x16xQiXU9-Kw4R7QH6JUmU9rNgjoc8axfobU4K0fepbuzZo=',
+      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_A_S_UserArgs_MT64x16xoibvcO4b_AacfEnfBAm8EOSdzCRwAEt5ykOUAQ7CV-A=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_addrstore_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_addrstore_char.ambr
index 2cd8ee199087..d597a4fd3595 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_addrstore_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_addrstore_char.ambr
@@ -2,11 +2,11 @@
 # name: test_r3_addrstore_gfx942_golden
   list([
     dict({
-      'basename': 'Cijk_Alik_Bljk_F8NSS_BH_SABV_SAV_UserArgs_MT128x5N13Z-g1GXe67Ir9KWCsSwQKtmhW_BmmlBQKPsExGCE=',
+      'basename': 'Cijk_Alik_Bljk_F8NSS_BH_SABV_SAV_UserArgs_MT128xQIJ47iwV0NEW-zv1uMo9e4TWqAtnHXYXaezMCqiihBQ=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_F8NSS_BH_SABV_SAV_UserArgs_MT256xqtJ3dcRQt26eeUfjcQFKYi-dFaSJVLY964Nf6__J7tw=',
+      'basename': 'Cijk_Alik_Bljk_F8NSS_BH_SABV_SAV_UserArgs_MT256xenLQ3p54yPcO1JeE5CS3X4gImMrWK0jdWp1OfySc8fo=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_datamover_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_datamover_char.ambr
index 6b00a991fb38..cf10d198d9a5 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_datamover_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_datamover_char.ambr
@@ -2,11 +2,11 @@
 # name: test_r3_datamover_gfx1250_golden
   list([
     dict({
-      'basename': 'Cijk_Alik_Bljk_F4BS_BH_UserArgs_MT16x16x128_MI16_QobKpeqHxanvL7qwfLO2v0ata2Bsw5nva5aJLS5RhY=',
+      'basename': 'Cijk_Alik_Bljk_F4BS_BH_UserArgs_MT16x16x128_MI16MVzLZOqm-oDhwdblMUu9axemtO1vW31Ug1---EVAxT4=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_F4BS_BH_UserArgs_MT32x32x128_MI167MEFMPkl0yanICd8aYGkQ-qZqG20dGGqi59GoZn4JvU=',
+      'basename': 'Cijk_Alik_Bljk_F4BS_BH_UserArgs_MT32x32x128_MI16dmV4LHVjZgzfC3xHWPtz7ckoJEtYPaKw8vtI2DbVaZQ=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_globalwrite_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_globalwrite_char.ambr
index 53fd9401db18..44152bab836d 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_globalwrite_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_globalwrite_char.ambr
@@ -2,11 +2,11 @@
 # name: test_r3_globalwrite_gfx942_golden
   list([
     dict({
-      'basename': 'Cijk_Alik_Bljk_F8NBS_BH_SABV_SAV_UserArgs_MT64x6G2VsKSdxS2FftTrXCGTzEQ291zPMixhpOFDE5dBdV6I=',
+      'basename': 'Cijk_Alik_Bljk_F8NBS_BH_SABV_SAV_UserArgs_MT64x6gKZiD60fBT_DGiACz-WnXGSFKpnkRqwElW8nhIQkDmo=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_F8NBS_BH_SABV_SAV_UserArgs_MT64x6ole293Uk2j2i6sNNZagIOoYJ0Pda6ZsUtbWD56YNsIo=',
+      'basename': 'Cijk_Alik_Bljk_F8NBS_BH_SABV_SAV_UserArgs_MT64x6xFm5Mr--fnK7on3rPgUni0DPIwWR_n1RNICi0KaEyPc=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_gsu_on_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_gsu_on_char.ambr
index 09f26df7612d..8cd9b20e23de 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_gsu_on_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_gsu_on_char.ambr
@@ -2,11 +2,11 @@
 # name: test_r3_gsu_on_golden
   list([
     dict({
-      'basename': 'Cijk_Ailk_Bljk_S_B_UserArgs_MT16x16x32_MI16x16x1LWOHt42XBgtzeyeJ5wmGxPDPFAUi7Qug-lCA8UxPzrQ=',
+      'basename': 'Cijk_Ailk_Bljk_S_B_UserArgs_MT16x16x32_MI16x16x1EQ5cU_Yh1QEkJ6nnbhNUhz7zESlzWA0zequT1j4dwGg=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bljk_S_B_UserArgs_MT16x16x32_MI16x16x1XItsdpT3EmyYBajibXSdoD3A2ZSiJQRDpIFM3lmKYwQ=',
+      'basename': 'Cijk_Ailk_Bljk_S_B_UserArgs_MT16x16x32_MI16x16x1pQOquPz4tDlSTww7MnItGkhfIVZnzt1HdCIZA15tq8A=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_kwafeat_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_kwafeat_char.ambr
index 733a2e937e06..dc516d06db45 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_kwafeat_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_kwafeat_char.ambr
@@ -2,11 +2,11 @@
 # name: test_r3_kwafeat_gfx942_golden
   list([
     dict({
-      'basename': 'Cijk_Ailk_Bjlk_S_B_UserArgs_MT128x128x16_MI16x16HFa3ermAzGoH7SKOpP_QUDezickLkOw4Z1yHm8e7s74=',
+      'basename': 'Cijk_Ailk_Bjlk_S_B_UserArgs_MT128x128x16_MI16x16-gVgHuwPYh1yhH4Nn5TKKlY9vpb-n9d3VceeD8SeGb0=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bjlk_S_B_UserArgs_MT128x128x16_MI16x16tGuRm6G38JZzdeXGcIiY0l_TSQJjZ_-4EaSQouJfAg4=',
+      'basename': 'Cijk_Ailk_Bjlk_S_B_UserArgs_MT128x128x16_MI16x16y6Z4pe8nS2xEEonDe3cVeIpE4NDvM0ZkwiAUUWXUayI=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_kwfeat_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_kwfeat_char.ambr
index af1ceb4dd56b..4cb571dfe9f1 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_kwfeat_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_kwfeat_char.ambr
@@ -2,11 +2,11 @@
 # name: test_r3_kwfeat_gfx942_golden
   list([
     dict({
-      'basename': 'Cijk_Alik_Bljk_BBS_BH_UserArgs_MT64x16x32_MI16x1O91NX_HfZlYNYwF0vawiAKRdSw3YdYy4AIcL0G2BvIY=',
+      'basename': 'Cijk_Alik_Bljk_BBS_BH_UserArgs_MT64x16x32_MI16x1INwonYWsDX3BcdTJinphJoKELwQVjRxr577t1cC4414=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_BBS_BH_UserArgs_MT64x16x32_MI16x1OVC7wdVuDf7TvEnmyXPf75NtJ-Dhpf2QSCxfhwAYSlQ=',
+      'basename': 'Cijk_Alik_Bljk_BBS_BH_UserArgs_MT64x16x32_MI16x1qunh74giuLyNl0pFhl1Oh8ETnoLBnJVoZXrExdxnx4Y=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_lra_tr_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_lra_tr_char.ambr
index e4395032e66b..801da51f62f9 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_lra_tr_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_lra_tr_char.ambr
@@ -2,11 +2,11 @@
 # name: test_r3_lra_tr_gfx950_golden
   list([
     dict({
-      'basename': 'Cijk_Ailk_Bljk_BBS_BH_UserArgs_MT128x128x32_MI168Z1tQHhatXbhtOsiOfrOYel0MqQd8XzRl-hTp03ojqE=',
+      'basename': 'Cijk_Ailk_Bljk_BBS_BH_UserArgs_MT128x128x32_MI165LlQlE1FeU-Kf6qaP8WgHtINB1SYIQpohYlIN-PzMaQ=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bljk_BBS_BH_UserArgs_MT128x128x32_MI16DEs-MWvRdz38Ac2z1IsQG8KyJhN6TISikzJQTNv41sk=',
+      'basename': 'Cijk_Ailk_Bljk_BBS_BH_UserArgs_MT128x128x32_MI16xQgGUiSxhs_nKahOvJa7Q0oApPouackoTOl0IxpcqIM=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_streamk_gfx1250_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_streamk_gfx1250_char.ambr
index 61799447ef6e..6d2cae7bbeb7 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_streamk_gfx1250_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_streamk_gfx1250_char.ambr
@@ -2,7 +2,7 @@
 # name: test_r3_streamk_gfx1250_golden
   list([
     dict({
-      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgNg-IV7t7EnYLL6mzebH3P5CbTswCA1Faj4hopzB-L_A=',
+      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgWuFPSmyxgZZS9lNLqQt2ml5kLoTKRLL3JkGmxXB1SBo=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_streamk_tdmsplit_gfx1250_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_streamk_tdmsplit_gfx1250_char.ambr
index 36d436eadb57..8464cbc6e995 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_streamk_tdmsplit_gfx1250_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_streamk_tdmsplit_gfx1250_char.ambr
@@ -2,7 +2,7 @@
 # name: test_r3_streamk_tdmsplit_gfx1250_golden
   list([
     dict({
-      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgZMQJ-hHyLz5yAyZk2P0zHciaExSDGnjyuEM24FkFLYA=',
+      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgRTb8_28SDU2S5aKZHEQM10z4b86hnOEYR4I_Cfu-3L8=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_streamk_xccm_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_streamk_xccm_char.ambr
index 7033d7ccdaec..19b66fd3bc0c 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_streamk_xccm_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_streamk_xccm_char.ambr
@@ -2,7 +2,7 @@
 # name: test_r3_streamk_xccm_gfx942_golden
   list([
     dict({
-      'basename': 'Cijk_Ailk_Bjlk_S_B_UserArgs_MT128x128x16_MI16x16Q4G6v0XQqtOnBCAqj33I-M22jIOY2tBQTLmh_bnsnKk=',
+      'basename': 'Cijk_Ailk_Bjlk_S_B_UserArgs_MT128x128x16_MI16x1655JwJfajhdjLcqlG8O5yyKtv7Vo7U3waWsFP3JMjqDI=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_subtile_lr_fp8_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_subtile_lr_fp8_char.ambr
index 477166f9ddbd..80226755a310 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_subtile_lr_fp8_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r3_subtile_lr_fp8_char.ambr
@@ -2,11 +2,11 @@
 # name: test_r3_subtile_lr_fp8_gfx950_golden
   list([
     dict({
-      'basename': 'Cijk_Alik_Bljk_F8SS_BH_UserArgs_MT128x128x256_MIndzHBWrUpa1mrfxnuueCTD8RSkQZxEnufkSndRoAdk8=',
+      'basename': 'Cijk_Alik_Bljk_F8SS_BH_UserArgs_MT128x128x256_MISJhfkGOM_LzznPGqHAFxwR60zHaJ_0QtubMCaumHZ8c=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_F8SS_BH_UserArgs_MT64x64x256_MI16FNWdVmc-73giYtPdH0fEFcwh6DjLB-UimxsZRpWHxfw=',
+      'basename': 'Cijk_Alik_Bljk_F8SS_BH_UserArgs_MT64x64x256_MI16VQE8vablR9BeR8bmyMsBQHlC6NGzvTg3bMpjT5hkvvY=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r4_asmaddr2_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r4_asmaddr2_char.ambr
index 0310a75063c2..02a68abb1d94 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r4_asmaddr2_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r4_asmaddr2_char.ambr
@@ -2,7 +2,7 @@
 # name: test_r4_asmaddr2_bf16_srvw_golden
   list([
     dict({
-      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_UserArgs_MT128x128x169Vh33ssk2RJEhKf6yNJ6mR6BOmNfDkSAUWs4GoU42Rw=',
+      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_UserArgs_MT128x128x16eGS1iaEQJ4szAijJQCYuiFxFTu0jdVE69jXnU8Bt5vI=',
       'err': 0,
     }),
   ])
@@ -10,7 +10,7 @@
 # name: test_r4_asmaddr2_fp32_golden
   list([
     dict({
-      'basename': 'Cijk_Ailk_Bjlk_S_B_UserArgs_MT64x256x16_MI16x16xjuZ9vKhEVpDsEC0tTerq_meJoK9mJLBTGOrzJvZ1W34=',
+      'basename': 'Cijk_Ailk_Bjlk_S_B_UserArgs_MT64x256x16_MI16x16xkuMUYv2yigRdN5gHIIWq4sexgk1WYvk1bT2yjVIjsKU=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r4_fp8mx_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r4_fp8mx_char.ambr
index 6e6eedc2149b..aabe454730fb 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r4_fp8mx_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r4_fp8mx_char.ambr
@@ -2,19 +2,19 @@
 # name: test_r4_fp8mx_gfx950_golden
   list([
     dict({
-      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArg5fK5RlSID_JXUB0WJxoYxBriKchsUACsajF5qCorDA0=',
+      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgBKesUaERJZl2UlrK_jE4_-aYzAWh4GcZRmqB8bCbKSM=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArg7ZH5DE3bVklqyGgK2QLWz9t5UupkabgTwEytC6U-WcQ=',
+      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgDC4-hn4Yw1xJFpeDLR49CUEEbgSRuR_HT8tjMNTgR0U=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgGCPDHiRtHcUrtT1JepolcGEuOGpATC_EPGyOtTJBJc8=',
+      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgEy6hm1-XqozC380jVSU_1K38VrWHZ1EqVmxjLzs3gFU=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgKr746cebXktPduqg9CdaWHK3SvbWQgLQ6q12tM1HAoA=',
+      'basename': 'Cijk_Alik_Bljk_F8SS_MXAE8B32_MXBE8B32_BH_UserArgsbNVNE2e7gtd4XtkphOKJJreOWvWUwB3_d1LhdVbv-E=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r4_shiftvector2_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r4_shiftvector2_char.ambr
index 1cb1787352ab..87c31b48533e 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r4_shiftvector2_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_r4_shiftvector2_char.ambr
@@ -2,15 +2,15 @@
 # name: test_r4_shiftvector2_golden
   list([
     dict({
-      'basename': 'Cijk_Ailk_Bjlk_S_B_UserArgs_MT128x64x16_SN_LDSB0d7xBIkLX0by-V53pr0Jiycp7CDqzRspnvNlabFO5C8U=',
+      'basename': 'Cijk_Ailk_Bjlk_S_B_UserArgs_MT128x64x16_SN_LDSB0cpTLakS5pBsOjvIqONKoxEnFWj2OStRESZi_W5GNUfg=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bjlk_S_B_UserArgs_MT64x128x16_SN_LDSB07aNnyJnxYoKp-C8NRKI92yaYFRqHP86lOc8cR6Ee5lA=',
+      'basename': 'Cijk_Ailk_Bjlk_S_B_UserArgs_MT64x128x16_SN_LDSB07NEbxnIKWhJFBTer4G8mAo7M9KXd8yyUWSgZ8sxSiUQ=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bjlk_S_B_UserArgs_MT64x64x16_SN_LDSB0_7uIfWIee_-LORKiqC-JJiQSvHpXF0UO1njh7h_2GFws=',
+      'basename': 'Cijk_Ailk_Bjlk_S_B_UserArgs_MT64x64x16_SN_LDSB0_kpznmTOlwvXU9I81jXTwAm4PIxF-Mi5iN2LWzSoJvfo=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx90a_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx90a_char.ambr
index dd9341e189ba..3f59c2ac1560 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx90a_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx90a_char.ambr
@@ -2,11 +2,11 @@
 # name: test_gfx90a_dominant_golden
   list([
     dict({
-      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_UserArgs_MT128x3RI_1NMp92IvBLvchYFl1Jfrp8mjF-qu7u3Iy7Ry5_QM=',
+      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_UserArgs_MT128x37-tIwpmdZHAGw0AyBwa_4lSbssae5-6aDXuzUjFNvQw=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_UserArgs_MT64x16LwdbZud9SuTLQVaDG5xPNktPoWfn4pdzE-sn0UXbdnE=',
+      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_UserArgs_MT64x1675NDkm6_odNVaIyOaPHLKxi2jCf2tcAWWUf9L3RFlcY=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx90a_db_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx90a_db_char.ambr
index ca9dc3359a3d..81d2deb3de79 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx90a_db_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx90a_db_char.ambr
@@ -2,7 +2,7 @@
 # name: test_gfx90a_db_golden
   list([
     dict({
-      'basename': 'Cijk_Ailk_Bjlk_D_B_UserArgs_MT4x16x16_MI4x4x4_SNmZ4sQjK4k_nmpHn_CcIfhd5nPd1OjDShAyAkErLoLvY=',
+      'basename': 'Cijk_Ailk_Bjlk_D_B_UserArgs_MT4x16x16_MI4x4x4_SNJQ-k7QK3Ia0NAZ7Q3M2Q-GF69hquEHAkmSLDYYWUe70=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx90a_hhs_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx90a_hhs_char.ambr
index 39d3282b738f..89522778412e 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx90a_hhs_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx90a_hhs_char.ambr
@@ -2,11 +2,11 @@
 # name: test_gfx90a_hhs_golden
   list([
     dict({
-      'basename': 'Cijk_Alik_Bljk_HHS_BH_Bias_A_S_UserArgs_MT64x16x2Is2ktKjOrVjH0RGMLAjmZgas5pGBjZWs-tdEVQm5Lc=',
+      'basename': 'Cijk_Alik_Bljk_HHS_BH_Bias_A_S_UserArgs_MT64x16xOFXXVKPIih86zgkczLJRBouzGer_dP-rzqR86gsmuZY=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_HHS_BH_Bias_A_S_UserArgs_MT64x16xtYotte_HoaUlOCcQg3EqOL47xArl9YGzuC7_6ydTzuU=',
+      'basename': 'Cijk_Alik_Bljk_HHS_BH_Bias_A_S_UserArgs_MT64x16xmNyA8myNxwNjZoiy9CjyxlkCiEXByi4SU9NwZRSs6D0=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx942_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx942_char.ambr
index 1c8407890508..8565028fbaf2 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx942_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx942_char.ambr
@@ -2,19 +2,19 @@
 # name: test_gfx942_dominant_golden
   list([
     dict({
-      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_A_S_UserArgs_MT128x325SRMO2il20gdWixbCawqC7P7tnEq4ZNHifaNJ6NWMBc=',
+      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_A_S_UserArgs_MT128x32JACyUcuoh7sO0G15BhuQ7QxWtUeQyT2mRkqgfb25unw=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_A_S_UserArgs_MT128x32Aq1Om9swigCGGiJAHW60efd6GTPeJOucetGIbI2_LIw=',
+      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_A_S_UserArgs_MT128x32YV6GjpOOowqr42UTSBdmG2lpwky91eJFdVcC8h7Xzd4=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_A_S_UserArgs_MT64x16xYaayYOrcrM4BuYyIzkeD75leFpuQUNGwafuO_5kIixc=',
+      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_A_S_UserArgs_MT64x16x2MLCeKH-GsFV_p4XATG8MNz8Dj1fok2bc84PKj53Lo8=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_A_S_UserArgs_MT64x16xv21MMjZ0uaPycUD4xS5fWzxW5mFOoq_EWwFIAW6woL4=',
+      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_A_S_UserArgs_MT64x16xS0STepP-bOL5dlH3cxvgGmDg84QioFXlwfulJSNM-z0=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx942_db_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx942_db_char.ambr
index 12c12fbd8d73..6946dec4fea9 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx942_db_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx942_db_char.ambr
@@ -2,7 +2,7 @@
 # name: test_gfx942_db_golden
   list([
     dict({
-      'basename': 'Cijk_Ailk_Bljk_D_B_UserArgs_MT16x16x8_MI16x16x1_YzGWc21SpxPo2ZEG5Ipv3rGTez62l8haI-v8lj3j1s8=',
+      'basename': 'Cijk_Ailk_Bljk_D_B_UserArgs_MT16x16x8_MI16x16x1_91h-AU75nF_9TmVXcwdcv40iUkX6ioBwBr6EeADVV4E=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx942_f8n_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx942_f8n_char.ambr
index 2eff85658a8b..95d214864605 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx942_f8n_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx942_f8n_char.ambr
@@ -2,35 +2,35 @@
 # name: test_gfx942_f8n_golden
   list([
     dict({
-      'basename': 'Cijk_Ailk_Bljk_F8NHS_BH_Bias_S_HA_S_SAB_SCD_SAV_729DuPqd3fQX45vbGAJtLPzHnT89KeMIbWwA14-iGEc=',
+      'basename': 'Cijk_Ailk_Bljk_F8NHS_BH_Bias_S_HA_S_SAB_SCD_SAV_6AafIojyN1RWBly_ZS1zlt297tFr1dJ8Ak3NQdk8sFU=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bljk_F8NHS_BH_Bias_S_HA_S_SAB_SCD_SAV_DKqKX09o8w273UkUJ4ErrfZRNcdy8hHftqkBMzHyZ7c=',
+      'basename': 'Cijk_Ailk_Bljk_F8NHS_BH_Bias_S_HA_S_SAB_SCD_SAV_EPq325jQgkjqx4a3ok1cDPZAAhJwNIxL0LJU9xRMY9s=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bljk_F8NHS_BH_Bias_S_HA_S_SAB_SCD_SAV_DMO3UlKd4l5MANtrOTjO97uKlAUrgIOE6TfetJvgSXE=',
+      'basename': 'Cijk_Ailk_Bljk_F8NHS_BH_Bias_S_HA_S_SAB_SCD_SAV_T7OUaMERQ99vrK5uciRa5VG3IXw5sdf_yYnUNSd2F-8=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bljk_F8NHS_BH_Bias_S_HA_S_SAB_SCD_SAV_EPsdfffWPkkbNo7oJrqHLvGOCejHyizMMGf1whufybQ=',
+      'basename': 'Cijk_Ailk_Bljk_F8NHS_BH_Bias_S_HA_S_SAB_SCD_SAV_VJDdd1ypH_zN9Qq-Jgk8_yNu5tPWU-uXzX_uuF5s4c0=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bljk_F8NHS_BH_Bias_S_HA_S_SAB_SCD_SAV_Hcsox9Ly57YNCddTdXbWslXPkLOZBQRjGErFrO3yxqI=',
+      'basename': 'Cijk_Ailk_Bljk_F8NHS_BH_Bias_S_HA_S_SAB_SCD_SAV_jRH0IDqthedQC6xSNT2M9j8HdyXkX-INXd2K-C6KVtk=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bljk_F8NHS_BH_Bias_S_HA_S_SAB_SCD_SAV_cj1EXNJfR_GVUCg4CbJBD3KaUCcco0V84kuZwigOXkE=',
+      'basename': 'Cijk_Ailk_Bljk_F8NHS_BH_Bias_S_HA_S_SAB_SCD_SAV_nFFRkxleXObhLzS6uESFkCgPGGLyI35n0jbz4pm13Zo=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bljk_F8NHS_BH_Bias_S_HA_S_SAB_SCD_SAV_mNy-dif8fOJMyT8UyBv0MZ9ffb_VxaUHeJstA5TOY2w=',
+      'basename': 'Cijk_Ailk_Bljk_F8NHS_BH_Bias_S_HA_S_SAB_SCD_SAV_rkLSH1FyGS3FyWC5GRNYcYf2hNZErEsjpJC88MhMQ1Y=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bljk_F8NHS_BH_Bias_S_HA_S_SAB_SCD_SAV_sCaojU4a7MnaGAQDShncPaG1hk-xCTJRrs1nq-pon9Y=',
+      'basename': 'Cijk_Ailk_Bljk_F8NHS_BH_Bias_S_HA_S_SAB_SCD_SAV_trK8Ljxks7OYw1DHyfYRPZs2bQYLXWE-Zqkjuvh1hCw=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx942_gg_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx942_gg_char.ambr
index 71ff7df90b22..b8fc0ec3044c 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx942_gg_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx942_gg_char.ambr
@@ -2,19 +2,19 @@
 # name: test_gfx942_gg_golden
   list([
     dict({
-      'basename': 'Cijk_Ailk_Bljk_HHS_BH_Bias_HA_S_UserArgs_MT128x15XQ3ldzb0Gj-JW5fpgaDvnFoXZXwC2jJ7ts_WYAaBbQ=',
+      'basename': 'Cijk_Ailk_Bljk_HHS_BH_Bias_HA_S_UserArgs_MT128x1BC6OjTae_lC3rScmocFhMirnREwfuS1GJy4hQUgXSBI=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bljk_HHS_BH_Bias_HA_S_UserArgs_MT128x1V4sU5hUzI65ueBxJPbD47cU_-4EXy4NBLjAW7H7Otbo=',
+      'basename': 'Cijk_Ailk_Bljk_HHS_BH_Bias_HA_S_UserArgs_MT128x1CXp_S3IN4snsMOMnvr_loNsuVpSxv3GBaPc0qjBvQ78=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bljk_HHS_BH_Bias_HA_S_UserArgs_MT128x1_7ieeaTpQBNiPeMi-Mmr6aDHQ9_ZOuuDJHHV4vIc9p4=',
+      'basename': 'Cijk_Ailk_Bljk_HHS_BH_Bias_HA_S_UserArgs_MT128x1EKV7ia6IDkmUVEVp9chtGnKy88z9lKljeKIiRdhQcEg=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Ailk_Bljk_HHS_BH_Bias_HA_S_UserArgs_MT128x1i-roPN0Pi4r3SueZWPAt8qCeHxB1PrsHFMbUv4NVgeg=',
+      'basename': 'Cijk_Ailk_Bljk_HHS_BH_Bias_HA_S_UserArgs_MT128x1gr7hsWrI_VbbidC-wW-BMZj4XlOeqfZOKciqRYcgf44=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx942_grad_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx942_grad_char.ambr
index b22b76517415..c0d6576397cf 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx942_grad_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx942_grad_char.ambr
@@ -2,7 +2,7 @@
 # name: test_gfx942_grad_golden
   list([
     dict({
-      'basename': 'Cijk_Ailk_Bljk_HHS_BH_Bias_D_GradH_A_S_UserArgs_bvPVyGWSUXAhqGpJcSGBO6JsinYZfW1vqlVdVXfvcQg=',
+      'basename': 'Cijk_Ailk_Bljk_HHS_BH_Bias_D_GradH_A_S_UserArgs_adSntHj-AzzReDb5lrZrIaar1W37eOiohg3tkwrqQaY=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx942_hss_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx942_hss_char.ambr
index e9b4bd9a83d8..2174dea6ac5c 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx942_hss_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx942_hss_char.ambr
@@ -2,7 +2,7 @@
 # name: test_gfx942_hss_golden
   list([
     dict({
-      'basename': 'Cijk_Alik_Bljk_HSS_BH_Bias_S_UserArgs_MT64x16x32hciFhDwVu3m43HLDu0ClEUX4IPgtLGxL39OuGxCb3Bw=',
+      'basename': 'Cijk_Alik_Bljk_HSS_BH_Bias_S_UserArgs_MT64x16x32E2Tx7-uPnt75Jaaf4M9tkY0tSJLlJiBrqsrm9GuvfY0=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx950_bbs_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx950_bbs_char.ambr
index 0f2c03a7db9a..bfe5328dc514 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx950_bbs_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx950_bbs_char.ambr
@@ -2,11 +2,11 @@
 # name: test_gfx950_bbs_golden
   list([
     dict({
-      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_UserArgs_MT64x16J1hjuComZBn8ynD3nMST543ZRWQnrP_CrC6IMMrzFgI=',
+      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_UserArgs_MT64x16ASZp0_s6XqvyWB8q007SdZuQZucGQQxMxSN8SQEYzYQ=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_UserArgs_MT64x16YbxrsWR0Ay1gm4YdA--tYN3aj5AK4N1xZpWdnrwMB_A=',
+      'basename': 'Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_UserArgs_MT64x16U_Da3xMKbb9sYFyv53pLDkg1iEPjqLxZnpSOVdcf9L8=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx950_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx950_char.ambr
index 8a3cf6b291eb..a08c1f4f656f 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx950_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx950_char.ambr
@@ -2,11 +2,11 @@
 # name: test_gfx950_dominant_golden
   list([
     dict({
-      'basename': 'Cijk_Alik_Bljk_F8HS_BH_UserArgs_MT64x64x128_MI16uzY40g8cX-K5IzxejKKFdTHUZHafhvKfvDnG5po25gs=',
+      'basename': 'Cijk_Alik_Bljk_F8HS_BH_UserArgs_MT64x64x128_MI16jEThGENwCZmn5D6ev1iiWTGTQmCuXKkGHLSuBvp96cM=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_F8HS_BH_UserArgs_MT64x64x128_MI32oTCAel08qrNV6VuUfKehTfxoGKTVp1mGhWKtp99WfJs=',
+      'basename': 'Cijk_Alik_Bljk_F8HS_BH_UserArgs_MT64x64x128_MI32sPWxWOfmdwFZBD7FFVUyGsKdTkuar_j5UwxNKKGSa-I=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx950_dtl_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx950_dtl_char.ambr
index af5d1f4c1c7b..7d2048991ce5 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx950_dtl_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx950_dtl_char.ambr
@@ -2,19 +2,19 @@
 # name: test_gfx950_dtl_golden
   list([
     dict({
-      'basename': 'Cijk_Alik_Bljk_S_B_UserArgs_MT16x16x32_MI16x16x1Encfa4zmXKJwzQl-I7wbG4q6Lt8UwWv8a4nbk80MEp4=',
+      'basename': 'Cijk_Alik_Bljk_S_B_UserArgs_MT16x16x32_MI16x16x186fGnMiOJOZM8q--0jWbysszjiVbNdIV2L2Y59EPK9s=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_S_B_UserArgs_MT16x16x32_MI16x16x1om-GzEoROYXMrJI9TVtTT-hp34pZ191zA_LxjgTyAiI=',
+      'basename': 'Cijk_Alik_Bljk_S_B_UserArgs_MT16x16x32_MI16x16x1MV-KOM1e9xO6AR3BJvu7alz3okgOhFbkNQfqNrgFPB0=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_S_B_UserArgs_MT16x16x32_MI16x16x1yUVl3pG_BN6ydUuVhUIAwi_Ngr0Lpngzr6TV7B-l9i4=',
+      'basename': 'Cijk_Alik_Bljk_S_B_UserArgs_MT16x16x32_MI16x16x1hSpuCxYU4YvlrH3JfUjGFTwAu0MxBiO-W5P10nWZlME=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_S_B_UserArgs_MT16x16x32_MI16x16x1zsrR2ngHKYLdYmsOB_LJvY7Mq62qKsMYBcX1IiBugf8=',
+      'basename': 'Cijk_Alik_Bljk_S_B_UserArgs_MT16x16x32_MI16x16x1vUj5siBerz1V-n-y2U11BtvjZcVP1qRie7sEn-yFSa0=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx950_hhs_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx950_hhs_char.ambr
index 2caffef2dcc6..5a88f6c35f79 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx950_hhs_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx950_hhs_char.ambr
@@ -2,11 +2,11 @@
 # name: test_gfx950_hhs_golden
   list([
     dict({
-      'basename': 'Cijk_Alik_Bljk_HHS_BH_Bias_A_S_UserArgs_MT64x16x7BctTdfI8G4uFQcCxzhrjixoYFmomu5ucHxfsPSK7nk=',
+      'basename': 'Cijk_Alik_Bljk_HHS_BH_Bias_A_S_UserArgs_MT64x16xfvmkG0Ic0vzeZ9PO4yL33kvH5qy7uX9KFi-6vYp6Wlg=',
       'err': 0,
     }),
     dict({
-      'basename': 'Cijk_Alik_Bljk_HHS_BH_Bias_A_S_UserArgs_MT64x16xGH3sI80LSP-Cpep4TPFdVpOMCXF4ibeBDv9BesP4SoE=',
+      'basename': 'Cijk_Alik_Bljk_HHS_BH_Bias_A_S_UserArgs_MT64x16xnPsGTDHruXmPToVpazJ7uBK5iwgHInnZ1XT8-qUJH7w=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx950_hss_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx950_hss_char.ambr
index b3809749b56d..e72ce0e63f86 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx950_hss_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx950_hss_char.ambr
@@ -2,7 +2,7 @@
 # name: test_gfx950_hss_golden
   list([
     dict({
-      'basename': 'Cijk_Ailk_Bljk_SH_HHS_BH_Bias_UserArgs_MT64x64x3Lt5IQfaC412_wChmBpw4I3DVdFOVwVYYJl6UhZBfzMI=',
+      'basename': 'Cijk_Ailk_Bljk_SH_HHS_BH_Bias_UserArgs_MT64x64x3N98T9YzMrAryUGTjE3rqFc5JxdR8Vrkqq11myu7GDTY=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx950_i8gsu_char.ambr b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx950_i8gsu_char.ambr
index 0743a37bcdb6..726f4d78dc45 100644
--- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx950_i8gsu_char.ambr
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/__snapshots__/test_seed_gfx950_i8gsu_char.ambr
@@ -2,7 +2,7 @@
 # name: test_gfx950_i8gsu_golden
   list([
     dict({
-      'basename': 'Cijk_Ailk_Bljk_I8I8S_BH_HA_S_UserArgs_MT16x16x64VxwuKLGtgmdNEtFSkXCEAwB3jkItdaKyzyjLVM8Bhkc=',
+      'basename': 'Cijk_Ailk_Bljk_I8I8S_BH_HA_S_UserArgs_MT16x16x64Qc4RQc8rRG1Dk7an9I9c-t8d9WZRjbiS1lokFBjP8Co=',
       'err': 0,
     }),
   ])
diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_streamk_work_stealing.py b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_streamk_work_stealing.py
new file mode 100644
index 000000000000..445ed1fde4ec
--- /dev/null
+++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_streamk_work_stealing.py
@@ -0,0 +1,527 @@
+# Copyright © Advanced Micro Devices, Inc., or its affiliates.
+# SPDX-License-Identifier: MIT
+"""Unit tests for the single-hop next-neighbor StreamK work-stealing codegen.
+
+These tests assert that the work-stealing assembly is emitted by the helper
+methods on the ``StreamK`` base class, and -- crucially -- that those helpers
+are only ever reached behind the codegen-time ``StreamKWorkStealing`` toggle.
+They import rocisa instructions and inspect emitted modules rather than matching
+source text; the toggle gating and the Solution-level validation are verified by
+executing the *real* source (via the AST) so the assertions track the actual code.
+
+Emission contract (single-hop next-neighbor + sticky-home + static auto-reset):
+  * The steal fires on an empty home fetch with no neighbor structural-extra guard.
+  * The steal & home atomic bounds are the predecessor-inclusive self-reset value.
+  * A per-WG sticky-empty SGPR gates the home fetch.
+  * kernelEnd emits no explicit reset; per-queue counters auto-reset.
+"""
+
+import ast
+import inspect
+import textwrap
+import types
+
+import pytest
+
+# Prime the component registry before StreamK imports (avoids circular import).
+from Tensile.KernelWriterAssembly import KernelWriterAssembly  # noqa: F401
+
+from rocisa.code import Module, Label
+from rocisa.instruction import (
+    SAddU32,
+    SAndB32,
+    SAtomicInc,
+    SBarrier,
+    SCBranchSCC1,
+    SCmpGeU32,
+    SCmpLtU32,
+    SMovB32,
+    SSubU32,
+)
+
+from Tensile.Common.ValidParameters import validParameters
+from Tensile.Components.StreamK import (
+    StreamK,
+    StreamKDynamic,
+    StreamKHybrid,
+    streamKVariantClass,
+)
+from Tensile.SolutionStructs import Solution
+from Tensile.SolutionStructs.Utilities import reject
+
+
+# ---------------------------------------------------------------------------
+# Fakes: just enough of a "writer" for the standalone helper methods.
+#
+# The helpers only touch ``writer.sgprPool`` (checkOut / checkOutAligned /
+# checkIn) and emit rocisa instructions via free functions (sgpr/vgpr), so a
+# tiny pool that hands out monotonically increasing register indices is all
+# that is required -- no KernelWriter, no GPU.
+# ---------------------------------------------------------------------------
+class _FakeSgprPool:
+    def __init__(self, start: int = 100):
+        self._next = start
+
+    def checkOut(self, n: int, name: str = "", *args, **kwargs) -> int:
+        reg = self._next
+        self._next += n
+        return reg
+
+    def checkOutAligned(self, n: int, align: int, name: str = "", *args, **kwargs) -> int:
+        if self._next % align:
+            self._next += align - (self._next % align)
+        reg = self._next
+        self._next += n
+        return reg
+
+    def checkIn(self, *args, **kwargs):
+        return None
+
+
+class _FakeWriter:
+    def __init__(self, numXCD: int = 8, cacheLineBytes: int = 128):
+        self.sgprPool = _FakeSgprPool()
+        # gfx942/gfx950 mirror origami get_default_num_xcds == 8 and origami
+        # get_default_cache_line_bytes == 128; the helpers read the per-arch
+        # queue count from writer.states.archCaps["NumXCD"] and the per-queue
+        # counter stride from writer.states.archCaps["CacheLineBytes"].
+        self.states = types.SimpleNamespace(
+            archCaps={"NumXCD": numXCD, "CacheLineBytes": cacheLineBytes})
+
+
+def _mk_label(base: str) -> Label:
+    return Label(base, "")
+
+
+# gfx942/gfx950 both map to 8 queues in the per-arch lookup; the helpers key
+# their fast-mask constants off ``kernel["ISA"]``.
+_WS_KERNEL = {"ISA": (9, 4, 0)}
+
+
+def _stream_k_instance(streamk: int) -> StreamK:
+    """A concrete StreamK variant (helpers live on the base class)."""
+    return streamKVariantClass(streamk)()
+
+
+def _imm_in(inst, value: int) -> bool:
+    """True if *inst* carries *value* as an immediate operand.
+
+    rocisa renders immediates inconsistently -- ints passed straight through
+    print as decimal ("7"), while values passed as ``hex(...)`` print as
+    "0x..." -- so normalise every param through ``int(p, 0)`` and compare.
+    """
+    for p in inst.getParams():
+        try:
+            if int(str(p), 0) == value:
+                return True
+        except (TypeError, ValueError):
+            continue
+    return False
+
+
+def _flat(module: Module) -> list:
+    return list(module.flatitems())
+
+
+# ---------------------------------------------------------------------------
+# AST helpers: read the *real* StreamK / Solution source and reason about it.
+# ---------------------------------------------------------------------------
+def _const_slice(subscript: ast.Subscript):
+    s = subscript.slice
+    if isinstance(s, ast.Constant):
+        return s.value
+    return None
+
+
+def _is_subscript_on(node, name: str, key: str) -> bool:
+    return (
+        isinstance(node, ast.Subscript)
+        and isinstance(node.value, ast.Name)
+        and node.value.id == name
+        and _const_slice(node) == key
+    )
+
+
+def _ws_guarded_calls(func) -> set:
+    """Names of ``self.streamKWorkStealing*`` calls that sit inside an
+    ``if kernel["StreamKWorkStealing"]:`` block in *func* (recursing into
+    nested closures)."""
+    tree = ast.parse(textwrap.dedent(inspect.getsource(func)))
+    guarded = set()
+    for node in ast.walk(tree):
+        if isinstance(node, ast.If) and _is_subscript_on(
+            node.test, "kernel", "StreamKWorkStealing"
+        ):
+            for sub in ast.walk(node):
+                if isinstance(sub, ast.Call) and isinstance(sub.func, ast.Attribute):
+                    if sub.func.attr.startswith("streamKWorkStealing"):
+                        guarded.add(sub.func.attr)
+    return guarded
+
+
+def _all_ws_calls(func) -> set:
+    """Every ``self.streamKWorkStealing*`` call in *func*, guarded or not."""
+    tree = ast.parse(textwrap.dedent(inspect.getsource(func)))
+    calls = set()
+    for node in ast.walk(tree):
+        if isinstance(node, ast.Call) and isinstance(node.func, ast.Attribute):
+            if node.func.attr.startswith("streamKWorkStealing"):
+                calls.add(node.func.attr)
+    return calls
+
+
+def _source_of(func) -> str:
+    return inspect.getsource(func)
+
+
+def _extract_ws_validation():
+    """Compile the *real* ``if state["StreamKWorkStealing"]:`` block out of
+    ``Solution.assignDerivedParameters`` into a standalone callable so the
+    actual rejection logic can be exercised without a full Solution state.
+
+    The block reads ``isaInfoMap[isa].asmCaps`` (for the HasSAtomic / gfx1250
+    gate), so ``isa`` and ``isaInfoMap`` are threaded in as parameters -- the
+    real function derives ``isa = tuple(state["ISA"])`` and receives
+    ``isaInfoMap`` as an argument."""
+    tree = ast.parse(
+        textwrap.dedent(inspect.getsource(Solution.assignDerivedParameters))
+    )
+    target = None
+    for node in ast.walk(tree):
+        if isinstance(node, ast.If) and _is_subscript_on(
+            node.test, "state", "StreamKWorkStealing"
+        ):
+            target = node
+            break
+    assert target is not None, "could not find StreamKWorkStealing validation block"
+
+    func = ast.FunctionDef(
+        name="_validate",
+        args=ast.arguments(
+            posonlyargs=[],
+            args=[
+                ast.arg("state"),
+                ast.arg("printRejectionReason"),
+                ast.arg("reject"),
+                ast.arg("isa"),
+                ast.arg("isaInfoMap"),
+            ],
+            vararg=None,
+            kwonlyargs=[],
+            kw_defaults=[],
+            kwarg=None,
+            defaults=[],
+        ),
+        body=[target],
+        decorator_list=[],
+        returns=None,
+        type_params=[],
+    )
+    mod = ast.Module(body=[func], type_ignores=[])
+    ast.fix_missing_locations(mod)
+    ns: dict = {}
+    exec(compile(mod, "<ws-validation>", "exec"), ns)
+    return ns["_validate"]
+
+
+# gfx1250 (12,5,0) has no scalar atomics; gfx942/gfx950 do. Mirror just the
+# HasSAtomic capability the work-stealing validation reads.
+_NO_SATOMIC_ISAS = {(12, 5, 0)}
+
+
+def _fake_isa_info_map(isa):
+    isa = tuple(isa)
+    asm_caps = {"HasSAtomic": isa not in _NO_SATOMIC_ISAS}
+    return {isa: types.SimpleNamespace(asmCaps=asm_caps)}
+
+
+# ===========================================================================
+# 1. ValidParameters: the codegen-time toggle exists and is boolean.
+# ===========================================================================
+class TestValidParameters:
+    def test_work_stealing_param_exists(self):
+        assert "StreamKWorkStealing" in validParameters
+
+    def test_work_stealing_param_is_zero_one(self):
+        assert validParameters["StreamKWorkStealing"] == [0, 1]
+
+
+# ===========================================================================
+# 2. The work-stealing helper methods exist on the StreamK base class.
+# ===========================================================================
+class TestHelperMethodsExist:
+    @pytest.mark.parametrize(
+        "name",
+        [
+            "streamKWorkStealingHomeBound",
+            "streamKWorkStealingSteal",
+        ],
+    )
+    def test_method_is_defined_on_base(self, name):
+        assert callable(getattr(StreamK, name))
+
+
+# ===========================================================================
+# 3a. Home auto-reset bound: the predecessor's workgroup count is folded into
+#     the atomic_inc bound (NOT disabled with 0xFFFFFFFF).
+# ===========================================================================
+class TestHomeBoundEmission:
+    def _emit(self):
+        sk = _stream_k_instance(4)
+        writer = _FakeWriter()
+        module = Module("home-bound")
+        sBound = writer.sgprPool.checkOut(1, "bound")
+        sQueueIdx = writer.sgprPool.checkOut(1, "queueIdx")
+        sk.streamKWorkStealingHomeBound(
+            writer, module, _WS_KERNEL, sBound, sQueueIdx, "skGrid"
+        )
+        return _flat(module)
+
+    def test_computes_predecessor_index(self):
+        items = self._emit()
+        # p = (q - 1) & 0x7
+        assert any(isinstance(i, SSubU32) and _imm_in(i, 1) for i in items), (
+            "expected q-1 to reach the predecessor queue"
+        )
+        assert any(isinstance(i, SAndB32) and _imm_in(i, 0x7) for i in items), (
+            "expected (q-1) & 0x7 wrap for the predecessor index"
+        )
+
+    def test_adds_predecessor_workgroups_to_bound(self):
+        items = self._emit()
+        # The predecessor structural share plus the final fold-in add.
+        assert sum(isinstance(i, SAddU32) for i in items) >= 2, (
+            "expected the W_(q-1) structural add and the bound fold-in"
+        )
+
+    def test_does_not_disable_auto_reset(self):
+        items = self._emit()
+        assert not any(
+            isinstance(i, SMovB32) and _imm_in(i, 0xFFFFFFFF) for i in items
+        ), "the home bound must be a finite predecessor-inclusive value"
+
+
+# ===========================================================================
+# 3b. Next-neighbor steal: one unconditional (past home-empty) NEXT-neighbor steal with
+#     the static predecessor-inclusive auto-reset bound. No remainder/extra
+#     guards.
+# ===========================================================================
+class TestStealEmission:
+    def _emit(self):
+        sk = _stream_k_instance(4)
+        writer = _FakeWriter()
+        module = Module("steal")
+        sQueueIdx = writer.sgprPool.checkOut(1, "queueIdx")
+        sWorkItemIdx = writer.sgprPool.checkOut(1, "workItemIdx")
+        sk.streamKWorkStealingSteal(
+            writer, module, _WS_KERNEL, sQueueIdx, sWorkItemIdx, "skGrid", _mk_label
+        )
+        return _flat(module)
+
+    def test_neighbor_walk_is_plus_one_then_wrap(self):
+        items = self._emit()
+        # +1 to advance to the next neighbor ...
+        assert any(
+            isinstance(i, SAddU32) and _imm_in(i, 1) for i in items
+        ), "expected +1 advance to the next queue"
+        # ... wrapped within the 8 queues (single-hop next-neighbor) via & 0x7.
+        assert any(
+            isinstance(i, SAndB32) and _imm_in(i, 0x7) for i in items
+        ), "expected (queueIdx+1) & 0x7 wrap"
+
+    def test_no_neighbor_extra_guard(self):
+        # The next-neighbor steal always fires on an empty home fetch: the old
+        # ">= remainder / neighbor-has-no-structural-extra" guard is removed.
+        items = self._emit()
+        assert not any(isinstance(i, SCmpGeU32) for i in items), (
+            "the next-neighbor steal must NOT guard on a neighbor structural extra"
+        )
+
+    def test_exactly_one_atomic_increment(self):
+        items = self._emit()
+        atomics = [i for i in items if isinstance(i, SAtomicInc)]
+        assert len(atomics) == 1, "the steal must emit exactly one atomic"
+
+    def test_atomic_bound_is_predecessor_inclusive_not_all_ones(self):
+        items = self._emit()
+        # The static bound tiles_s + W_s + W_q - 1 ends in a -1 (SSubU32 imm 1)
+        # and must never load the 0xFFFFFFFF disable-auto-reset sentinel.
+        assert not any(
+            isinstance(i, SMovB32) and _imm_in(i, 0xFFFFFFFF) for i in items
+        ), "the stolen atomic must use a finite self-reset bound, not 0xFFFFFFFF"
+        assert any(
+            isinstance(i, SSubU32) and _imm_in(i, 1) for i in items
+        ), "expected the '- 1' that forms the atomic_inc auto-reset bound"
+
+    def test_guards_on_valid_home_fetch(self):
+        # A valid home fetch (index < TotalItems) must short-circuit the steal.
+        items = self._emit()
+        assert any(isinstance(i, SCmpLtU32) for i in items)
+        assert any(isinstance(i, SCBranchSCC1) for i in items)
+
+    def test_no_reset_barrier(self):
+        items = self._emit()
+        assert not any(isinstance(i, SBarrier) for i in items), (
+            "the steal helper must not emit a reset barrier"
+        )
+
+
+# ===========================================================================
+# 3c. Sticky-home: the home fetch is gated by a persistent StreamKStickyEmpty
+#     SGPR, and the flag is latched on an empty home fetch. Verified against
+#     the real graWorkGroup source.
+# ===========================================================================
+class TestStickyHomeGate:
+    @pytest.mark.parametrize(
+        "func", [StreamKDynamic.graWorkGroup, StreamKHybrid.graWorkGroup]
+    )
+    def test_home_fetch_is_gated_by_sticky_flag(self, func):
+        src = _source_of(func)
+        assert "StreamKStickyEmpty" in src, (
+            "the home fetch must be gated by the persistent sticky-empty SGPR"
+        )
+        # A steal-only skip label proves the home s_atomic_inc is bypassed once
+        # the WG has gone sticky.
+        assert "SK_StealOnly" in src, (
+            "sticky WGs must skip the home fetch and steal only"
+        )
+
+    @pytest.mark.parametrize(
+        "func", [StreamKDynamic.graWorkGroup, StreamKHybrid.graWorkGroup]
+    )
+    def test_steal_passes_grid_sgpr(self, func):
+        # The steal now needs the mode-appropriate grid SGPR for its bound.
+        src = _source_of(func)
+        grid = "SKGrid" if func is StreamKHybrid.graWorkGroup else "skGrid"
+        assert "streamKWorkStealingSteal" in src
+        assert grid in src
+
+
+# ===========================================================================
+# 3d. kernelEnd emits no work-stealing reset and no completion counter.
+# ===========================================================================
+class TestNoExplicitReset:
+    @pytest.mark.parametrize("func", [StreamKDynamic.kernelEnd, StreamKHybrid.kernelEnd])
+    def test_kernelend_has_no_ws_calls(self, func):
+        assert _all_ws_calls(func) == set(), (
+            "kernelEnd must not emit any work-stealing reset"
+        )
+
+    @pytest.mark.parametrize("func", [StreamKDynamic.kernelEnd, StreamKHybrid.kernelEnd])
+    def test_kernelend_has_no_completion_counter(self, func):
+        src = _source_of(func)
+        assert "0x80" not in src
+        assert "completion" not in src.lower() or "no explicit" in src.lower()
+
+
+# ===========================================================================
+# 5. Per-architecture queue-count lookup (C1): fast masking requires a
+#    power-of-two queue count; the count is read from the per-arch capability
+#    writer.states.archCaps["NumXCD"] (the codegen mirror of origami
+#    get_default_num_xcds).
+# ===========================================================================
+class TestQueueConstants:
+    @pytest.mark.parametrize("isa", [(9, 4, 0), (9, 5, 0)])
+    def test_supported_arches_use_eight_power_of_two_queues(self, isa):
+        sk = _stream_k_instance(4)
+        # gfx942/gfx950 mirror origami get_default_num_xcds == 8 and
+        # get_default_cache_line_bytes == 128, so the constants tuple is exactly
+        # (numQueues=8, mask=7, log2=3, cacheLineLog2=log2(128)=7).
+        writer = _FakeWriter(numXCD=8, cacheLineBytes=128)
+        assert sk._wsQueueConstants(writer, {"ISA": isa}) == (8, 7, 3, 7)
+
+    def test_non_power_of_two_queue_count_asserts(self):
+        # The shift/AND fast masking is only valid for a power-of-two queue
+        # count; a non-power-of-two NumXCD cap must trip the guard assert.
+        sk = _stream_k_instance(4)
+        writer = _FakeWriter(numXCD=6)
+        with pytest.raises(AssertionError):
+            sk._wsQueueConstants(writer, {"ISA": (9, 9, 0)})
+
+
+# ===========================================================================
+# 3e. Absence-by-toggle: the helpers are only reached behind the
+#     ``kernel["StreamKWorkStealing"]`` gate at every callsite. Verified
+#     against the real source so "off" provably emits nothing extra.
+# ===========================================================================
+class TestCallsitesAreToggleGated:
+    def test_sk4_grawg_steal_calls_are_all_gated(self):
+        guarded = _ws_guarded_calls(StreamKDynamic.graWorkGroup)
+        allcalls = _all_ws_calls(StreamKDynamic.graWorkGroup)
+        assert {"streamKWorkStealingHomeBound", "streamKWorkStealingSteal"} <= guarded
+        # Nothing slips through ungated.
+        assert allcalls == guarded
+
+    def test_sk5_grawg_steal_calls_are_all_gated(self):
+        guarded = _ws_guarded_calls(StreamKHybrid.graWorkGroup)
+        allcalls = _all_ws_calls(StreamKHybrid.graWorkGroup)
+        assert {"streamKWorkStealingHomeBound", "streamKWorkStealingSteal"} <= guarded
+        assert allcalls == guarded
+
+
+# ===========================================================================
+# 4. Solution validation: the real rejection logic from
+#    assignDerivedParameters, executed in isolation.
+# ===========================================================================
+class TestSolutionValidation:
+    def setup_method(self):
+        self.validate = _extract_ws_validation()
+
+    def _run(self, *, streamk, atomic, work_stealing=1, debug_streamk=0, isa=(9, 4, 2)):
+        isa = tuple(isa)
+        state = {
+            "StreamKWorkStealing": work_stealing,
+            "StreamK": streamk,
+            "StreamKAtomic": atomic,
+            "DebugStreamK": debug_streamk,
+            "ISA": isa,
+        }
+        self.validate(state, False, reject, isa, _fake_isa_info_map(isa))
+        return state
+
+    @pytest.mark.parametrize("streamk", [0, 1, 2, 3])
+    def test_rejected_when_streamk_not_4_or_5(self, streamk):
+        state = self._run(streamk=streamk, atomic=0)
+        assert state["Valid"] is False
+
+    @pytest.mark.parametrize("streamk", [4, 5])
+    def test_accepted_for_dynamic_and_hybrid_without_atomic(self, streamk):
+        state = self._run(streamk=streamk, atomic=0)
+        assert state.get("Valid", True) is True
+
+    @pytest.mark.parametrize("streamk", [4, 5])
+    def test_rejected_with_atomic(self, streamk):
+        state = self._run(streamk=streamk, atomic=1)
+        assert state["Valid"] is False
+
+    @pytest.mark.parametrize("debug", [1, 2, 3])
+    def test_rejected_with_debug_streamk(self, debug):
+        # DebugStreamK overrides can break the W_q>=1 precondition the per-queue
+        # auto-reset relies on, so the combination is rejected.
+        state = self._run(streamk=4, atomic=0, debug_streamk=debug)
+        assert state["Valid"] is False
+
+    def test_off_toggle_is_inert_even_for_bad_combo(self):
+        # With the toggle off the guard must not fire, even for a combination
+        # that would otherwise be rejected.
+        state = self._run(streamk=3, atomic=1, work_stealing=0, debug_streamk=3)
+        assert "Valid" not in state
+
+    @pytest.mark.parametrize("streamk", [4, 5])
+    def test_rejected_on_gfx1250_no_scalar_atomic(self, streamk):
+        # gfx1250 lacks scalar atomics (HasSAtomic=false); the steal path emits
+        # s_atomic_inc with no vector fallback, so work stealing is rejected.
+        state = self._run(streamk=streamk, atomic=0, isa=(12, 5, 0))
+        assert state["Valid"] is False
+
+    @pytest.mark.parametrize("isa", [(9, 4, 2), (9, 5, 0)])
+    @pytest.mark.parametrize("streamk", [4, 5])
+    def test_accepted_on_scalar_atomic_arches(self, isa, streamk):
+        # gfx942/gfx950 have scalar atomics, so work stealing stays accepted.
+        state = self._run(streamk=streamk, atomic=0, isa=isa)
+        assert state.get("Valid", True) is True
+
+    def test_off_toggle_is_inert_on_gfx1250(self):
+        # StreamKWorkStealing=0 on gfx1250 must not fire the reject.
+        state = self._run(streamk=4, atomic=0, work_stealing=0, isa=(12, 5, 0))
+        assert "Valid" not in state
diff --git a/projects/hipblaslt/tensilelite/include/Tensile/ContractionSolution.hpp b/projects/hipblaslt/tensilelite/include/Tensile/ContractionSolution.hpp
index 2ef52162110c..27c049c56fda 100644
--- a/projects/hipblaslt/tensilelite/include/Tensile/ContractionSolution.hpp
+++ b/projects/hipblaslt/tensilelite/include/Tensile/ContractionSolution.hpp
@@ -382,6 +382,17 @@ namespace TensileLite
         // the launch grid and the packed args can never disagree.
         bool                 streamK5EffectiveDynamic(Problem const&  problem,
                                                       Hardware const& hardware) const;
+        // Selection-time predicate for the StreamK dynamic-queue / work-stealing
+        // path. The SK4 and dynamic sub-path of SK5 kernels hardcode a
+        // power-of-two per-XCD queue count and mask indices with (Q-1); that fast
+        // masking is only valid when the device exposes a power-of-two number of
+        // XCDs. Returns false (and warns once) when this solution would take the
+        // dynamic-queue path but the hardware's NUM_XCD is not a power of two
+        // (e.g. MI300A = 6), so the solution is EXCLUDED from selection rather
+        // than silently degraded to tree reduction. All other solutions return
+        // true. Wired into softwarePredicate() (SolutionLibrary.hpp).
+        bool                 streamKDynamicQueueSupported(Problem const&  problem,
+                                                          Hardware const& hardware) const;
         size_t               partialTileSize(size_t skGrid) const;
 
         static float computeGranularity(float x);
diff --git a/projects/hipblaslt/tensilelite/include/Tensile/SolutionLibrary.hpp b/projects/hipblaslt/tensilelite/include/Tensile/SolutionLibrary.hpp
index 689785b7b612..1e94054251aa 100644
--- a/projects/hipblaslt/tensilelite/include/Tensile/SolutionLibrary.hpp
+++ b/projects/hipblaslt/tensilelite/include/Tensile/SolutionLibrary.hpp
@@ -80,7 +80,13 @@ namespace TensileLite
         switch(searchType)
         {
         case SolutionLibrarySearchType::DEFAULT:
-            return (*solutions.problemPredicate)(problem) && (*solutions.taskPredicate)(task);
+            // streamKDynamicQueueSupported() excludes StreamK dynamic-queue /
+            // work-stealing solutions (SK4 and the dynamic sub-path of SK5) on
+            // devices whose XCD count is not a power of two, warning the user
+            // once. This is reject-and-continue: selection falls through to
+            // another (SK3-static / non-StreamK) solution for the GEMM.
+            return (*solutions.problemPredicate)(problem) && (*solutions.taskPredicate)(task)
+                   && solutions.streamKDynamicQueueSupported(problem, hardware);
             break;
         case SolutionLibrarySearchType::GEMM_TYPE_ONLY:
             return isGemmTypeSame(solutions, problem);
diff --git a/projects/hipblaslt/tensilelite/rocisa/rocisa/include/hardware_caps.hpp b/projects/hipblaslt/tensilelite/rocisa/rocisa/include/hardware_caps.hpp
index a6d5f122e516..16331fba8197 100644
--- a/projects/hipblaslt/tensilelite/rocisa/rocisa/include/hardware_caps.hpp
+++ b/projects/hipblaslt/tensilelite/rocisa/rocisa/include/hardware_caps.hpp
@@ -606,6 +606,18 @@ inline std::map<std::string, int> initArchCaps(const IsaVersion& isaVersion)
     rv["LDSBankCount"] = 64;
     rv["LDSBankWidth"] = 4; // bytes per bank
 
+    // Per-XCD work-queue count baked into StreamK dynamic-queue kernels. Single
+    // codegen-side mirror of origami get_default_num_xcds(): gfx942/gfx950 bake
+    // 8 (the MI300X value), every other arch 1. gfx942 covers BOTH MI300X (8
+    // XCDs) and MI300A (6 XCDs), which codegen cannot tell apart, so it always
+    // bakes 8; the host guard rejects a device whose runtime NUM_XCD != this
+    // baked value (so MI300A's 6 is excluded at runtime). Power-of-two keeps the
+    // StreamK queue masking (AND/shift) valid.
+    rv["NumXCD"] = checkInList(isaVersion, {{9, 4, 2}, {9, 5, 0}}) ? 8 : 1;
+
+    // Per-queue counter stride = L2 cache-line size (uniform 128B on supported archs).
+    rv["CacheLineBytes"] = 128;
+
     return rv;
 }
 
diff --git a/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp b/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp
index a6782f85c4fb..660b453a13c5 100644
--- a/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp
+++ b/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp
@@ -43,6 +43,7 @@
 #include <cmath>
 #include <cstddef>
 #include <cstdlib>
+#include <mutex>
 #include <random>
 
 #include <sstream>
@@ -56,6 +57,100 @@
 
 namespace TensileLite
 {
+    namespace
+    {
+        // The dynamic-queue StreamK kernels (SK4 and the SK4 sub-path of SK5)
+        // bake a fixed power-of-two per-XCD queue count for fast index masking.
+        // Codegen derives it from the arch's XCD count (StreamK.py
+        // _wsQueueConstants / archCaps["NumXCD"], mirroring origami
+        // get_default_num_xcds); the host reads the SAME origami value here so
+        // codegen and the runtime guard stay in lockstep. Returns 0 when the
+        // architecture cannot be determined (guard treats that as unsupported).
+        inline size_t streamKBakedQueueCount(Hardware const& hardware)
+        {
+            auto const* hipAMDGPU = dynamic_cast<hip::HipAMDGPU const*>(&hardware);
+            if(hipAMDGPU == nullptr || hipAMDGPU->analyticalHardware == nullptr)
+                return 0;
+            try
+            {
+                return origami::hardware_t::get_default_num_xcds(
+                    hipAMDGPU->analyticalHardware->arch);
+            }
+            catch(std::exception const&)
+            {
+                // origami throws for architectures without a hardcoded default
+                // XCD count; treat that as "cannot determine" (0 == unsupported)
+                // rather than propagating the exception through solution
+                // selection.
+                return 0;
+            }
+        }
+
+        // Per-XCD counter stride (bytes) for the dynamic-queue work-queue
+        // region. Set equal to the hardware L2 cache-line size so each per-XCD
+        // atomic counter occupies its own line (no false sharing). Sourced from
+        // origami (hardware_t::get_default_cache_line_bytes) -- the SAME value
+        // the codegen mirrors via rocisa archCaps["CacheLineBytes"]
+        // (StreamK.py _wsQueueConstants). Host (origami) and codegen (archCaps)
+        // strides are two mirrors of the one origami cache-line size, so the
+        // workspace the host reserves matches the layout the kernel addresses.
+        // Returns 0 when the architecture cannot be determined.
+        inline size_t streamKPerQueueStrideBytes(Hardware const& hardware)
+        {
+            auto const* hipAMDGPU = dynamic_cast<hip::HipAMDGPU const*>(&hardware);
+            if(hipAMDGPU == nullptr || hipAMDGPU->analyticalHardware == nullptr)
+                return 0;
+            return origami::hardware_t::get_default_cache_line_bytes(
+                hipAMDGPU->analyticalHardware->arch);
+        }
+
+        // The dynamic-queue fetch / work stealing is only correct when the
+        // device's runtime NUM_XCD is a power of two AND equals the baked
+        // per-XCD queue count. Returns true (UNSUPPORTED) when the hardware is
+        // unknown (not a HipAMDGPU, missing analytical hardware, or no baked
+        // per-XCD queue count), when NUM_XCD is 0, not a power of two, or
+        // NUM_XCD != baked (e.g. MI300A's 6 XCDs, or a 4-XCD partition of an
+        // 8-XCD gfx942). Unknown hardware is treated as UNSUPPORTED: the
+        // dynamic-queue solution is then excluded from selection and a
+        // non-dynamic-queue solution serves the GEMM, rather than staying
+        // selectable while the per-XCD counter workspace is sized with an
+        // unknown (0) queue count (which would under-allocate). Kept isolated
+        // here so it stays trivially unit-testable (see CuCount_test.cpp).
+        inline bool streamKDynamicQueueUnsupported(Hardware const& hardware)
+        {
+            auto const* hipAMDGPU = dynamic_cast<hip::HipAMDGPU const*>(&hardware);
+            if(hipAMDGPU == nullptr || hipAMDGPU->analyticalHardware == nullptr)
+                return true;
+            size_t baked  = streamKBakedQueueCount(hardware);
+            size_t numXCD = hipAMDGPU->analyticalHardware->NUM_XCD;
+            return baked == 0 || numXCD == 0 || (numXCD & (numXCD - 1)) != 0
+                   || numXCD != baked;
+        }
+
+        // Emit a single, user-visible warning (not once-per-call spam) when a
+        // StreamK dynamic-queue / work-stealing solution is excluded from
+        // selection because the device's XCD count does not match the compiled
+        // per-XCD queue count. This is what surfaces the reject to the user
+        // instead of silently degrading to tree reduction.
+        void warnStreamKDynamicQueueUnsupportedOnce(Hardware const& hardware)
+        {
+            static std::once_flag warnedFlag;
+            std::call_once(warnedFlag, [&]() {
+                size_t      numXCD    = 0;
+                auto const* hipAMDGPU = dynamic_cast<hip::HipAMDGPU const*>(&hardware);
+                if(hipAMDGPU != nullptr && hipAMDGPU->analyticalHardware != nullptr)
+                    numXCD = hipAMDGPU->analyticalHardware->NUM_XCD;
+                size_t baked = streamKBakedQueueCount(hardware);
+                std::cerr << "hipBLASLt Warning: StreamK dynamic-queue (work-stealing) solutions "
+                             "require the device's XCD count to be a power of two and to equal the "
+                             "compiled per-XCD queue count; this device reports NUM_XCD="
+                          << numXCD << " with a compiled per-XCD queue count of " << baked
+                          << ", so those solutions are excluded from selection and a "
+                             "non-work-stealing solution will be used instead.\n";
+            });
+        }
+    }
+
     enum class KERNELARGTYPE
     {
         NORMAL   = 0,
@@ -3218,6 +3313,40 @@ namespace TensileLite
             const bool effectiveDynamic = (sizeMapping.streamK == 5)
                                               ? streamK5EffectiveDynamic(problem, hardware)
                                               : false;
+            // Defensive: dynamic-queue / work-stealing StreamK solutions are
+            // excluded from selection on devices whose runtime XCD count is not
+            // a power of two or does not equal the baked per-XCD queue count
+            // (see streamKDynamicQueueSupported() wired into softwarePredicate).
+            // The normal path therefore never reaches solve() for such a
+            // solution; a different (SK3-static / non-StreamK) solution serves
+            // the GEMM instead. If we DO get here it means the software
+            // predicate was bypassed (e.g. an explicit select-by-index), so
+            // reject EXPLICITLY rather than silently running the fixed-mask
+            // kernel with a mismatched queue count (which would corrupt
+            // results).
+            const bool dynamicQueuePath
+                = (sizeMapping.streamK == 4)
+                  || (sizeMapping.streamK == 5 && effectiveDynamic);
+            if(dynamicQueuePath && streamKDynamicQueueUnsupported(hardware))
+            {
+                warnStreamKDynamicQueueUnsupportedOnce(hardware);
+                // Fail EARLY -- before workspace sizing / kernel-arg packing --
+                // when NUM_XCD is unknown (baked queue count == 0, e.g. missing
+                // analyticalHardware). Sizing the per-XCD counter region with a
+                // 0 queue count would under-allocate the workspace the kernel
+                // writes; reject with an actionable message instead.
+                if(streamKBakedQueueCount(hardware) == 0)
+                    throw std::runtime_error(
+                        "hipBLASLt Error: StreamK dynamic-queue (work-stealing) requires a known "
+                        "NUM_XCD (analyticalHardware unavailable); refusing to size the per-XCD "
+                        "counter workspace with an unknown queue count. "
+                        "Select a non-work-stealing solution instead.");
+                throw std::runtime_error(
+                    "hipBLASLt Error: StreamK dynamic-queue (work-stealing) solution selected on a "
+                    "device whose XCD count is not a power of two or does not equal the compiled "
+                    "per-XCD queue count; this kernel is unsupported here. "
+                    "Select a non-work-stealing solution instead.");
+            }
             if(sizeMapping.streamK == 4)
                 sk.reduction = origami::reduction_t::tree;
             else if(sizeMapping.streamK == 5)
@@ -3241,12 +3370,20 @@ namespace TensileLite
             {
                 size_t idealWorkspace = partialTileSize(sk.grid);
                 // SK4 and SK5-dynamic need the per-XCD work-queue region; SK5-static
-                // sizes like standalone SK3.
-                if(sizeMapping.streamK == 4
-                   || (sizeMapping.streamK == 5 && effectiveDynamic))
-                    idealWorkspace += 256 * 8;
+                // sizes like standalone SK3. The region is sized as (per-queue
+                // stride) * (baked per-XCD queue count), both sourced from origami:
+                // the stride is the L2 cache-line size (get_default_cache_line_bytes,
+                // 128 B for gfx942/gfx950) so each counter owns its line (no false
+                // sharing), and the queue count is the per-arch XCD count. The
+                // acceptance guard requires runtime NUM_XCD == baked, so this equals
+                // cacheLineBytes * NUM_XCD for every device that reaches here.
+                if(dynamicQueuePath)
+                    idealWorkspace
+                        += streamKPerQueueStrideBytes(hardware) * streamKBakedQueueCount(hardware);
                 // If given workspace is less than ideal, we can fall back to DP mode
-                // Performance will likely be lower, but the kernel can run if workspace is unavailable
+                // Performance will likely be lower, but the kernel can run if workspace is unavailable.
+                // (The non-power-of-two XCD case is handled earlier by explicit
+                // rejection, not by a silent fall back to tree reduction.)
                 if(idealWorkspace > problem.workspaceSize())
                 {
                     sk.reduction = origami::reduction_t::tree;
@@ -3672,9 +3809,19 @@ namespace TensileLite
                 else if(skGrid > 0 && (tiles % skGrid != 0 && !streamKDP && !forceDPOnly))
                 {
                     size_t idealWorkspace = partialTileSize(skGrid);
+                    // Reserve the per-XCD work-queue region for the dynamic-queue
+                    // path. Sized as (per-queue stride) * (baked per-XCD queue
+                    // count), both from origami: the stride is the L2 cache-line
+                    // size (get_default_cache_line_bytes, 128 B for gfx942/gfx950)
+                    // so each counter owns its line (no false sharing), and the
+                    // queue count is the per-arch XCD count (e.g. 8 for
+                    // gfx942/gfx950). This may slightly over-report on a device
+                    // that falls back to tree reduction (e.g. MI300A), which is
+                    // safe (never under-sized).
                     if(sizeMapping.streamK == 4
                        || (sizeMapping.streamK == 5 && effectiveDynamic))
-                        idealWorkspace += 256 * 8;
+                        idealWorkspace
+                            += streamKPerQueueStrideBytes(hardware) * streamKBakedQueueCount(hardware);
                     // If given workspace is less than ideal, we can fall back to DP mode
                     // Performance will likely be lower, but the kernel can run if workspace is unavailable
                     if(idealWorkspace <= problem.workspaceSize())
@@ -3967,6 +4114,43 @@ namespace TensileLite
         return effectiveDynamic;
     }
 
+    bool ContractionSolution::streamKDynamicQueueSupported(Problem const&  problem,
+                                                           Hardware const& hardware) const
+    {
+        // Only StreamK solutions can ever take the dynamic-queue / work-stealing
+        // path; everything else (SK3-static, non-StreamK) is always selectable.
+        if(sizeMapping.streamK != 4 && sizeMapping.streamK != 5)
+            return true;
+
+        // Fast/common path: on hardware whose runtime XCD count is a power of
+        // two AND equals the baked per-XCD queue count, the fixed queue masking
+        // is valid, so nothing is excluded. Unknown hardware (missing analytical
+        // info / no baked count) is treated as UNSUPPORTED by the predicate, so
+        // it falls through to the reject-and-continue path below rather than
+        // being kept. Checked before streamK5EffectiveDynamic() so the mainline
+        // gfx942(MI300X)/gfx950 path stays cheap.
+        if(!streamKDynamicQueueUnsupported(hardware))
+            return true;
+
+        // Runtime XCD count is not a power of two or does not equal the baked
+        // per-XCD queue count. Only the dynamic-queue sub-path is affected: an
+        // SK5 solution that resolves to the static (SK3) sub-path for this
+        // problem stays valid and selectable.
+        const bool dynamicQueue
+            = (sizeMapping.streamK == 4)
+              || (sizeMapping.streamK == 5 && streamK5EffectiveDynamic(problem, hardware));
+        if(!dynamicQueue)
+            return true;
+
+        // Reject-and-continue: exclude this dynamic-queue / work-stealing
+        // solution from selection (return false) and warn the user ONCE so they
+        // are informed rather than silently degraded to tree reduction. Because
+        // this is a selection-time predicate, other solutions (SK3-static,
+        // non-StreamK) remain available to serve the GEMM.
+        warnStreamKDynamicQueueUnsupportedOnce(hardware);
+        return false;
+    }
+
     namespace
     {
         size_t getSKGridImpl(ContractionSolution const& self,
diff --git a/projects/hipblaslt/tensilelite/tests/CuCount_test.cpp b/projects/hipblaslt/tensilelite/tests/CuCount_test.cpp
index 5ab444dceb68..ba9796c9c5d7 100644
--- a/projects/hipblaslt/tensilelite/tests/CuCount_test.cpp
+++ b/projects/hipblaslt/tensilelite/tests/CuCount_test.cpp
@@ -697,8 +697,9 @@ TEST(StreamK5WorkspaceRegressionTest, QueryAndLaunchAgreeForDynamicMode)
     size_t ws = env.solution.requiredWorkspaceSize(problem, env.device);
     EXPECT_GT(ws, 0u) << "Dynamic mode with partial tiles must request workspace";
 
-    // The workspace must be at least partialTileSize; the +2048 queue
-    // region is included by both query and launch so they agree.
+    // The workspace must be at least partialTileSize; the per-XCD work-queue
+    // region (cacheLineBytes * numXCD = 128 * 8 = 1024 B on gfx942/gfx950) is
+    // included by both query and launch so they agree.
     EXPECT_GE(ws, env.solution.partialTileSize(grid))
         << "Workspace must cover at least partialTileSize(grid)";
 }
@@ -723,7 +724,7 @@ TEST(StreamK5WorkspaceRegressionTest, StaticModeOmitsQueueRegion)
     size_t ws = env.solution.requiredWorkspaceSize(problem, env.device);
 
     // Static (SK3) path does not use the work-queue, so workspace
-    // should be exactly partialTileSize — no +2048.
+    // should be exactly partialTileSize — no per-XCD work-queue region.
     EXPECT_EQ(ws, env.solution.partialTileSize(grid))
         << "OFF workspace must equal partialTileSize(staticGrid)";
 }
@@ -804,3 +805,200 @@ TEST(Sk3Sk5OffPartition512Test, NativeSk3MatchesSk5OffHostPack)
         EXPECT_NE(sk3Pack.grid, sk5OnPack.grid)
             << "512^3 static path oversubscribes; dynamic path should not match";
 }
+
+// ===========================================================================
+// StreamKDynamicQueueXcdGateTest -- MI300A (NUM_XCD=6) reject-and-continue.
+//
+// SK4 / SK5-dynamic work-stealing kernels bake a fixed power-of-two per-XCD
+// queue count (8 for gfx942/gfx950) and mask indices with (Q-1), so they are
+// valid only when the device's runtime NUM_XCD equals that baked count. MI300A
+// reports gfx942 (baked 8) but has 6 XCDs; a mismatched partition (e.g. a 4-XCD
+// slice of an 8-XCD gfx942) is likewise rejected. The host excludes such a
+// solution from selection (streamKDynamicQueueSupported wired into
+// softwarePredicate) and warns once instead of silently degrading. The
+// production predicates live in a .cpp anonymous namespace, so -- like
+// computeStreamKHostPack above -- this test mirrors them over a hip::HipAMDGPU
+// mock (6 -> reject, 4 -> reject, 8 -> allow, unknown -> allow). Not run on
+// real MI300A silicon.
+// ===========================================================================
+namespace
+{
+    // Mirror of the anonymous-namespace helper in ContractionSolution.cpp
+    // (streamKBakedQueueCount): the baked per-XCD queue count comes from
+    // origami's per-arch XCD count. Kept in lockstep with the production code.
+    inline size_t streamKBakedQueueCountRef(Hardware const& hardware)
+    {
+        auto const* hipAMDGPU = dynamic_cast<hip::HipAMDGPU const*>(&hardware);
+        if(hipAMDGPU == nullptr || hipAMDGPU->analyticalHardware == nullptr)
+            return 0;
+        try
+        {
+            return origami::hardware_t::get_default_num_xcds(
+                hipAMDGPU->analyticalHardware->arch);
+        }
+        catch(std::exception const&)
+        {
+            return 0;
+        }
+    }
+
+    // Byte-for-byte mirror of the anonymous-namespace numeric predicate in
+    // ContractionSolution.cpp (streamKDynamicQueueUnsupported). Kept in lockstep
+    // with the production code; if that predicate changes, update this too.
+    // Unknown hardware (not a HipAMDGPU, no analytical hardware, or no baked
+    // per-XCD queue count) is treated as UNSUPPORTED (returns true).
+    inline bool streamKDynamicQueueUnsupportedRef(Hardware const& hardware)
+    {
+        auto const* hipAMDGPU = dynamic_cast<hip::HipAMDGPU const*>(&hardware);
+        if(hipAMDGPU == nullptr || hipAMDGPU->analyticalHardware == nullptr)
+            return true;
+        size_t baked  = streamKBakedQueueCountRef(hardware);
+        size_t numXCD = hipAMDGPU->analyticalHardware->NUM_XCD;
+        return baked == 0 || numXCD == 0 || (numXCD & (numXCD - 1)) != 0
+               || numXCD != baked;
+    }
+
+    // Mirror of ContractionSolution::streamKDynamicQueueSupported(). Returns
+    // true when the solution is SELECTABLE, false when it must be EXCLUDED
+    // (dynamic-queue / work-stealing on a non-power-of-two XCD device). streamK
+    // is sizeMapping.streamK; effectiveDynamic is the SK5 sub-mode result
+    // (ignored for streamK != 5). Kept in lockstep with the production member.
+    inline bool streamKDynamicQueueSupportedRef(int             streamK,
+                                                bool            effectiveDynamic,
+                                                Hardware const& hardware)
+    {
+        if(streamK != 4 && streamK != 5)
+            return true;
+        if(!streamKDynamicQueueUnsupportedRef(hardware))
+            return true;
+        const bool dynamicQueue = (streamK == 4) || (streamK == 5 && effectiveDynamic);
+        return !dynamicQueue; // dynamic-queue on non-pow2 XCD -> excluded
+    }
+
+    // gfx942 analytical hardware with a caller-chosen XCD count. NUM_XCD is the
+    // 5th positional arg of origami::hardware_t (see origami/hardware.hpp).
+    origami::hardware_t makeGfx942HardwareWithXcd(size_t numXCD)
+    {
+        using arch_t = origami::hardware_t::architecture_t;
+        return origami::hardware_t(arch_t::gfx942,
+                                   304, // N_CU (MI300X SPX)
+                                   163840,
+                                   262144,
+                                   numXCD,
+                                   1.0,
+                                   1.0,
+                                   1.0,
+                                   4000000,
+                                   1.2,
+                                   1,
+                                   std::make_tuple(0.0, 0.008, 0.0));
+    }
+
+    hip::HipAMDGPU makeGfx942DeviceWithXcd(size_t numXCD)
+    {
+        hip::HipAMDGPU device;
+        device.processor          = AMDGPU::Processor::gfx942;
+        device.computeUnitCount   = 304;
+        device.deviceName         = "test-gfx942-xcd";
+        device.analyticalHardware = std::make_shared<origami::hardware_t>(
+            makeGfx942HardwareWithXcd(numXCD));
+        return device;
+    }
+} // namespace
+
+TEST(StreamKDynamicQueueXcdGateTest, RejectsMi300aSixXcd)
+{
+    hip::HipAMDGPU mi300a = makeGfx942DeviceWithXcd(6);
+    Hardware const& hw    = mi300a;
+    EXPECT_TRUE(streamKDynamicQueueUnsupportedRef(hw))
+        << "MI300A (NUM_XCD=6, not a power of two) must flag the dynamic-queue "
+           "work-stealing path as unsupported";
+}
+
+TEST(StreamKDynamicQueueXcdGateTest, AllowsMi300xEightXcd)
+{
+    hip::HipAMDGPU mi300x = makeGfx942DeviceWithXcd(8);
+    Hardware const& hw    = mi300x;
+    EXPECT_FALSE(streamKDynamicQueueUnsupportedRef(hw))
+        << "MI300X (NUM_XCD=8, power of two) must keep the dynamic-queue path";
+}
+
+TEST(StreamKDynamicQueueXcdGateTest, RejectsGfx942FourXcdPowerOfTwoButMismatched)
+{
+    // A 4-XCD partition (e.g. a CPX-style slice) of an 8-XCD gfx942: 4 IS a
+    // power of two, but the kernel bakes Q=8, so runtime NUM_XCD (4) != baked
+    // (8) and the fixed Q=8 masking would mis-map queues. Must be rejected.
+    hip::HipAMDGPU  gfx942Cpx = makeGfx942DeviceWithXcd(4);
+    Hardware const& hw        = gfx942Cpx;
+    EXPECT_EQ(streamKBakedQueueCountRef(hw), 8u)
+        << "gfx942 must bake origami's per-arch XCD count (8)";
+    EXPECT_TRUE(streamKDynamicQueueUnsupportedRef(hw))
+        << "gfx942 with NUM_XCD=4 (power of two but != baked 8) must be rejected";
+}
+
+TEST(StreamKDynamicQueueXcdGateTest, AllowsGfx950EightXcd)
+{
+    // gfx950 (local MI355X) analytical hardware advertises 8 XCDs.
+    hip::HipAMDGPU gfx950   = makeHipDeviceWithAnalytical(makeGfx950AnalyticalHardware());
+    Hardware const& hw      = gfx950;
+    EXPECT_FALSE(streamKDynamicQueueUnsupportedRef(hw))
+        << "gfx950 (NUM_XCD=8) must keep the dynamic-queue work-stealing path";
+}
+
+TEST(StreamKDynamicQueueXcdGateTest, MissingAnalyticalHardwareIsUnsupported)
+{
+    // Unknown hardware (null analyticalHardware -> unknown NUM_XCD / baked
+    // queue count == 0) must be treated as UNSUPPORTED so the dynamic-queue
+    // solution is excluded from selection and a non-dynamic-queue solution
+    // serves the GEMM, rather than staying selectable while the per-XCD counter
+    // workspace is sized with an unknown (0) queue count (under-allocation).
+    hip::HipAMDGPU noAnalytical;
+    noAnalytical.processor     = AMDGPU::Processor::gfx942;
+    noAnalytical.deviceName    = "test-gfx942-no-analytical";
+    Hardware const& hwNoAnalyt = noAnalytical;
+    ASSERT_EQ(noAnalytical.analyticalHardware, nullptr);
+    EXPECT_TRUE(streamKDynamicQueueUnsupportedRef(hwNoAnalyt))
+        << "Missing analyticalHardware (unknown NUM_XCD) must be treated as unsupported";
+    // And the selection predicate must therefore EXCLUDE the dynamic-queue
+    // solution (SK4) while keeping non-dynamic-queue solutions selectable.
+    EXPECT_FALSE(streamKDynamicQueueSupportedRef(4, /*effectiveDynamic=*/false, hwNoAnalyt))
+        << "SK4 work-stealing solution must be excluded when NUM_XCD is unknown";
+    EXPECT_TRUE(streamKDynamicQueueSupportedRef(3, /*effectiveDynamic=*/false, hwNoAnalyt))
+        << "SK3-static solution must remain selectable when NUM_XCD is unknown";
+}
+
+// Selection-predicate contract: on MI300A (6 XCD) the dynamic-queue solution is
+// EXCLUDED from selection (supported == false) so a different solution serves
+// the GEMM, while on MI300X (8 XCD) the identical solution stays selectable.
+TEST(StreamKDynamicQueueXcdGateTest, ExcludesDynamicQueueSolutionOnMi300a)
+{
+    hip::HipAMDGPU  mi300a = makeGfx942DeviceWithXcd(6);
+    hip::HipAMDGPU  mi300x = makeGfx942DeviceWithXcd(8);
+    Hardware const& hwA    = mi300a;
+    Hardware const& hwX    = mi300x;
+
+    // SK4 is always dynamic-queue.
+    EXPECT_FALSE(streamKDynamicQueueSupportedRef(4, /*effectiveDynamic=*/false, hwA))
+        << "SK4 work-stealing solution must be excluded from selection on MI300A";
+    EXPECT_TRUE(streamKDynamicQueueSupportedRef(4, /*effectiveDynamic=*/false, hwX))
+        << "SK4 work-stealing solution must remain selectable on MI300X";
+
+    // SK5 only takes the dynamic-queue path when it resolves to the dynamic
+    // sub-mode; the static (SK3) sub-mode stays selectable even on MI300A.
+    EXPECT_FALSE(streamKDynamicQueueSupportedRef(5, /*effectiveDynamic=*/true, hwA))
+        << "SK5-dynamic must be excluded from selection on MI300A";
+    EXPECT_TRUE(streamKDynamicQueueSupportedRef(5, /*effectiveDynamic=*/false, hwA))
+        << "SK5-static (SK3 sub-path) must remain selectable on MI300A";
+}
+
+// Non-dynamic-queue solutions must never be excluded, so the GEMM still runs.
+TEST(StreamKDynamicQueueXcdGateTest, KeepsNonDynamicQueueSolutionsOnMi300a)
+{
+    hip::HipAMDGPU  mi300a = makeGfx942DeviceWithXcd(6);
+    Hardware const& hwA    = mi300a;
+
+    EXPECT_TRUE(streamKDynamicQueueSupportedRef(0, /*effectiveDynamic=*/false, hwA))
+        << "Non-StreamK solution must remain selectable on MI300A";
+    EXPECT_TRUE(streamKDynamicQueueSupportedRef(3, /*effectiveDynamic=*/false, hwA))
+        << "SK3-static solution must remain selectable on MI300A";
+}
diff --git a/shared/origami/include/origami/hardware.hpp b/shared/origami/include/origami/hardware.hpp
index 6646f7e9e5c1..f78d8051f0ca 100644
--- a/shared/origami/include/origami/hardware.hpp
+++ b/shared/origami/include/origami/hardware.hpp
@@ -757,6 +757,17 @@ class ORIGAMI_EXPORT hardware_t {
    */
   static size_t get_default_num_xcds(architecture_t arch);
 
+  /**
+   * @brief Get the default L2 cache-line size (in bytes) for an architecture.
+   *
+   * Returns the per-arch L2 cache-line size used for the StreamK per-queue
+   * counter stride (currently uniform 128 B across supported archs).
+   *
+   * @param arch Architecture enum value
+   * @return L2 cache-line size in bytes
+   */
+  static size_t get_default_cache_line_bytes(architecture_t arch);
+
   /**
    * @brief Check if the hardware described by properties is supported.
    *
diff --git a/shared/origami/src/origami/hardware.cpp b/shared/origami/src/origami/hardware.cpp
index 10d7e560ca4d..d93926d5c38e 100644
--- a/shared/origami/src/origami/hardware.cpp
+++ b/shared/origami/src/origami/hardware.cpp
@@ -213,6 +213,11 @@ size_t hardware_t::get_default_num_xcds(architecture_t arch) {
   }
 }
 
+size_t hardware_t::get_default_cache_line_bytes(architecture_t /*arch*/) {
+  // Per-arch L2 cache-line size, currently uniform 128 B across supported archs.
+  return 128;
+}
+
 void hardware_t::print() const {
   std::cout << "================== Hardware Configuration ==================\n";
   std::cout << "Number of CUs (N_CU)      : " << N_CU << "\n";