diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 70b2e08c833a0..a5aa3f03afe14 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/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(); @@ -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; diff --git a/clang/test/SemaSYCL/undefined-functor.cpp b/clang/test/SemaSYCL/undefined-functor.cpp new file mode 100644 index 0000000000000..496987a3953d0 --- /dev/null +++ b/clang/test/SemaSYCL/undefined-functor.cpp @@ -0,0 +1,40 @@ +// 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 {} +}; + +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(FunctorWithCallOpDefined{}); + }); + +}