From 4555f4468e3077569f2eb93d3f0c81e1161d829c Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Tue, 10 Mar 2020 19:17:10 +0300 Subject: [PATCH 01/26] [SYCL] Support intel::reqd_work_group_size (fix #6) Signed-off-by: Aleksander Fadeev --- .../SemaSYCL/intel-reqd-work-group-size.cpp | 122 ++++++++++++++++++ 1 file changed, 122 insertions(+) create mode 100644 clang/test/SemaSYCL/intel-reqd-work-group-size.cpp diff --git a/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp b/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp new file mode 100644 index 0000000000000..40f2bfe12a64b --- /dev/null +++ b/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp @@ -0,0 +1,122 @@ +// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -verify -DTRIGGER_ERROR %s +// RUN: %clang_cc1 -fsycl-is-device -ast-dump %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-host -fsyntax-only -verify %s + +#ifndef __SYCL_DEVICE_ONLY__ +// expected-no-diagnostics +class Functor { +public: + [[intel::reqd_work_group_size(4, 1, 1)]] void operator()() {} + +}; + +template +void kernel(Func kernelFunc) { + kernelFunc(); +} + +void bar() { + Functor f; + kernel(f); +} +#else +[[intel::reqd_work_group_size(4, 1, 1)]] void f4x1x1() {} // expected-note {{conflicting attribute is here}} +// expected-note@-1 {{conflicting attribute is here}} +[[intel::reqd_work_group_size(32, 1, 1)]] void f32x1x1() {} // expected-note {{conflicting attribute is here}} + +[[intel::reqd_work_group_size(16, 1, 1)]] void f16x1x1() {} // expected-note {{conflicting attribute is here}} +[[intel::reqd_work_group_size(16, 16, 1)]] void f16x16x1() {} // expected-note {{conflicting attribute is here}} + +[[intel::reqd_work_group_size(32, 32, 1)]] void f32x32x1() {} // expected-note {{conflicting attribute is here}} +[[intel::reqd_work_group_size(32, 32, 32)]] void f32x32x32() {} // expected-note {{conflicting attribute is here}} + +class Functor16 { +public: + [[intel::reqd_work_group_size(16, 1)]] void operator()() {} +}; + +class Functor16x16x16 { +public: + [[intel::reqd_work_group_size(16, 16, 16)]] void operator()() {} +}; + +class Functor8 { // expected-error {{conflicting attributes applied to a SYCL kernel}} +public: + [[intel::reqd_work_group_size(8, 1, 1)]] void operator()() { // expected-note {{conflicting attribute is here}} + f4x1x1(); + } +}; + +class Functor { +public: + void operator()() { + f4x1x1(); + } +}; + +class FunctorAttr { +public: + __attribute__((reqd_work_group_size(128, 128, 128))) void operator()() {} +}; + +template +__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { + kernelFunc(); +} + +void bar() { + Functor16 f16; + kernel(f16); + + Functor f; + kernel(f); + + Functor16x16x16 f16x16x16; + kernel(f16x16x16); + + FunctorAttr fattr; + kernel(fattr); + + kernel([]() [[intel::reqd_work_group_size(32, 32, 32)]] { + f32x32x32(); + }); + + +#ifdef TRIGGER_ERROR + Functor8 f8; + kernel(f8); + + kernel([]() { // expected-error {{conflicting attributes applied to a SYCL kernel}} + f4x1x1(); + f32x1x1(); + }); + + kernel([]() { // expected-error {{conflicting attributes applied to a SYCL kernel}} + f16x1x1(); + f16x16x1(); + }); + + kernel([]() { // expected-error {{conflicting attributes applied to a SYCL kernel}} + f32x32x32(); + f32x32x1(); + }); + + // expected-error@+1 {{expected variable name or 'this' in lambda capture list}} + kernel([[intel::reqd_work_group_size(32, 32, 32)]] []() { + f32x32x32(); + }); + +#endif +} + +// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name1 +// CHECK: SYCLIntelReqdWorkGroupSizeAttr {{.*}} 1 1 16 +// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name2 +// CHECK: SYCLIntelReqdWorkGroupSizeAttr {{.*}} 1 1 4 +// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name3 +// CHECK: SYCLIntelReqdWorkGroupSizeAttr {{.*}} 16 16 16 +// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name4 +// CHECK: ReqdWorkGroupSizeAttr {{.*}} 128 128 128 +// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name5 +// CHECK: SYCLIntelReqdWorkGroupSizeAttr {{.*}} 32 32 32 +#endif // __SYCL_DEVICE_ONLY__ From c704d1a45153f7511cfce629ee8b33ff8db09f47 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Wed, 11 Mar 2020 12:41:03 +0300 Subject: [PATCH 02/26] [SYCL] Support intel::reqd_work_group_size (fix #7) Signed-off-by: Aleksander Fadeev --- clang/include/clang/Basic/Attr.td | 8 +++++ .../include/clang/Basic/AttributeCommonInfo.h | 1 + clang/lib/CodeGen/CodeGenFunction.cpp | 10 ++++++ clang/lib/Sema/SemaDeclAttr.cpp | 35 +++++++++++++++---- clang/lib/Sema/SemaSYCL.cpp | 19 ++++++++++ ...a-attribute-supported-attributes-list.test | 1 + 6 files changed, 67 insertions(+), 7 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 088d90f505f9a..ecbb28e2bdb18 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -2422,6 +2422,14 @@ def ReqdWorkGroupSize : InheritableAttr { let Documentation = [Undocumented]; } +def SYCLIntelReqdWorkGroupSize : InheritableAttr { + let Spellings = [CXX11<"intel","reqd_work_group_size">]; + let Args = [IntArgument<"XDim">, DefaultIntArgument<"YDim", 1>, + DefaultIntArgument<"ZDim", 1>]; + let Subjects = SubjectList<[Function], ErrorDiag>; + let Documentation = [Undocumented]; +} + def WorkGroupSizeHint : InheritableAttr { // Does not have a [[]] spelling because it is an OpenCL-related attribute. let Spellings = [GNU<"work_group_size_hint">]; diff --git a/clang/include/clang/Basic/AttributeCommonInfo.h b/clang/include/clang/Basic/AttributeCommonInfo.h index 2fbaa8dfc22fa..3895784dea50e 100644 --- a/clang/include/clang/Basic/AttributeCommonInfo.h +++ b/clang/include/clang/Basic/AttributeCommonInfo.h @@ -160,6 +160,7 @@ class AttributeCommonInfo { auto ParsedAttr = getParsedKind(); if (ParsedAttr == AT_SYCLIntelKernelArgsRestrict || (ParsedAttr == AT_ReqdWorkGroupSize && isCXX11Attribute()) || + (ParsedAttr == AT_SYCLIntelReqdWorkGroupSize && isCXX11Attribute()) || (ParsedAttr == AT_IntelReqdSubGroupSize && isCXX11Attribute()) || ParsedAttr == AT_SYCLIntelNumSimdWorkItems || ParsedAttr == AT_SYCLIntelMaxWorkGroupSize || diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index c2925b28c8d45..11066db4f5cae 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -644,6 +644,16 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD, Fn->setMetadata("reqd_work_group_size", llvm::MDNode::get(Context, AttrMDArgs)); } + if (const SYCLIntelReqdWorkGroupSizeAttr *A = + FD->getAttr()) { + llvm::Metadata *AttrMDArgs[] = { + llvm::ConstantAsMetadata::get(Builder.getInt32(A->getXDim())), + llvm::ConstantAsMetadata::get(Builder.getInt32(A->getYDim())), + llvm::ConstantAsMetadata::get(Builder.getInt32(A->getZDim()))}; + Fn->setMetadata("reqd_work_group_size", + llvm::MDNode::get(Context, AttrMDArgs)); + } + if (const IntelReqdSubGroupSizeAttr *A = FD->getAttr()) { llvm::Metadata *AttrMDArgs[] = { diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index f71b8327224aa..70526e5486ccb 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -2889,6 +2889,9 @@ static bool checkWorkGroupSizeValues(Sema &S, Decl *D, const ParsedAttr &Attr, if (const auto *A = D->getAttr()) Result &= checkZeroDim(A, A->getXDim(), A->getYDim(), A->getZDim()); return Result; + if (const auto *A = D->getAttr()) + Result &= checkZeroDim(A, A->getXDim(), A->getYDim(), A->getZDim()); + return Result; } if (const auto *A = D->getAttr()) @@ -2912,9 +2915,18 @@ static bool checkWorkGroupSizeValues(Sema &S, Decl *D, const ParsedAttr &Attr, Result &= false; } } + if (const auto *A = D->getAttr()) { + if (!(WGSize[0] >= A->getXDim() && WGSize[1] >= A->getYDim() && + WGSize[2] >= A->getZDim())) { + S.Diag(Attr.getLoc(), diag::err_conflicting_sycl_function_attributes) + << Attr << A->getSpelling(); + Result &= false; + } + } return Result; } +#include // Handles reqd_work_group_size, work_group_size_hint and max_work_group_size template static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { @@ -2922,15 +2934,18 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { return; uint32_t WGSize[3]; - for (unsigned i = 0; i < 3; ++i) { - const Expr *E = AL.getArgAsExpr(i); - if (!checkUInt32Argument(S, AL, E, WGSize[i], i, + if (std::is_same::value) { + WGSize[1] = SYCLIntelReqdWorkGroupSizeAttr::DefaultYDim; + WGSize[2] = SYCLIntelReqdWorkGroupSizeAttr::DefaultZDim; + } + if (i < AL.getNumArgs()) + if (!checkUInt32Argument(S, AL, AL.getArgAsExpr(i), WGSize[i], i, /*StrictlyUnsigned=*/true)) return; - if (WGSize[i] == 0) { - S.Diag(AL.getLoc(), diag::err_attribute_argument_is_zero) - << AL << E->getSourceRange(); - return; + if (WGSize[i] == 0) { + S.Diag(AL.getLoc(), diag::err_attribute_argument_is_zero) + << AL << E->getSourceRange(); + return; } } @@ -7660,6 +7675,9 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, case ParsedAttr::AT_ReqdWorkGroupSize: handleWorkGroupSize(S, D, AL); break; + case ParsedAttr::AT_SYCLIntelReqdWorkGroupSize: + handleWorkGroupSize(S, D, AL); + break; case ParsedAttr::AT_SYCLIntelMaxWorkGroupSize: handleWorkGroupSize(S, D, AL); break; @@ -8162,6 +8180,9 @@ void Sema::ProcessDeclAttributeList(Scope *S, Decl *D, // diag::err_attribute_wrong_decl_type + ExpectedKernelFunction. Diag(D->getLocation(), diag::err_opencl_kernel_attr) << A; D->setInvalidDecl(); + } else if (const auto *A = D->getAttr()) { + Diag(D->getLocation(), diag::err_opencl_kernel_attr) << A; + D->setInvalidDecl(); } else if (const auto *A = D->getAttr()) { Diag(D->getLocation(), diag::err_opencl_kernel_attr) << A; D->setInvalidDecl(); diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index b31bc2ec7cdde..c0d0b95cf32bb 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -410,6 +410,8 @@ class MarkDeviceFunction : public RecursiveASTVisitor { Attrs.insert(A); if (auto *A = FD->getAttr()) Attrs.insert(A); + if (auto *A = FD->getAttr()) + Attrs.insert(A); // Allow the following kernel attributes only on lambda functions and // function objects that are called directly from a kernel (i.e. the one // passed to the parallel_for function). For all other cases, @@ -1370,6 +1372,23 @@ void Sema::MarkDevice(void) { } break; } + case attr::Kind::SYCLIntelReqdWorkGroupSize: { + auto *Attr = cast(A); + if (auto *Existing = SYCLKernel->getAttr()) { + if (Existing->getXDim() != Attr->getXDim() || + Existing->getYDim() != Attr->getYDim() || + Existing->getZDim() != Attr->getZDim()) { + Diag(SYCLKernel->getLocation(), + diag::err_conflicting_sycl_kernel_attributes); + Diag(Existing->getLocation(), diag::note_conflicting_attribute); + Diag(Attr->getLocation(), diag::note_conflicting_attribute); + SYCLKernel->setInvalidDecl(); + } + } else { + SYCLKernel->addAttr(A); + } + break; + } case attr::Kind::SYCLIntelKernelArgsRestrict: case attr::Kind::SYCLIntelNumSimdWorkItems: case attr::Kind::SYCLIntelMaxGlobalWorkDim: diff --git a/clang/test/Misc/pragma-attribute-supported-attributes-list.test b/clang/test/Misc/pragma-attribute-supported-attributes-list.test index 2a196f48973b2..282014396e86b 100644 --- a/clang/test/Misc/pragma-attribute-supported-attributes-list.test +++ b/clang/test/Misc/pragma-attribute-supported-attributes-list.test @@ -140,6 +140,7 @@ // CHECK-NEXT: SYCLDevice (SubjectMatchRule_function) // CHECK-NEXT: SYCLDeviceIndirectlyCallable (SubjectMatchRule_function) // CHECK-NEXT: SYCLIntelKernelArgsRestrict (SubjectMatchRule_function) +// CHECK-NEXT: SYCLIntelReqdWorkGroupSize (SubjectMatchRule_function) // CHECK-NEXT: ScopedLockable (SubjectMatchRule_record) // CHECK-NEXT: Section (SubjectMatchRule_function, SubjectMatchRule_variable_is_global, SubjectMatchRule_objc_method, SubjectMatchRule_objc_property) // CHECK-NEXT: SetTypestate (SubjectMatchRule_function_is_member) From d6e357550eeebfe629c579588f1fe79f68f0d1c2 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Wed, 11 Mar 2020 12:45:45 +0300 Subject: [PATCH 03/26] [SYCL] Support intel::reqd_work_group_size (fix #7.2) Signed-off-by: Aleksander Fadeev --- clang/lib/Sema/SemaDeclAttr.cpp | 15 ++++++++------- 1 file changed, 8 insertions(+), 7 deletions(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 70526e5486ccb..1bc6eea296ca4 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -2938,14 +2938,15 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { WGSize[1] = SYCLIntelReqdWorkGroupSizeAttr::DefaultYDim; WGSize[2] = SYCLIntelReqdWorkGroupSizeAttr::DefaultZDim; } - if (i < AL.getNumArgs()) - if (!checkUInt32Argument(S, AL, AL.getArgAsExpr(i), WGSize[i], i, - /*StrictlyUnsigned=*/true)) + for (unsigned i = 0; i < 3; ++i) { + if (i < AL.getNumArgs()) + if (!checkUInt32Argument(S, AL, AL.getArgAsExpr(i), WGSize[i], i, + /*StrictlyUnsigned=*/true)) + return; + if (WGSize[i] == 0) { + S.Diag(AL.getLoc(), diag::err_attribute_argument_is_zero) + << AL << E->getSourceRange(); return; - if (WGSize[i] == 0) { - S.Diag(AL.getLoc(), diag::err_attribute_argument_is_zero) - << AL << E->getSourceRange(); - return; } } From 18b0e94a8328fceb3c6d0cd79855256bfc1c5453 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Wed, 11 Mar 2020 13:21:45 +0300 Subject: [PATCH 04/26] [SYCL] Support intel::reqd_work_group_size (fix #7.3) Signed-off-by: Aleksander Fadeev --- clang/lib/Sema/SemaDeclAttr.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 1bc6eea296ca4..f89e3cae487b5 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -2945,7 +2945,7 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { return; if (WGSize[i] == 0) { S.Diag(AL.getLoc(), diag::err_attribute_argument_is_zero) - << AL << E->getSourceRange(); + << AL << AL.getArgAsExpr(i)->getSourceRange(); return; } } From 6a313bfc7f6938cdf2ff48ca04e305821823c0ba Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Wed, 11 Mar 2020 16:51:30 +0300 Subject: [PATCH 05/26] [SYCL] Support intel::reqd_work_group_size (fix #7.4) Signed-off-by: Aleksander Fadeev --- clang/lib/Sema/SemaDeclAttr.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index f89e3cae487b5..c7659ec1fa8f8 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -36,6 +36,7 @@ #include "llvm/ADT/STLExtras.h" #include "llvm/ADT/StringExtras.h" #include "llvm/Support/MathExtras.h" +#include using namespace clang; using namespace sema; @@ -2926,7 +2927,7 @@ static bool checkWorkGroupSizeValues(Sema &S, Decl *D, const ParsedAttr &Attr, return Result; } -#include +//#include // Handles reqd_work_group_size, work_group_size_hint and max_work_group_size template static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { From 995352eca68a31a5d689ce01509d8f41fcc3888c Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Thu, 12 Mar 2020 18:10:40 +0300 Subject: [PATCH 06/26] [SYCL] Support intel::reqd_work_group_size (fix #8) Signed-off-by: Aleksander Fadeev --- clang/include/clang/Basic/Attr.td | 13 ++------- .../include/clang/Basic/AttributeCommonInfo.h | 1 - clang/lib/CodeGen/CodeGenFunction.cpp | 10 ------- clang/lib/Sema/SemaDeclAttr.cpp | 29 +++++-------------- clang/lib/Sema/SemaSYCL.cpp | 19 ------------ .../SemaSYCL/intel-reqd-work-group-size.cpp | 17 +++++++---- 6 files changed, 22 insertions(+), 67 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index ecbb28e2bdb18..77c8a1a64df90 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -2415,16 +2415,9 @@ def NoDeref : TypeAttr { def ReqdWorkGroupSize : InheritableAttr { let Spellings = [GNU<"reqd_work_group_size">, - CXX11<"cl","reqd_work_group_size">]; - let Args = [UnsignedArgument<"XDim">, UnsignedArgument<"YDim">, - UnsignedArgument<"ZDim">]; - let Subjects = SubjectList<[Function], ErrorDiag>; - let Documentation = [Undocumented]; -} - -def SYCLIntelReqdWorkGroupSize : InheritableAttr { - let Spellings = [CXX11<"intel","reqd_work_group_size">]; - let Args = [IntArgument<"XDim">, DefaultIntArgument<"YDim", 1>, + CXX11<"cl","reqd_work_group_size">, + CXX11<"intel","reqd_work_group_size">]; + let Args = [UnsignedArgument<"XDim">, DefaultIntArgument<"YDim", 1>, DefaultIntArgument<"ZDim", 1>]; let Subjects = SubjectList<[Function], ErrorDiag>; let Documentation = [Undocumented]; diff --git a/clang/include/clang/Basic/AttributeCommonInfo.h b/clang/include/clang/Basic/AttributeCommonInfo.h index 3895784dea50e..2fbaa8dfc22fa 100644 --- a/clang/include/clang/Basic/AttributeCommonInfo.h +++ b/clang/include/clang/Basic/AttributeCommonInfo.h @@ -160,7 +160,6 @@ class AttributeCommonInfo { auto ParsedAttr = getParsedKind(); if (ParsedAttr == AT_SYCLIntelKernelArgsRestrict || (ParsedAttr == AT_ReqdWorkGroupSize && isCXX11Attribute()) || - (ParsedAttr == AT_SYCLIntelReqdWorkGroupSize && isCXX11Attribute()) || (ParsedAttr == AT_IntelReqdSubGroupSize && isCXX11Attribute()) || ParsedAttr == AT_SYCLIntelNumSimdWorkItems || ParsedAttr == AT_SYCLIntelMaxWorkGroupSize || diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index 11066db4f5cae..c2925b28c8d45 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -644,16 +644,6 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD, Fn->setMetadata("reqd_work_group_size", llvm::MDNode::get(Context, AttrMDArgs)); } - if (const SYCLIntelReqdWorkGroupSizeAttr *A = - FD->getAttr()) { - llvm::Metadata *AttrMDArgs[] = { - llvm::ConstantAsMetadata::get(Builder.getInt32(A->getXDim())), - llvm::ConstantAsMetadata::get(Builder.getInt32(A->getYDim())), - llvm::ConstantAsMetadata::get(Builder.getInt32(A->getZDim()))}; - Fn->setMetadata("reqd_work_group_size", - llvm::MDNode::get(Context, AttrMDArgs)); - } - if (const IntelReqdSubGroupSizeAttr *A = FD->getAttr()) { llvm::Metadata *AttrMDArgs[] = { diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index c7659ec1fa8f8..32639b720924c 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -2890,9 +2890,6 @@ static bool checkWorkGroupSizeValues(Sema &S, Decl *D, const ParsedAttr &Attr, if (const auto *A = D->getAttr()) Result &= checkZeroDim(A, A->getXDim(), A->getYDim(), A->getZDim()); return Result; - if (const auto *A = D->getAttr()) - Result &= checkZeroDim(A, A->getXDim(), A->getYDim(), A->getZDim()); - return Result; } if (const auto *A = D->getAttr()) @@ -2916,14 +2913,6 @@ static bool checkWorkGroupSizeValues(Sema &S, Decl *D, const ParsedAttr &Attr, Result &= false; } } - if (const auto *A = D->getAttr()) { - if (!(WGSize[0] >= A->getXDim() && WGSize[1] >= A->getYDim() && - WGSize[2] >= A->getZDim())) { - S.Diag(Attr.getLoc(), diag::err_conflicting_sycl_function_attributes) - << Attr << A->getSpelling(); - Result &= false; - } - } return Result; } @@ -2935,13 +2924,15 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { return; uint32_t WGSize[3]; - if (std::is_same::value) { - WGSize[1] = SYCLIntelReqdWorkGroupSizeAttr::DefaultYDim; - WGSize[2] = SYCLIntelReqdWorkGroupSizeAttr::DefaultZDim; - } + if (AL.getNormalizedFullName() == "intel::reqd_work_group_size") { + WGSize[1] = ReqdWorkGroupSizeAttr::DefaultYDim; + WGSize[2] = ReqdWorkGroupSizeAttr::DefaultZDim; + } else if (!checkAttributeNumArgs(S, AL, 3)) + return; + for (unsigned i = 0; i < 3; ++i) { if (i < AL.getNumArgs()) - if (!checkUInt32Argument(S, AL, AL.getArgAsExpr(i), WGSize[i], i, + if (!checkUInt32Argument(S, AL, AL.getArgAsExpr(i), WGSize[i], i, /*StrictlyUnsigned=*/true)) return; if (WGSize[i] == 0) { @@ -7677,9 +7668,6 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, case ParsedAttr::AT_ReqdWorkGroupSize: handleWorkGroupSize(S, D, AL); break; - case ParsedAttr::AT_SYCLIntelReqdWorkGroupSize: - handleWorkGroupSize(S, D, AL); - break; case ParsedAttr::AT_SYCLIntelMaxWorkGroupSize: handleWorkGroupSize(S, D, AL); break; @@ -8182,9 +8170,6 @@ void Sema::ProcessDeclAttributeList(Scope *S, Decl *D, // diag::err_attribute_wrong_decl_type + ExpectedKernelFunction. Diag(D->getLocation(), diag::err_opencl_kernel_attr) << A; D->setInvalidDecl(); - } else if (const auto *A = D->getAttr()) { - Diag(D->getLocation(), diag::err_opencl_kernel_attr) << A; - D->setInvalidDecl(); } else if (const auto *A = D->getAttr()) { Diag(D->getLocation(), diag::err_opencl_kernel_attr) << A; D->setInvalidDecl(); diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index c0d0b95cf32bb..b31bc2ec7cdde 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -410,8 +410,6 @@ class MarkDeviceFunction : public RecursiveASTVisitor { Attrs.insert(A); if (auto *A = FD->getAttr()) Attrs.insert(A); - if (auto *A = FD->getAttr()) - Attrs.insert(A); // Allow the following kernel attributes only on lambda functions and // function objects that are called directly from a kernel (i.e. the one // passed to the parallel_for function). For all other cases, @@ -1372,23 +1370,6 @@ void Sema::MarkDevice(void) { } break; } - case attr::Kind::SYCLIntelReqdWorkGroupSize: { - auto *Attr = cast(A); - if (auto *Existing = SYCLKernel->getAttr()) { - if (Existing->getXDim() != Attr->getXDim() || - Existing->getYDim() != Attr->getYDim() || - Existing->getZDim() != Attr->getZDim()) { - Diag(SYCLKernel->getLocation(), - diag::err_conflicting_sycl_kernel_attributes); - Diag(Existing->getLocation(), diag::note_conflicting_attribute); - Diag(Attr->getLocation(), diag::note_conflicting_attribute); - SYCLKernel->setInvalidDecl(); - } - } else { - SYCLKernel->addAttr(A); - } - break; - } case attr::Kind::SYCLIntelKernelArgsRestrict: case attr::Kind::SYCLIntelNumSimdWorkItems: case attr::Kind::SYCLIntelMaxGlobalWorkDim: diff --git a/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp b/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp index 40f2bfe12a64b..f62fa0b8cb281 100644 --- a/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp +++ b/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp @@ -30,9 +30,16 @@ void bar() { [[intel::reqd_work_group_size(32, 32, 1)]] void f32x32x1() {} // expected-note {{conflicting attribute is here}} [[intel::reqd_work_group_size(32, 32, 32)]] void f32x32x32() {} // expected-note {{conflicting attribute is here}} +#ifdef TRIGGER_ERROR +class Functor32 { +public: + [[cl::reqd_work_group_size(32)]] void operator()() {} // expected-error {{'reqd_work_group_size' attribute requires exactly 3 arguments}} +}; +#endif + class Functor16 { public: - [[intel::reqd_work_group_size(16, 1)]] void operator()() {} + [[intel::reqd_work_group_size(16)]] void operator()() {} }; class Functor16x16x16 { @@ -110,13 +117,13 @@ void bar() { } // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name1 -// CHECK: SYCLIntelReqdWorkGroupSizeAttr {{.*}} 1 1 16 +// CHECK: ReqdWorkGroupSizeAttr {{.*}} 1 1 16 // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name2 -// CHECK: SYCLIntelReqdWorkGroupSizeAttr {{.*}} 1 1 4 +// CHECK: ReqdWorkGroupSizeAttr {{.*}} 1 1 4 // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name3 -// CHECK: SYCLIntelReqdWorkGroupSizeAttr {{.*}} 16 16 16 +// CHECK: ReqdWorkGroupSizeAttr {{.*}} 16 16 16 // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name4 // CHECK: ReqdWorkGroupSizeAttr {{.*}} 128 128 128 // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name5 -// CHECK: SYCLIntelReqdWorkGroupSizeAttr {{.*}} 32 32 32 +// CHECK: ReqdWorkGroupSizeAttr {{.*}} 32 32 32 #endif // __SYCL_DEVICE_ONLY__ From 421217acd9046b03a8452eda3938a46a04480568 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Thu, 12 Mar 2020 18:14:03 +0300 Subject: [PATCH 07/26] [SYCL] Support intel::reqd_work_group_size (fix #8.2) Signed-off-by: Aleksander Fadeev --- clang/lib/Sema/SemaDeclAttr.cpp | 2 -- clang/test/Misc/pragma-attribute-supported-attributes-list.test | 1 - 2 files changed, 3 deletions(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 32639b720924c..51ed63d612636 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -36,7 +36,6 @@ #include "llvm/ADT/STLExtras.h" #include "llvm/ADT/StringExtras.h" #include "llvm/Support/MathExtras.h" -#include using namespace clang; using namespace sema; @@ -2916,7 +2915,6 @@ static bool checkWorkGroupSizeValues(Sema &S, Decl *D, const ParsedAttr &Attr, return Result; } -//#include // Handles reqd_work_group_size, work_group_size_hint and max_work_group_size template static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { diff --git a/clang/test/Misc/pragma-attribute-supported-attributes-list.test b/clang/test/Misc/pragma-attribute-supported-attributes-list.test index 282014396e86b..2a196f48973b2 100644 --- a/clang/test/Misc/pragma-attribute-supported-attributes-list.test +++ b/clang/test/Misc/pragma-attribute-supported-attributes-list.test @@ -140,7 +140,6 @@ // CHECK-NEXT: SYCLDevice (SubjectMatchRule_function) // CHECK-NEXT: SYCLDeviceIndirectlyCallable (SubjectMatchRule_function) // CHECK-NEXT: SYCLIntelKernelArgsRestrict (SubjectMatchRule_function) -// CHECK-NEXT: SYCLIntelReqdWorkGroupSize (SubjectMatchRule_function) // CHECK-NEXT: ScopedLockable (SubjectMatchRule_record) // CHECK-NEXT: Section (SubjectMatchRule_function, SubjectMatchRule_variable_is_global, SubjectMatchRule_objc_method, SubjectMatchRule_objc_property) // CHECK-NEXT: SetTypestate (SubjectMatchRule_function_is_member) From f9c58228ac0eebaf22a363f8413f6de2808c9a61 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Fri, 13 Mar 2020 15:26:16 +0300 Subject: [PATCH 08/26] [SYCL] Support intel::reqd_work_group_size (fix #8.3) Signed-off-by: Aleksander Fadeev --- clang/include/clang/Basic/Attr.td | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 77c8a1a64df90..2347efffcab17 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -2417,7 +2417,7 @@ def ReqdWorkGroupSize : InheritableAttr { let Spellings = [GNU<"reqd_work_group_size">, CXX11<"cl","reqd_work_group_size">, CXX11<"intel","reqd_work_group_size">]; - let Args = [UnsignedArgument<"XDim">, DefaultIntArgument<"YDim", 1>, + let Args = [IntArgument<"XDim">, DefaultIntArgument<"YDim", 1>, DefaultIntArgument<"ZDim", 1>]; let Subjects = SubjectList<[Function], ErrorDiag>; let Documentation = [Undocumented]; From eb5c79741290d36b4eef777b64b07e7c27881644 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Fri, 13 Mar 2020 16:55:28 +0300 Subject: [PATCH 09/26] [SYCL] Support intel::reqd_work_group_size (fix #9) Signed-off-by: Aleksander Fadeev --- clang/include/clang/Basic/Attr.td | 4 +++- clang/include/clang/Basic/AttrDocs.td | 11 +++++++++++ 2 files changed, 14 insertions(+), 1 deletion(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 2347efffcab17..e1008ba2964fa 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -2413,6 +2413,8 @@ def NoDeref : TypeAttr { let Documentation = [NoDerefDocs]; } +//Default arguments in ReqWorkGroupSize can be used only with +//intel::reqd_work_group_size spelling. def ReqdWorkGroupSize : InheritableAttr { let Spellings = [GNU<"reqd_work_group_size">, CXX11<"cl","reqd_work_group_size">, @@ -2420,7 +2422,7 @@ def ReqdWorkGroupSize : InheritableAttr { let Args = [IntArgument<"XDim">, DefaultIntArgument<"YDim", 1>, DefaultIntArgument<"ZDim", 1>]; let Subjects = SubjectList<[Function], ErrorDiag>; - let Documentation = [Undocumented]; + let Documentation = [ReqdWorkGroupSizeAttrDocs]; } def WorkGroupSizeHint : InheritableAttr { diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 910c6b2bee61b..44e4e359856ec 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -1983,6 +1983,17 @@ device kernel, the attribute is ignored and it is not propagated to a kernel. }]; } +def ReqdWorkGroupSizeAttrDocs : Documentation { + let Category = DocCatFunction; + let Heading = "cl::reqd_work_group_size, intel::reqd_work_group_size"; + let Content = [{ +Applies to a device function/lambda function. Indicates the dimensions +of a work group. Values must be positive integers. Default arguments +equaling to one can be used for second and third dimensions if +intel::reqd_work_group_size spelling was applied. + }]; +} + def SYCLIntelMaxWorkGroupSizeAttrDocs : Documentation { let Category = DocCatFunction; let Heading = "max_work_group_size (IntelFPGA)"; From 9c1c13033c199f7befd218def9bfb3449f2e9cbc Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Mon, 16 Mar 2020 13:35:05 +0300 Subject: [PATCH 10/26] [SYCL] Support intel::reqd_work_group_size (fix #9.2) Signed-off-by: Aleksander Fadeev --- clang/lib/Sema/SemaDeclAttr.cpp | 2 +- clang/test/SemaSYCL/intel-reqd-work-group-size.cpp | 1 - 2 files changed, 1 insertion(+), 2 deletions(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 51ed63d612636..499f4ec56271d 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -2930,7 +2930,7 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { for (unsigned i = 0; i < 3; ++i) { if (i < AL.getNumArgs()) - if (!checkUInt32Argument(S, AL, AL.getArgAsExpr(i), WGSize[i], i, + if (!checkUInt32Argument(S, AL, AL.getArgAsExpr(i), WGSize[i], i, /*StrictlyUnsigned=*/true)) return; if (WGSize[i] == 0) { diff --git a/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp b/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp index f62fa0b8cb281..92f3b2db87e77 100644 --- a/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp +++ b/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp @@ -7,7 +7,6 @@ class Functor { public: [[intel::reqd_work_group_size(4, 1, 1)]] void operator()() {} - }; template From a8b1121ab5e7289729037fb144efa01473ae0da7 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Mon, 16 Mar 2020 13:46:05 +0300 Subject: [PATCH 11/26] [SYCL] Support intel::reqd_work_group_size (fix #9.3) Signed-off-by: Aleksander Fadeev --- clang/test/SemaSYCL/intel-reqd-work-group-size.cpp | 9 ++++----- 1 file changed, 4 insertions(+), 5 deletions(-) diff --git a/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp b/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp index 92f3b2db87e77..75a282c590485 100644 --- a/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp +++ b/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp @@ -23,10 +23,10 @@ void bar() { // expected-note@-1 {{conflicting attribute is here}} [[intel::reqd_work_group_size(32, 1, 1)]] void f32x1x1() {} // expected-note {{conflicting attribute is here}} -[[intel::reqd_work_group_size(16, 1, 1)]] void f16x1x1() {} // expected-note {{conflicting attribute is here}} +[[intel::reqd_work_group_size(16, 1, 1)]] void f16x1x1() {} // expected-note {{conflicting attribute is here}} [[intel::reqd_work_group_size(16, 16, 1)]] void f16x16x1() {} // expected-note {{conflicting attribute is here}} -[[intel::reqd_work_group_size(32, 32, 1)]] void f32x32x1() {} // expected-note {{conflicting attribute is here}} +[[intel::reqd_work_group_size(32, 32, 1)]] void f32x32x1() {} // expected-note {{conflicting attribute is here}} [[intel::reqd_work_group_size(32, 32, 32)]] void f32x32x32() {} // expected-note {{conflicting attribute is here}} #ifdef TRIGGER_ERROR @@ -84,10 +84,9 @@ void bar() { kernel(fattr); kernel([]() [[intel::reqd_work_group_size(32, 32, 32)]] { - f32x32x32(); + f32x32x32(); }); - #ifdef TRIGGER_ERROR Functor8 f8; kernel(f8); @@ -108,7 +107,7 @@ void bar() { }); // expected-error@+1 {{expected variable name or 'this' in lambda capture list}} - kernel([[intel::reqd_work_group_size(32, 32, 32)]] []() { + kernel([[intel::reqd_work_group_size(32, 32, 32)]][]() { f32x32x32(); }); From 93a022d98cf80aea39a7b9eb8dd43e3ebb91fa7c Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Mon, 16 Mar 2020 18:20:25 +0300 Subject: [PATCH 12/26] [SYCL] Support intel::reqd_work_group_size (fix #10) Signed-off-by: Aleksander Fadeev --- clang/lib/Sema/SemaDeclAttr.cpp | 2 +- clang/test/SemaSYCL/intel-reqd-work-group-size.cpp | 9 +++++++++ 2 files changed, 10 insertions(+), 1 deletion(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 499f4ec56271d..16d44f86c3354 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -2922,7 +2922,7 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { return; uint32_t WGSize[3]; - if (AL.getNormalizedFullName() == "intel::reqd_work_group_size") { + if (AL.getAttributeSpellingListIndex() == 2) { WGSize[1] = ReqdWorkGroupSizeAttr::DefaultYDim; WGSize[2] = ReqdWorkGroupSizeAttr::DefaultZDim; } else if (!checkAttributeNumArgs(S, AL, 3)) diff --git a/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp b/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp index 75a282c590485..1476d273a0c51 100644 --- a/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp +++ b/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp @@ -34,6 +34,10 @@ class Functor32 { public: [[cl::reqd_work_group_size(32)]] void operator()() {} // expected-error {{'reqd_work_group_size' attribute requires exactly 3 arguments}} }; +class Functor33 { +public: + [[intel::reqd_work_group_size(32, 1, -1)]] void operator()() {} // expected-error {{'reqd_work_group_size' attribute requires a non-negative integral compile time constant expression}} +}; #endif class Functor16 { @@ -41,6 +45,11 @@ class Functor16 { [[intel::reqd_work_group_size(16)]] void operator()() {} }; +class Functor64 { +public: + [[intel::reqd_work_group_size(64, 64)]] void operator()() {} +}; + class Functor16x16x16 { public: [[intel::reqd_work_group_size(16, 16, 16)]] void operator()() {} From fb7e14eb90a7e904cd828480e1bf7622db72fefc Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Mon, 16 Mar 2020 18:25:17 +0300 Subject: [PATCH 13/26] [SYCL] Support intel::reqd_work_group_size (fix #10.2) Signed-off-by: Aleksander Fadeev --- .../test/SemaSYCL/intel-reqd-work-group-size.cpp | 16 ++++++++-------- 1 file changed, 8 insertions(+), 8 deletions(-) diff --git a/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp b/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp index 1476d273a0c51..046293b889927 100644 --- a/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp +++ b/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp @@ -6,7 +6,7 @@ // expected-no-diagnostics class Functor { public: - [[intel::reqd_work_group_size(4, 1, 1)]] void operator()() {} + [[intel::reqd_work_group_size(4)]] void operator()() {} }; template @@ -19,14 +19,14 @@ void bar() { kernel(f); } #else -[[intel::reqd_work_group_size(4, 1, 1)]] void f4x1x1() {} // expected-note {{conflicting attribute is here}} +[[intel::reqd_work_group_size(4)]] void f4x1x1() {} // expected-note {{conflicting attribute is here}} // expected-note@-1 {{conflicting attribute is here}} -[[intel::reqd_work_group_size(32, 1, 1)]] void f32x1x1() {} // expected-note {{conflicting attribute is here}} +[[intel::reqd_work_group_size(32)]] void f32x1x1() {} // expected-note {{conflicting attribute is here}} -[[intel::reqd_work_group_size(16, 1, 1)]] void f16x1x1() {} // expected-note {{conflicting attribute is here}} -[[intel::reqd_work_group_size(16, 16, 1)]] void f16x16x1() {} // expected-note {{conflicting attribute is here}} +[[intel::reqd_work_group_size(16)]] void f16x1x1() {} // expected-note {{conflicting attribute is here}} +[[intel::reqd_work_group_size(16)]] void f16x16x1() {} // expected-note {{conflicting attribute is here}} -[[intel::reqd_work_group_size(32, 32, 1)]] void f32x32x1() {} // expected-note {{conflicting attribute is here}} +[[intel::reqd_work_group_size(32, 32)]] void f32x32x1() {} // expected-note {{conflicting attribute is here}} [[intel::reqd_work_group_size(32, 32, 32)]] void f32x32x32() {} // expected-note {{conflicting attribute is here}} #ifdef TRIGGER_ERROR @@ -36,7 +36,7 @@ class Functor32 { }; class Functor33 { public: - [[intel::reqd_work_group_size(32, 1, -1)]] void operator()() {} // expected-error {{'reqd_work_group_size' attribute requires a non-negative integral compile time constant expression}} + [[intel::reqd_work_group_size(32, -4)]] void operator()() {} // expected-error {{'reqd_work_group_size' attribute requires a non-negative integral compile time constant expression}} }; #endif @@ -57,7 +57,7 @@ class Functor16x16x16 { class Functor8 { // expected-error {{conflicting attributes applied to a SYCL kernel}} public: - [[intel::reqd_work_group_size(8, 1, 1)]] void operator()() { // expected-note {{conflicting attribute is here}} + [[intel::reqd_work_group_size(8)]] void operator()() { // expected-note {{conflicting attribute is here}} f4x1x1(); } }; From 8a3f1b80fc9b43467792e1724d13c2de0d45304e Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Tue, 17 Mar 2020 17:53:47 +0300 Subject: [PATCH 14/26] The test commit, didn' have practical benefit Signed-off-by: Aleksander Fadeev --- clang/lib/Sema/SemaDeclAttr.cpp | 2 ++ clang/test/SemaSYCL/intel-reqd-work-group-size.cpp | 2 +- 2 files changed, 3 insertions(+), 1 deletion(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 16d44f86c3354..53039ff18e0a7 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -2923,6 +2923,8 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { uint32_t WGSize[3]; if (AL.getAttributeSpellingListIndex() == 2) { + WorkGroupAttr *Ex = D->getAttr + std::string str(Ex->getSpelling()); WGSize[1] = ReqdWorkGroupSizeAttr::DefaultYDim; WGSize[2] = ReqdWorkGroupSizeAttr::DefaultZDim; } else if (!checkAttributeNumArgs(S, AL, 3)) diff --git a/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp b/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp index 046293b889927..999314408cd52 100644 --- a/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp +++ b/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp @@ -6,7 +6,7 @@ // expected-no-diagnostics class Functor { public: - [[intel::reqd_work_group_size(4)]] void operator()() {} + [[intel::reqd_work_group_size(4)]] [[intel::reqd_work_group_size(4)]] void operator()() {} }; template From dc0eaac2aaf1a4d4a9a1b80be9287ad7a556677e Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Tue, 17 Mar 2020 22:23:31 +0300 Subject: [PATCH 15/26] Implementing DefaultUnsignedArgument class Signed-off-by: Aleksander Fadeev --- clang/include/clang/Basic/Attr.td | 16 ++++++++++------ clang/lib/Sema/SemaDeclAttr.cpp | 4 +--- .../test/SemaSYCL/intel-reqd-work-group-size.cpp | 2 +- clang/utils/TableGen/ClangAttrEmitter.cpp | 3 +++ 4 files changed, 15 insertions(+), 10 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index e1008ba2964fa..00549976d9514 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -236,6 +236,10 @@ class DefaultIntArgument : IntArgument { int Default = default; } +class DefaultUnsignedArgument : UnsignedArgument { + int Default = default; +} + // This argument is more complex, it includes the enumerator type name, // a list of strings to accept, and a list of enumerators to map them to. class EnumArgument values, @@ -2413,14 +2417,14 @@ def NoDeref : TypeAttr { let Documentation = [NoDerefDocs]; } -//Default arguments in ReqWorkGroupSize can be used only with -//intel::reqd_work_group_size spelling. +// Default arguments in ReqWorkGroupSize can be used only with +// intel::reqd_work_group_size spelling. def ReqdWorkGroupSize : InheritableAttr { let Spellings = [GNU<"reqd_work_group_size">, - CXX11<"cl","reqd_work_group_size">, - CXX11<"intel","reqd_work_group_size">]; - let Args = [IntArgument<"XDim">, DefaultIntArgument<"YDim", 1>, - DefaultIntArgument<"ZDim", 1>]; + CXX11<"intel","reqd_work_group_size">, + CXX11<"cl","reqd_work_group_size">]; + let Args = [UnsignedArgument<"XDim">, DefaultUnsignedArgument<"YDim", 1>, + DefaultUnsignedArgument<"ZDim", 1>]; let Subjects = SubjectList<[Function], ErrorDiag>; let Documentation = [ReqdWorkGroupSizeAttrDocs]; } diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 53039ff18e0a7..499f4ec56271d 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -2922,9 +2922,7 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { return; uint32_t WGSize[3]; - if (AL.getAttributeSpellingListIndex() == 2) { - WorkGroupAttr *Ex = D->getAttr - std::string str(Ex->getSpelling()); + if (AL.getNormalizedFullName() == "intel::reqd_work_group_size") { WGSize[1] = ReqdWorkGroupSizeAttr::DefaultYDim; WGSize[2] = ReqdWorkGroupSizeAttr::DefaultZDim; } else if (!checkAttributeNumArgs(S, AL, 3)) diff --git a/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp b/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp index 999314408cd52..046293b889927 100644 --- a/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp +++ b/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp @@ -6,7 +6,7 @@ // expected-no-diagnostics class Functor { public: - [[intel::reqd_work_group_size(4)]] [[intel::reqd_work_group_size(4)]] void operator()() {} + [[intel::reqd_work_group_size(4)]] void operator()() {} }; template diff --git a/clang/utils/TableGen/ClangAttrEmitter.cpp b/clang/utils/TableGen/ClangAttrEmitter.cpp index ce95d78d7db03..f314036f7e946 100644 --- a/clang/utils/TableGen/ClangAttrEmitter.cpp +++ b/clang/utils/TableGen/ClangAttrEmitter.cpp @@ -1304,6 +1304,9 @@ createArgument(const Record &Arg, StringRef Attr, else if (ArgName == "DefaultIntArgument") Ptr = std::make_unique( Arg, Attr, "int", Arg.getValueAsInt("Default")); + else if (ArgName == "DefaultUnsignedArgument") + Ptr = std::make_unique( + Arg, Attr, "unsigned", Arg.getValueAsInt("Default")); else if (ArgName == "IntArgument") Ptr = std::make_unique(Arg, Attr, "int"); else if (ArgName == "StringArgument") From 04a4c835c85fc2e3a63f67eb83ece45dac7598b0 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Wed, 18 Mar 2020 14:54:31 +0300 Subject: [PATCH 16/26] Impelmention getAttributeSpellingListIndex() Signed-off-by: Aleksander Fadeev --- clang/lib/Sema/SemaDeclAttr.cpp | 4 +++- clang/test/SemaSYCL/intel-reqd-work-group-size.cpp | 2 +- 2 files changed, 4 insertions(+), 2 deletions(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 499f4ec56271d..157ca53004ed5 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -2922,7 +2922,9 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { return; uint32_t WGSize[3]; - if (AL.getNormalizedFullName() == "intel::reqd_work_group_size") { + if (AL.getKind() == ParsedAttr::AT_ReqdWorkGroupSize && + AL.getAttributeSpellingListIndex() == + ReqdWorkGroupSizeAttr::CXX11_intel_reqd_work_group_size) { WGSize[1] = ReqdWorkGroupSizeAttr::DefaultYDim; WGSize[2] = ReqdWorkGroupSizeAttr::DefaultZDim; } else if (!checkAttributeNumArgs(S, AL, 3)) diff --git a/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp b/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp index 046293b889927..1691ecc45298b 100644 --- a/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp +++ b/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp @@ -24,7 +24,7 @@ void bar() { [[intel::reqd_work_group_size(32)]] void f32x1x1() {} // expected-note {{conflicting attribute is here}} [[intel::reqd_work_group_size(16)]] void f16x1x1() {} // expected-note {{conflicting attribute is here}} -[[intel::reqd_work_group_size(16)]] void f16x16x1() {} // expected-note {{conflicting attribute is here}} +[[intel::reqd_work_group_size(16, 16)]] void f16x16x1() {} // expected-note {{conflicting attribute is here}} [[intel::reqd_work_group_size(32, 32)]] void f32x32x1() {} // expected-note {{conflicting attribute is here}} [[intel::reqd_work_group_size(32, 32, 32)]] void f32x32x32() {} // expected-note {{conflicting attribute is here}} From fc108050f94a754c2453db448322580ab868f8dd Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Wed, 18 Mar 2020 16:10:48 +0300 Subject: [PATCH 17/26] Resolving cinflict ClangAttrEmitter.cpp Signed-off-by: Aleksander Fadeev --- clang/utils/TableGen/ClangAttrEmitter.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/utils/TableGen/ClangAttrEmitter.cpp b/clang/utils/TableGen/ClangAttrEmitter.cpp index f314036f7e946..504ededa36edf 100644 --- a/clang/utils/TableGen/ClangAttrEmitter.cpp +++ b/clang/utils/TableGen/ClangAttrEmitter.cpp @@ -2332,7 +2332,7 @@ static void emitAttributes(RecordKeeper &Records, raw_ostream &OS, SemanticSpellingMap SemanticToSyntacticMap; std::string SpellingEnum; - if (!ElideSpelling) + if (Spellings.size() > 1) SpellingEnum = CreateSemanticSpellings(Spellings, SemanticToSyntacticMap); if (Header) OS << SpellingEnum; From 971eaa469f444e60f5750e9dbff4b2e91506e7fc Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Thu, 19 Mar 2020 00:58:35 +0300 Subject: [PATCH 18/26] The AttrDocs.td changing and other little fixes Signed-off-by: Aleksander Fadeev --- clang/include/clang/Basic/AttrDocs.td | 1 + clang/lib/Sema/SemaDeclAttr.cpp | 3 +-- clang/test/SemaSYCL/intel-reqd-work-group-size.cpp | 6 +++--- 3 files changed, 5 insertions(+), 5 deletions(-) diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 44e4e359856ec..2c3cfa9c1469f 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -1991,6 +1991,7 @@ Applies to a device function/lambda function. Indicates the dimensions of a work group. Values must be positive integers. Default arguments equaling to one can be used for second and third dimensions if intel::reqd_work_group_size spelling was applied. +This attribute is also valid for OpenCL C. }]; } diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 157ca53004ed5..aadd7fed2ac2d 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -2931,8 +2931,7 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { return; for (unsigned i = 0; i < 3; ++i) { - if (i < AL.getNumArgs()) - if (!checkUInt32Argument(S, AL, AL.getArgAsExpr(i), WGSize[i], i, + if (i < AL.getNumArgs() && !checkUInt32Argument(S, AL, AL.getArgAsExpr(i), WGSize[i], i, /*StrictlyUnsigned=*/true)) return; if (WGSize[i] == 0) { diff --git a/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp b/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp index 1691ecc45298b..128da59fa6d4c 100644 --- a/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp +++ b/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp @@ -1,6 +1,6 @@ -// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -verify -DTRIGGER_ERROR %s -// RUN: %clang_cc1 -fsycl-is-device -ast-dump %s | FileCheck %s -// RUN: %clang_cc1 -fsycl-is-host -fsyntax-only -verify %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -verify -DTRIGGER_ERROR %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -ast-dump %s | FileCheck %s +// RUN: %clang_cc1 -fsycl -fsycl-is-host -fsyntax-only -verify %s #ifndef __SYCL_DEVICE_ONLY__ // expected-no-diagnostics From 65fc0bae4e581e91ed99ebe60dcf089bed1ca144 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Thu, 19 Mar 2020 01:17:58 +0300 Subject: [PATCH 19/26] Formatting Signed-off-by: Aleksander Fadeev --- clang/lib/Sema/SemaDeclAttr.cpp | 7 ++++--- clang/test/SemaSYCL/intel-reqd-work-group-size.cpp | 4 ++-- 2 files changed, 6 insertions(+), 5 deletions(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index aadd7fed2ac2d..b6394311b74ab 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -2931,9 +2931,10 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { return; for (unsigned i = 0; i < 3; ++i) { - if (i < AL.getNumArgs() && !checkUInt32Argument(S, AL, AL.getArgAsExpr(i), WGSize[i], i, - /*StrictlyUnsigned=*/true)) - return; + if (i < AL.getNumArgs() && + !checkUInt32Argument(S, AL, AL.getArgAsExpr(i), WGSize[i], i, + /*StrictlyUnsigned=*/true)) + return; if (WGSize[i] == 0) { S.Diag(AL.getLoc(), diag::err_attribute_argument_is_zero) << AL << AL.getArgAsExpr(i)->getSourceRange(); diff --git a/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp b/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp index 128da59fa6d4c..be970cda66e18 100644 --- a/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp +++ b/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp @@ -23,10 +23,10 @@ void bar() { // expected-note@-1 {{conflicting attribute is here}} [[intel::reqd_work_group_size(32)]] void f32x1x1() {} // expected-note {{conflicting attribute is here}} -[[intel::reqd_work_group_size(16)]] void f16x1x1() {} // expected-note {{conflicting attribute is here}} +[[intel::reqd_work_group_size(16)]] void f16x1x1() {} // expected-note {{conflicting attribute is here}} [[intel::reqd_work_group_size(16, 16)]] void f16x16x1() {} // expected-note {{conflicting attribute is here}} -[[intel::reqd_work_group_size(32, 32)]] void f32x32x1() {} // expected-note {{conflicting attribute is here}} +[[intel::reqd_work_group_size(32, 32)]] void f32x32x1() {} // expected-note {{conflicting attribute is here}} [[intel::reqd_work_group_size(32, 32, 32)]] void f32x32x32() {} // expected-note {{conflicting attribute is here}} #ifdef TRIGGER_ERROR From 2776fe50b6ff2a40b1f26ffed407f0fdb5d1985d Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Thu, 19 Mar 2020 17:29:30 +0300 Subject: [PATCH 20/26] AttrDocs.td and formatting Signed-off-by: Aleksander Fadeev --- clang/include/clang/Basic/AttrDocs.td | 28 +++++++++++++++---- .../SemaSYCL/intel-reqd-work-group-size.cpp | 4 +-- clang/utils/TableGen/ClangAttrEmitter.cpp | 6 ++-- 3 files changed, 28 insertions(+), 10 deletions(-) diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 2c3cfa9c1469f..11686f8fc976b 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -1987,11 +1987,29 @@ def ReqdWorkGroupSizeAttrDocs : Documentation { let Category = DocCatFunction; let Heading = "cl::reqd_work_group_size, intel::reqd_work_group_size"; let Content = [{ -Applies to a device function/lambda function. Indicates the dimensions -of a work group. Values must be positive integers. Default arguments -equaling to one can be used for second and third dimensions if -intel::reqd_work_group_size spelling was applied. -This attribute is also valid for OpenCL C. +This attribute is documented by both OpenCL and SYCL standards +and allows to specify exact *local_work_size* which must be used as +argument to **clEnqueueNDRangeKernel** (in OpenCL) or to +**parallel_for** in SYCL. This allows the compiler to optimize the +generated code appropriately for this kernel. + +While semantic of this attribute is the same between OpenCL and SYCL, +spelling is a bit different: + +SYCL 1.2.1 describes ``[[cl::reqd_work_group_size(X, Y, Z)]]`` spelling: this +attribute is legal on device functions and their specification is propagated +down to any caller of those device functions, such that the kernel attributes +are the sum of all the kernel attributes of all device functions called. +See section 6.7 Attributes for more details. + +As Intel extension, ``[[intel::reqd_work_group_size(X, Y, Z)]]`` spelling is allowed +which features optional arguments `Y` and `Z`, which simplifies its usage if +only 1- or 2-dimensional ND-range is assumed by a programmer. These arguments +defaults to ``1``. + +In OpenCL C, this attribute is available in GNU spelling +(``__attribute__((reqd_work_group_size(X, Y, Z)))``), see, for example, section +6.7.2 Optional Attribute Qualifiers of OpenCL 1.2 specification }]; } diff --git a/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp b/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp index be970cda66e18..361b3d083c0ef 100644 --- a/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp +++ b/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp @@ -38,7 +38,7 @@ class Functor33 { public: [[intel::reqd_work_group_size(32, -4)]] void operator()() {} // expected-error {{'reqd_work_group_size' attribute requires a non-negative integral compile time constant expression}} }; -#endif +#endif // TRIGGER_ERROR class Functor16 { public: @@ -120,7 +120,7 @@ void bar() { f32x32x32(); }); -#endif +#endif // TRIGGER_ERROR } // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name1 diff --git a/clang/utils/TableGen/ClangAttrEmitter.cpp b/clang/utils/TableGen/ClangAttrEmitter.cpp index 504ededa36edf..37594e7b679c9 100644 --- a/clang/utils/TableGen/ClangAttrEmitter.cpp +++ b/clang/utils/TableGen/ClangAttrEmitter.cpp @@ -1304,9 +1304,6 @@ createArgument(const Record &Arg, StringRef Attr, else if (ArgName == "DefaultIntArgument") Ptr = std::make_unique( Arg, Attr, "int", Arg.getValueAsInt("Default")); - else if (ArgName == "DefaultUnsignedArgument") - Ptr = std::make_unique( - Arg, Attr, "unsigned", Arg.getValueAsInt("Default")); else if (ArgName == "IntArgument") Ptr = std::make_unique(Arg, Attr, "int"); else if (ArgName == "StringArgument") @@ -1315,6 +1312,9 @@ createArgument(const Record &Arg, StringRef Attr, Ptr = std::make_unique(Arg, Attr); else if (ArgName == "UnsignedArgument") Ptr = std::make_unique(Arg, Attr, "unsigned"); + else if (ArgName == "DefaultUnsignedArgument") + Ptr = std::make_unique(Arg, Attr, "unsigned", + Arg.getValueAsInt("Default")); else if (ArgName == "VariadicUnsignedArgument") Ptr = std::make_unique(Arg, Attr, "unsigned"); else if (ArgName == "VariadicStringArgument") From 545cb9f4d5543de712161a191d6c768b53abfaa3 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Thu, 19 Mar 2020 19:34:14 +0300 Subject: [PATCH 21/26] adding GNU in AttrDocs.td Signed-off-by: Aleksander Fadeev --- clang/include/clang/Basic/AttrDocs.td | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 11686f8fc976b..5fe898640f50f 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -1985,7 +1985,8 @@ device kernel, the attribute is ignored and it is not propagated to a kernel. def ReqdWorkGroupSizeAttrDocs : Documentation { let Category = DocCatFunction; - let Heading = "cl::reqd_work_group_size, intel::reqd_work_group_size"; + let Heading = "GNU, CXX11"; let Content = [{ This attribute is documented by both OpenCL and SYCL standards and allows to specify exact *local_work_size* which must be used as From 60541fd018fc9d870ff297ebcfc6e78a6af42534 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Fri, 20 Mar 2020 10:51:47 +0300 Subject: [PATCH 22/26] Docs.td modifying Signed-off-by: Aleksander Fadeev --- clang/include/clang/Basic/AttrDocs.td | 12 ++++++------ clang/test/SemaSYCL/intel-reqd-work-group-size.cpp | 4 ++-- 2 files changed, 8 insertions(+), 8 deletions(-) diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 5fe898640f50f..2cd6f84941dc8 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -1985,22 +1985,22 @@ device kernel, the attribute is ignored and it is not propagated to a kernel. def ReqdWorkGroupSizeAttrDocs : Documentation { let Category = DocCatFunction; - let Heading = "GNU, CXX11"; + let Heading = "__attribute__((reqd_work_group_size)), [[cl::reqd_work_group_size]], + [[intel::reqd_work_group_size]]"; let Content = [{ This attribute is documented by both OpenCL and SYCL standards and allows to specify exact *local_work_size* which must be used as argument to **clEnqueueNDRangeKernel** (in OpenCL) or to **parallel_for** in SYCL. This allows the compiler to optimize the -generated code appropriately for this kernel. +generated code appropriately for the kernel to which attribute is applied. While semantic of this attribute is the same between OpenCL and SYCL, spelling is a bit different: SYCL 1.2.1 describes ``[[cl::reqd_work_group_size(X, Y, Z)]]`` spelling: this -attribute is legal on device functions and their specification is propagated -down to any caller of those device functions, such that the kernel attributes -are the sum of all the kernel attributes of all device functions called. +attribute is legal on device functions and is propagated down to any caller of +those device functions, such that the kernel attributes are the sum of all +attributes of all device functions called in this kernel. See section 6.7 Attributes for more details. As Intel extension, ``[[intel::reqd_work_group_size(X, Y, Z)]]`` spelling is allowed diff --git a/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp b/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp index 361b3d083c0ef..432b981c58a85 100644 --- a/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp +++ b/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp @@ -23,10 +23,10 @@ void bar() { // expected-note@-1 {{conflicting attribute is here}} [[intel::reqd_work_group_size(32)]] void f32x1x1() {} // expected-note {{conflicting attribute is here}} -[[intel::reqd_work_group_size(16)]] void f16x1x1() {} // expected-note {{conflicting attribute is here}} +[[intel::reqd_work_group_size(16)]] void f16x1x1() {} // expected-note {{conflicting attribute is here}} [[intel::reqd_work_group_size(16, 16)]] void f16x16x1() {} // expected-note {{conflicting attribute is here}} -[[intel::reqd_work_group_size(32, 32)]] void f32x32x1() {} // expected-note {{conflicting attribute is here}} +[[intel::reqd_work_group_size(32, 32)]] void f32x32x1() {} // expected-note {{conflicting attribute is here}} [[intel::reqd_work_group_size(32, 32, 32)]] void f32x32x32() {} // expected-note {{conflicting attribute is here}} #ifdef TRIGGER_ERROR From 50e397d9ad3190d6232e56fe922a25e054cb793e Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Fri, 20 Mar 2020 14:48:51 +0300 Subject: [PATCH 23/26] fix AttrDocs.td heading Signed-off-by: Aleksander Fadeev --- clang/include/clang/Basic/AttrDocs.td | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 2cd6f84941dc8..4048e3800b799 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -1985,8 +1985,7 @@ device kernel, the attribute is ignored and it is not propagated to a kernel. def ReqdWorkGroupSizeAttrDocs : Documentation { let Category = DocCatFunction; - let Heading = "__attribute__((reqd_work_group_size)), [[cl::reqd_work_group_size]], - [[intel::reqd_work_group_size]]"; + let Heading = "reqd_work_group_size" let Content = [{ This attribute is documented by both OpenCL and SYCL standards and allows to specify exact *local_work_size* which must be used as From cf200c4ade022fdc711bc2d90c53e891582e5e93 Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Fri, 20 Mar 2020 16:20:05 +0300 Subject: [PATCH 24/26] Minifix AttrDocs Signed-off-by: Aleksander Fadeev --- clang/include/clang/Basic/AttrDocs.td | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 4048e3800b799..83cf7c5be7fb6 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -1985,7 +1985,7 @@ device kernel, the attribute is ignored and it is not propagated to a kernel. def ReqdWorkGroupSizeAttrDocs : Documentation { let Category = DocCatFunction; - let Heading = "reqd_work_group_size" + let Heading = "reqd_work_group_size"; let Content = [{ This attribute is documented by both OpenCL and SYCL standards and allows to specify exact *local_work_size* which must be used as From 23e0bbbafb38380412cbbc2c1ebf4f91129abefd Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Mon, 23 Mar 2020 15:16:29 +0300 Subject: [PATCH 25/26] Formatting Signed-off-by: Aleksander Fadeev --- clang/test/SemaSYCL/intel-reqd-work-group-size.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp b/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp index 432b981c58a85..361b3d083c0ef 100644 --- a/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp +++ b/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp @@ -23,10 +23,10 @@ void bar() { // expected-note@-1 {{conflicting attribute is here}} [[intel::reqd_work_group_size(32)]] void f32x1x1() {} // expected-note {{conflicting attribute is here}} -[[intel::reqd_work_group_size(16)]] void f16x1x1() {} // expected-note {{conflicting attribute is here}} +[[intel::reqd_work_group_size(16)]] void f16x1x1() {} // expected-note {{conflicting attribute is here}} [[intel::reqd_work_group_size(16, 16)]] void f16x16x1() {} // expected-note {{conflicting attribute is here}} -[[intel::reqd_work_group_size(32, 32)]] void f32x32x1() {} // expected-note {{conflicting attribute is here}} +[[intel::reqd_work_group_size(32, 32)]] void f32x32x1() {} // expected-note {{conflicting attribute is here}} [[intel::reqd_work_group_size(32, 32, 32)]] void f32x32x32() {} // expected-note {{conflicting attribute is here}} #ifdef TRIGGER_ERROR From 6d2bce7e171d6e2c7c3acf9b336709d7691a6f9a Mon Sep 17 00:00:00 2001 From: Aleksander Fadeev Date: Mon, 23 Mar 2020 16:38:29 +0300 Subject: [PATCH 26/26] AttrDocs.td fix Signed-off-by: Aleksander Fadeev --- clang/include/clang/Basic/AttrDocs.td | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 83cf7c5be7fb6..23bb435017d13 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -2003,13 +2003,13 @@ attributes of all device functions called in this kernel. See section 6.7 Attributes for more details. As Intel extension, ``[[intel::reqd_work_group_size(X, Y, Z)]]`` spelling is allowed -which features optional arguments `Y` and `Z`, which simplifies its usage if +which features optional arguments `Y` and `Z`, those simplifies its usage if only 1- or 2-dimensional ND-range is assumed by a programmer. These arguments defaults to ``1``. In OpenCL C, this attribute is available in GNU spelling -(``__attribute__((reqd_work_group_size(X, Y, Z)))``), see, for example, section -6.7.2 Optional Attribute Qualifiers of OpenCL 1.2 specification +(``__attribute__((reqd_work_group_size(X, Y, Z)))``), see section +6.7.2 Optional Attribute Qualifiers of OpenCL 1.2 specification for details. }]; }