diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index cb91cc1707b23..84c9ed1188166 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 {} }; @@ -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 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. @@ -2714,14 +2718,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 {} }; diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 2d40ea3103286..2ccdce6cbd484 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -337,6 +337,14 @@ 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/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}}