[SYCL][SPIR-V] Don't use 'print' to get JointMatrix typename (#7054)

Author: MrSidims

clang/lib/CodeGen/CodeGenTypes.cpp

@@ -80,19 +80,30 @@ void CodeGenTypes::addRecordTypeName(const RecordDecl *RD,
                 OS << "i" << TTy->getIntegerBitWidth();
                 break;
               }
-            } else if (TTy->isBFloatTy())
+            } else if (TTy->isHalfTy()) {
+              OS << "half";
+            } else if (TTy->isFloatTy()) {
+              OS << "float";
+            } else if (TTy->isDoubleTy()) {
+              OS << "double";
+            } else if (TTy->isBFloatTy()) {
               OS << "bfloat16";
-            else if (TTy->isStructTy()) {
+            } else if (TTy->isStructTy()) {
               StringRef LlvmTyName = TTy->getStructName();
-              // Emit half/bfloat16 for sycl[::*]::{half,bfloat16}
+              // Emit half/bfloat16/tf32 for sycl[::*]::{half,bfloat16,tf32}
               if (LlvmTyName.startswith("class.sycl::") ||
                   LlvmTyName.startswith("class.__sycl_internal::"))
                 LlvmTyName = LlvmTyName.rsplit("::").second;
+              if (LlvmTyName != "half" && LlvmTyName != "bfloat16" &&
+                  LlvmTyName != "tf32")
+                llvm_unreachable("Wrong matrix base type!");
               OS << LlvmTyName;
-            } else
-              TTy->print(OS, false, true);
-          } else if (TemplateArg.getKind() == TemplateArgument::Integral)
+            } else {
+              llvm_unreachable("Wrong matrix base type!");
+            }
+          } else if (TemplateArg.getKind() == TemplateArgument::Integral) {
             OS << TemplateArg.getAsIntegral();
+          }
         }
         Ty->setName(OS.str());
         return;

clang/test/CodeGenSYCL/matrix.cpp

@@ -21,6 +21,7 @@ void f3(__spv::__spirv_JointMatrixINTEL<char, 10, 2, 0, 0> *matrix) {}
 namespace sycl {
   class half {};
   class bfloat16 {};
+  class tf32 {};
 }
 typedef sycl::half my_half;
 
@@ -32,3 +33,9 @@ void f5(__spv::__spirv_JointMatrixINTEL<sycl::bfloat16, 10, 2, 0, 0> *matrix) {}
 
 // CHECK: @_Z2f6{{.*}}(%spirv.JointMatrixINTEL._i128_10_2_0_0
 void f6(__spv::__spirv_JointMatrixINTEL<_BitInt(128), 10, 2, 0, 0> *matrix) {}
+
+// CHECK: @_Z2f7{{.*}}(%spirv.JointMatrixINTEL._tf32_10_2_0_0
+void f7(__spv::__spirv_JointMatrixINTEL<sycl::tf32, 10, 2, 0, 0> *matrix) {}
+
+// CHECK: @_Z2f8{{.*}}(%spirv.JointMatrixINTEL._double_5_10_0_1
+void f8(__spv::__spirv_JointMatrixINTEL<double, 5, 10, 0, 1> *matrix) {}