Skip to content

Fix MIArchVgpr=1 + LSU=2 + SK and fix grouped gemm solution selection #345

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
AlexBrownAMD merged 3 commits into from
Jun 24, 2025
Merged

Fix MIArchVgpr=1 + LSU=2 + SK and fix grouped gemm solution selection #345

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
99e5772
Fix MIArchVgpr=1 + LSU=2 + SK and fix grouped gemm solution selection
AlexBrownAMD Jun 23, 2025
c2f6613
Limit launch grid by supplied workspace size
AlexBrownAMD Jun 23, 2025
2c4bf23
Merge branch 'develop' into FixBugs
AlexBrownAMD Jun 23, 2025
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
2 changes: 1 addition & 1 deletion projects/hipblaslt/tensilelite/Tensile/Components/StreamK.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1388,7 +1388,7 @@ def fixupBatch(self, writer, kernel, ss, batchIdx, edge, gwvw, \
# AccVgpr write
# if kernel.enabledSetPrioSplitLDS:
# kStr += inst("s_setprio", "0", "")
if codeAccVgprWrite is not None:
if codeAccVgprWrite is not None and kernel["LocalSplitU"] == 1:
regsPerScalar = writer.states.bpeCinternal // writer.states.bpr # register per scalar
# loop over store instructions within one batch
for elementIdx in range(0, len(batchElements)):
Expand Down
14 changes: 10 additions & 4 deletions projects/hipblaslt/tensilelite/Tensile/Source/lib/source/ContractionSolution.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2875,10 +2875,13 @@ namespace TensileLite
}
const bool streamKDP = Debug::Instance().useStreamKDataParrallel();
auto tiles = problem.getNumTiles(sizeMapping, gsu);
size_t skGrid = getSKGrid(problem, hardware, tiles);
// Get space required for partial tiles
if(tiles % skGrid != 0 && !streamKDP)
size += partialTileSize(skGrid);
if(tiles > 0) // Grouped GEMM reports 0 tiles
{
size_t skGrid = getSKGrid(problem, hardware, tiles);
// Get space required for partial tiles
if(skGrid > 0 && tiles % skGrid != 0 && !streamKDP)
size += partialTileSize(skGrid);
}
}
else
{
Expand Down Expand Up @@ -3189,6 +3192,9 @@ namespace TensileLite
}
}
}

if (tiles % skGrid != 0 && partialTileSize(skGrid) > problem.workspaceSize())
skGrid = tiles;
return skGrid;
}
// Limit the CUs Stream-K is launched on either max or the specified,
Expand Down
Loading