change mfmaIter from kStr to Code.Module
@@ -5734,25 +5734,26 @@ class KernelWriterAssembly(KernelWriter):
##############################################################################
def mfmaIter(self, kernel, m, innerUnroll):
- # TODO: should return Code.Module or string?
- kStr = ""
+ imod = Code.Module("mfmaCode")
if kernel["ProblemType"]["DataType"].isBFloat16():
- kStr += "s_nop 1\n" # a must; for VALU packing writes to be consumed by matrix instruction
+ # a must; for VALU packing writes to be consumed by matrix instruction
# this can be removed once the vregs are directly loaded via ds_read
- for b in range(0, self.numColInsts):
- for a in range(0, self.numRowInsts):
+ imod.addInst("s_nop ","1","VALU packing writes to be consumed by matrix instruction")
+
for iui in range(0, innerUnroll):
+ for a in range(0, self.numRowInsts):
+ for b in range(0, self.numColInsts):
accIdx = b * self.numRowInsts + a
accStart = accIdx * self.destAgprs
accEnd = accStart + self.destAgprs - 1
aStr = "v[%s+%u]" % ("vgprValuA_X%u_I%u" % (m, iui), a)
bStr = "v[%s+%u]" % ("vgprValuB_X%u_I%u" % (m, iui), b)
- kStr += "v_mfma_f32_%ux%ux%u%s a[%u:%u], %s, %s, a[%u:%u]%s" \
- % (kernel["MatrixInstM"], kernel["MatrixInstN"], kernel["MatrixInstK"], kernel["ProblemType"]["DataType"].toNameAbbrev(), accStart, accEnd, aStr, bStr, accStart, accEnd, self.endLine)
+ imod.addText("v_mfma_f32_%ux%ux%u%s a[%u:%u], %s, %s, a[%u:%u]%s" \
+ % (kernel["MatrixInstM"], kernel["MatrixInstN"], kernel["MatrixInstK"], kernel["ProblemType"]["DataType"].toNameAbbrev(), accStart, accEnd, aStr, bStr, accStart, accEnd, self.endLine))
# TODO Broadcast code
#kStr += "v_mfma_f32_%ux%ux%uf32 a[%u:%u], %s, %s, a[%u:%u] blgp:%u%s" \
# % (kernel["MatrixInstM"], kernel["MatrixInstN"], kernel["MatrixInstK"], accStart, accEnd, bStr, aStr, accStart, accEnd, a, self.endLine)
- return kStr
+ return imod
##############################################################################
# MAC Iteration
