Skip to content

Revert "fix(tensilelite): use the 1024-byte LDS granularity on gfx11 … #11672

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
tony-davis merged 1 commit into from
Sep 3, 2026
Merged

Revert "fix(tensilelite): use the 1024-byte LDS granularity on gfx11 … #11672

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
15 changes: 4 additions & 11 deletions projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -245,8 +245,7 @@ def getVgprOccupancy(self, numThreads, vgprs, doubleVgpr=False):
def getOccupancy(self, numThreads, vgprs, sgprs, ldsSize, accvgprs=0, doubleVgpr=False):

deviceLdsSize = self.states.archCaps["DeviceLDS"]
ldsLimitedOccupancy = self.getLdsLimitedOccupancy(
deviceLdsSize, ldsSize, self.states.archCaps["LdsGranularity"])
ldsLimitedOccupancy = self.getLdsLimitedOccupancy(deviceLdsSize, ldsSize)

if not doubleVgpr:
vgprLimitedOccupancy = self.getVgprOccupancy(numThreads, vgprs, doubleVgpr)
Expand Down Expand Up @@ -280,20 +279,14 @@ def getVgpr(vgpr, doubleVgpr):
return lastVgprs, initOccupancy

@staticmethod
def getLdsLimitedOccupancy(deviceLdsSize, ldsSize, granularity):
"""Max. number of workgroups that can run on the same CU due to LDS size

granularity is the archCaps["LdsGranularity"] allocation granule: a workgroup
occupies a whole number of granules, so the rounding can cost a workgroup at a
boundary.
"""
def getLdsLimitedOccupancy(deviceLdsSize, ldsSize):
if ldsSize == 0:
# No LDS usage: LDS is not the binding constraint.
# Return a large sentinel so other limits (VGPR, wave cap) win in min().
return deviceLdsSize // granularity
return deviceLdsSize // 256
# As ldsSize gets large, rounding might push us slightly higher than deviceLdsSize.
# Clamp at deviceLdsSize
ldsSize = min(ldsSize + granularity - 1, deviceLdsSize) & ~(granularity - 1)
ldsSize = min(ldsSize + 255, deviceLdsSize) & 0xffffff00 # 256-byte granularity

ldsLimitedOccupancy = deviceLdsSize//ldsSize
return ldsLimitedOccupancy
Expand Down
4 changes: 2 additions & 2 deletions projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_occupancy.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -210,7 +210,7 @@ def test_gfx11_1024_vgpr_parts(isa):
)
def test_lds_limited_occupancy_gfx950(lds_bytes, expected_occ):
"""getLdsLimitedOccupancy on gfx950 with 163840-byte device LDS."""
occ = KernelWriterAssembly.getLdsLimitedOccupancy(163840, lds_bytes, 256)
occ = KernelWriterAssembly.getLdsLimitedOccupancy(163840, lds_bytes)
assert occ == expected_occ, (
f"LDS={lds_bytes} B: expected {expected_occ} blocks/CU, got {occ}"
)
Expand Down Expand Up @@ -339,7 +339,7 @@ def test_lds_limited_occupancy_matches_hip_oracle(desc, numRegs, staticLDS,
2 — matching HIP — but the code-gen VGPR pool over-estimate produces stored occ=1.
"""
device_lds = 163840 # gfx950 160 KB
lds_occ = KernelWriterAssembly.getLdsLimitedOccupancy(device_lds, staticLDS, 256)
lds_occ = KernelWriterAssembly.getLdsLimitedOccupancy(device_lds, staticLDS)

if staticLDS > device_lds // 2:
# LDS-limited: the LDS formula alone reproduces the HIP result.
Expand Down
3 changes: 0 additions & 3 deletions projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_occupancy_buildtime.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -442,9 +442,6 @@ def _make_max_vgpr_writer(arch_acc_unified=True, vgpr_size=256, agpr_size=256, s
"ArchAccUnifiedRegs": int(arch_acc_unified),
"MaxWavesPerSimd": 8 if arch_acc_unified else 10,
"DeviceLDS": 163840,
# Mirrors the real cap: gfx11 (the non-unified-AccVgpr case here)
# allocates LDS in 1024-byte granules, CDNA in 256-byte ones.
"LdsGranularity": 256 if arch_acc_unified else 1024,
},
regCaps={"MaxVgpr": 256, "PhysicalMaxSgpr": 800},
doubleVgpr=arch_acc_unified,
Expand Down
5 changes: 0 additions & 5 deletions projects/hipblaslt/tensilelite/rocisa/rocisa/include/hardware_caps.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -564,11 +564,6 @@ inline std::map<std::string, int> initArchCaps(const IsaVersion& isaVersion)
deviceLDS = 327680;

rv["DeviceLDS"] = deviceLDS;
// LDS allocation granule: a workgroup occupies a whole number of granules, so this
// rounding can cost a resident workgroup at a boundary. gfx11 allocates 1024 bytes
// at a time, other targets 256. TODO: gfx10/gfx12 might also need a different
// granularity.
rv["LdsGranularity"] = isaVersion[0] == 11 ? 1024 : 256;
rv["CMPXWritesSGPR"] = checkNotInList(isaVersion[0], {10, 11, 12});
rv["HasWave32"] = checkInList(isaVersion[0], {10, 11, 12});
rv["HasSchedMode"] = checkInList(isaVersion[0], {12});
Expand Down
Loading