[clang] Diagnose usages of static methods to define SYCL free function kernels (#18761)

Static class method support will be added after 2025.3 due to
impossibility of its implementation with the current approach using
integration header. This PR adds diagnostic messages for both static and
non-static class methods

Author: dklochkov-emb

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -9434,8 +9434,6 @@ def err_deleted_inherited_ctor_use : Error<
 def note_called_by : Note<"called by %0">;
 def err_kern_type_not_void_return : Error<
   "kernel function type %0 must have void return type">;
-def err_kern_is_nonstatic_method : Error<
-  "kernel function %0 must be a free function or static member function">;
 def err_config_scalar_return : Error<
   "CUDA special function '%0' must have scalar return type">;
 def err_kern_call_not_global_function : Error<
@@ -12930,6 +12928,9 @@ def err_free_function_return_type : Error<
   "SYCL free function kernel should have return type 'void'">;
 def err_free_function_first_occurrence_missing_attr: Error<
   "the first occurrence of SYCL kernel free function should be declared with 'sycl-nd-range-kernel' or 'sycl-single-task-kernel' compile time properties">;
+def err_free_function_class_method : Error<
+  "%select{static |}0class method cannot be used to define a SYCL kernel free function kernel">;
+
 
 // SYCL kernel entry point diagnostics
 def err_sycl_entry_point_invalid : Error<

clang/lib/Sema/SemaDeclAttr.cpp

@@ -5137,8 +5137,8 @@ static void handleGlobalAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
   }
   if (const auto *Method = dyn_cast<CXXMethodDecl>(FD)) {
     if (Method->isInstance()) {
-      S.Diag(Method->getBeginLoc(), diag::err_kern_is_nonstatic_method)
-          << Method;
+      S.Diag(Method->getBeginLoc(), diag::err_free_function_class_method)
+          << /*non static*/ 1 << Method;
       return;
     }
     S.Diag(Method->getBeginLoc(), diag::warn_kern_is_method) << Method;

clang/lib/Sema/SemaSYCL.cpp

@@ -5835,6 +5835,13 @@ static bool CheckFreeFunctionDiagnostics(Sema &S, FunctionDecl *FD) {
     return S.Diag(FD->getLocation(), diag::err_free_function_return_type);
   }
 
+  if (const auto *MD = llvm::dyn_cast<CXXMethodDecl>(FD))
+    // The free function extension specification usual methods to be used to
+    // define a free function kernel. We also disallow static methods because we
+    // use integration header.
+    return S.Diag(FD->getLocation(), diag::err_free_function_class_method)
+           << !MD->isStatic() << MD->getSourceRange();
+
   for (ParmVarDecl *Param : FD->parameters()) {
     if (Param->hasDefaultArg()) {
       return S.Diag(Param->getLocation(),

clang/test/SemaCUDA/bad-attributes.cu

@@ -51,7 +51,7 @@ __global__ __host__ void z12();  // expected-error {{attributes are not compatib
 // expected-note@-1 {{conflicting attribute is here}}
 
 struct S {
-  __global__ void foo() {};  // expected-error {{must be a free function or static member function}}
+  __global__ void foo() {};  // expected-error {{class method cannot be used to define a SYCL kernel free function kernel}}
   __global__ static void bar(); // expected-warning {{kernel function 'bar' is a member function}}
   // Although this is implicitly inline, we shouldn't warn.
   __global__ static void baz() {}; // expected-warning {{kernel function 'baz' is a member function}}

clang/test/SemaCUDA/lambda.cu

@@ -49,7 +49,7 @@ public:
 
 int main(void) {
   auto lambda_kernel = [&]__global__(){};
-  // com-error@-1 {{kernel function 'operator()' must be a free function or static member function}}
+  // com-error@-1 {{class method cannot be used to define a SYCL kernel free function kernel}}
 
   int b;
   [&](){ hd(b); }();

clang/test/SemaSYCL/free_function_negative.cpp

@@ -69,3 +69,52 @@ singleTaskKernelReturnType(int Value) {
 [[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]] int 
 ndRangeKernelReturnType(int Value) {
 }
+
+class TestClass {
+public:
+
+// expected-error@+2 {{class method cannot be used to define a SYCL kernel free function kernel}}
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]] void 
+ndRangeKernelMethod(int Value) {
+}
+
+// expected-error@+2 {{class method cannot be used to define a SYCL kernel free function kernel}}
+[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 2)]] void 
+singleTaskKernelMethod(int Value) {
+}
+
+// expected-error@+2 {{static class method cannot be used to define a SYCL kernel free function kernel}}
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]]
+static void StaticndRangeKernelMethod(int Value) {
+}
+
+// expected-error@+2 {{static class method cannot be used to define a SYCL kernel free function kernel}}
+[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 2)]] 
+static void StaticsingleTaskKernelMethod(int Value) {
+}
+
+};
+
+struct TestStruct {
+
+// expected-error@+2 {{class method cannot be used to define a SYCL kernel free function kernel}}
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]] void 
+ndRangeKernelMethod(int Value) {
+}
+
+// expected-error@+2 {{class method cannot be used to define a SYCL kernel free function kernel}}
+[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 2)]] void 
+singleTaskKernelMethod(int Value) {
+}
+
+// expected-error@+2 {{static class method cannot be used to define a SYCL kernel free function kernel}}
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]]
+static void StaticndRangeKernelMethod(int Value) {
+}
+
+// expected-error@+2 {{static class method cannot be used to define a SYCL kernel free function kernel}}
+[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 2)]] 
+static void StaticsingleTaskKernelMethod(int Value) {
+}
+
+};

sycl/test/extensions/free_function_kernels/free_function_kernels_diagnostics.cpp

@@ -27,17 +27,6 @@ void singleTaskKernelVariadic(...) {}
 SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::nd_range_kernel<1>)
 void ndRangeKernelVariadic(...) {}
 
-class DummyClass {
-public:
-  // Diagnostic for these violations of the restrictions haven't been
-  // implemented yet.
-  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::single_task_kernel)
-  void singleTaskKernelNonStaticMemberFunc(int *Ptr) {}
-
-  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::nd_range_kernel<2>)
-  void ndRangeKernelNonStaticMemberFunc(float *Ptr) {}
-};
-
 // expected-error@+2 {{SYCL free function kernel should have return type 'void'}}
 SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclexp::single_task_kernel)
 float singleTaskKernelNonVoid() { return 0.0F; }