From ebe888ac6f2c0d18ad5c3d6ea03ae370ec9539a9 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Fri, 11 Jun 2021 18:00:59 -0700 Subject: [PATCH 1/5] [SYCL] Deprecate [[intel::reqd_work_group_size]] attribute spelling Since we have SYCL 2020 replacement for `intel::reqd_work_group_size` - sycl::reqd_work_group_size, it makes sense do deprecate intel one. Signed-off-by: Soumi Manna --- clang/include/clang/Basic/AttrDocs.td | 4 ++ clang/lib/Sema/SemaDeclAttr.cpp | 9 ++++ .../CodeGenSYCL/sycl-multi-kernel-attr.cpp | 6 +-- .../intel-max-global-work-dim-device.cpp | 16 +++---- .../intel-reqd-work-group-size-device.cpp | 32 +++++++------- .../intel-reqd-work-group-size-host.cpp | 8 ++-- .../SemaSYCL/num_simd_work_items_device.cpp | 42 ++++++++++--------- .../SemaSYCL/reqd-work-group-size-device.cpp | 8 ++-- ...ice-intel-max-global-work-dim-template.cpp | 10 ++--- ...ce-intel-reqd-work-group-size-template.cpp | 10 ++--- ...cl-device-num_simd_work_items-template.cpp | 6 +-- 11 files changed, 84 insertions(+), 67 deletions(-) diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index cb91cc1707b23..7fc4559ad7f1a 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -2659,6 +2659,10 @@ 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 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 2d40ea3103286..d1403ae8267ab 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -337,6 +337,15 @@ void Sema::CheckDeprecatedSYCLAttributeSpelling(const ParsedAttr &A, return; } + // Deprecate [[intel::reqd_work_group_size]] attribute spelling in favor + // of SYCL 2020 attribute spelling [[sycl::reqd_work_group_size]]. + if (A.hasScope() && A.getScopeName()->isStr("intel")) { + if (A.getKind() == ParsedAttr::AT_ReqdWorkGroupSize) { + 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/CodeGenSYCL/sycl-multi-kernel-attr.cpp b/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp index 3c7c80016cc23..fd4e014d5825a 100644 --- a/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp +++ b/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp @@ -12,17 +12,17 @@ class Functor { class Functor1 { public: - [[intel::reqd_sub_group_size(2), intel::reqd_work_group_size(64, 32, 32)]] void operator()() const {} + [[intel::reqd_sub_group_size(2), sycl::reqd_work_group_size(64, 32, 32)]] void operator()() const {} }; template class Functor2 { public: - [[intel::reqd_work_group_size(SIZE, SIZE1, SIZE2)]] void operator()() const {} + [[sycl::reqd_work_group_size(SIZE, SIZE1, SIZE2)]] void operator()() const {} }; template -[[intel::reqd_work_group_size(N, N1, N2)]] void func() {} +[[sycl::reqd_work_group_size(N, N1, N2)]] void func() {} int main() { q.submit([&](handler &h) { 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 29cc7f084e2d4..a93096acec29e 100644 --- a/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp +++ b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp @@ -67,7 +67,7 @@ struct TRIFuncObjGood2 { }; struct TRIFuncObjGood3 { - [[intel::reqd_work_group_size(1)]] + [[sycl::reqd_work_group_size(1)]] [[intel::max_global_work_dim(0)]] void operator()() const {} }; @@ -85,7 +85,7 @@ struct TRIFuncObjGood5 { }; struct TRIFuncObjGood6 { - [[intel::reqd_work_group_size(4, 1, 1)]] + [[sycl::reqd_work_group_size(4, 1, 1)]] [[intel::max_global_work_dim(3)]] void operator()() const {} }; @@ -119,7 +119,7 @@ void TRIFuncObjBad::operator()() const {} // attributes when merging, so the test compiles without // any diagnostic when it shouldn't. struct TRIFuncObjBad1 { - [[intel::reqd_work_group_size(4, 4, 4)]] void + [[sycl::reqd_work_group_size(4, 4, 4)]] void operator()() const; }; @@ -166,7 +166,9 @@ 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'}} + [[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?}} [[intel::max_global_work_dim(0)]] void operator()() const {} }; @@ -184,7 +186,7 @@ struct TRIFuncObjBad8 { operator()() const; }; -[[intel::reqd_work_group_size(4, 4, 4)]] // expected-error{{all 'reqd_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}} +[[sycl::reqd_work_group_size(4, 4, 4)]] // expected-error{{all 'reqd_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}} void TRIFuncObjBad8::operator()() const {} struct TRIFuncObjBad9 { @@ -201,7 +203,7 @@ void TRIFuncObjBad9::operator()() const {} struct TRIFuncObjBad10 { // expected-error@+2{{all 'reqd_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}} // expected-warning@+1{{implicit conversion changes signedness: 'int' to 'unsigned long long'}} - [[intel::reqd_work_group_size(-4, 1)]] + [[sycl::reqd_work_group_size(-4, 1)]] [[intel::max_global_work_dim(0)]] void operator()() const {} }; @@ -219,7 +221,7 @@ struct TRIFuncObjBad12 { }; struct TRIFuncObjBad13 { - [[intel::reqd_work_group_size(4)]] + [[sycl::reqd_work_group_size(4)]] [[intel::max_global_work_dim(-2)]] // expected-error{{'max_global_work_dim' attribute requires integer constant between 0 and 3 inclusive}} void operator()() const {} }; 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 74449a3777fa8..52c9cdfe9a6f5 100644 --- a/clang/test/SemaSYCL/intel-reqd-work-group-size-device.cpp +++ b/clang/test/SemaSYCL/intel-reqd-work-group-size-device.cpp @@ -10,7 +10,7 @@ queue q; // expected-no-diagnostics class Functor { public: - [[intel::reqd_work_group_size(4)]] void operator()() const {} + [[sycl::reqd_work_group_size(4)]] void operator()() const {} }; void bar() { @@ -21,15 +21,14 @@ void bar() { } #else -[[intel::reqd_work_group_size(4)]] void f4x1x1() {} // expected-note {{conflicting attribute is here}} +[[sycl::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)]] void f32x1x1() {} // expected-note {{conflicting attribute is here}} +[[sycl::reqd_work_group_size(32)]] void f32x1x1() {} // expected-note {{conflicting attribute is here}} +[[sycl::reqd_work_group_size(16)]] void f16x1x1() {} // expected-note {{conflicting attribute is here}} +[[sycl::reqd_work_group_size(16, 16)]] 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, 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}} +[[sycl::reqd_work_group_size(32, 32)]] void f32x32x1() {} // expected-note {{conflicting attribute is here}} +[[sycl::reqd_work_group_size(32, 32, 32)]] void f32x32x32() {} // expected-note {{conflicting attribute is here}} #ifdef TRIGGER_ERROR class Functor32 { @@ -43,33 +42,34 @@ class Functor32 { class Functor33 { public: // expected-warning@+1{{implicit conversion changes signedness: 'int' to 'unsigned long long'}} - [[intel::reqd_work_group_size(32, -4)]] void operator()() const {} + [[sycl::reqd_work_group_size(32, -4)]] void operator()() const {} }; class Functor30 { public: // expected-warning@+1 2{{implicit conversion changes signedness: 'int' to 'unsigned long long'}} - [[intel::reqd_work_group_size(30, -30, -30)]] void operator()() const {} + [[sycl::reqd_work_group_size(30, -30, -30)]] void operator()() const {} }; class Functor16 { public: - [[intel::reqd_work_group_size(16)]] void operator()() const {} + [[sycl::reqd_work_group_size(16)]] void operator()() const {} }; class Functor64 { public: - [[intel::reqd_work_group_size(64, 64)]] void operator()() const {} + [[sycl::reqd_work_group_size(64, 64)]] void operator()() const {} }; class Functor16x16x16 { public: - [[intel::reqd_work_group_size(16, 16, 16)]] void operator()() const {} + [[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?}} }; class Functor8 { // expected-error {{conflicting attributes applied to a SYCL kernel}} public: - [[intel::reqd_work_group_size(8)]] void operator()() const { // expected-note {{conflicting attribute is here}} + [[sycl::reqd_work_group_size(8)]] void operator()() const { // expected-note {{conflicting attribute is here}} f4x1x1(); } }; @@ -107,7 +107,7 @@ int main() { Functor30 f30; h.single_task(f30); - h.single_task([]() [[intel::reqd_work_group_size(32, 32, 32)]] { + h.single_task([]() [[sycl::reqd_work_group_size(32, 32, 32)]] { f32x32x32(); }); #ifdef TRIGGER_ERROR @@ -130,7 +130,7 @@ int main() { }); // expected-error@+1 {{expected variable name or 'this' in lambda capture list}} - h.single_task([[intel::reqd_work_group_size(32, 32, 32)]][]() { + h.single_task([[sycl::reqd_work_group_size(32, 32, 32)]][]() { f32x32x32(); }); diff --git a/clang/test/SemaSYCL/intel-reqd-work-group-size-host.cpp b/clang/test/SemaSYCL/intel-reqd-work-group-size-host.cpp index cd3655dd60f7f..f9d98ef6be877 100644 --- a/clang/test/SemaSYCL/intel-reqd-work-group-size-host.cpp +++ b/clang/test/SemaSYCL/intel-reqd-work-group-size-host.cpp @@ -1,13 +1,13 @@ // RUN: %clang_cc1 -fsycl-is-host -Wno-sycl-2017-compat -fsyntax-only -verify %s // expected-no-diagnostics -[[intel::reqd_work_group_size(4)]] void f4x1x1() {} +[[sycl::reqd_work_group_size(4)]] void f4x1x1() {} -[[intel::reqd_work_group_size(16)]] void f16x1x1() {} +[[sycl::reqd_work_group_size(16)]] void f16x1x1() {} -[[intel::reqd_work_group_size(32, 32, 32)]] void f32x32x32() {} +[[sycl::reqd_work_group_size(32, 32, 32)]] void f32x32x32() {} class Functor64 { public: - [[intel::reqd_work_group_size(64, 64, 64)]] void operator()() const {} + [[sycl::reqd_work_group_size(64, 64, 64)]] void operator()() const {} }; diff --git a/clang/test/SemaSYCL/num_simd_work_items_device.cpp b/clang/test/SemaSYCL/num_simd_work_items_device.cpp index bfa01d8b782cf..0841662f9b7bb 100644 --- a/clang/test/SemaSYCL/num_simd_work_items_device.cpp +++ b/clang/test/SemaSYCL/num_simd_work_items_device.cpp @@ -48,13 +48,13 @@ struct FuncObj { // can be evenly divided by the [[intel::num_simd_work_items()]] attribute. struct TRIFuncObjBad1 { [[intel::num_simd_work_items(3)]] // 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(3, 6, 5)]] // expected-note{{conflicting attribute is here}} + [[sycl::reqd_work_group_size(3, 6, 5)]] // expected-note{{conflicting attribute is here}} void operator()() const {} }; struct TRIFuncObjBad2 { - [[intel::reqd_work_group_size(3, 6, 5)]] // expected-note{{conflicting attribute is here}} + [[sycl::reqd_work_group_size(3, 6, 5)]] // expected-note{{conflicting attribute is here}} [[intel::num_simd_work_items(3)]] // expected-error{{'num_simd_work_items' attribute must evenly divide the work-group size for the 'reqd_work_group_size' attribute}} void operator()() const {} @@ -70,7 +70,7 @@ struct TRIFuncObjBad2 { // values. This will prevent to redeclare the function with a different dimensionality. struct TRIFuncObjBad3 { [[intel::num_simd_work_items(3)]] // 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(3)]] //expected-note{{conflicting attribute is here}} + [[sycl::reqd_work_group_size(3)]] //expected-note{{conflicting attribute is here}} void operator()() const {} }; @@ -82,7 +82,7 @@ struct TRIFuncObjBad3 { // (one, two, and three argument) can be used instead of assuming default // values. This will prevent to redeclare the function with a different dimensionality. struct TRIFuncObjBad4 { - [[intel::reqd_work_group_size(3)]] // expected-note{{conflicting attribute is here}} + [[sycl::reqd_work_group_size(3)]] // expected-note{{conflicting attribute is here}} [[intel::num_simd_work_items(3)]] // expected-error{{'num_simd_work_items' attribute must evenly divide the work-group size for the 'reqd_work_group_size' attribute}} void operator()() const {} @@ -96,7 +96,7 @@ struct TRIFuncObjBad4 { // values. This will prevent to redeclare the function with a different dimensionality. struct TRIFuncObjBad5 { [[intel::num_simd_work_items(4)]] // 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, 64)]] // expected-note{{conflicting attribute is here}} + [[sycl::reqd_work_group_size(4, 64)]] // expected-note{{conflicting attribute is here}} void operator()() const {} }; @@ -108,7 +108,7 @@ struct TRIFuncObjBad5 { // (one, two, and three argument) can be used instead of assuming default // values. This will prevent to redeclare the function with a different dimensionality. struct TRIFuncObjBad6 { - [[intel::reqd_work_group_size(4, 64)]] // expected-note{{conflicting attribute is here}} + [[sycl::reqd_work_group_size(4, 64)]] // expected-note{{conflicting attribute is here}} [[intel::num_simd_work_items(4)]] // expected-error{{'num_simd_work_items' attribute must evenly divide the work-group size for the 'reqd_work_group_size' attribute}} void operator()() const {} @@ -129,9 +129,11 @@ 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}} -[[intel::reqd_work_group_size(4, 2, 3)]] void func1(); // expected-note{{conflicting attribute is here}} +[[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}} +[[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?}} [[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}} @@ -149,62 +151,62 @@ __attribute__((reqd_work_group_size(4, 2, 5))) void func5(); // expected-note{{c // Tests for incorrect argument values for Intel FPGA num_simd_work_items and reqd_work_group_size function attributes struct TRIFuncObjBad9 { - [[intel::reqd_work_group_size(5, 5, 5)]] + [[sycl::reqd_work_group_size(5, 5, 5)]] [[intel::num_simd_work_items(0)]] // expected-error{{'num_simd_work_items' attribute requires a positive integral compile time constant expression}} void operator()() const {} }; struct TRIFuncObjBad10 { [[intel::num_simd_work_items(0)]] // expected-error{{'num_simd_work_items' attribute requires a positive integral compile time constant expression}} - [[intel::reqd_work_group_size(5, 5, 5)]] void + [[sycl::reqd_work_group_size(5, 5, 5)]] void operator()() const {} }; struct TRIFuncObjBad11 { [[intel::num_simd_work_items(3.f)]] // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'float'}} - [[intel::reqd_work_group_size(64, 64, 64)]] + [[sycl::reqd_work_group_size(64, 64, 64)]] void operator()() const {} }; struct TRIFuncObjBad12 { - [[intel::reqd_work_group_size(64, 64, 64)]] + [[sycl::reqd_work_group_size(64, 64, 64)]] [[intel::num_simd_work_items(3.f)]] // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'float'}} void operator()() const {} }; struct TRIFuncObjBad13 { - [[intel::reqd_work_group_size(0)]] // expected-error{{'reqd_work_group_size' attribute must be greater than 0}} + [[sycl::reqd_work_group_size(0)]] // expected-error{{'reqd_work_group_size' attribute must be greater than 0}} [[intel::num_simd_work_items(0)]] // expected-error{{'num_simd_work_items' attribute requires a positive integral compile time constant expression}} void operator()() const {} }; struct TRIFuncObjBad14 { [[intel::num_simd_work_items(0)]] // expected-error{{'num_simd_work_items' attribute requires a positive integral compile time constant expression}} - [[intel::reqd_work_group_size(0)]] // expected-error{{'reqd_work_group_size' attribute must be greater than 0}} + [[sycl::reqd_work_group_size(0)]] // expected-error{{'reqd_work_group_size' attribute must be greater than 0}} void operator()() const {} }; struct TRIFuncObjBad15 { [[intel::num_simd_work_items(3.f)]] // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'float'}} - [[intel::reqd_work_group_size(3.f)]] // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'float'}} + [[sycl::reqd_work_group_size(3.f)]] // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'float'}} void operator()() const {} }; struct TRIFuncObjBad16 { - [[intel::reqd_work_group_size(3.f)]] // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'float'}} + [[sycl::reqd_work_group_size(3.f)]] // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'float'}} [[intel::num_simd_work_items(3.f)]] // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'float'}} void operator()() const {} }; struct TRIFuncObjBad17 { [[intel::num_simd_work_items(3)]] - [[intel::reqd_work_group_size(3, 3, 3.f)]] // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'float'}} + [[sycl::reqd_work_group_size(3, 3, 3.f)]] // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'float'}} void operator()() const {} }; struct TRIFuncObjBad18 { [[intel::num_simd_work_items(-1)]] // expected-error{{'num_simd_work_items' attribute requires a positive integral compile time constant expression}} - [[intel::reqd_work_group_size(-1)]] // expected-warning{{implicit conversion changes signedness: 'int' to 'unsigned long long'}} + [[sycl::reqd_work_group_size(-1)]] // expected-warning{{implicit conversion changes signedness: 'int' to 'unsigned long long'}} void operator()() const {} }; @@ -216,12 +218,12 @@ struct TRIFuncObjBad18 { // [[intel::num_simd_work_items()]] attribute. struct TRIFuncObjGood1 { [[intel::num_simd_work_items(4)]] - [[intel::reqd_work_group_size(3, 64, 4)]] void + [[sycl::reqd_work_group_size(3, 64, 4)]] void operator()() const {} }; struct TRIFuncObjGood2 { - [[intel::reqd_work_group_size(3, 64, 4)]] + [[sycl::reqd_work_group_size(3, 64, 4)]] [[intel::num_simd_work_items(4)]] void operator()() const {} }; diff --git a/clang/test/SemaSYCL/reqd-work-group-size-device.cpp b/clang/test/SemaSYCL/reqd-work-group-size-device.cpp index c737afe30ef50..cda6a368be58c 100644 --- a/clang/test/SemaSYCL/reqd-work-group-size-device.cpp +++ b/clang/test/SemaSYCL/reqd-work-group-size-device.cpp @@ -17,7 +17,7 @@ queue q; [[sycl::reqd_work_group_size(32, 32, 32)]] void f32x32x32() {} // expected-note {{conflicting attribute is here}} // No diagnostic because the attributes are synonyms with identical behavior. -[[intel::reqd_work_group_size(4, 4, 4)]] void four(); +[[sycl::reqd_work_group_size(4, 4, 4)]] void four(); [[sycl::reqd_work_group_size(4, 4, 4)]] void four(); // OK // Same for the default values. @@ -27,11 +27,11 @@ queue q; // kernel, so the one, two, and three arg forms of the attribute are actually // *different* attributes. This means that you should not be able to redeclare // the function with a different dimensionality. -[[intel::reqd_work_group_size(4)]] void four_again(); +[[sycl::reqd_work_group_size(4)]] void four_again(); [[sycl::reqd_work_group_size(4)]] void four_again(); // OK -[[intel::reqd_work_group_size(4, 1)]] void four_again(); // OK [[sycl::reqd_work_group_size(4, 1)]] void four_again(); // OK -[[intel::reqd_work_group_size(4, 1, 1)]] void four_again(); // OK +[[sycl::reqd_work_group_size(4, 1)]] void four_again(); // OK +[[sycl::reqd_work_group_size(4, 1, 1)]] void four_again(); // OK [[sycl::reqd_work_group_size(4, 1, 1)]] void four_again(); // OK // The GNU and [[cl::reqd_work_group_size]] spellings are deprecated in SYCL diff --git a/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp b/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp index 41bb2ce80dde5..52f441a496983 100644 --- a/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp +++ b/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp @@ -79,17 +79,17 @@ template [[intel::max_global_work_dim(0)]] void func5(); template -[[intel::reqd_work_group_size(N)]] void func6(); // expected-error {{all 'reqd_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}} +[[sycl::reqd_work_group_size(N)]] void func6(); // expected-error {{all 'reqd_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}} template [[intel::max_global_work_dim(0)]] void func6(); template -[[intel::reqd_work_group_size(N, N)]] void func7(); // expected-error {{all 'reqd_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}} +[[sycl::reqd_work_group_size(N, N)]] void func7(); // expected-error {{all 'reqd_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}} template [[intel::max_global_work_dim(0)]] void func7(); template -[[intel::reqd_work_group_size(N, N, N)]] void func8(); // expected-error {{all 'reqd_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}} +[[sycl::reqd_work_group_size(N, N, N)]] void func8(); // expected-error {{all 'reqd_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}} template [[intel::max_global_work_dim(0)]] void func8(); @@ -99,12 +99,12 @@ template [[intel::max_global_work_dim(0)]] void func9(); template -[[intel::reqd_work_group_size(N)]] void func10(); +[[sycl::reqd_work_group_size(N)]] void func10(); template [[intel::max_global_work_dim(0)]] void func10(); template -[[intel::reqd_work_group_size(N, N)]] void func11(); +[[sycl::reqd_work_group_size(N, N)]] void func11(); template [[intel::max_global_work_dim(0)]] void func11(); 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 8338899af5f1f..680d66014264c 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 @@ -7,7 +7,7 @@ template // expected-error@+1 3{{integral constant expression must have integral or unscoped enumeration type, not 'S'}} -[[intel::reqd_work_group_size(Ty{}, Ty1{}, Ty2{})]] void func() {} +[[sycl::reqd_work_group_size(Ty{}, Ty1{}, Ty2{})]] void func() {} struct S {}; void var() { @@ -20,17 +20,17 @@ void var() { int foo(); // expected-error@+2 3{{expression is not an integral constant expression}} // expected-note@+1 3{{non-constexpr function 'foo' cannot be used in a constant expression}} -[[intel::reqd_work_group_size(foo() + 12, foo() + 12, foo() + 12)]] void func1(); +[[sycl::reqd_work_group_size(foo() + 12, foo() + 12, foo() + 12)]] void func1(); // Test that checks expression is a constant expression. constexpr int bar() { return 0; } -[[intel::reqd_work_group_size(bar() + 12, bar() + 12, bar() + 12)]] void func2(); // OK +[[sycl::reqd_work_group_size(bar() + 12, bar() + 12, bar() + 12)]] void func2(); // OK // Test that checks template parameter support on member function of class template. template class KernelFunctor { public: - [[intel::reqd_work_group_size(SIZE, SIZE1, SIZE2)]] void operator()() {} + [[sycl::reqd_work_group_size(SIZE, SIZE1, SIZE2)]] void operator()() {} }; int main() { @@ -59,7 +59,7 @@ int main() { // Test that checks template parameter support on function. template -[[intel::reqd_work_group_size(N, N1, N2)]] void func3() {} +[[sycl::reqd_work_group_size(N, N1, N2)]] void func3() {} int check() { func3<8, 8, 8>(); diff --git a/clang/test/SemaSYCL/sycl-device-num_simd_work_items-template.cpp b/clang/test/SemaSYCL/sycl-device-num_simd_work_items-template.cpp index 61d867d4a1013..77698887ec8df 100644 --- a/clang/test/SemaSYCL/sycl-device-num_simd_work_items-template.cpp +++ b/clang/test/SemaSYCL/sycl-device-num_simd_work_items-template.cpp @@ -106,17 +106,17 @@ template [[intel::num_simd_work_items(N)]] void func7(); // expected-error{{'num_simd_work_items' attribute must evenly divide the work-group size for the 'reqd_work_group_size' attribute}} template -[[intel::reqd_work_group_size(N, N, N)]] void func8(); // expected-note{{conflicting attribute is here}} +[[sycl::reqd_work_group_size(N, N, N)]] void func8(); // expected-note{{conflicting attribute is here}} template [[intel::num_simd_work_items(3)]] void func8(); // expected-error{{'num_simd_work_items' attribute must evenly divide the work-group size for the 'reqd_work_group_size' attribute}} template -[[intel::reqd_work_group_size(X, Y, Z)]] void func9(); // expected-note{{conflicting attribute is here}} +[[sycl::reqd_work_group_size(X, Y, Z)]] void func9(); // expected-note{{conflicting attribute is here}} template [[intel::num_simd_work_items(N)]] void func9(); // expected-error{{'num_simd_work_items' attribute must evenly divide the work-group size for the 'reqd_work_group_size' attribute}} template -[[intel::reqd_work_group_size(X, Y, Z)]] void func10(); // expected-note{{conflicting attribute is here}} +[[sycl::reqd_work_group_size(X, Y, Z)]] void func10(); // expected-note{{conflicting attribute is here}} template [[intel::num_simd_work_items(3)]] void func10(); // expected-error{{'num_simd_work_items' attribute must evenly divide the work-group size for the 'reqd_work_group_size' attribute}} From 20f8b0305eb4396efc4e77f938cf2e4618f4e3f1 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Fri, 11 Jun 2021 18:17:28 -0700 Subject: [PATCH 2/5] update Doc and fix wrong merge conflicts Signed-off-by: Soumi Manna --- clang/include/clang/Basic/AttrDocs.td | 56 +++++++++++++-------------- 1 file changed, 28 insertions(+), 28 deletions(-) diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 7fc4559ad7f1a..1857639d14456 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -2526,14 +2526,14 @@ of the kernel the attribute is applied to. // argument to the attribute is the one which increments fastest. struct func { [[intel::num_simd_work_items(4)]] - [[intel::reqd_work_group_size(7, 4, 64)]] + [[sycl::reqd_work_group_size(7, 4, 64)]] void operator()() const {} }; // Note, '8' is evenly divisible by '8'; in SYCL, the last // argument to the attribute is the one which increments fastest. struct bar { - [[intel::reqd_work_group_size(1, 1, 8)]] + [[sycl::reqd_work_group_size(1, 1, 8)]] [[intel::num_simd_work_items(8)]] void operator()() const {} }; @@ -2671,30 +2671,6 @@ In OpenCL C, this attribute is available with the GNU spelling __kernel __attribute__((reqd_work_group_size(8, 16, 32))) void test() {} - }]; -} - -def WorkGroupSizeHintAttrDocs : Documentation { - let Category = DocCatFunction; - let Content = [{ -Applies to a device function/lambda function. Hint to the compiler on the work- -group size most likely to be used when launching the kernel at runtime. Each -argument must be a nonzero integral constant expression whose value is -unsigned. The number of dimensional values defined provide additional -information to the compiler on the dimensionality most likely to be used when -launching the kernel at runtime. - -The GNU spelling is deprecated in SYCL mode. - -.. code-block:: c++ - - [[sycl::work_group_size_hint(4, 4, 4)]] void foo() {} - - class Foo { - public: - [[sycl::work_group_size_hint(2, 2, 2)]] void operator()() const {} - }; - The arguments to ``reqd_work_group_size`` are ordered based on which index increments the fastest. In OpenCL, the first argument is the index that increments the fastest, and in SYCL, the last argument is the index that @@ -2718,14 +2694,14 @@ in the ``reqd_work_group_size`` attribute. // argument to the attribute is the one which increments fastest. struct func { [[intel::num_simd_work_items(4)]] - [[intel::reqd_work_group_size(7, 4, 64)]] + [[sycl::reqd_work_group_size(7, 4, 64)]] void operator()() const {} }; // Note, '8' is evenly divisible by '8'; in SYCL, the last // argument to the attribute is the one which increments fastest. struct bar { - [[intel::reqd_work_group_size(1, 1, 8)]] + [[sycl::reqd_work_group_size(1, 1, 8)]] [[intel::num_simd_work_items(8)]] void operator()() const {} }; @@ -2769,6 +2745,30 @@ in the ``reqd_work_group_size`` attribute. }]; } +def WorkGroupSizeHintAttrDocs : Documentation { + let Category = DocCatFunction; + let Content = [{ +Applies to a device function/lambda function. Hint to the compiler on the work- +group size most likely to be used when launching the kernel at runtime. Each +argument must be a nonzero integral constant expression whose value is +unsigned. The number of dimensional values defined provide additional +information to the compiler on the dimensionality most likely to be used when +launching the kernel at runtime. + +The GNU spelling is deprecated in SYCL mode. + +.. code-block:: c++ + + [[sycl::work_group_size_hint(4, 4, 4)]] void foo() {} + + class Foo { + public: + [[sycl::work_group_size_hint(2, 2, 2)]] void operator()() const {} + }; + + }]; +} + def SYCLIntelMaxWorkGroupSizeAttrDocs : Documentation { let Category = DocCatFunction; let Heading = "intel::max_work_group_size"; From 86041256464ca3711c55d2c7bdc783f5ccc61032 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Mon, 14 Jun 2021 06:17:10 -0700 Subject: [PATCH 3/5] address review comments Signed-off-by: Soumi Manna --- clang/include/clang/Basic/AttrDocs.td | 2 +- clang/lib/Sema/SemaDeclAttr.cpp | 11 +++++------ 2 files changed, 6 insertions(+), 7 deletions(-) diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 1857639d14456..9b93645034f5d 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -2660,7 +2660,7 @@ 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 SYCL 2020 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 diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index d1403ae8267ab..58f0090bce8fe 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -338,12 +338,11 @@ void Sema::CheckDeprecatedSYCLAttributeSpelling(const ParsedAttr &A, } // Deprecate [[intel::reqd_work_group_size]] attribute spelling in favor - // of SYCL 2020 attribute spelling [[sycl::reqd_work_group_size]]. - if (A.hasScope() && A.getScopeName()->isStr("intel")) { - if (A.getKind() == ParsedAttr::AT_ReqdWorkGroupSize) { - DiagnoseDeprecatedAttribute(A, "sycl", NewName); - return; - } + // 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. From 4e265e3641b5194e65f5b4c6ce9ffbf97d009ee0 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Mon, 14 Jun 2021 06:21:47 -0700 Subject: [PATCH 4/5] fix format issues Signed-off-by: Soumi Manna --- clang/lib/Sema/SemaDeclAttr.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 58f0090bce8fe..2ccdce6cbd484 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -339,8 +339,8 @@ void Sema::CheckDeprecatedSYCLAttributeSpelling(const ParsedAttr &A, // 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")) { + if (A.getKind() == ParsedAttr::AT_ReqdWorkGroupSize && A.hasScope() && + A.getScopeName()->isStr("intel")) { DiagnoseDeprecatedAttribute(A, "sycl", NewName); return; } From f8371ab66ac9ee17970aede04a8f107558659c94 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Mon, 14 Jun 2021 06:33:27 -0700 Subject: [PATCH 5/5] remove NFC change from doc Signed-off-by: Soumi Manna --- clang/include/clang/Basic/AttrDocs.td | 48 +++++++++++++-------------- 1 file changed, 24 insertions(+), 24 deletions(-) diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 9b93645034f5d..84c9ed1188166 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -2671,6 +2671,30 @@ In OpenCL C, this attribute is available with the GNU spelling __kernel __attribute__((reqd_work_group_size(8, 16, 32))) void test() {} + }]; +} + +def WorkGroupSizeHintAttrDocs : Documentation { + let Category = DocCatFunction; + let Content = [{ +Applies to a device function/lambda function. Hint to the compiler on the work- +group size most likely to be used when launching the kernel at runtime. Each +argument must be a nonzero integral constant expression whose value is +unsigned. The number of dimensional values defined provide additional +information to the compiler on the dimensionality most likely to be used when +launching the kernel at runtime. + +The GNU spelling is deprecated in SYCL mode. + +.. code-block:: c++ + + [[sycl::work_group_size_hint(4, 4, 4)]] void foo() {} + + class Foo { + public: + [[sycl::work_group_size_hint(2, 2, 2)]] void operator()() const {} + }; + The arguments to ``reqd_work_group_size`` are ordered based on which index increments the fastest. In OpenCL, the first argument is the index that increments the fastest, and in SYCL, the last argument is the index that @@ -2745,30 +2769,6 @@ in the ``reqd_work_group_size`` attribute. }]; } -def WorkGroupSizeHintAttrDocs : Documentation { - let Category = DocCatFunction; - let Content = [{ -Applies to a device function/lambda function. Hint to the compiler on the work- -group size most likely to be used when launching the kernel at runtime. Each -argument must be a nonzero integral constant expression whose value is -unsigned. The number of dimensional values defined provide additional -information to the compiler on the dimensionality most likely to be used when -launching the kernel at runtime. - -The GNU spelling is deprecated in SYCL mode. - -.. code-block:: c++ - - [[sycl::work_group_size_hint(4, 4, 4)]] void foo() {} - - class Foo { - public: - [[sycl::work_group_size_hint(2, 2, 2)]] void operator()() const {} - }; - - }]; -} - def SYCLIntelMaxWorkGroupSizeAttrDocs : Documentation { let Category = DocCatFunction; let Heading = "intel::max_work_group_size";