From 3a136cb42dd42b86fde001eef6c88a3b782ed085 Mon Sep 17 00:00:00 2001 From: "Sidorov, Dmitry" Date: Fri, 30 Sep 2022 06:51:33 -0700 Subject: [PATCH 1/3] [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 --- clang/lib/CodeGen/CodeGenTypes.cpp | 21 +++++++++++++++------ clang/test/CodeGenSYCL/matrix.cpp | 17 ++++++++++------- 2 files changed, 25 insertions(+), 13 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenTypes.cpp b/clang/lib/CodeGen/CodeGenTypes.cpp index f29fb445c31d..79228bab2983 100644 --- a/clang/lib/CodeGen/CodeGenTypes.cpp +++ b/clang/lib/CodeGen/CodeGenTypes.cpp @@ -57,11 +57,14 @@ void CodeGenTypes::addRecordTypeName(const RecordDecl *RD, if (auto TemplateDecl = dyn_cast(RD)) { ArrayRef 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; diff --git a/clang/test/CodeGenSYCL/matrix.cpp b/clang/test/CodeGenSYCL/matrix.cpp index a36151859051..1c630b8007af 100644 --- a/clang/test/CodeGenSYCL/matrix.cpp +++ b/clang/test/CodeGenSYCL/matrix.cpp @@ -5,18 +5,18 @@ #include namespace __spv { - template + template struct __spirv_JointMatrixINTEL; } // CHECK: @_Z2f1{{.*}}(%spirv.JointMatrixINTEL._float_5_10_0_1 -void f1(__spv::__spirv_JointMatrixINTEL *matrix) {} +void f1(__spv::__spirv_JointMatrixINTEL *matrix) {} // CHECK: @_Z2f2{{.*}}(%spirv.JointMatrixINTEL._long_10_2_0_0 -void f2(__spv::__spirv_JointMatrixINTEL *matrix) {} +void f2(__spv::__spirv_JointMatrixINTEL *matrix) {} // CHECK: @_Z2f3{{.*}}(%spirv.JointMatrixINTEL._char_10_2_0_0 -void f3(__spv::__spirv_JointMatrixINTEL *matrix) {} +void f3(__spv::__spirv_JointMatrixINTEL *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 *matrix) {} +void f4(__spv::__spirv_JointMatrixINTEL *matrix) {} // CHECK: @_Z2f5{{.*}}(%spirv.JointMatrixINTEL._bfloat16_10_2_0_0 -void f5(__spv::__spirv_JointMatrixINTEL *matrix) {} +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) {} +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 *matrix) {} From 1405c9d072229294e2c9b9f49fe073091b8d021c Mon Sep 17 00:00:00 2001 From: "Sidorov, Dmitry" Date: Fri, 30 Sep 2022 08:07:55 -0700 Subject: [PATCH 2/3] Add min number of template parameters check Signed-off-by: Sidorov, Dmitry --- clang/lib/CodeGen/CodeGenTypes.cpp | 7 +++++-- 1 file changed, 5 insertions(+), 2 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenTypes.cpp b/clang/lib/CodeGen/CodeGenTypes.cpp index 79228bab2983..0a1dfed8df2b 100644 --- a/clang/lib/CodeGen/CodeGenTypes.cpp +++ b/clang/lib/CodeGen/CodeGenTypes.cpp @@ -57,11 +57,14 @@ void CodeGenTypes::addRecordTypeName(const RecordDecl *RD, if (auto TemplateDecl = dyn_cast(RD)) { ArrayRef TemplateArgs = TemplateDecl->getTemplateArgs().asArray(); + [[maybe_unused]] constexpr size_t MinMatrixParameter = 5; constexpr size_t MaxMatrixParameter = 6; - assert(TemplateArgs.size() <= MaxMatrixParameter && + const size_t TemplateArgsSize = TemplateArgs.size(); + assert((TemplateArgsSize >= MinMatrixParameter && + TemplateArgsSize <= MaxMatrixParameter) && "Too many template parameters for JointMatrixINTEL type"); OS << "spirv.JointMatrixINTEL."; - for (size_t I = 0; I != TemplateArgs.size(); ++I) { + for (size_t I = 0; I != TemplateArgsSize; ++I) { if (TemplateArgs[I].getKind() == TemplateArgument::Type) { OS << "_"; llvm::Type *TTy = ConvertType(TemplateArgs[I].getAsType()); From 141524c74c48324d0e30bc7d639d66d3f0e8a3b5 Mon Sep 17 00:00:00 2001 From: "Sidorov, Dmitry" Date: Sun, 2 Oct 2022 18:27:32 -0700 Subject: [PATCH 3/3] Remove min template params Signed-off-by: Sidorov, Dmitry --- clang/lib/CodeGen/CodeGenTypes.cpp | 16 ++++++++-------- 1 file changed, 8 insertions(+), 8 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenTypes.cpp b/clang/lib/CodeGen/CodeGenTypes.cpp index 0a1dfed8df2b..712b39e7a877 100644 --- a/clang/lib/CodeGen/CodeGenTypes.cpp +++ b/clang/lib/CodeGen/CodeGenTypes.cpp @@ -57,12 +57,10 @@ void CodeGenTypes::addRecordTypeName(const RecordDecl *RD, if (auto TemplateDecl = dyn_cast(RD)) { ArrayRef TemplateArgs = TemplateDecl->getTemplateArgs().asArray(); - [[maybe_unused]] constexpr size_t MinMatrixParameter = 5; - constexpr size_t MaxMatrixParameter = 6; + constexpr size_t NumOfMatrixParameters = 6; const size_t TemplateArgsSize = TemplateArgs.size(); - assert((TemplateArgsSize >= MinMatrixParameter && - TemplateArgsSize <= MaxMatrixParameter) && - "Too many template parameters for JointMatrixINTEL type"); + assert(TemplateArgsSize == NumOfMatrixParameters && + "Incorrect number of template parameters for JointMatrixINTEL"); OS << "spirv.JointMatrixINTEL."; for (size_t I = 0; I != TemplateArgsSize; ++I) { if (TemplateArgs[I].getKind() == TemplateArgument::Type) { @@ -99,10 +97,12 @@ void CodeGenTypes::addRecordTypeName(const RecordDecl *RD, TTy->print(OS, false, true); } 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 + // Last template parameter of __spirv_JointMatrixINTEL 'Use' is + // optional in SPIR-V, so If it has 'Unnecessary' value - skip it. + // MatrixUse::Unnecessary defined as '3' in spirv_types.hpp. constexpr size_t Unnecessary = 3; - if (!(I == MaxMatrixParameter && IntTemplateParam == Unnecessary)) + if (!(I == NumOfMatrixParameters && + IntTemplateParam == Unnecessary)) OS << "_" << IntTemplateParam; } }