From b9529a66b7b6cd337575c83b89071ac8872f76e6 Mon Sep 17 00:00:00 2001 From: "Sidorov, Dmitry" Date: Tue, 14 Feb 2023 07:47:29 -0800 Subject: [PATCH 1/5] [SYCL] Represent JointMatrixINTEL type as extension type Signed-off-by: Sidorov, Dmitry --- clang/lib/CodeGen/CodeGenTypes.cpp | 144 +++++++++++++++++------------ clang/lib/CodeGen/CodeGenTypes.h | 8 ++ clang/test/CodeGenSYCL/matrix.cpp | 35 +++---- 3 files changed, 111 insertions(+), 76 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenTypes.cpp b/clang/lib/CodeGen/CodeGenTypes.cpp index bf8e4006e1d69..ff0e560f62050 100644 --- a/clang/lib/CodeGen/CodeGenTypes.cpp +++ b/clang/lib/CodeGen/CodeGenTypes.cpp @@ -51,65 +51,6 @@ void CodeGenTypes::addRecordTypeName(const RecordDecl *RD, StringRef suffix) { SmallString<256> TypeName; llvm::raw_svector_ostream OS(TypeName); - // If RD is spirv_JointMatrixINTEL type, mangle differently. - if (CGM.getTriple().isSPIRV() || CGM.getTriple().isSPIR()) { - if (RD->getQualifiedNameAsString() == "__spv::__spirv_JointMatrixINTEL") { - if (auto TemplateDecl = dyn_cast(RD)) { - ArrayRef TemplateArgs = - TemplateDecl->getTemplateArgs().asArray(); - OS << "spirv.JointMatrixINTEL."; - for (auto &TemplateArg : TemplateArgs) { - OS << "_"; - if (TemplateArg.getKind() == TemplateArgument::Type) { - llvm::Type *TTy = ConvertType(TemplateArg.getAsType()); - if (TTy->isIntegerTy()) { - switch (TTy->getIntegerBitWidth()) { - case 8: - OS << "char"; - break; - case 16: - OS << "short"; - break; - case 32: - OS << "int"; - break; - case 64: - OS << "long"; - break; - default: - OS << "i" << TTy->getIntegerBitWidth(); - break; - } - } 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()) { - StringRef LlvmTyName = TTy->getStructName(); - // 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 { - llvm_unreachable("Wrong matrix base type!"); - } - } else if (TemplateArg.getKind() == TemplateArgument::Integral) { - OS << TemplateArg.getAsIntegral(); - } - } - Ty->setName(OS.str()); - return; - } - } - } OS << RD->getKindName() << '.'; // FIXME: We probably want to make more tweaks to the printing policy. For @@ -460,6 +401,79 @@ llvm::Type *CodeGenTypes::ConvertFunctionTypeInternal(QualType QFT) { return ResultType; } +template +llvm::Type* getJointMatrixINTELExtType(llvm::Type *CompTy, + ArrayRef TemplateArgs, + const unsigned Val = 0) { + // TODO: we should actually have exactly 5 template parameters: 1 for + // type and 4 for type parameters. But in previous version of the SPIR-V + // spec we have Layout matrix type parameter, that was later removed. + // Once we update to the newest version of the spec - this should be updated. + assert((TemplateArgs.size() == 5 || TemplateArgs.size() == 6) && + "Wrong JointMatrixINTEL template parameters number"); + // This is required to represent optional Optional + // 'Component Type Interpretation' parameter + using ParamsType = + typename std::conditional, + SmallVector>::type; + ParamsType Params; + if constexpr (NeedTypeInterpret) + Params = { 0, 0, 0, 0, 0, Val}; + else + Params = { 0, 0, 0, 0, 0}; + for (size_t I = 1; I != TemplateArgs.size(); ++I) { + assert(TemplateArgs[I].getKind() == + TemplateArgument::Integral && + "Wrong JointMatrixINTEL template parameter"); + Params[I - 1] = TemplateArgs[I].getAsIntegral().getExtValue(); + } + return llvm::TargetExtType::get( + CompTy->getContext(), "spirv.JointMatrixINTEL", {CompTy}, Params); +} + +/// ConvertSYCLJointMatrixINTELType - Convert SYCL joint_matrix type +/// which is represented as a pointer to a structure to LLVM extension type +/// with the parameters that follow SPIR-V JointMatrixINTEL type. +/// The expected representation is: +/// target("spirv.JointMatrixINTEL", %element_type, %rows%, %cols%, %scope%, +/// %use%, (optional) %element_type_interpretation%) +llvm::Type *CodeGenTypes::ConvertSYCLJointMatrixINTELType(RecordDecl *RD) { + auto *TemplateDecl = cast(RD); + ArrayRef TemplateArgs = + TemplateDecl->getTemplateArgs().asArray(); + assert(TemplateArgs[0].getKind() == TemplateArgument::Type && + "1st JointMatrixINTEL template parameter must be type"); + llvm::Type *CompTy = ConvertType(TemplateArgs[0].getAsType()); + + // Per JointMatrixINTEL spec the type can have an Optional + // 'Component Type Interpretation' parameter. We should emit it in case + // if on SYCL level joint matrix accepts 'bfloat16' or 'tf32' objects as + // matrix's components. Yet bfloat16 should be represented as 'int16' and + // 'tf32' as 'float' types. + if (CompTy->isStructTy()) { + StringRef LlvmTyName = CompTy->getStructName(); + // Emit half/int16/float for sycl[::*]::{half,bfloat16,tf32} + if (LlvmTyName.startswith("class.sycl::") || + LlvmTyName.startswith("class.__sycl_internal::")) + LlvmTyName = LlvmTyName.rsplit("::").second; + if (LlvmTyName == "half") { + CompTy = llvm::Type::getHalfTy(getLLVMContext()); + return getJointMatrixINTELExtType(CompTy, TemplateArgs); + } else if (LlvmTyName == "tf32") { + CompTy = llvm::Type::getHalfTy(getLLVMContext()); + // 'tf32' interpretation is mapped to '0' + return getJointMatrixINTELExtType(CompTy, TemplateArgs, 0); + } else if (LlvmTyName == "bfloat16") { + CompTy = llvm::Type::getInt16Ty(getLLVMContext()); + // 'bfloat16' interpretation is mapped to '1' + return getJointMatrixINTELExtType(CompTy, TemplateArgs, 1); + } else { + llvm_unreachable("Wrong matrix base type!"); + } + } + return getJointMatrixINTELExtType(CompTy, TemplateArgs); +} + /// ConvertType - Convert the specified type to its LLVM form. llvm::Type *CodeGenTypes::ConvertType(QualType T) { T = Context.getCanonicalType(T); @@ -745,6 +759,18 @@ llvm::Type *CodeGenTypes::ConvertType(QualType T) { llvm::Type *PointeeType = ConvertTypeForMem(ETy); if (PointeeType->isVoidTy()) PointeeType = llvm::Type::getInt8Ty(getLLVMContext()); + if (CGM.getTriple().isSPIRV() || CGM.getTriple().isSPIR()) { + const Type* ClangETy = ETy.getTypePtrOrNull(); + if (ClangETy && ClangETy->isStructureOrClassType()) { + RecordDecl *RD = ClangETy->getAsCXXRecordDecl(); + if (RD->getQualifiedNameAsString() == + "__spv::__spirv_JointMatrixINTEL") { + ResultType = ConvertSYCLJointMatrixINTELType(RD); + break; + } + } + } + unsigned AS = getTargetAddressSpace(ETy); ResultType = llvm::PointerType::get(PointeeType, AS); break; diff --git a/clang/lib/CodeGen/CodeGenTypes.h b/clang/lib/CodeGen/CodeGenTypes.h index e76fda95513f6..3f198b2a3de1a 100644 --- a/clang/lib/CodeGen/CodeGenTypes.h +++ b/clang/lib/CodeGen/CodeGenTypes.h @@ -133,6 +133,14 @@ class CodeGenTypes { /// memory representation is usually i8 or i32, depending on the target. llvm::Type *ConvertTypeForMem(QualType T, bool ForBitField = false); + /// ConvertSYCLJointMatrixINTELType - Convert SYCL joint_matrix type + /// which is represented as a pointer to a structure to LLVM extension type + /// with the parameters that follow SPIR-V JointMatrixINTEL type. + /// The expected representation is: + /// target("spirv.JointMatrixINTEL", %element_type, %rows%, %cols%, %scope%, + /// %use%, (optional) %element_type_interpretation%) + llvm::Type *ConvertSYCLJointMatrixINTELType(RecordDecl *RD); + /// GetFunctionType - Get the LLVM function type for \arg Info. llvm::FunctionType *GetFunctionType(const CGFunctionInfo &Info); diff --git a/clang/test/CodeGenSYCL/matrix.cpp b/clang/test/CodeGenSYCL/matrix.cpp index 69469811047fd..ee70e7a909c0b 100644 --- a/clang/test/CodeGenSYCL/matrix.cpp +++ b/clang/test/CodeGenSYCL/matrix.cpp @@ -5,18 +5,19 @@ #include namespace __spv { - template + template struct __spirv_JointMatrixINTEL; } -// CHECK: @_Z2f1{{.*}}(%spirv.JointMatrixINTEL._float_5_10_0_1 -void f1(__spv::__spirv_JointMatrixINTEL *matrix) {} +target("spirv.JointMatrixINTEL", i16, 16, 16, 2, 3, 1, 1) +// CHECK: @_Z2f1{{.*}}(target("spirv.JointMatrixINTEL", float, 5, 10, 0, 1, 0) +void f1(__spv::__spirv_JointMatrixINTEL *matrix) {} -// CHECK: @_Z2f2{{.*}}(%spirv.JointMatrixINTEL._long_10_2_0_0 -void f2(__spv::__spirv_JointMatrixINTEL *matrix) {} +// CHECK: @_Z2f2{{.*}}(target("spirv.JointMatrixINTEL", i64, 10, 2, 0, 0, 0) +void f2(__spv::__spirv_JointMatrixINTEL *matrix) {} -// CHECK: @_Z2f3{{.*}}(%spirv.JointMatrixINTEL._char_10_2_0_0 -void f3(__spv::__spirv_JointMatrixINTEL *matrix) {} +// CHECK: @_Z2f3{{.*}}(target("spirv.JointMatrixINTEL", i8, 10, 2, 0, 0, 0) +void f3(__spv::__spirv_JointMatrixINTEL *matrix) {} namespace sycl { class half {}; @@ -25,17 +26,17 @@ namespace sycl { } typedef sycl::half my_half; -// CHECK: @_Z2f4{{.*}}(%spirv.JointMatrixINTEL._half_10_2_0_0 -void f4(__spv::__spirv_JointMatrixINTEL *matrix) {} +// CHECK: @_Z2f4{{.*}}(target("spirv.JointMatrixINTEL", half, 10, 2, 0, 0, 0) +void f4(__spv::__spirv_JointMatrixINTEL *matrix) {} -// CHECK: @_Z2f5{{.*}}(%spirv.JointMatrixINTEL._bfloat16_10_2_0_0 -void f5(__spv::__spirv_JointMatrixINTEL *matrix) {} +// CHECK: @_Z2f5{{.*}}(target("spirv.JointMatrixINTEL", i16, 10, 2, 0, 0, 0, 1) +void f5(__spv::__spirv_JointMatrixINTEL *matrix) {} -// CHECK: @_Z2f6{{.*}}(%spirv.JointMatrixINTEL._i128_10_2_0_0 -void f6(__spv::__spirv_JointMatrixINTEL<_BitInt(128), 10, 2, 0, 0> *matrix) {} +// CHECK: @_Z2f6{{.*}}(target("spirv.JointMatrixINTEL", i128, 10, 2, 0, 0, 0) +void f6(__spv::__spirv_JointMatrixINTEL<_BitInt(128), 10, 2, 0, 0, 0> *matrix) {} -// CHECK: @_Z2f7{{.*}}(%spirv.JointMatrixINTEL._tf32_10_2_0_0 -void f7(__spv::__spirv_JointMatrixINTEL *matrix) {} +// CHECK: @_Z2f7{{.*}}(target("spirv.JointMatrixINTEL", float, 10, 2, 0, 0, 0, 0) +void f7(__spv::__spirv_JointMatrixINTEL *matrix) {} -// CHECK: @_Z2f8{{.*}}(%spirv.JointMatrixINTEL._double_5_10_0_1 -void f8(__spv::__spirv_JointMatrixINTEL *matrix) {} +// CHECK: @_Z2f8{{.*}}(target("spirv.JointMatrixINTEL", double, 5, 10, 0, 1, 0) +void f8(__spv::__spirv_JointMatrixINTEL *matrix) {} From 366c3e38385402ea78da3f60c1a0a76dfc6a0a33 Mon Sep 17 00:00:00 2001 From: "Sidorov, Dmitry" Date: Fri, 31 Mar 2023 05:46:17 -0700 Subject: [PATCH 2/5] Apply clang-format Signed-off-by: Sidorov, Dmitry --- clang/lib/CodeGen/CodeGenTypes.cpp | 24 +++++++++++------------- 1 file changed, 11 insertions(+), 13 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenTypes.cpp b/clang/lib/CodeGen/CodeGenTypes.cpp index ff0e560f62050..ea9a39ffdf78d 100644 --- a/clang/lib/CodeGen/CodeGenTypes.cpp +++ b/clang/lib/CodeGen/CodeGenTypes.cpp @@ -402,7 +402,7 @@ llvm::Type *CodeGenTypes::ConvertFunctionTypeInternal(QualType QFT) { } template -llvm::Type* getJointMatrixINTELExtType(llvm::Type *CompTy, +llvm::Type *getJointMatrixINTELExtType(llvm::Type *CompTy, ArrayRef TemplateArgs, const unsigned Val = 0) { // TODO: we should actually have exactly 5 template parameters: 1 for @@ -418,17 +418,16 @@ llvm::Type* getJointMatrixINTELExtType(llvm::Type *CompTy, SmallVector>::type; ParamsType Params; if constexpr (NeedTypeInterpret) - Params = { 0, 0, 0, 0, 0, Val}; + Params = {0, 0, 0, 0, 0, Val}; else - Params = { 0, 0, 0, 0, 0}; + Params = {0, 0, 0, 0, 0}; for (size_t I = 1; I != TemplateArgs.size(); ++I) { - assert(TemplateArgs[I].getKind() == - TemplateArgument::Integral && + assert(TemplateArgs[I].getKind() == TemplateArgument::Integral && "Wrong JointMatrixINTEL template parameter"); Params[I - 1] = TemplateArgs[I].getAsIntegral().getExtValue(); } - return llvm::TargetExtType::get( - CompTy->getContext(), "spirv.JointMatrixINTEL", {CompTy}, Params); + return llvm::TargetExtType::get(CompTy->getContext(), + "spirv.JointMatrixINTEL", {CompTy}, Params); } /// ConvertSYCLJointMatrixINTELType - Convert SYCL joint_matrix type @@ -442,7 +441,7 @@ llvm::Type *CodeGenTypes::ConvertSYCLJointMatrixINTELType(RecordDecl *RD) { ArrayRef TemplateArgs = TemplateDecl->getTemplateArgs().asArray(); assert(TemplateArgs[0].getKind() == TemplateArgument::Type && - "1st JointMatrixINTEL template parameter must be type"); + "1st JointMatrixINTEL template parameter must be type"); llvm::Type *CompTy = ConvertType(TemplateArgs[0].getAsType()); // Per JointMatrixINTEL spec the type can have an Optional @@ -760,14 +759,13 @@ llvm::Type *CodeGenTypes::ConvertType(QualType T) { if (PointeeType->isVoidTy()) PointeeType = llvm::Type::getInt8Ty(getLLVMContext()); if (CGM.getTriple().isSPIRV() || CGM.getTriple().isSPIR()) { - const Type* ClangETy = ETy.getTypePtrOrNull(); + const Type *ClangETy = ETy.getTypePtrOrNull(); if (ClangETy && ClangETy->isStructureOrClassType()) { - RecordDecl *RD = ClangETy->getAsCXXRecordDecl(); - if (RD->getQualifiedNameAsString() == - "__spv::__spirv_JointMatrixINTEL") { + RecordDecl *RD = ClangETy->getAsCXXRecordDecl(); + if (RD->getQualifiedNameAsString() == "__spv::__spirv_JointMatrixINTEL") { ResultType = ConvertSYCLJointMatrixINTELType(RD); break; - } + } } } From 5d1630460ca190377ba47d1254d6f2146b908d1e Mon Sep 17 00:00:00 2001 From: "Sidorov, Dmitry" Date: Fri, 31 Mar 2023 08:39:57 -0700 Subject: [PATCH 3/5] Fix issue and test Signed-off-by: Sidorov, Dmitry --- clang/lib/CodeGen/CodeGenTypes.cpp | 2 +- clang/test/CodeGenSYCL/matrix.cpp | 1 - sycl/test/matrix/legacy/matrix-int8-test.cpp | 6 +++--- sycl/test/matrix/matrix-int8-test.cpp | 6 +++--- 4 files changed, 7 insertions(+), 8 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenTypes.cpp b/clang/lib/CodeGen/CodeGenTypes.cpp index ea9a39ffdf78d..cb8d48a9dddbe 100644 --- a/clang/lib/CodeGen/CodeGenTypes.cpp +++ b/clang/lib/CodeGen/CodeGenTypes.cpp @@ -459,7 +459,7 @@ llvm::Type *CodeGenTypes::ConvertSYCLJointMatrixINTELType(RecordDecl *RD) { CompTy = llvm::Type::getHalfTy(getLLVMContext()); return getJointMatrixINTELExtType(CompTy, TemplateArgs); } else if (LlvmTyName == "tf32") { - CompTy = llvm::Type::getHalfTy(getLLVMContext()); + CompTy = llvm::Type::getFloatTy(getLLVMContext()); // 'tf32' interpretation is mapped to '0' return getJointMatrixINTELExtType(CompTy, TemplateArgs, 0); } else if (LlvmTyName == "bfloat16") { diff --git a/clang/test/CodeGenSYCL/matrix.cpp b/clang/test/CodeGenSYCL/matrix.cpp index ee70e7a909c0b..b2c0c51adba6e 100644 --- a/clang/test/CodeGenSYCL/matrix.cpp +++ b/clang/test/CodeGenSYCL/matrix.cpp @@ -9,7 +9,6 @@ namespace __spv { struct __spirv_JointMatrixINTEL; } -target("spirv.JointMatrixINTEL", i16, 16, 16, 2, 3, 1, 1) // CHECK: @_Z2f1{{.*}}(target("spirv.JointMatrixINTEL", float, 5, 10, 0, 1, 0) void f1(__spv::__spirv_JointMatrixINTEL *matrix) {} diff --git a/sycl/test/matrix/legacy/matrix-int8-test.cpp b/sycl/test/matrix/legacy/matrix-int8-test.cpp index a0c2edb62c2f1..852c877b46fc4 100644 --- a/sycl/test/matrix/legacy/matrix-int8-test.cpp +++ b/sycl/test/matrix/legacy/matrix-int8-test.cpp @@ -1,8 +1,8 @@ // RUN: %clangxx -fsycl -fsycl-device-only -O2 -DSYCL_EXT_ONEAPI_MATRIX_VERSION=1 -S -emit-llvm -o - %s | FileCheck %s -// CHECK-DAG: %spirv.JointMatrixINTEL._char_12_48_0_3 = type opaque -// CHECK-DAG: %spirv.JointMatrixINTEL._int_12_12_0_3 = type opaque -// CHECK-DAG: %spirv.JointMatrixINTEL._char_48_12_3_3 = type opaque +// CHECK-DAG: target("spirv.JointMatrixINTEL", i8, 12, 48, 0, 3, 0) +// CHECK-DAG: target("spirv.JointMatrixINTEL", i32, 12, 12, 0, 3, 0) +// CHECK-DAG: target("spirv.JointMatrixINTEL", i8, 48, 12, 3, 3, 0) #include #include diff --git a/sycl/test/matrix/matrix-int8-test.cpp b/sycl/test/matrix/matrix-int8-test.cpp index de8721bca3b09..99f60423ca212 100644 --- a/sycl/test/matrix/matrix-int8-test.cpp +++ b/sycl/test/matrix/matrix-int8-test.cpp @@ -1,8 +1,8 @@ // RUN: %clangxx -fsycl -fsycl-device-only -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 -O2 -S -emit-llvm -o - %s | FileCheck %s -// CHECK-DAG: %spirv.JointMatrixINTEL._char_12_48_0_3_0 = type opaque -// CHECK-DAG: %spirv.JointMatrixINTEL._int_12_12_3_3_2 = type opaque -// CHECK-DAG: %spirv.JointMatrixINTEL._char_48_12_2_3_1 = type opaque +// CHECK-DAG: target("spirv.JointMatrixINTEL", i8, 12, 48, 0, 3, 0) +// CHECK-DAG: target("spirv.JointMatrixINTEL", i32, 12, 12, 3, 3, 2) +// CHECK-DAG: target("spirv.JointMatrixINTEL", i8, 48, 12, 2, 3, 1) #include #include From 3edb1bdd773a8999a2ddea90215336b654c6f50f Mon Sep 17 00:00:00 2001 From: "Sidorov, Dmitry" Date: Mon, 3 Apr 2023 06:00:54 -0700 Subject: [PATCH 4/5] check for RD Signed-off-by: Sidorov, Dmitry --- clang/lib/CodeGen/CodeGenTypes.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/clang/lib/CodeGen/CodeGenTypes.cpp b/clang/lib/CodeGen/CodeGenTypes.cpp index cb8d48a9dddbe..230255134a357 100644 --- a/clang/lib/CodeGen/CodeGenTypes.cpp +++ b/clang/lib/CodeGen/CodeGenTypes.cpp @@ -762,7 +762,8 @@ llvm::Type *CodeGenTypes::ConvertType(QualType T) { const Type *ClangETy = ETy.getTypePtrOrNull(); if (ClangETy && ClangETy->isStructureOrClassType()) { RecordDecl *RD = ClangETy->getAsCXXRecordDecl(); - if (RD->getQualifiedNameAsString() == "__spv::__spirv_JointMatrixINTEL") { + if (RD && + RD->getQualifiedNameAsString() == "__spv::__spirv_JointMatrixINTEL") { ResultType = ConvertSYCLJointMatrixINTELType(RD); break; } From 8d29f3c3521b87fb109a56f6343a5c111806d845 Mon Sep 17 00:00:00 2001 From: "Sidorov, Dmitry" Date: Thu, 6 Apr 2023 08:33:49 -0700 Subject: [PATCH 5/5] Typos fixes Signed-off-by: Sidorov, Dmitry --- clang/lib/CodeGen/CodeGenTypes.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenTypes.cpp b/clang/lib/CodeGen/CodeGenTypes.cpp index 230255134a357..ed063a36e5705 100644 --- a/clang/lib/CodeGen/CodeGenTypes.cpp +++ b/clang/lib/CodeGen/CodeGenTypes.cpp @@ -411,8 +411,8 @@ llvm::Type *getJointMatrixINTELExtType(llvm::Type *CompTy, // Once we update to the newest version of the spec - this should be updated. assert((TemplateArgs.size() == 5 || TemplateArgs.size() == 6) && "Wrong JointMatrixINTEL template parameters number"); - // This is required to represent optional Optional - // 'Component Type Interpretation' parameter + // This is required to represent optional 'Component Type Interpretation' + // parameter using ParamsType = typename std::conditional, SmallVector>::type; @@ -444,10 +444,10 @@ llvm::Type *CodeGenTypes::ConvertSYCLJointMatrixINTELType(RecordDecl *RD) { "1st JointMatrixINTEL template parameter must be type"); llvm::Type *CompTy = ConvertType(TemplateArgs[0].getAsType()); - // Per JointMatrixINTEL spec the type can have an Optional + // Per JointMatrixINTEL spec the type can have an optional // 'Component Type Interpretation' parameter. We should emit it in case // if on SYCL level joint matrix accepts 'bfloat16' or 'tf32' objects as - // matrix's components. Yet bfloat16 should be represented as 'int16' and + // matrix's components. Yet 'bfloat16' should be represented as 'int16' and // 'tf32' as 'float' types. if (CompTy->isStructTy()) { StringRef LlvmTyName = CompTy->getStructName();