-
Notifications
You must be signed in to change notification settings - Fork 137
Enable mbsk spmm #2119
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.
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
Enable mbsk spmm #2119
Conversation
vin-huang
commented
May 24, 2025
- Added requiredSynchornizerSizes()
- Enable MBSK for SPMM
- Add MBSK test for SPMM
d3cb107
to
94935bb
Compare
The Azure CI was failed due to cannot find proper device. Hip error: 'no ROCm-capable device is detected'(100) at /agent/_work/1/s/library/src/amd_detail/hipblaslt.cpp:171 |
90eaf6b
to
c1c8201
Compare
…r size for the MBSK kernel.
* Add refineOccupancy to set the proper occupancy limit. * Allocated the required SGPRs needed by storeRemapAddStore() before calculating occupancy.
c1c8201
to
c189be1
Compare
@@ -867,14 +867,19 @@ def GSUSynccodegen(self, writer, kernel, tmpVgpr, tmpVgprSize, tmpVgprDynamic, b | |||
vmcnt = SyncloadedData = SyncloadedData -1 | |||
module.add(SWaitCnt(lgkmcnt=lgkmcnt, vmcnt=vmcnt, vscnt=vscnt, comment="(wait for buffer ready)")) | |||
|
|||
if ((gwvw % 2) == 1): | |||
if kernel["ProblemType"]["DataType"].isInt8() or kernel["ProblemType"]["DataType"].isInt32(): |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
should we check kernel["ProblemType"]["ComputeDataType"].isInt32() ?
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
In the original case, which didn't consider the output datatype is int32 from mfma/smfma instruction.
So, here I use kernel["ProblemType"]["DataType"] to identify the output datatype.