call amd_mixed_dot in ROCm 2.2
@@ -167,8 +167,7 @@ def writeSolutionsAndKernels(outputPath, problemTypes, solutions, kernels, kerne
kernelHeaderFile.write("#include \"KernelHeader.h\"\n")
kernelHeaderFile.write("\n\n")
kernelHeaderFile.write("__device__ inline int GenDot4(int a, int b, int c) { \n")
- kernelHeaderFile.write(" typedef struct { int c0:8,c1:8,c2:8,c3:8; } C4I8;\n")
- kernelHeaderFile.write(" typedef union { int32_t i; C4I8 z; } PkInt8x4;\n")
+ kernelHeaderFile.write(" typedef union { int32_t i; char4 z; } PkInt8x4;\n")
kernelHeaderFile.write(" PkInt8x4 va, vb; va.i = a; vb.i = b;\n")
kernelHeaderFile.write("//if (__oclc_ISA_version == 906)\n")
@@ -177,7 +176,7 @@ def writeSolutionsAndKernels(outputPath, problemTypes, solutions, kernels, kerne
kernelHeaderFile.write("//}\n")
kernelHeaderFile.write("//else\n")
kernelHeaderFile.write("//{\n")
- kernelHeaderFile.write(" return c + (vb.z.c3*va.z.c3 + vb.z.c2*va.z.c2 + vb.z.c1*va.z.c1 + vb.z.c0*va.z.c0); }\n")
+ kernelHeaderFile.write(" return amd_mixed_dot(va.z, vb.z, c, true); }\n")
kernelHeaderFile.write("//}\n")
kernelHeaderFile.write("\n\n")
else:
