[SYCL][SPIR-V] Drop Unnecessary Matrix Use parameter

Last parameter 'Use' is optional in SPIR-V, but is not optional
in DPCPP headers. If it has Unnecessary value - skip it

Signed-off-by: Sidorov, Dmitry <dmitry.sidorov@intel.com>

Author: MrSidims

clang/lib/CodeGen/CodeGenTypes.cpp

@@ -57,11 +57,14 @@ void CodeGenTypes::addRecordTypeName(const RecordDecl *RD,
       if (auto TemplateDecl = dyn_cast<ClassTemplateSpecializationDecl>(RD)) {
         ArrayRef<TemplateArgument> TemplateArgs =
             TemplateDecl->getTemplateArgs().asArray();
+        constexpr size_t MaxMatrixParameter = 6;
+        assert(TemplateArgs.size() <= MaxMatrixParameter &&
+               "Too many template parameters for JointMatrixINTEL type");
         OS << "spirv.JointMatrixINTEL.";
-        for (auto &TemplateArg : TemplateArgs) {
-          OS << "_";
-          if (TemplateArg.getKind() == TemplateArgument::Type) {
-            llvm::Type *TTy = ConvertType(TemplateArg.getAsType());
+        for (size_t I = 0; I != TemplateArgs.size(); ++I) {
+          if (TemplateArgs[I].getKind() == TemplateArgument::Type) {
+            OS << "_";
+            llvm::Type *TTy = ConvertType(TemplateArgs[I].getAsType());
             if (TTy->isIntegerTy()) {
               switch (TTy->getIntegerBitWidth()) {
               case 8:
@@ -91,8 +94,14 @@ void CodeGenTypes::addRecordTypeName(const RecordDecl *RD,
               OS << LlvmTyName;
             } else
               TTy->print(OS, false, true);
-          } else if (TemplateArg.getKind() == TemplateArgument::Integral)
-            OS << TemplateArg.getAsIntegral();
+          } else if (TemplateArgs[I].getKind() == TemplateArgument::Integral) {
+            const auto IntTemplateParam = TemplateArgs[I].getAsIntegral();
+            // Last parameter 'Use' is optional in SPIR-V, but is not optional
+            // in DPCPP headers. If it has Unnecessary value - skip it
+            constexpr size_t Unnecessary = 3;
+            if (!(I == MaxMatrixParameter && IntTemplateParam == Unnecessary))
+              OS << "_" << IntTemplateParam;
+          }
         }
         Ty->setName(OS.str());
         return;

clang/test/CodeGenSYCL/matrix.cpp

@@ -5,18 +5,18 @@
 #include <stdint.h>
 
 namespace __spv {
-  template <typename T, size_t R, size_t C, uint32_t U, uint32_t S>
+  template <typename T, size_t R, size_t C, uint32_t L, uint32_t S, uint32_t U>
   struct __spirv_JointMatrixINTEL;
 }
 
 // CHECK: @_Z2f1{{.*}}(%spirv.JointMatrixINTEL._float_5_10_0_1
-void f1(__spv::__spirv_JointMatrixINTEL<float, 5, 10, 0, 1> *matrix) {}
+void f1(__spv::__spirv_JointMatrixINTEL<float, 5, 10, 0, 1, 3> *matrix) {}
 
 // CHECK: @_Z2f2{{.*}}(%spirv.JointMatrixINTEL._long_10_2_0_0
-void f2(__spv::__spirv_JointMatrixINTEL<uint64_t, 10, 2, 0, 0> *matrix) {}
+void f2(__spv::__spirv_JointMatrixINTEL<uint64_t, 10, 2, 0, 0, 3> *matrix) {}
 
 // CHECK: @_Z2f3{{.*}}(%spirv.JointMatrixINTEL._char_10_2_0_0
-void f3(__spv::__spirv_JointMatrixINTEL<char, 10, 2, 0, 0> *matrix) {}
+void f3(__spv::__spirv_JointMatrixINTEL<char, 10, 2, 0, 0, 3> *matrix) {}
 
 namespace sycl {
   class half {};
@@ -25,10 +25,13 @@ namespace sycl {
 typedef sycl::half my_half;
 
 // CHECK: @_Z2f4{{.*}}(%spirv.JointMatrixINTEL._half_10_2_0_0
-void f4(__spv::__spirv_JointMatrixINTEL<my_half, 10, 2, 0, 0> *matrix) {}
+void f4(__spv::__spirv_JointMatrixINTEL<my_half, 10, 2, 0, 0, 3> *matrix) {}
 
 // CHECK: @_Z2f5{{.*}}(%spirv.JointMatrixINTEL._bfloat16_10_2_0_0
-void f5(__spv::__spirv_JointMatrixINTEL<sycl::bfloat16, 10, 2, 0, 0> *matrix) {}
+void f5(__spv::__spirv_JointMatrixINTEL<sycl::bfloat16, 10, 2, 0, 0, 3> *matrix) {}
 
 // CHECK: @_Z2f6{{.*}}(%spirv.JointMatrixINTEL._i128_10_2_0_0
-void f6(__spv::__spirv_JointMatrixINTEL<_BitInt(128), 10, 2, 0, 0> *matrix) {}
+void f6(__spv::__spirv_JointMatrixINTEL<_BitInt(128), 10, 2, 0, 0, 3> *matrix) {}
+
+// CHECK: @_Z2f7{{.*}}(%spirv.JointMatrixINTEL._char_10_2_0_0_1
+void f7(__spv::__spirv_JointMatrixINTEL<char, 10, 2, 0, 0, 1> *matrix) {}