Skip to content

[hipblaslt] Manual revert KRingShift #7443

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
yenong-amd merged 4 commits into from
May 15, 2026
Merged

[hipblaslt] Manual revert KRingShift #7443

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
bc9a194
[hipblaslt] Revert BAddrInterleave/KRingShift feature (PR #3679)
yenong-amd May 14, 2026
38d55f3
[hipblaslt] Strip BAddrInterleave/KRingShift keys from library YAML f…
yenong-amd May 14, 2026
7a6f866
"Fix skipGroTileOffsetsLoop flag"
yenong-amd May 14, 2026
03bbe8d
Merge branch 'develop' into users/yenong-amd/revert_kringshift
yenong-amd May 14, 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
6 changes: 3 additions & 3 deletions ...asm_full/gfx950/gfx950/Equality/gfx950_Cijk_Alik_Bljk_BBS_BH_BiasSB_HAS_SAV_UserArgs.yaml
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -375886,7 +375886,7 @@
AssertSummationElementMultiple: 1
AssignedDerivedParameters: true
AssignedProblemIndependentDerivedParameters: true
BAddrInterleave: true
BAddrInterleave: false
BaseName: Cijk_Alik_Bljk_BBS_BH_BiasSB_HAS_SAV_UserArgs_MT128x192x128_MI1n-ZRL8mxl6uWAXsQRkqoULpPWXdDeqIc8yirhD35fpI=
BufferLoad: true
BufferStore: true
Expand Down Expand Up @@ -375932,7 +375932,7 @@
InterleaveAlpha: 0
InternalSupportParams: {KernArgsVersion: 2, SupportCustomStaggerU: true, SupportCustomWGM: true,
SupportUserGSU: false, UseUniversalArgs: true}
KRingShift: true
KRingShift: false
Kernel: true
KernelLanguage: Assembly
KernelNameMin: Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x192x128_MI16x16x1_SN_LDSB1_AFC0_AG0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_DTVMXSA0_DTVMXSB0_DTVSM0_DPLBn1_EPS0_ELFLR0_EMLLn1_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA2048_LBSPPB256_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_3_MO40_MGRIPMn1_NTn1_NTA0_NTB0_NTC1_NTD4_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SGROBn1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW8_SK3_SKFTR0_SKXCCM0_SGRO0_TDMI0_TDMS0_TIN0_TLDS2_TLDSM1_ULSGRO0_USL1_USLMX0_UIOFGRO0_UPLRP0_USFGRO0_USI0_VSn1_VWA8_VWB1_WSGRA0_WSGRB0_WS64_WG16_16_1
Expand Down Expand Up @@ -376058,7 +376058,7 @@
SourceSwap: true
SpaceFillingAlgo: []
StaggerU: 16
StaggerUMapping: 1
StaggerUMapping: 0
StaggerUStride: 256
StorePriorityOpt: false
StoreRemapVectorWidth: 0
Expand Down
6 changes: 3 additions & 3 deletions ...l/gfx950/gfx950_id75a3/Equality/gfx950_Cijk_Alik_Bljk_BBS_BH_BiasSB_HAS_SAV_UserArgs.yaml
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -371587,7 +371587,7 @@
AssertSummationElementMultiple: 1
AssignedDerivedParameters: true
AssignedProblemIndependentDerivedParameters: true
BAddrInterleave: true
BAddrInterleave: false
BaseName: Cijk_Alik_Bljk_BBS_BH_BiasSB_HAS_SAV_UserArgs_MT128x192x128_MI1n-ZRL8mxl6uWAXsQRkqoULpPWXdDeqIc8yirhD35fpI=
BufferLoad: true
BufferStore: true
Expand Down Expand Up @@ -371633,7 +371633,7 @@
InterleaveAlpha: 0
InternalSupportParams: {KernArgsVersion: 2, SupportCustomStaggerU: true, SupportCustomWGM: true,
SupportUserGSU: false, UseUniversalArgs: true}
KRingShift: true
KRingShift: false
Kernel: true
KernelLanguage: Assembly
KernelNameMin: Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x192x128_MI16x16x1_SN_LDSB1_AFC0_AG0_AFEM1_AFEM1_ASEM1_CLR0_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_DTVMXSA0_DTVMXSB0_DTVSM0_DPLBn1_EPS0_ELFLR0_EMLLn1_FDSI0_GRPM1_GRVWA8_GRVWB8_GSU0_GSUAMB_GLS0_ISA950_IU1_K1_LDSTI0_LBSPPA2048_LBSPPB256_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_3_MO40_MGRIPMn1_NTn1_NTA0_NTB0_NTC1_NTD4_NTM0_NEPBS0_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SGROBn1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW8_SK3_SKFTR0_SKXCCM0_SGRO0_TDMI0_TIN0_TLDS2_TLDSM1_ULSGRO0_USL1_UIOFGRO0_UPLRP0_USFGRO0_VSn1_VWA8_VWB1_WSGRA0_WSGRB0_WS64_WG16_16_1
Expand Down Expand Up @@ -371759,7 +371759,7 @@
SourceSwap: true
SpaceFillingAlgo: []
StaggerU: 16
StaggerUMapping: 1
StaggerUMapping: 0
StaggerUStride: 256
StorePriorityOpt: false
StoreRemapVectorWidth: 0
Expand Down
10 changes: 0 additions & 10 deletions projects/hipblaslt/tensilelite/Tensile/Common/GlobalParameters.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -420,23 +420,13 @@
{"DirectToVgprMXSA": [False]},
{"DirectToVgprMXSB": [False]},
{"DirectToVgprSparseMetadata": [False]},
# Restricted address remap features (default off unless explicitly enabled in the solution config):
{"BAddrInterleave": [False]},
{"KRingShift": [False]},
{"DirectToLds": [0]},
{"UseSubtileImpl": [False]},
{"UseSgprForGRO": [-1]},
{"UseInstOffsetForGRO": [0]},
{"AssertSummationElementMultiple": [1]},
{"AssertFree0ElementMultiple": [1]},
{"AssertFree1ElementMultiple": [1]},
# Address-interleave restriction (default disabled):
# When >0, the solution requires tiles1=(SizeJ/MT1) to have lowbit(tiles1)>1 (i.e. G>1),
{"AssertFree1DivByMT1LowbitGT1": [0]},
# KRingShift wrap restriction (default disabled):
# Encodes a runtime predicate that ensures (k + KRingShift) does not wrap in main loop
# (wrap is allowed only in tail loop where codegen applies the correction).
{"AssertKRingShiftTailWrapOnly": [0]},
{"AssertAIGreaterThanEqual": [-1]},
{"AssertAILessThanEqual": [-1]},
{"StaggerU": [32]}, # recommend [0,32]
Expand Down
17 changes: 0 additions & 17 deletions projects/hipblaslt/tensilelite/Tensile/Common/ValidParameters.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -376,14 +376,6 @@ def makeValidMatrixInstructions():
"DirectToVgprMXSA": [False, True],
"DirectToVgprMXSB": [False, True],
"DirectToVgprSparseMetadata": [False, True],
# B address interleave (restricted): non-contiguous tile columns for TN/NN-like B (TLUB == False),
# with runtime G chosen as the largest power-of-two factor of (N/MT1), capped by LVCB.
# Requires SizeJ % MT1 == 0 at runtime; otherwise falls back to original mapping.
"BAddrInterleave": [False, True],
# K ring-shift (restricted): apply a per-WG shift along the summation (K) dimension so that
# the B-side base K address for each workgroup is cacheline-aligned/congruent, while preserving
# correctness via tail-loop ring wrap. Intended for TN/NN-like B (TLUB == False).
"KRingShift": [False, True],
# Attempt to load directly from global memory into LDS.
# Assembly only
# Requires BufferLoad, assembler support for lds modifier on buffer
Expand Down Expand Up @@ -500,15 +492,6 @@ def makeValidMatrixInstructions():
# 1 indicates no assertion (since all sizes are multiples of 1)
"AssertFree1ElementMultiple": [1, 2, 4, 8, 16, 32],
# Address-interleave restriction:
Comment thread
yenong-amd marked this conversation as resolved.
Outdated
# If >0, require tiles1=(Free1Size / MT1) to have lowbit(tiles1)>1 (i.e. G>1).
# This matches the kernel's initBInterleaveG logic:
# - require Free1Size % MT1 == 0
# - compute lowbit(tiles1)
# - enable only if min(lowbit, LVCB) > 1
"AssertFree1DivByMT1LowbitGT1": -1,
# KRingShift wrap restriction (packed integer; see Solution.py for encoding):
# If >0, require any (k + KRingShift) wrap to occur only in tail loop (no main-loop wrap).
"AssertKRingShiftTailWrapOnly": -1,
# Assertions that require arithmetic intensity to be specified value.
# Arithmetic intensity measures the ratio of computation to memory bandwidth required for a problem.
# These predicates can be used to adjust solution selection compute-bound or memory-bound problems.
Expand Down
10 changes: 0 additions & 10 deletions projects/hipblaslt/tensilelite/Tensile/Contractions.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -466,16 +466,6 @@ def FromOriginalKeyPair(cls, pair):
if key == "AssertAILessThanEqual":
return cls("AILessThanEqual", value=value) if value > 0 else None

# Address-interleave restriction:
# Require tiles1 = Free1Size / MT1 to be a power-of-two (and divisible).
if key == "AssertFree1DivByMT1LowbitGT1":
return cls("Free1SizeDivByValueLowbitGT1", index=0, value=value) if value > 0 else None

# KRingShift wrap restriction (packed value; see Solution.py):
# Require that any (k + KRingShift) wrap occurs only in tail loop (no main-loop wrap).
if key == "AssertKRingShiftTailWrapOnly":
return cls("KRingShiftTailWrapOnly", index=-1, value=value) if value > 0 else None

if key.endswith('Multiple'):
if value == 1:
return None
Expand Down
126 changes: 22 additions & 104 deletions projects/hipblaslt/tensilelite/Tensile/KernelWriter.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5276,116 +5276,40 @@ def kernelBody( self, kernel, tensorParametersA, tensorParametersB ):
if tailLoopOpt2nd and (globalReadMode2nd == 2):
module.add(self.doTailLoopOpt(kernel, tensorParameters2nd))
else:
# Keep per-tensor tail branching for tc2 when tc1 uses tailLoopOpt.
if kernel["KRingShift"] and kernel["BufferLoad"] and tc2 in ("A", "B"):
labelNoKRS = Label(self.labels.getNameInc(f"KRS_tail_noop_{tc2}"), "")
labelDoneKRS = Label(self.labels.getNameInc(f"KRS_tail_done_{tc2}"), "")
labelNoKRS.comment = f"KRS: tail no-KRS path for {tc2} (sgprKRingShift==0)"
labelDoneKRS.comment = f"KRS: tail KRS branch join for {tc2}"
module.add(SCmpEQU32(src0=sgpr("KRingShift"), src1=0, comment="KRS: sgprKRingShift==0 ?"))
module.add(SCBranchSCC1(labelName=labelNoKRS.getLabelName(), comment="KRS: take no-KRS tail loads"))
module.add(self.globalReadDo(kernel, globalReadMode2nd, tensorParameters2nd, krTailForceDisable=False))
module.add(SBranch(labelName=labelDoneKRS.getLabelName(), comment="KRS: skip no-KRS tail loads"))
module.add(labelNoKRS)
module.add(self.globalReadDo(kernel, globalReadMode2nd, tensorParameters2nd, krTailForceDisable=True))
module.add(labelDoneKRS)
else:
module.add(self.globalReadDo(kernel, globalReadMode2nd, tensorParameters2nd))
module.add(self.globalReadDo(kernel, globalReadMode2nd, tensorParameters2nd))
else:
# KRS: If both tail global-read blocks (A/B) are eligible for KRS, do ONE runtime branch and
# share ONE set of labels for A/B. When sgprKRingShift==0, force both tail blocks down the
# original "load-only" path (no KRS_TAIL_OFFSET_* at all).
# skip wait for DTL if global load 1st is DTL
skip2ndWaitForDtl = kernel["DirectToLds%s"%tc1]
krsTailBranchable1 = kernel["KRingShift"] and kernel["BufferLoad"] and tc1 in ("A", "B")
krsTailBranchable2 = kernel["KRingShift"] and kernel["BufferLoad"] and tc2 in ("A", "B") \
and not (tailLoopOpt2nd and (globalReadMode2nd == 2))
if krsTailBranchable1 and krsTailBranchable2:
labelNoKRS = Label(self.labels.getNameInc("KRS_tail_noop_AB"), "")
labelDoneKRS = Label(self.labels.getNameInc("KRS_tail_done_AB"), "")
labelNoKRS.comment = "KRS: tail no-KRS path for A/B (sgprKRingShift==0)"
labelDoneKRS.comment = "KRS: tail KRS branch join for A/B"

module.add(SCmpEQU32(src0=sgpr("KRingShift"), src1=0, comment="KRS: sgprKRingShift==0 ?"))
module.add(SCBranchSCC1(labelName=labelNoKRS.getLabelName(), comment="KRS: take no-KRS tail loads (A+B)"))

# KRS-enabled path: A then B
module.add(self.globalReadDo(kernel, globalReadMode1st, tensorParameters1st, krTailForceDisable=False))
module.addComment1("Update M0 for DTLDS")
moduleTmp = self.directToLdsM0Update(kernel, 2, tensorParameters2nd, skip2ndWaitForDtl)
module.add(replaceHolder(moduleTmp, 0))
module.addComment1("Tail global read %s"%tc2)
module.add(self.globalReadDo(kernel, globalReadMode2nd, tensorParameters2nd, krTailForceDisable=False))
module.add(SBranch(labelName=labelDoneKRS.getLabelName(), comment="KRS: skip no-KRS tail loads (A+B)"))
module.add(self.globalReadDo(kernel, globalReadMode1st, tensorParameters1st))

# no-KRS path: A then B (load-only)
module.add(labelNoKRS)
module.add(self.globalReadDo(kernel, globalReadMode1st, tensorParameters1st, krTailForceDisable=True))
if "MX" in tensorParameters1st:
module.addComment1("Update M0 for DTLDS")
moduleTmp = self.directToLdsM0Update(kernel, 2, tensorParameters2nd, skip2ndWaitForDtl)
moduleTmp = self.directToLdsM0Update(kernel, 1, tensorParameters1st["MX"], skipWait=skip2ndWaitForDtl)
module.add(replaceHolder(moduleTmp, 0))
module.addComment1("Tail global read %s"%tc2)
module.add(self.globalReadDo(kernel, globalReadMode2nd, tensorParameters2nd, krTailForceDisable=True))
module.add(labelDoneKRS)
else:
# Fallback: keep per-tensor tail branching.
if krsTailBranchable1:
labelNoKRS = Label(self.labels.getNameInc(f"KRS_tail_noop_{tc1}"), "")
labelDoneKRS = Label(self.labels.getNameInc(f"KRS_tail_done_{tc1}"), "")
labelNoKRS.comment = f"KRS: tail no-KRS path for {tc1} (sgprKRingShift==0)"
labelDoneKRS.comment = f"KRS: tail KRS branch join for {tc1}"
module.add(SCmpEQU32(src0=sgpr("KRingShift"), src1=0, comment="KRS: sgprKRingShift==0 ?"))
module.add(SCBranchSCC1(labelName=labelNoKRS.getLabelName(), comment="KRS: take no-KRS tail loads"))
module.add(self.globalReadDo(kernel, globalReadMode1st, tensorParameters1st, krTailForceDisable=False))
module.add(SBranch(labelName=labelDoneKRS.getLabelName(), comment="KRS: skip no-KRS tail loads"))
module.add(labelNoKRS)
module.add(self.globalReadDo(kernel, globalReadMode1st, tensorParameters1st, krTailForceDisable=True))
module.add(labelDoneKRS)
module.addComment1("Tail global read MXS%s"%tc1)
if tailLoopOpt1st and (globalReadMode1st == 2):
module.add(self.doTailLoopOpt(kernel, tensorParameters1st["MX"]))
else:
module.add(self.globalReadDo(kernel, globalReadMode1st, tensorParameters1st))

#TODO: To handle KRS with MX
if "MX" in tensorParameters1st:
module.addComment1("Update M0 for DTLDS")
moduleTmp = self.directToLdsM0Update(kernel, 1, tensorParameters1st["MX"], skipWait=skip2ndWaitForDtl)
module.add(replaceHolder(moduleTmp, 0))
module.addComment1("Tail global read MXS%s"%tc1)
if tailLoopOpt1st and (globalReadMode1st == 2):
module.add(self.doTailLoopOpt(kernel, tensorParameters1st["MX"]))
else:
module.add(self.globalReadDo(kernel, globalReadMode1st, tensorParameters1st["MX"]))

if "MX" in tensorParameters2nd:
module.addComment1("Update M0 for DTLDS")
moduleTmp = self.directToLdsM0Update(kernel, 1, tensorParameters2nd["MX"], skipWait=skip2ndWaitForDtl)
module.add(replaceHolder(moduleTmp, 0))
module.addComment1("Tail global read MXS%s"%tc2)
if tailLoopOpt2nd and (globalReadMode2nd == 2):
module.add(self.doTailLoopOpt(kernel, tensorParameters2nd["MX"]))
else:
module.add(self.globalReadDo(kernel, globalReadMode2nd, tensorParameters2nd["MX"]))
module.add(self.globalReadDo(kernel, globalReadMode1st, tensorParameters1st["MX"]))

if "MX" in tensorParameters2nd:
module.addComment1("Update M0 for DTLDS")
moduleTmp = self.directToLdsM0Update(kernel, 2, tensorParameters2nd, skip2ndWaitForDtl)
moduleTmp = self.directToLdsM0Update(kernel, 1, tensorParameters2nd["MX"], skipWait=skip2ndWaitForDtl)
module.add(replaceHolder(moduleTmp, 0))
module.addComment1("Tail global read %s"%tc2)
module.addComment1("Tail global read MXS%s"%tc2)
if tailLoopOpt2nd and (globalReadMode2nd == 2):
module.add(self.doTailLoopOpt(kernel, tensorParameters2nd))
module.add(self.doTailLoopOpt(kernel, tensorParameters2nd["MX"]))
else:
if kernel["KRingShift"] and kernel["BufferLoad"] and tc2 in ("A", "B"):
labelNoKRS = Label(self.labels.getNameInc(f"KRS_tail_noop_{tc2}"), "")
labelDoneKRS = Label(self.labels.getNameInc(f"KRS_tail_done_{tc2}"), "")
labelNoKRS.comment = f"KRS: tail no-KRS path for {tc2} (sgprKRingShift==0)"
labelDoneKRS.comment = f"KRS: tail KRS branch join for {tc2}"
module.add(SCmpEQU32(src0=sgpr("KRingShift"), src1=0, comment="KRS: sgprKRingShift==0 ?"))
module.add(SCBranchSCC1(labelName=labelNoKRS.getLabelName(), comment="KRS: take no-KRS tail loads"))
module.add(self.globalReadDo(kernel, globalReadMode2nd, tensorParameters2nd, krTailForceDisable=False))
module.add(SBranch(labelName=labelDoneKRS.getLabelName(), comment="KRS: skip no-KRS tail loads"))
module.add(labelNoKRS)
module.add(self.globalReadDo(kernel, globalReadMode2nd, tensorParameters2nd, krTailForceDisable=True))
module.add(labelDoneKRS)
else:
module.add(self.globalReadDo(kernel, globalReadMode2nd, tensorParameters2nd))
module.add(self.globalReadDo(kernel, globalReadMode2nd, tensorParameters2nd["MX"]))

module.addComment1("Update M0 for DTLDS")
moduleTmp = self.directToLdsM0Update(kernel, 2, tensorParameters2nd, skip2ndWaitForDtl)
module.add(replaceHolder(moduleTmp, 0))
module.addComment1("Tail global read %s"%tc2)
if tailLoopOpt2nd and (globalReadMode2nd == 2):
module.add(self.doTailLoopOpt(kernel, tensorParameters2nd))
else:
module.add(self.globalReadDo(kernel, globalReadMode2nd, tensorParameters2nd))

doA = False
doB = False
Expand Down Expand Up @@ -8415,12 +8339,6 @@ def vgprAllocationImplSubtile():
if kernel["StreamK"] and kernel["StreamKAtomic"] == 0:
self.addSgprVarToPool("SrdWS")

if kernel["BAddrInterleave"]:
self.defineSgpr("BInterleaveG", 1)
self.addSgprVarToPool("BInterleaveG")
if kernel["KRingShift"]:
self.defineSgpr("KRingShift", 1)
self.addSgprVarToPool("KRingShift")
# gfx1250 frees the SK constant SGPRs later in moveStreamKConstantsToVgpr
# after their values have been copied to VGPRs. Freeing them here would let
# temp allocs clobber kernel arguments before they are copied.
Expand Down
Loading