[SYCL] Fix crash caused by functor without call operator as Kernel

clang/lib/Sema/SemaSYCL.cpp

@@ -976,6 +976,28 @@ static QualType GetSYCLKernelObjectType(const FunctionDecl *KernelCaller) {
   return KernelParamTy.getUnqualifiedType();
 }
 
+static CXXMethodDecl *getOperatorParens(const CXXRecordDecl *Rec) {
+  for (auto *MD : Rec->methods()) {
+    if (MD->getOverloadedOperator() == OO_Call)
+      return MD;
+  }
+  return nullptr;
+}
+
+// Fetch the associated call operator of the kernel object
+// (of either the lambda or the function object).
+static CXXMethodDecl *
+GetCallOperatorOfKernelObject(const CXXRecordDecl *KernelObjType) {
+  CXXMethodDecl *CallOperator = nullptr;
+  if (!KernelObjType)
+    return CallOperator;
+  if (KernelObjType->isLambda())
+    CallOperator = KernelObjType->getLambdaCallOperator();
+  else
+    CallOperator = getOperatorParens(KernelObjType);
+  return CallOperator;
+}
+
 /// Creates a kernel parameter descriptor
 /// \param Src  field declaration to construct name from
 /// \param Ty   the desired parameter type
@@ -2399,14 +2421,6 @@ class SyclOptReportCreator : public SyclKernelFieldHandler {
   }
 };
 
-static CXXMethodDecl *getOperatorParens(const CXXRecordDecl *Rec) {
-  for (auto *MD : Rec->methods()) {
-    if (MD->getOverloadedOperator() == OO_Call)
-      return MD;
-  }
-  return nullptr;
-}
-
 static bool isESIMDKernelType(const CXXRecordDecl *KernelObjType) {
   const CXXMethodDecl *OpParens = getOperatorParens(KernelObjType);
   return (OpParens != nullptr) && OpParens->hasAttr<SYCLSimdAttr>();
@@ -2494,13 +2508,10 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
 
     // Fetch the kernel object and the associated call operator
     // (of either the lambda or the function object).
-    CXXRecordDecl *KernelObj =
+    const CXXRecordDecl *KernelObj =
         GetSYCLKernelObjectType(KernelCallerFunc)->getAsCXXRecordDecl();
-    CXXMethodDecl *WGLambdaFn = nullptr;
-    if (KernelObj->isLambda())
-      WGLambdaFn = KernelObj->getLambdaCallOperator();
-    else
-      WGLambdaFn = getOperatorParens(KernelObj);
+    CXXMethodDecl *WGLambdaFn = GetCallOperatorOfKernelObject(KernelObj);
+
     assert(WGLambdaFn && "non callable object is passed as kernel obj");
     // Mark the function that it "works" in a work group scope:
     // NOTE: In case of parallel_for_work_item the marker call itself is
@@ -3031,7 +3042,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
 static bool IsSYCLUnnamedKernel(Sema &SemaRef, const FunctionDecl *FD) {
   if (!SemaRef.getLangOpts().SYCLUnnamedLambda)
     return false;
-  QualType FunctorTy = GetSYCLKernelObjectType(FD);
+  const QualType FunctorTy = GetSYCLKernelObjectType(FD);
   QualType TmplArgTy = calculateKernelNameType(SemaRef.Context, FD);
   return SemaRef.Context.hasSameType(FunctorTy, TmplArgTy);
 }
@@ -3457,7 +3468,7 @@ void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc,
   const CXXRecordDecl *KernelObj =
       GetSYCLKernelObjectType(KernelFunc)->getAsCXXRecordDecl();
 
-  if (!KernelObj) {
+  if (!GetCallOperatorOfKernelObject(KernelObj)) {
     Diag(Args[0]->getExprLoc(), diag::err_sycl_kernel_not_function_object);
     KernelFunc->setInvalidDecl();
     return;

clang/test/SemaSYCL/undefined-functor.cpp

@@ -0,0 +1,50 @@
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -verify -fsyntax-only %s
+// This test checks that an error is thrown when a functor without a call operator defined is used as a kernel.
+
+#include "sycl.hpp"
+
+using namespace sycl;
+queue q;
+
+struct FunctorWithoutCallOperator; // expected-note {{forward declaration of 'FunctorWithoutCallOperator'}}
+
+struct StructDefined {
+  int x;
+};
+
+class FunctorWithCallOpDefined {
+  int x;
+  public:
+  void operator()() const {}
+};
+
+class FunctorWithCallOpDeclared {
+  int x;
+  public:
+  void operator()() const;
+};
+
+int main() {
+
+  q.submit([&](sycl::handler &cgh) {
+    // expected-error@#KernelSingleTask {{kernel parameter must be a lambda or function object}}
+    // expected-error@+2 {{invalid use of incomplete type 'FunctorWithoutCallOperator'}}
+    // expected-note@+1 {{in instantiation of function template specialization}}
+    cgh.single_task(FunctorWithoutCallOperator{});
+  });
+
+  q.submit([&](sycl::handler &cgh) {
+    // expected-error@#KernelSingleTask {{kernel parameter must be a lambda or function object}}
+    // expected-note@+1 {{in instantiation of function template specialization}}
+    cgh.single_task(StructDefined{});
+  });
+
+  q.submit([&](sycl::handler &cgh) {
+    cgh.single_task(FunctorWithCallOpDeclared{});
+  });
+
+  q.submit([&](sycl::handler &cgh) {
+    cgh.single_task(FunctorWithCallOpDefined{});
+  });
+
+}