Skip to content
Merged
Show file tree
Hide file tree
Changes from 45 commits
Commits
Show all changes
53 commits
Select commit Hold shift + click to select a range
bb2b526
[hipblaslt][tensilelite] Single-hop next-neighbor StreamK work stealing
jaopaulolc Jun 15, 2026
948a5a1
[hipblaslt][tensilelite] Tests for single-hop StreamK work stealing
jaopaulolc Jun 15, 2026
70324cb
Merge branch 'develop' into users/jolabega/dyn-streamk-queue-work-ste…
jaopaulolc Jun 18, 2026
b3453ff
Merge branch 'develop' into users/jolabega/dyn-streamk-queue-work-ste…
jaopaulolc Jun 19, 2026
b47579a
Merge branch 'develop' into users/jolabega/dyn-streamk-queue-work-ste…
jaopaulolc Jun 23, 2026
a251b3a
Merge branch 'develop' into users/jolabega/dyn-streamk-queue-work-ste…
jaopaulolc Jun 24, 2026
327120e
[tensilelite] Fix parameter types in StreamK work-stealing test YAMLs
jaopaulolc Jun 24, 2026
73a4f85
Merge remote-tracking branch 'origin/develop' into users/jolabega/dyn…
jaopaulolc Jun 24, 2026
fd43325
Merge branch 'develop' into users/jolabega/dyn-streamk-queue-work-ste…
jaopaulolc Jun 24, 2026
189649e
[tensilelite] Regenerate characterization snapshots for StreamKWorkSt…
jaopaulolc Jun 24, 2026
a4821d8
Merge branch 'develop' into users/jolabega/dyn-streamk-queue-work-ste…
jaopaulolc Jun 24, 2026
d09e44a
[tensilelite] Reset StreamK work-stealing counters on deferred-store …
jaopaulolc Jun 25, 2026
c457dd2
Remove MergeFiles
Alex-Vasile Jun 25, 2026
1796e31
Further fixes
Alex-Vasile Jun 25, 2026
e56cb24
[tensilelite] Drop host-side StreamK WS synchronizer reset on this br…
jaopaulolc Jun 25, 2026
3bd3d11
Merge branch 'develop' into users/jolabega/dyn-streamk-queue-work-ste…
jaopaulolc Jun 25, 2026
a8d3123
Merge remote-tracking branch 'origin/develop' into users/jolabega/dyn…
jaopaulolc Jun 26, 2026
7c9cccc
Merge remote-tracking branch 'origin/develop' into users/jolabega/dyn…
jaopaulolc Jul 6, 2026
8971162
[tensilelite] Document StreamKWorkStealing required-param golden chan…
jaopaulolc Jul 6, 2026
dfd1ff5
Merge branch 'develop' into users/jolabega/dyn-streamk-queue-work-ste…
jaopaulolc Jul 6, 2026
cf49645
Merge branch 'develop' into users/jolabega/dyn-streamk-queue-work-ste…
jaopaulolc Jul 6, 2026
1504c79
Merge remote-tracking branch 'origin/develop' into users/jolabega/dyn…
jaopaulolc Jul 7, 2026
87e202c
Merge branch 'develop' into users/jolabega/dyn-streamk-queue-work-ste…
jaopaulolc Jul 7, 2026
f6561c8
Merge branch 'develop' into users/jolabega/dyn-streamk-queue-work-ste…
jaopaulolc Jul 7, 2026
29941d4
Merge branch 'develop' into users/jolabega/dyn-streamk-queue-work-ste…
jaopaulolc Jul 8, 2026
e25ac2e
Merge branch 'develop' into users/jolabega/dyn-streamk-queue-work-ste…
jaopaulolc Jul 9, 2026
19775e0
Merge branch 'develop' into users/jolabega/dyn-streamk-queue-work-ste…
jaopaulolc Jul 10, 2026
9e67bbd
feat(tensilelite): gate StreamK work stealing to power-of-two XCD counts
jaopaulolc Jul 13, 2026
eaf679d
feat(tensilelite): reset-free single-hop next-neighbor StreamK work s…
jaopaulolc Jul 13, 2026
ae1f803
Merge branch 'develop' into users/jolabega/dyn-streamk-queue-work-ste…
jaopaulolc Jul 13, 2026
541e93d
Merge remote-tracking branch 'origin/develop' into users/jolabega/dyn…
jaopaulolc Jul 13, 2026
91bc727
feat(tensilelite): derive StreamK work-queue count from per-arch XCD …
jaopaulolc Jul 13, 2026
1ad17c6
Merge branch 'develop' into users/jolabega/dyn-streamk-queue-work-ste…
jaopaulolc Jul 13, 2026
79924af
refactor(tensilelite): tidy StreamK work-stealing comments and derive…
jaopaulolc Jul 13, 2026
f64b3fe
refactor(tensilelite): set StreamK queue stride to the origami cache-…
jaopaulolc Jul 13, 2026
4ece523
Merge branch 'develop' into users/jolabega/dyn-streamk-queue-work-ste…
jaopaulolc Jul 14, 2026
e5934aa
refactor(tensilelite): trim StreamK work-stealing comments, docs, and…
jaopaulolc Jul 14, 2026
f9569ca
fix(tensilelite): reject StreamK dynamic-queue/work-stealing on unkno…
jaopaulolc Jul 14, 2026
09d0417
Merge branch 'develop' into users/jolabega/dyn-streamk-queue-work-ste…
jaopaulolc Jul 14, 2026
725af13
Merge remote-tracking branch 'origin/develop' into users/jolabega/dyn…
jaopaulolc Jul 17, 2026
4cf435d
Regenerate characterization .ambr snapshots after merging develop
jaopaulolc Jul 17, 2026
61c460b
Regenerate test_r3_streamk_tdmsplit_gfx1250 golden after merging develop
jaopaulolc Jul 17, 2026
a6bc8b3
Merge branch 'develop' into users/jolabega/dyn-streamk-queue-work-ste…
jaopaulolc Jul 17, 2026
bf9d27b
Merge branch 'develop' into users/jolabega/dyn-streamk-queue-work-ste…
jaopaulolc Jul 20, 2026
df0564c
Merge branch 'develop' into users/jolabega/dyn-streamk-queue-work-ste…
jaopaulolc Jul 20, 2026
94e0669
Merge branch 'develop' into users/jolabega/dyn-streamk-queue-work-ste…
jaopaulolc Jul 21, 2026
9d8750c
Merge branch 'develop' into users/jolabega/dyn-streamk-queue-work-ste…
jaopaulolc Jul 21, 2026
7b0ef09
Merge remote-tracking branch 'origin/develop' into users/jolabega/dyn…
jaopaulolc Jul 22, 2026
d23e85c
Regenerate characterization .ambr snapshots after merging develop
jaopaulolc Jul 22, 2026
11ec315
Merge branch 'develop' into users/jolabega/dyn-streamk-queue-work-ste…
jaopaulolc Jul 22, 2026
8a1b2ef
Merge remote-tracking branch 'origin/develop' into users/jolabega/dyn…
jaopaulolc Jul 24, 2026
684ef65
Regenerate characterization .ambr snapshots after merging develop
jaopaulolc Jul 24, 2026
a52fa99
Merge branch 'develop' into users/jolabega/dyn-streamk-queue-work-ste…
jaopaulolc Jul 24, 2026
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
Original file line number Diff line number Diff line change
Expand Up @@ -568,6 +568,7 @@
{"StreamK": [0]},
{"StreamKForceDPOnly": [0]},
{"StreamKAtomic": [0]},
{"StreamKWorkStealing": [0]},
{"StreamKXCCMapping": [0]},
{"StreamKFixupTreeReduction": [0]},
{"DebugStreamK": [0]},
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -122,6 +122,7 @@ def getRequiredParametersMin() -> set:
'StoreVectorWidth',
'StreamK',
'StreamKForceDPOnly',
'StreamKWorkStealing',
'StreamKXCCMapping',
'StreamKFixupTreeReduction',
'SuppressNoLoadLoop',
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -816,6 +816,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
Expand Down
266 changes: 241 additions & 25 deletions projects/hipblaslt/tensilelite/Tensile/Components/StreamK.py

