[SYCL] do not emit diagnostics if class with virtual method is used as argument type

Author: dklochkov-emb

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -12970,7 +12970,7 @@ def err_free_function_first_occurrence_missing_attr: Error<
 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 of kernel free function can not %select{be virtually derived |have virtual methods}1">;
+  "argument type '%0' virtually inherited from base class is not supported as a SYCL kernel free function kernel arguments">;
 
 
 // SYCL kernel entry point diagnostics

clang/lib/Sema/SemaSYCL.cpp

@@ -1574,6 +1574,26 @@ class KernelObjVisitor {
     else if (ParamTy->isStructureOrClassType()) {
       if (KP_FOR_EACH(handleStructType, Param, ParamTy)) {
         CXXRecordDecl *RD = ParamTy->getAsCXXRecordDecl();
+        llvm::SmallVector<const CXXRecordDecl *, 8> WorkList;
+        llvm::SmallPtrSet<const CXXRecordDecl *, 8> Visited;
+        if (RD)
+          WorkList.push_back(RD);
+        while (!WorkList.empty()) {
+          const CXXRecordDecl *Cur = WorkList.pop_back_val();
+          if (!Cur || !Visited.insert(Cur).second)
+            continue;
+          for (const auto &Base : Cur->bases()) {
+            if (Base.isVirtual()) {
+              SemaSYCLRef.SemaRef.Diag(Param->getLocation(),
+                                       diag::err_free_function_virtual_arg)
+                  << Cur->getNameAsString() << Param->getSourceRange();
+              return;
+            }
+            const CXXRecordDecl *BaseDecl = Base.getType()->getAsCXXRecordDecl();
+            if (BaseDecl)
+              WorkList.push_back(BaseDecl);
+          }
+        }
         visitRecord(nullptr, Param, RD, ParamTy, Handlers...);
       }
     } else if (ParamTy->isUnionType())
@@ -5848,44 +5868,6 @@ void SemaSYCL::MarkDevices() {
     checkSYCLAddIRAttributesFunctionAttrConflicts(T.GetSYCLKernel());
   }
 }
-
-static bool hasVirtuals(Sema &S, ParmVarDecl *Param) {
-  const clang::CXXRecordDecl *ParamType =
-      Param->getType()->getAsCXXRecordDecl();
-  if (!ParamType || !ParamType->isThisDeclarationADefinition())
-    return false;
-
-  llvm::SmallVector<const clang::CXXRecordDecl *, 8> WorkList;
-  llvm::SmallPtrSet<const clang::CXXRecordDecl *, 8> Visited;
-  WorkList.push_back(ParamType);
-
-  while (!WorkList.empty()) {
-    const clang::CXXRecordDecl *Cur = WorkList.pop_back_val();
-    if (!Visited.insert(Cur).second)
-      continue;
-
-    // Check for virtual bases
-    for (const clang::CXXBaseSpecifier &Base : Cur->bases()) {
-      if (Base.isVirtual())
-        return S.Diag(Param->getLocation(), diag::err_free_function_virtual_arg)
-               << ParamType->getNameAsString() << 0 << Param->getSourceRange();
-      const clang::CXXRecordDecl *BaseDecl =
-          Base.getType()->getAsCXXRecordDecl();
-      if (BaseDecl && BaseDecl->isThisDeclarationADefinition())
-        WorkList.push_back(BaseDecl);
-    }
-
-    // Check for virtual member functions
-    for (const auto *Method : Cur->methods()) {
-      if (Method->isVirtual()) {
-        return S.Diag(Param->getLocation(), diag::err_free_function_virtual_arg)
-               << ParamType->getNameAsString() << 1 << Param->getSourceRange();
-      }
-    }
-  }
-  return false;
-}
-
 static bool CheckFreeFunctionDiagnostics(Sema &S, const FunctionDecl *FD) {
   if (FD->isVariadic()) {
     return S.Diag(FD->getLocation(), diag::err_free_function_variadic_args);
@@ -5908,8 +5890,6 @@ static bool CheckFreeFunctionDiagnostics(Sema &S, const FunctionDecl *FD) {
                     diag::err_free_function_with_default_arg)
              << Param->getSourceRange();
     }
-    if (hasVirtuals(S, Param))
-      return true;
   }
   return false;
 }

clang/test/SemaSYCL/free_function_negative.cpp

@@ -119,23 +119,24 @@ static void StaticsingleTaskKernelMethod(int Value) {
 
 };
 
+
 class Base {};
 class Derived : virtual public Base {};
 
-// expected-error@+2 {{argument type Derived of kernel free function can not be virtually derived}}
+// expected-error@+2 2 {{argument type 'Derived' virtually inherited from base class is not supported as a SYCL kernel free function kernel arguments}}
 [[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 2)]] 
 void VirtualInheritArg(Derived Value) {
 }
 
-// expected-error@+2 {{argument type Derived of kernel free function can not be virtually derived}}
+// expected-error@+2 4 {{argument type 'Derived' virtually inherited from base class is not supported as a SYCL kernel free function kernel arguments}}
 [[__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 {
 };
 
-// expected-error@+2 {{argument type Derived1 of kernel free function can not be virtually derived}}
+// expected-error@+2 2 {{argument type 'Derived' virtually inherited from base class is not supported as a SYCL kernel free function kernel arguments}}
 [[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 2)]] 
 void VirtualInheritArg2(Derived1 Value) {
 }
@@ -144,55 +145,8 @@ class Base1 {};
 class Derived2 : public Base1, public virtual Base {
 };
 
-// expected-error@+2 {{argument type Derived2 of kernel free function can not be virtually derived}}
+// expected-error@+2 2 {{argument type 'Derived2' virtually inherited from base class is not supported as a SYCL kernel free function kernel arguments}}
 [[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 2)]] 
 void VirtualInheritArg3(Derived2 Value) {
 }
 
-struct Derived3 : virtual public Base {
-};
-
-// expected-error@+2 {{argument type Derived3 of kernel free function can not be virtually derived}}
-[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 2)]] 
-void VirtualInheritArg4(Derived3 Value) {
-}
-
-class Class {
-    public:
-    virtual void virtualMethod() {}
-};
-
-// expected-error@+2 {{argument type Class of kernel free function can not have virtual methods}}
-[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 2)]] 
-void VirtualMethodArg1(Class Value) {
-}
-
-class Class1: public Class {
-    public:
-    void NewMethod() {}
-};
-
-// expected-error@+2 {{argument type Class1 of kernel free function can not have virtual methods}}
-[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 2)]] 
-void VirtualMethodArg2(Class1 Value) {
-}
-
-class Class2: Base1 {
-    public:
-    virtual void virtualMethod() {}
-};
-
-// expected-error@+2 {{argument type Class2 of kernel free function can not have virtual methods}}
-[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]] 
-void VirtualMethodArg2(Class2 Value) {
-}
-
-class Class3 {
-    public:
-    virtual ~Class3() {}
-};
-
-// expected-error@+2 {{argument type Class3 of kernel free function can not have virtual methods}}
-[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]] 
-void VirtualMethodArg3(Class3 Value) {
-}

sycl/test-e2e/FreeFunctionKernels/virtual_methods.cpp

@@ -0,0 +1,96 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+#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;
+}
+
+static 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>();
+  const 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;
+}