[SYCL] Add warning for attributes applied to non-kernel functions (#15154)

According to [5.8.1. Kernel
attributes](https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:kernel.attributes),
kernel attributes can only be applied to SYCL kernel functions, but not
to a regular device functions.

Based on this information, this pull request introduces a warning
diagnostic whenever specified attributes are applied to regular device
functions.

Author: dyniols

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -12358,6 +12358,8 @@ def err_bit_cast_type_size_mismatch : Error<
   "__builtin_bit_cast source size does not equal destination size (%0 vs %1)">;
 
 // SYCL-specific diagnostics
+def warn_sycl_incorrect_use_attribute_non_kernel_function : Warning<
+  "%0 attribute can only be applied to a SYCL kernel function">, InGroup<SyclStrict>;
 def warn_sycl_kernel_num_of_template_params : Warning<
   "'sycl_kernel' attribute only applies to a function template with at least"
   " two template parameters">, InGroup<IgnoredAttributes>;

clang/include/clang/Sema/SemaSYCL.h

@@ -259,6 +259,8 @@ class SemaSYCL : public SemaBase {
   // useful notes that shows where the kernel was called.
   bool DiagnosingSYCLKernel = false;
 
+  llvm::DenseSet<const FunctionDecl *> SYCLKernelFunctions;
+
 public:
   SemaSYCL(Sema &S);
 
@@ -300,6 +302,10 @@ class SemaSYCL : public SemaBase {
   void addSyclDeviceDecl(Decl *d) { SyclDeviceDecls.insert(d); }
   llvm::SetVector<Decl *> &syclDeviceDecls() { return SyclDeviceDecls; }
 
+  void addSYCLKernelFunction(const FunctionDecl *FD) {
+    SYCLKernelFunctions.insert(FD);
+  }
+
   /// Lazily creates and returns SYCL integration header instance.
   SYCLIntegrationHeader &getSyclIntegrationHeader() {
     if (SyclIntHeader == nullptr)
@@ -375,6 +381,8 @@ class SemaSYCL : public SemaBase {
                                    SourceLocation Loc,
                                    DeviceDiagnosticReason Reason);
 
+  void performSYCLDelayedAttributesAnalaysis(const FunctionDecl *FD);
+
   /// Tells whether given variable is a SYCL explicit SIMD extension's "private
   /// global" variable - global variable in the private address space.
   bool isSYCLEsimdPrivateGlobal(VarDecl *VDecl) {

clang/lib/Sema/Sema.cpp

@@ -1909,6 +1909,10 @@ class DeferredDiagnosticsEmitter
   void checkFunc(SourceLocation Loc, FunctionDecl *FD) {
     auto &Done = DoneMap[InOMPDeviceContext > 0 ? 1 : 0];
     FunctionDecl *Caller = UsePath.empty() ? nullptr : UsePath.back();
+
+    if (!Caller && S.LangOpts.SYCLIsDevice)
+      S.SYCL().performSYCLDelayedAttributesAnalaysis(FD);
+
     if ((!ShouldEmitRootNode && !S.getLangOpts().OpenMP && !Caller) ||
         S.shouldIgnoreInHostDeviceCheck(FD) || InUsePath.count(FD))
       return;

clang/lib/Sema/SemaSYCL.cpp

@@ -882,6 +882,8 @@ class SingleDeviceFunctionTracker {
         // having a kernel lambda with a lambda call inside of it.
         KernelBody = CurrentDecl;
       }
+      if (KernelBody)
+        Parent.SemaSYCLRef.addSYCLKernelFunction(KernelBody);
     }
 
     // Recurse.
@@ -6852,3 +6854,19 @@ ExprResult SemaSYCL::ActOnUniqueStableNameExpr(SourceLocation OpLoc,
 
   return BuildUniqueStableNameExpr(OpLoc, LParen, RParen, TSI);
 }
+
+void SemaSYCL::performSYCLDelayedAttributesAnalaysis(const FunctionDecl *FD) {
+  if (SYCLKernelFunctions.contains(FD))
+    return;
+
+  for (const auto *KernelAttr : std::vector<AttributeCommonInfo *>{
+           FD->getAttr<SYCLReqdWorkGroupSizeAttr>(),
+           FD->getAttr<IntelReqdSubGroupSizeAttr>(),
+           FD->getAttr<SYCLWorkGroupSizeHintAttr>(),
+           FD->getAttr<VecTypeHintAttr>()}) {
+    if (KernelAttr)
+      Diag(KernelAttr->getLoc(),
+           diag::warn_sycl_incorrect_use_attribute_non_kernel_function)
+          << KernelAttr;
+  }
+}

clang/test/SemaSYCL/check-work-group-size-hint-device.cpp

@@ -17,7 +17,8 @@
 
 // Produce a conflicting attribute warning when the args are different.
 [[sycl::work_group_size_hint(4, 1, 1)]] void f3();    // expected-note {{previous attribute is here}}
-[[sycl::work_group_size_hint(1, 1, 32)]] void f3() {} // expected-warning {{attribute 'work_group_size_hint' is already applied with different arguments}}
+[[sycl::work_group_size_hint(1, 1, 32)]] void f3() {} // expected-warning {{attribute 'work_group_size_hint' is already applied with different arguments}} \
+// expected-warning {{'work_group_size_hint' attribute can only be applied to a SYCL kernel function}}
 
 // 1 and 2 dim versions
 [[sycl::work_group_size_hint(2)]] void f4();    // ok
@@ -70,10 +71,13 @@ void instantiate() {
   f8<0>(); // expected-note {{in instantiation}}
 #endif
 
+  // expected-warning@#f9prev {{'work_group_size_hint' attribute can only be applied to a SYCL kernel function}}
   f9<1, 1, 1>(); // OK, args are the same on the redecl.
 
   // expected-warning@#f9 {{attribute 'work_group_size_hint' is already applied with different arguments}}
   // expected-note@#f9prev {{previous attribute is here}}
+  // expected-warning@#f9prev {{'work_group_size_hint' attribute can only be applied to a SYCL kernel function}}
+
   f9<1, 2, 3>(); // expected-note {{in instantiation}}
 }
 
@@ -97,14 +101,14 @@ class Functor16x2x1 {
 
 class Functor4x4x4 {
 public:
-  [[sycl::work_group_size_hint(4, 4, 4)]] void operator()() const {};
+  [[sycl::work_group_size_hint(4, 4, 4)]] void operator()() const {}; // expected-warning {{'work_group_size_hint' attribute can only be applied to a SYCL kernel function}}
 };
 
 // Checking whether propagation of the attribute happens or not, according to the SYCL version.
 #if defined(EXPECT_PROP) // if attribute is propagated, then we expect errors here
 void f8x8x8(){};
 #else // otherwise no error
-[[sycl::work_group_size_hint(8, 8, 8)]] void f8x8x8(){};
+[[sycl::work_group_size_hint(8, 8, 8)]] void f8x8x8(){}; // expected-warning {{'work_group_size_hint' attribute can only be applied to a SYCL kernel function}}
 #endif
 class FunctorNoProp {
 public:

clang/test/SemaSYCL/device_has.cpp

@@ -17,7 +17,6 @@ enum class aspect {
 
 [[sycl::device_has("123")]] void func1() {}                          // expected-error{{'device_has' attribute argument is invalid; argument must be device aspect of type sycl::aspect}}
 [[sycl::device_has(fake_cl::sycl::aspect::aspect1)]] void func2() {} // expected-error{{'device_has' attribute argument is invalid; argument must be device aspect of type sycl::aspect}}
-
 [[sycl::device_has(sycl::aspect::cpu)]] void func3();   // expected-note{{previous attribute is here}}
 [[sycl::device_has(sycl::aspect::gpu)]] void func3() {} // expected-warning{{attribute 'device_has' is already applied}}
 

clang/test/SemaSYCL/intel-max-work-group-size.cpp

@@ -70,15 +70,15 @@ void instantiate() {
 // a declaration along with [[sycl::reqd_work_group_size(X1, Y1, Z1)]]
 // attribute, check to see if values of reqd_work_group_size arguments are
 // equal or less than values coming from max_work_group_size attribute.
-[[sycl::reqd_work_group_size(64, 64, 64)]] // expected-note {{conflicting attribute is here}}
+[[sycl::reqd_work_group_size(64, 64, 64)]] // expected-note {{conflicting attribute is here}} // expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
 [[intel::max_work_group_size(64, 16, 64)]] // expected-error {{'max_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}}
 void
 f9() {}
 
 [[intel::max_work_group_size(4, 4, 4)]] void f10();
 [[sycl::reqd_work_group_size(2, 2, 2)]] void f10(); // OK
 
-[[sycl::reqd_work_group_size(2, 2, 2)]] [[intel::max_work_group_size(4, 4, 4)]] void f11() {} // OK
+[[sycl::reqd_work_group_size(2, 2, 2)]] [[intel::max_work_group_size(4, 4, 4)]] void f11() {} // expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
 
 [[sycl::reqd_work_group_size(64, 64, 64)]] void f12(); // expected-note {{conflicting attribute is here}}
 [[intel::max_work_group_size(16, 16, 16)]] void f12(); // expected-error {{'max_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}}
@@ -91,14 +91,14 @@ f13() {}
 [[sycl::reqd_work_group_size(64, 64, 64)]] void f14(); // expected-error{{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}}
 
 [[cl::reqd_work_group_size(1, 2, 3)]] // expected-warning {{attribute 'cl::reqd_work_group_size' is deprecated}} \
-		                      // expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}}
+		                      // expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}} // expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
 [[intel::max_work_group_size(1, 2, 3)]] void
 f15() {} // OK
 
 [[intel::max_work_group_size(2, 3, 7)]] void f16(); // expected-note {{conflicting attribute is here}}
 [[sycl::reqd_work_group_size(7, 3, 2)]] void f16(); // expected-error{{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}}
 
-[[intel::max_work_group_size(1, 2, 3)]] [[sycl::reqd_work_group_size(1, 2, 3)]] void f17(){}; // OK
+[[intel::max_work_group_size(1, 2, 3)]] [[sycl::reqd_work_group_size(1, 2, 3)]] void f17(){}; // expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
 
 [[sycl::reqd_work_group_size(16)]]            // expected-note {{conflicting attribute is here}}
 [[intel::max_work_group_size(16, 1, 1)]] void // expected-error {{'max_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}}

clang/test/SemaSYCL/reqd-sub-group-size.cpp

@@ -33,17 +33,19 @@ int main() {
   return 0;
 }
 [[intel::reqd_sub_group_size(16)]] SYCL_EXTERNAL void B();
-[[intel::reqd_sub_group_size(16)]] void A() {
+[[intel::reqd_sub_group_size(16)]] void A() // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}} 
+{ 
 }
 
-[[intel::reqd_sub_group_size(16)]] SYCL_EXTERNAL void B() {
+[[intel::reqd_sub_group_size(16)]] SYCL_EXTERNAL void B() { // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
   A();
 }
 // expected-note@+1 {{conflicting attribute is here}}
-[[intel::reqd_sub_group_size(2)]] void sg_size2() {}
+[[intel::reqd_sub_group_size(2)]] void sg_size2() {} // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
 
-// expected-note@+2 {{conflicting attribute is here}}
-// expected-error@+1 {{conflicting attributes applied to a SYCL kernel}}
+// expected-note@+3 {{conflicting attribute is here}}
+// expected-error@+2 {{conflicting attributes applied to a SYCL kernel}}
+// expected-warning@+1 {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
 [[intel::reqd_sub_group_size(4)]] __attribute__((sycl_device)) void sg_size4() {
   sg_size2();
 }
@@ -67,7 +69,7 @@ int main() {
 
 // No diagnostic is emitted because the arguments match.
 [[intel::reqd_sub_group_size(12)]] void same();
-[[intel::reqd_sub_group_size(12)]] void same() {} // OK
+[[intel::reqd_sub_group_size(12)]] void same() {} // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
 
 // No diagnostic because the attributes are synonyms with identical behavior.
 [[sycl::reqd_sub_group_size(12)]] void same(); // OK
@@ -117,10 +119,12 @@ int check() {
 
 // Test that checks template parameter support on function.
 template <int N>
-// expected-error@+1{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}}
+// expected-error@+2{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}}
+// expected-warning@+1 {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
 [[intel::reqd_sub_group_size(N)]] void func3() {}
 
 template <int N>
+// expected-warning@+1 {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
 [[intel::reqd_sub_group_size(4)]] void func4(); // expected-note {{previous attribute is here}}
 
 template <int N>

clang/test/SemaSYCL/reqd_work_group_size.cpp

@@ -22,28 +22,31 @@ class Functor30 {
 
 // Tests for 'reqd_work_group_size' attribute duplication.
 // No diagnostic is emitted because the arguments match. Duplicate attribute is silently ignored.
+// expected-warning@+1 {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
 [[sycl::reqd_work_group_size(6, 6, 6)]] [[sycl::reqd_work_group_size(6, 6, 6)]] void f2() {}
 
 // No diagnostic is emitted because the arguments match.
 [[sycl::reqd_work_group_size(32, 32, 32)]] void f3();
 [[sycl::reqd_work_group_size(32, 32, 32)]] void f3(); // OK
 
 // Produce a conflicting attribute warning when the args are different.
-[[sycl::reqd_work_group_size(6, 6, 6)]]         // expected-note {{previous attribute is here}}
+[[sycl::reqd_work_group_size(6, 6, 6)]]         // expected-note {{previous attribute is here}} // expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
 [[sycl::reqd_work_group_size(16, 16, 16)]] void // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}}
 f4() {}
 
 // Catch the easy case where the attributes are all specified at once with
 // different arguments.
 struct TRIFuncObjGood1 {
-  // expected-note@+2 {{previous attribute is here}}
-  // expected-error@+1 {{attribute 'reqd_work_group_size' is already applied with different arguments}}
+  // expected-note@+3 {{previous attribute is here}}
+  // expected-error@+2 {{attribute 'reqd_work_group_size' is already applied with different arguments}}
+  // expected-warning@+1 {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}} 
   [[sycl::reqd_work_group_size(64)]] [[sycl::reqd_work_group_size(128)]] void operator()() const {}
 };
 
 struct TRIFuncObjGood2 {
-  // expected-note@+2 {{previous attribute is here}}
-  // expected-error@+1 {{attribute 'reqd_work_group_size' is already applied with different arguments}}
+  // expected-note@+3 {{previous attribute is here}}
+  // expected-error@+2 {{attribute 'reqd_work_group_size' is already applied with different arguments}}
+  // expected-warning@+1 {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
   [[sycl::reqd_work_group_size(64, 64)]] [[sycl::reqd_work_group_size(128, 128)]] void operator()() const {}
 };
 
@@ -52,7 +55,8 @@ struct TRIFuncObjGood3 {
   operator()() const;
 };
 
-[[sycl::reqd_work_group_size(4, 4)]] // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}}
+[[sycl::reqd_work_group_size(4, 4)]] // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}} \
+// expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
 void
 TRIFuncObjGood3::operator()() const {}
 
@@ -73,7 +77,7 @@ class FunctorC {
 
 class Functor32 {
 public:
-  [[sycl::reqd_work_group_size(32, 1, 1)]]      // expected-note {{previous attribute is here}}
+  [[sycl::reqd_work_group_size(32, 1, 1)]]      // expected-note {{previous attribute is here}} // expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
   [[sycl::reqd_work_group_size(1, 1, 32)]] void // expected-error{{attribute 'reqd_work_group_size' is already applied with different arguments}}
   operator()() const {}
 };
@@ -105,16 +109,18 @@ void instantiate() {
   f7<1, 1, 1>(); // OK, args are the same on the redecl.
   // expected-error@#f7 {{attribute 'reqd_work_group_size' is already applied with different arguments}}
   // expected-note@#f7prev {{previous attribute is here}}
+  // expected-warning@#f7prev {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
+  // expected-warning@#f7prev {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
   f7<2, 2, 2>(); // expected-note {{in instantiation}}
 }
 
 // Tests for 'reqd_work_group_size' attribute duplication.
 
-[[sycl::reqd_work_group_size(8)]]            // expected-note {{previous attribute is here}}
+[[sycl::reqd_work_group_size(8)]]            // expected-note {{previous attribute is here}} // expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
 [[sycl::reqd_work_group_size(1, 1, 8)]] void // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}}
 f8(){};
 
-[[sycl::reqd_work_group_size(32, 32, 1)]]            // expected-note {{previous attribute is here}}
+[[sycl::reqd_work_group_size(32, 32, 1)]]            // expected-note {{previous attribute is here}} // expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
 [[sycl::reqd_work_group_size(32, 32)]] void f9() {}  // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}}
 
 // Test that template redeclarations also get diagnosed properly.
@@ -127,6 +133,8 @@ void test() {
   f10<64, 1, 1>(); // OK, args are the same on the redecl.
   // expected-error@#f10err {{attribute 'reqd_work_group_size' is already applied with different arguments}}
   // expected-note@#f10prev {{previous attribute is here}}
+  // expected-warning@#f10prev {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
+  // expected-warning@#f10prev {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
   f10<1, 1, 64>(); // expected-note {{in instantiation}}
 }
 
@@ -135,7 +143,8 @@ struct TRIFuncObjBad {
   operator()() const;
 };
 
-[[sycl::reqd_work_group_size(1, 1, 32)]] // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}}
+[[sycl::reqd_work_group_size(1, 1, 32)]] // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}} \
+// expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
 void
 TRIFuncObjBad::operator()() const {}
 
@@ -174,6 +183,7 @@ int main() {
   KernelFunctor<16, 1, 1>();
 }
 // Test that checks template parameter support on function.
+// expected-warning@+2 {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
 template <int N, int N1, int N2>
 [[sycl::reqd_work_group_size(N, N1, N2)]] void func3() {}
 

clang/test/SemaSYCL/sub-group-size.cpp

@@ -123,9 +123,11 @@ void calls_kernel_4() {
   sycl::kernel_single_task<class Kernel4>([]() { // #Kernel4
     // integer-error@#AttrFunc2{{kernel-called function must have a sub group size that matches the size specified for the kernel}}
     // integer-note@#Kernel4{{kernel declared here}}
+    // expected-warning@#AttrFunc2 {{'sub_group_size' attribute can only be applied to a SYCL kernel function}}
     AttrFunc2();
     // integer-error@#AttrExternalDefined2{{kernel-called function must have a sub group size that matches the size specified for the kernel}}
     // integer-note@#Kernel4{{kernel declared here}}
+    // expected-warning@#AttrExternalDefined2 {{'sub_group_size' attribute can only be applied to a SYCL kernel function}}
     AttrExternalDefined2();
     // integer-error@#AttrExternalNotDefined2{{kernel-called function must have a sub group size that matches the size specified for the kernel}}
     // integer-note@#Kernel4{{kernel declared here}}