From da9476559fe19d36f56fda228bf288b20d8b7009 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Mon, 30 Aug 2021 10:26:08 -0700 Subject: [PATCH 1/4] [SYCL] Remove deprecated intel::reqd_work_group_size attribute Signed-off-by: Soumi Manna --- clang/include/clang/Basic/Attr.td | 5 ++--- clang/include/clang/Basic/AttrDocs.td | 16 ++++------------ clang/lib/Sema/SemaDeclAttr.cpp | 8 -------- .../intel-max-global-work-dim-device.cpp | 4 +--- ...l-reqd-work-group-size-device-direct-prop.cpp | 3 +-- .../test/SemaSYCL/num_simd_work_items_device.cpp | 10 ++++------ ...evice-intel-reqd-work-group-size-template.cpp | 2 +- 7 files changed, 13 insertions(+), 35 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index a5df8351e8c8a..0c3d17f44564c 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -2985,11 +2985,10 @@ def NoDeref : TypeAttr { let Documentation = [NoDerefDocs]; } -// Default arguments can only be used with the intel::reqd_work_group_size or -// sycl::reqd_work_group_size spellings. +// Default arguments can only be used with the sycl::reqd_work_group_size +// spelling. def ReqdWorkGroupSize : InheritableAttr { let Spellings = [GNU<"reqd_work_group_size">, - CXX11<"intel", "reqd_work_group_size">, CXX11<"cl", "reqd_work_group_size">, CXX11<"sycl", "reqd_work_group_size">]; let Args = [ExprArgument<"XDim">, diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 5cb6f853910b1..f19ec7f150ddc 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -2673,10 +2673,10 @@ allows the Y and Z arguments to be optional. If not provided by the user, the value of Y and Z defaults to 1. See section 5.8.1 Kernel Attributes for more details. -In SYCL 1.2.1 mode, the ``intel::reqd_work_group_size``, -``cl::reqd_work_group_size``, and ``sycl::reqd_work_group_size`` attributes are -propagated from the function they are applied to onto the kernel which calls the -function. In SYCL 2020 mode, the attributes are not propagated to the kernel. +In SYCL 1.2.1 mode, the ``cl::reqd_work_group_size`` and +``sycl::reqd_work_group_size`` attributes are propagated from the function they +are applied to onto the kernel which calls the function. In SYCL 2020 mode, the +attributes are not propagated to the kernel. .. code-block:: c++ @@ -2700,14 +2700,6 @@ The ``[[cl::reqd_work_group_size(X, Y, Z)]]`` and ``__attribute__((reqd_work_group_size(X, Y, Z)))`` spellings are both deprecated in SYCL 2020. -As an Intel extension, the ``[[intel::reqd_work_group_size(X, Y, Z)]]`` -spelling is supported with the same semantics as the -``[[sycl::reqd_work_group_size(X, Y, Z)]]`` spelling. - -The ``[[intel::reqd_work_group_size(X, Y, Z)]]`` attribute spelling -is deprecated in favor of the SYCL 2020 attribute spelling -``[[sycl::reqd_work_group_size]]``. - In OpenCL C, this attribute is available with the GNU spelling (``__attribute__((reqd_work_group_size(X, Y, Z)))``), see section 6.7.2 Optional Attribute Qualifiers of OpenCL 1.2 specification for details. diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index bcbac73fd300c..0f945edfae820 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -338,14 +338,6 @@ void Sema::CheckDeprecatedSYCLAttributeSpelling(const ParsedAttr &A, return; } - // Deprecate [[intel::reqd_work_group_size]] attribute spelling in favor - // of the SYCL 2020 attribute spelling [[sycl::reqd_work_group_size]]. - if (A.getKind() == ParsedAttr::AT_ReqdWorkGroupSize && A.hasScope() && - A.getScopeName()->isStr("intel")) { - DiagnoseDeprecatedAttribute(A, "sycl", NewName); - return; - } - // All GNU-style spellings are deprecated in favor of a C++-style spelling. if (A.getSyntax() == ParsedAttr::AS_GNU) { // Note: we cannot suggest an automatic fix-it because GNU-style diff --git a/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp index 507001a889432..3dac19194380c 100644 --- a/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp +++ b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp @@ -166,9 +166,7 @@ struct TRIFuncObjBad5 { }; struct TRIFuncObjBad6 { - [[intel::reqd_work_group_size(4)]] // expected-error{{all 'reqd_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}} \ - // expected-warning {{attribute 'intel::reqd_work_group_size' is deprecated}} \ - // expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}} + [[sycl::reqd_work_group_size(4)]] // expected-error{{all 'reqd_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}} [[intel::max_global_work_dim(0)]] void operator()() const {} }; diff --git a/clang/test/SemaSYCL/intel-reqd-work-group-size-device-direct-prop.cpp b/clang/test/SemaSYCL/intel-reqd-work-group-size-device-direct-prop.cpp index 8909245dd2b19..1b3903849ce87 100644 --- a/clang/test/SemaSYCL/intel-reqd-work-group-size-device-direct-prop.cpp +++ b/clang/test/SemaSYCL/intel-reqd-work-group-size-device-direct-prop.cpp @@ -56,8 +56,7 @@ class Functor64 { class Functor16x16x16 { public: - [[intel::reqd_work_group_size(16, 16, 16)]] void operator()() const {} // expected-warning {{attribute 'intel::reqd_work_group_size' is deprecated}} \ - // expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}} + [[sycl::reqd_work_group_size(16, 16, 16)]] void operator()() const {} }; class FunctorAttr { diff --git a/clang/test/SemaSYCL/num_simd_work_items_device.cpp b/clang/test/SemaSYCL/num_simd_work_items_device.cpp index 858fe713e412b..48553b5fd0ac0 100644 --- a/clang/test/SemaSYCL/num_simd_work_items_device.cpp +++ b/clang/test/SemaSYCL/num_simd_work_items_device.cpp @@ -42,7 +42,7 @@ struct FuncObj { }; #ifdef TRIGGER_ERROR -// If the declaration has a [[intel::reqd_work_group_size]] +// If the declaration has a [[sycl::reqd_work_group_size]] // [[cl::reqd_work_group_size]] attribute, tests that check // if the work group size attribute argument (the last argument) // can be evenly divided by the [[intel::num_simd_work_items()]] attribute. @@ -60,7 +60,7 @@ struct TRIFuncObjBad2 { operator()() const {} }; -// Tests for the default values of [[intel::reqd_work_group_size()]]. +// Tests for the default values of [[sycl::reqd_work_group_size()]]. // FIXME: This should be accepted instead of error which turns out to be // an implementation bug that shouldn't be visible to the user as there @@ -131,9 +131,7 @@ struct TRIFuncObjBad8 { [[intel::num_simd_work_items(2)]] // expected-error{{'num_simd_work_items' attribute must evenly divide the work-group size for the 'reqd_work_group_size' attribute}} [[sycl::reqd_work_group_size(4, 2, 3)]] void func1(); // expected-note{{conflicting attribute is here}} -[[intel::reqd_work_group_size(4, 2, 3)]] // expected-note{{conflicting attribute is here}} \ - // expected-warning {{attribute 'intel::reqd_work_group_size' is deprecated}} \ - // expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}} +[[sycl::reqd_work_group_size(4, 2, 3)]] // expected-note{{conflicting attribute is here}} [[intel::num_simd_work_items(2)]] void func2(); // expected-error{{'num_simd_work_items' attribute must evenly divide the work-group size for the 'reqd_work_group_size' attribute}} [[intel::num_simd_work_items(2)]] // expected-error{{'num_simd_work_items' attribute must evenly divide the work-group size for the 'reqd_work_group_size' attribute}} @@ -211,7 +209,7 @@ struct TRIFuncObjBad18 { }; #endif // TRIGGER_ERROR -// If the declaration has a [[intel::reqd_work_group_size()]] +// If the declaration has a [[sycl::reqd_work_group_size()]] // or [[cl::reqd_work_group_size()]] or // __attribute__((reqd_work_group_size)) attribute, check to see // if the last argument can be evenly divided by the diff --git a/clang/test/SemaSYCL/sycl-device-intel-reqd-work-group-size-template.cpp b/clang/test/SemaSYCL/sycl-device-intel-reqd-work-group-size-template.cpp index 680d66014264c..12ffca8ab74f2 100644 --- a/clang/test/SemaSYCL/sycl-device-intel-reqd-work-group-size-template.cpp +++ b/clang/test/SemaSYCL/sycl-device-intel-reqd-work-group-size-template.cpp @@ -1,6 +1,6 @@ // RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -ast-dump -verify -pedantic %s | FileCheck %s -// Test that checks template parameter support for 'intel::reqd_work_group_size' attribute on sycl device. +// Test that checks template parameter support for 'sycl::reqd_work_group_size' attribute on sycl device. // Test that checks wrong function template instantiation and ensures that the type // is checked properly when instantiating from the template definition. From 6cb02f2fcf4decffebb9fbc6ea556b60c2904c0a Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Mon, 30 Aug 2021 12:17:16 -0700 Subject: [PATCH 2/4] update test with correct attribute spelling Signed-off-by: Soumi Manna --- sycl/test/kernel_param/attr.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test/kernel_param/attr.cpp b/sycl/test/kernel_param/attr.cpp index 1037d23f54686..104cf53b4bcf7 100755 --- a/sycl/test/kernel_param/attr.cpp +++ b/sycl/test/kernel_param/attr.cpp @@ -12,7 +12,7 @@ int main() { queue myQueue; myQueue.submit([&](handler &cgh) { cgh.parallel_for(Size, [=](item<1> ITEM) - [[intel::reqd_work_group_size(4)]]{}); + [[sycl::reqd_work_group_size(4)]]{}); }); } From 2dda79cb1e4cf9b97f7959a3da932682f2e2fb9b Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Tue, 31 Aug 2021 19:19:29 -0700 Subject: [PATCH 3/4] address review comments Signed-off-by: Soumi Manna --- clang/test/SemaSYCL/num_simd_work_items_device.cpp | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/clang/test/SemaSYCL/num_simd_work_items_device.cpp b/clang/test/SemaSYCL/num_simd_work_items_device.cpp index 48553b5fd0ac0..926829b93de2d 100644 --- a/clang/test/SemaSYCL/num_simd_work_items_device.cpp +++ b/clang/test/SemaSYCL/num_simd_work_items_device.cpp @@ -131,9 +131,12 @@ struct TRIFuncObjBad8 { [[intel::num_simd_work_items(2)]] // expected-error{{'num_simd_work_items' attribute must evenly divide the work-group size for the 'reqd_work_group_size' attribute}} [[sycl::reqd_work_group_size(4, 2, 3)]] void func1(); // expected-note{{conflicting attribute is here}} -[[sycl::reqd_work_group_size(4, 2, 3)]] // expected-note{{conflicting attribute is here}} +[[sycl::reqd_work_group_size(4, 2, 3)]] // expected-note{{conflicting attribute is here}} [[intel::num_simd_work_items(2)]] void func2(); // expected-error{{'num_simd_work_items' attribute must evenly divide the work-group size for the 'reqd_work_group_size' attribute}} +[[intel::reqd_work_group_size(4, 2, 9)]] // expected-warning{{unknown attribute 'reqd_work_group_size' ignored}} +[[intel::num_simd_work_items(2)]] void ignore(); + [[intel::num_simd_work_items(2)]] // expected-error{{'num_simd_work_items' attribute must evenly divide the work-group size for the 'reqd_work_group_size' attribute}} [[cl::reqd_work_group_size(4, 2, 3)]] void func3(); // expected-note{{conflicting attribute is here}} expected-warning {{attribute 'cl::reqd_work_group_size' is deprecated}} expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}} From 2a521c27c5afa3660049656f7e8a45bde60120c0 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Tue, 31 Aug 2021 19:37:47 -0700 Subject: [PATCH 4/4] update test with removed attribute spelling Signed-off-by: Soumi Manna --- clang/test/SemaSYCL/intel-reqd-work-group-size-device.cpp | 2 ++ clang/test/SemaSYCL/num_simd_work_items_device.cpp | 5 +---- 2 files changed, 3 insertions(+), 4 deletions(-) diff --git a/clang/test/SemaSYCL/intel-reqd-work-group-size-device.cpp b/clang/test/SemaSYCL/intel-reqd-work-group-size-device.cpp index 68d299d4d3a4f..a9ce1857d779b 100644 --- a/clang/test/SemaSYCL/intel-reqd-work-group-size-device.cpp +++ b/clang/test/SemaSYCL/intel-reqd-work-group-size-device.cpp @@ -38,6 +38,8 @@ class Functor33 { [[sycl::reqd_work_group_size(32, -4)]] void operator()() const {} }; +[[intel::reqd_work_group_size(4, 2, 9)]] void unknown() {} // expected-warning{{unknown attribute 'reqd_work_group_size' ignored}} + class Functor30 { public: // expected-warning@+1 2{{implicit conversion changes signedness: 'int' to 'unsigned long long'}} diff --git a/clang/test/SemaSYCL/num_simd_work_items_device.cpp b/clang/test/SemaSYCL/num_simd_work_items_device.cpp index 926829b93de2d..f37fcaf27f84d 100644 --- a/clang/test/SemaSYCL/num_simd_work_items_device.cpp +++ b/clang/test/SemaSYCL/num_simd_work_items_device.cpp @@ -43,7 +43,7 @@ struct FuncObj { #ifdef TRIGGER_ERROR // If the declaration has a [[sycl::reqd_work_group_size]] -// [[cl::reqd_work_group_size]] attribute, tests that check +// or [[cl::reqd_work_group_size]] attribute, tests that check // if the work group size attribute argument (the last argument) // can be evenly divided by the [[intel::num_simd_work_items()]] attribute. struct TRIFuncObjBad1 { @@ -134,9 +134,6 @@ struct TRIFuncObjBad8 { [[sycl::reqd_work_group_size(4, 2, 3)]] // expected-note{{conflicting attribute is here}} [[intel::num_simd_work_items(2)]] void func2(); // expected-error{{'num_simd_work_items' attribute must evenly divide the work-group size for the 'reqd_work_group_size' attribute}} -[[intel::reqd_work_group_size(4, 2, 9)]] // expected-warning{{unknown attribute 'reqd_work_group_size' ignored}} -[[intel::num_simd_work_items(2)]] void ignore(); - [[intel::num_simd_work_items(2)]] // expected-error{{'num_simd_work_items' attribute must evenly divide the work-group size for the 'reqd_work_group_size' attribute}} [[cl::reqd_work_group_size(4, 2, 3)]] void func3(); // expected-note{{conflicting attribute is here}} expected-warning {{attribute 'cl::reqd_work_group_size' is deprecated}} expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}}