Calculate number off accgrs required for requested instruction.
Generate requested type of instruction, need to fix number of instructions and vgpr addresses.
@@ -1515,7 +1515,7 @@ class KernelWriterAssembly(KernelWriter):
########################################
self.totalAgprs = 0
if "MatrixInstM" in kernel:
- self.totalAgprs = 32
+ self.totalAgprs = kernel["MatrixInstM"] * kernel["MatrixInstN"] * kernel["MatrixInstB"] // 64
########################################
# Register Pools
@@ -1878,6 +1878,15 @@ class KernelWriterAssembly(KernelWriter):
# single precision
elif kernel["ProblemType"]["DataType"].isSingle():
+ if self.version == (9,0,8):
+ a = 0 #todo fix vgpr addresses for a and b
+ b = 0 #todo
+ iui = 0 #todo
+ cStr = "v[%s+%u+%u*%u]" % ("vgprValuC", a, b, kernel["ThreadTile0"])
+ 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%uf32 a[0:%u], %s, %s, a[0:%u]%s" % (kernel["MatrixInstM"], kernel["MatrixInstN"], kernel["MatrixInstK"], self.totalAgprs - 1, aStr, bStr, self.totalAgprs - 1, self.endLine)
+ else:
for b in range(0, kernel["ThreadTile1"]):
for a in range(0, kernel["ThreadTile0"]):
for iui in range(0, innerUnroll):
