[SYCL] add clang diagnostic for illegal types of kernel free function arguments

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -12969,6 +12969,8 @@ 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">;
+def err_free_function_virtual_arg : Error<
+  "argument type '%0' virtually inherited from base class `%1` is not supported as a SYCL kernel argument">;
 
 
 // SYCL kernel entry point diagnostics

clang/lib/Sema/SemaSYCL.cpp

@@ -1257,12 +1257,19 @@ constructFreeFunctionKernelName(const FunctionDecl *FreeFunc,
     MC.mangleName(FreeFunc, Out);
     std::string MangledName(Out.str());
     size_t StartNums = MangledName.find_first_of("0123456789");
-    size_t EndNums = MangledName.find_first_not_of("0123456789", StartNums);
-    size_t NameLength =
-        std::stoi(MangledName.substr(StartNums, EndNums - StartNums));
-    size_t NewNameLength = 14 /*length of __sycl_kernel_*/ + NameLength;
-    NewName = MangledName.substr(0, StartNums) + std::to_string(NewNameLength) +
-              "__sycl_kernel_" + MangledName.substr(EndNums);
+    if (StartNums == std::string::npos) {
+      // Microsoft mangling name has template like ?FunctionName@@YAXH@Z
+      NewName =
+          MangledName.substr(0, 1) + "sycl_kernel_" + MangledName.substr(1);
+    } else {
+      size_t EndNums = MangledName.find_first_not_of("0123456789", StartNums);
+      size_t NameLength =
+          std::stoi(MangledName.substr(StartNums, EndNums - StartNums));
+      size_t NewNameLength = 14 /*length of __sycl_kernel_*/ + NameLength;
+      NewName = MangledName.substr(0, StartNums) +
+                std::to_string(NewNameLength) + "__sycl_kernel_" +
+                MangledName.substr(EndNums);
+    }
   }
   StableName = NewName;
   return {NewName, StableName};
@@ -1463,6 +1470,10 @@ class KernelObjVisitor {
       if (isSyclSpecialType(BaseTy, SemaSYCLRef))
         (void)std::initializer_list<int>{
             (Handlers.handleSyclSpecialType(Owner, Base, BaseTy), 0)...};
+      else if (Base.isVirtual())
+        SemaSYCLRef.SemaRef.Diag(Owner->getLocation(),
+                                 diag::err_free_function_virtual_arg)
+            << Owner->getNameAsString() << Base.getType().getAsString();
       else
         // For all other bases, visit the record
         visitRecord(Owner, Base, BaseTy->getAsCXXRecordDecl(), BaseTy,

clang/test/SemaSYCL/free_function_negative.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -verify=expected -fsycl-int-header=%t.h %s
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -verify=expected %s
 
 #include "sycl.hpp"
 
@@ -118,3 +118,33 @@ static void StaticsingleTaskKernelMethod(int Value) {
 }
 
 };
+
+
+class Base {};
+// expected-error@+1 8 {{argument type 'Derived' virtually inherited from base class `Base` is not supported as a SYCL kernel argument}}
+class Derived : virtual public Base {};
+
+[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 2)]] 
+void VirtualInheritArg(Derived Value) {
+}
+
+[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 2)]] 
+void VirtualInheritArg1(int a, Derived Value, float b, Derived Value1) {
+}
+
+class Derived1 : public Derived {
+};
+
+[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 2)]] 
+void VirtualInheritArg2(Derived1 Value) {
+}
+
+class Base1 {};
+
+// expected-error@+1 2 {{argument type 'Derived2' virtually inherited from base class `Base` is not supported as a SYCL kernel argumen}}
+class Derived2 : public Base1, public virtual Base {
+};
+
+[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 2)]] 
+void VirtualInheritArg3(Derived2 Value) {
+}

sycl/test-e2e/FreeFunctionKernels/virtual_methods.cpp

@@ -0,0 +1,101 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+/*
+ * Test to check class/struct type with virtual methods as SYCL free function
+ * kernel arguments.
+ */
+
+#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;
+static constexpr auto FFTestMark = "Free function Kernel Test:";
+static constexpr float offset = 1.1f;
+
+class Base {
+public:
+  virtual void virtual_method(float start) = 0;
+  virtual ~Base() = default;
+};
+
+class TestClass : public Base {
+  float data = 0.0f;
+
+public:
+  void virtual_method(float start) override {}
+
+  float calculate(float start, size_t id) {
+    return start + static_cast<float>(id) + data;
+  }
+
+  void setData(float value) { data = value; }
+};
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<2>))
+void func_range(TestClass *acc, float *ptr) {
+  size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
+  ptr[id] = acc->calculate(3.14f, id);
+}
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::single_task_kernel))
+void func_single(TestClass *acc, float *ptr) {
+  size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
+  ptr[id] = acc->calculate(3.14f, id);
+}
+
+int check_result(float *ptr) {
+  for (size_t i = 0; i < NUM; ++i) {
+    const float expected = 3.14f + static_cast<float>(i) + offset;
+    if (ptr[i] != expected)
+      return 1;
+  }
+  return 0;
+}
+
+int call_kernel_code(sycl::queue &q, sycl::kernel &kernel) {
+  float *ptr = sycl::malloc_shared<float>(NUM, q);
+  TestClass *obj = sycl::malloc_shared<TestClass>(1, q);
+  obj->setData(offset);
+
+  q.submit([&](sycl::handler &cgh) {
+     cgh.set_args(obj, ptr);
+     sycl::nd_range ndr{{NUM}, {WGSIZE}};
+     cgh.parallel_for(ndr, kernel);
+   }).wait();
+  int ret = check_result(ptr);
+  sycl::free(ptr, q);
+  sycl::free(obj, q);
+  return ret;
+}
+
+template <auto *Func>
+int test_arg_with_virtual_method(sycl::queue &q, sycl::context &ctxt,
+                                 std::string_view name) {
+  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>();
+  int ret = call_kernel_code(q, k_func);
+  if (ret != 0)
+    std::cerr << FFTestMark << name << " failed\n";
+  return ret;
+}
+
+int main() {
+  sycl::queue q;
+  sycl::context ctxt = q.get_context();
+  sycl::device dev = q.get_device();
+
+  int ret =
+      test_arg_with_virtual_method<func_range>(q, ctxt, "virtual_method_range");
+  ret |= test_arg_with_virtual_method<func_single>(q, ctxt,
+                                                   "virtual_method_single");
+  return ret;
+}