Large diffs are not rendered by default.

14 changes: 14 additions & 0 deletions projects/hipblaslt/tensilelite/Tensile/KernelWriter.py
Original file line number Diff line number Diff line change
Expand Up @@ -5250,6 +5250,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
Expand Down Expand Up @@ -9432,6 +9437,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:
Expand All @@ -9453,6 +9463,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:
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -1751,13 +1751,49 @@ 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
else:
# 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
Expand Down Expand Up @@ -2781,6 +2817,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"],
Expand Down
Original file line number Diff line number Diff line change
@@ -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]
Original file line number Diff line number Diff line change
@@ -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]
Original file line number Diff line number Diff line change
Expand Up @@ -346,4 +346,13 @@ No other equivalent mutants are accepted yet; widening rounds append their
accepted equivalents/pragmas here, each with its one-line reason.

## 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."
**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.
Original file line number Diff line number Diff line change
Expand Up @@ -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_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_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_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_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',
]),
})
# ---
Expand Down Expand Up @@ -55,7 +55,7 @@
'getitem_kernel_language': 'Assembly',
'iter_matches_keys': True,
'keys_is_list': True,
'len': 335,
'len': 336,
})
# ---
# name: test_solution_construction
Expand Down Expand Up @@ -293,6 +293,7 @@
'StreamKAtomic',
'StreamKFixupTreeReduction',
'StreamKForceDPOnly',
'StreamKWorkStealing',
'StreamKXCCMapping',
'SubGroup0',
'SubGroup1',
Expand Down Expand Up @@ -397,8 +398,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_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_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': 335,
'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_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_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': 336,
'problem_type': 'Cijk_Alik_Bjlk_HSS_BH_Bias_HA_S_SAV_UserArgs',
'stable': dict({
'DepthU': 32,
Expand Down
Loading
Loading