[SYCL] the first kernel function declaration should be added with attribute (#18405)

Docs [Defining a free function
kernel](https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/proposed/sycl_ext_oneapi_free_function_kernels.asciidoc#defining-a-free-function-kernel)
contains:
The property must appear on the first declaration of the function in the
translation unit. Redeclarations of the function may optionally be
decorated with the same property if the property argument is the same.
The effect is the same regardless of whether redeclarations are so
decorated.

This PR changes `isFreeFunction` to define if the first occurrence of
kernel free function has attribute or not. New clang diagnostics was
added to highlight that first occurrence does not have attribute but the
next has.

Author: dklochkov-emb

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -12871,7 +12871,8 @@ def err_free_function_variadic_args : Error<
   "free function kernel cannot be a variadic function">;
 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">;
 
 // SYCL kernel entry point diagnostics
 def err_sycl_entry_point_invalid : Error<

clang/lib/Sema/SemaSYCL.cpp

@@ -1159,14 +1159,33 @@ static target getAccessTarget(QualType FieldTy,
 }
 
 bool SemaSYCL::isFreeFunction(const FunctionDecl *FD) {
-  for (auto *IRAttr : FD->specific_attrs<SYCLAddIRAttributesFunctionAttr>()) {
-    SmallVector<std::pair<std::string, std::string>, 4> NameValuePairs =
-        IRAttr->getAttributeNameValuePairs(getASTContext());
-    for (const auto &NameValuePair : NameValuePairs) {
-      if (NameValuePair.first == "sycl-nd-range-kernel" ||
-          NameValuePair.first == "sycl-single-task-kernel") {
+  SourceLocation Loc = FD->getLocation();
+  bool NextDeclaredWithAttr = false;
+  for (FunctionDecl *Redecl : FD->redecls()) {
+    bool IsFreeFunctionAttr = false;
+    for (auto *IRAttr :
+         Redecl->specific_attrs<SYCLAddIRAttributesFunctionAttr>()) {
+      SmallVector<std::pair<std::string, std::string>, 4> NameValuePairs =
+          IRAttr->getAttributeNameValuePairs(getASTContext());
+      const auto it = std::find_if(
+          NameValuePairs.begin(), NameValuePairs.end(),
+          [](const auto &NameValuePair) {
+            return NameValuePair.first == "sycl-nd-range-kernel" ||
+                   NameValuePair.first == "sycl-single-task-kernel";
+          });
+      IsFreeFunctionAttr = it != NameValuePairs.end();
+    }
+    if (Redecl->isFirstDecl()) {
+      if (IsFreeFunctionAttr)
         return true;
+      if (NextDeclaredWithAttr) {
+        Diag(Loc, diag::err_free_function_first_occurrence_missing_attr);
+        Diag(Redecl->getLocation(), diag::note_previous_declaration);
+        return false;
       }
+    } else {
+      Loc = Redecl->getLocation();
+      NextDeclaredWithAttr = IsFreeFunctionAttr;
     }
   }
   return false;

clang/test/SemaSYCL/free_function_negative.cpp

@@ -10,6 +10,36 @@ foo(int start, ...) { // expected-error {{free function kernel cannot be a varia
 foo1(int start, ...) { // expected-error {{free function kernel cannot be a variadic function}}
 }
 
+// expected-note@+1 {{conflicting attribute is here}}
+[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 1)]] void
+foo2(int start);
+
+// expected-error@+1 {{attribute 'add_ir_attributes_function' is already applied with different arguments}}
+[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 2)]] void 
+foo2(int start) {
+}
+
+// expected-note@+1 {{previous declaration is here}}
+void foo3(int start, int *ptr);
+
+// expected-error@+2 {{the first occurrence of SYCL kernel free function should be declared with 'sycl-nd-range-kernel' or 'sycl-single-task-kernel' compile time properties}}
+[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 2)]] void 
+foo3(int start, int *ptr){}
+
+// expected-note@+1 {{previous declaration is here}}
+void foo4(float start, float *ptr);
+
+// expected-error@+2 {{the first occurrence of SYCL kernel free function should be declared with 'sycl-nd-range-kernel' or 'sycl-single-task-kernel' compile time properties}}
+[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 2)]] void 
+foo4(float start, float *ptr);
+
+[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 2)]] void 
+foo4(float start, float *ptr);
+
+[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 2)]] void 
+foo4(float start, float *ptr){}
+
+
 // expected-error@+2 {{a function with a default argument value cannot be used to define SYCL free function kernel}}
 [[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 2)]] void 
 singleTaskKernelDefaultValues(int Value = 1) {

sycl/test-e2e/FreeFunctionKernels/redeclaration.cpp

@@ -0,0 +1,102 @@
+// REQUIRES: aspect-usm_shared_allocations
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+// The name mangling for free function kernels currently does not work with PTX.
+// UNSUPPORTED: cuda, hip
+// UNSUPPORTED-INTENDED: Not implemented yet for Nvidia/AMD backends.
+
+#include <iostream>
+#include <sycl/detail/core.hpp>
+#include <sycl/ext/oneapi/free_function_queries.hpp>
+#include <sycl/kernel_bundle.hpp>
+#include <sycl/usm.hpp>
+
+namespace syclext = sycl::ext::oneapi;
+namespace syclexp = sycl::ext::oneapi::experimental;
+
+static constexpr size_t NUM = 1024;
+static constexpr size_t WGSIZE = 16;
+
+template <typename T> int check_result(T *ptr, T value) {
+  for (size_t i = 0; i < NUM; ++i) {
+    const T expected = value + static_cast<T>(i);
+    if (ptr[i] != expected) {
+      std::cout << "Kernel execution did not produce the expected result\n";
+      return 1;
+    }
+  }
+  return 0;
+}
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
+void free_func(int *ptr, int start);
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
+void free_func(int *ptr, int start);
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
+void free_func1(int *ptr, int start);
+
+void free_func1(int *ptr, int start);
+
+template <typename T>
+static int call_kernel_code(sycl::queue &q, sycl::kernel &kernel, T value) {
+  T *ptr = sycl::malloc_shared<T>(NUM, q);
+  q.submit([&](sycl::handler &cgh) {
+     if (value == 0)
+       cgh.set_args(ptr);
+     else
+       cgh.set_args(ptr, value);
+     sycl::nd_range ndr{{NUM}, {WGSIZE}};
+     cgh.parallel_for(ndr, kernel);
+   }).wait();
+  const int ret = check_result(ptr, value);
+  sycl::free(ptr, q);
+  return ret;
+}
+
+#define KERNEL_CODE(start, ptr, type)                                          \
+  size_t id =                                                                  \
+      syclext::this_work_item::get_nd_item<1>().get_global_linear_id();        \
+  ptr[id] = static_cast<type>(start) + static_cast<type>(id);
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
+void free_func(int *ptr, int start) { KERNEL_CODE(start, ptr, int); }
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
+void free_func1(int *ptr, int start) { KERNEL_CODE(start, ptr, int); }
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
+void free_func2(int *ptr, int start) { KERNEL_CODE(start, ptr, int); }
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
+void free_func2(float *ptr, float start) { KERNEL_CODE(start, ptr, float); }
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
+void free_func2(int *ptr) { KERNEL_CODE(0, ptr, int); }
+
+template <auto Func, typename T>
+int test_declarations(sycl::queue &q, sycl::context &ctxt, T value) {
+  auto exe_bndl =
+      syclexp::get_kernel_bundle<Func, sycl::bundle_state::executable>(ctxt);
+  sycl::kernel k_func = exe_bndl.template ext_oneapi_get_kernel<Func>();
+  return call_kernel_code<T>(q, k_func, value);
+}
+
+int main() {
+  sycl::queue q;
+  sycl::context ctxt = q.get_context();
+
+  int result{0};
+  result |= test_declarations<free_func, int>(q, ctxt, 3);
+  result |= test_declarations<free_func1, int>(q, ctxt, 3);
+  result |=
+      test_declarations<static_cast<void (*)(int *, int)>(free_func2), int>(
+          q, ctxt, 3);
+  result |= test_declarations<static_cast<void (*)(float *, float)>(free_func2),
+                              float>(q, ctxt, 3.14f);
+  result |= test_declarations<static_cast<void (*)(int *)>(free_func2), int>(
+      q, ctxt, 0);
+  return result;
+}