From b2bf16732a2a5128d61f0337279ae5a27d7e8bc8 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Tue, 18 Oct 2022 13:09:28 -0700 Subject: [PATCH 01/12] [SYCL] Fix crash caused by functor without call operator passed as Kernel. --- clang/lib/Sema/SemaSYCL.cpp | 8 +++++++- clang/test/SemaSYCL/undefined-functor.cpp | 18 ++++++++++++++++++ 2 files changed, 25 insertions(+), 1 deletion(-) create mode 100644 clang/test/SemaSYCL/undefined-functor.cpp diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 70b2e08c833a0..559599a8e217b 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -2412,6 +2412,11 @@ static bool isESIMDKernelType(const CXXRecordDecl *KernelObjType) { return (OpParens != nullptr) && OpParens->hasAttr(); } +static bool isSYCLKernelDefinedAsAFunctor(const CXXRecordDecl *KernelObjType) { + const CXXMethodDecl *OpParens = getOperatorParens(KernelObjType); + return (OpParens != nullptr); +} + class SyclKernelBodyCreator : public SyclKernelFieldHandler { SyclKernelDeclCreator &DeclCreator; llvm::SmallVector BodyStmts; @@ -3457,7 +3462,8 @@ void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc, const CXXRecordDecl *KernelObj = GetSYCLKernelObjectType(KernelFunc)->getAsCXXRecordDecl(); - if (!KernelObj) { + bool IsKernelAValidFunctor = isSYCLKernelDefinedAsAFunctor(KernelObj); + if (!KernelObj || !IsKernelAValidFunctor) { 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..35cfc0df9739a --- /dev/null +++ b/clang/test/SemaSYCL/undefined-functor.cpp @@ -0,0 +1,18 @@ +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -verify -fsyntax-only %s + +#include "sycl.hpp" + +using namespace sycl; +queue q; + +struct FunctorWithoutCallOperator; // expected-note {{forward declaration of 'FunctorWithoutCallOperator'}} + +int main() { + // expected-error@#KernelSingleTask {{kernel parameter must be a lambda or function object}} + q.submit([&](sycl::handler &cgh) { + // expected-error@+2 {{invalid use of incomplete type 'FunctorWithoutCallOperator'}} + // expected-note@+1 {{in instantiation of function template specialization}} + cgh.single_task(FunctorWithoutCallOperator{}); + }); + +} From e77b1348fdc8c6c81cfcb99e296fe0123ba0bdec Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Thu, 20 Oct 2022 20:48:46 -0700 Subject: [PATCH 02/12] Fix failing tests --- clang/lib/Sema/SemaSYCL.cpp | 16 +++++++++++----- 1 file changed, 11 insertions(+), 5 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 559599a8e217b..54084236020d5 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -2412,9 +2412,15 @@ static bool isESIMDKernelType(const CXXRecordDecl *KernelObjType) { return (OpParens != nullptr) && OpParens->hasAttr(); } -static bool isSYCLKernelDefinedAsAFunctor(const CXXRecordDecl *KernelObjType) { - const CXXMethodDecl *OpParens = getOperatorParens(KernelObjType); - return (OpParens != nullptr); +// Fetch the associated call operator of the kernel object +// (of either the lambda or the function object). +static bool IsCallOperatorDefined(const CXXRecordDecl *KernelObjType) { + const CXXMethodDecl *CallOperator = nullptr; + if (KernelObjType->isLambda()) + CallOperator = KernelObjType->getLambdaCallOperator(); + else + CallOperator = getOperatorParens(KernelObjType); + return (CallOperator != nullptr); } class SyclKernelBodyCreator : public SyclKernelFieldHandler { @@ -3462,8 +3468,8 @@ void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc, const CXXRecordDecl *KernelObj = GetSYCLKernelObjectType(KernelFunc)->getAsCXXRecordDecl(); - bool IsKernelAValidFunctor = isSYCLKernelDefinedAsAFunctor(KernelObj); - if (!KernelObj || !IsKernelAValidFunctor) { + bool IsKernelTypeValid = IsCallOperatorDefined(KernelObj); + if (!KernelObj || !IsKernelTypeValid) { Diag(Args[0]->getExprLoc(), diag::err_sycl_kernel_not_function_object); KernelFunc->setInvalidDecl(); return; From 6fd86bf6793ebc7deba20043daf594ce914928aa Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Thu, 20 Oct 2022 22:11:24 -0700 Subject: [PATCH 03/12] Fix failing test --- clang/lib/Sema/SemaSYCL.cpp | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 54084236020d5..90fa30ac21c32 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -2416,6 +2416,8 @@ static bool isESIMDKernelType(const CXXRecordDecl *KernelObjType) { // (of either the lambda or the function object). static bool IsCallOperatorDefined(const CXXRecordDecl *KernelObjType) { const CXXMethodDecl *CallOperator = nullptr; + if (!KernelObjType) + return false; if (KernelObjType->isLambda()) CallOperator = KernelObjType->getLambdaCallOperator(); else @@ -3469,7 +3471,7 @@ void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc, GetSYCLKernelObjectType(KernelFunc)->getAsCXXRecordDecl(); bool IsKernelTypeValid = IsCallOperatorDefined(KernelObj); - if (!KernelObj || !IsKernelTypeValid) { + if (!IsKernelTypeValid) { Diag(Args[0]->getExprLoc(), diag::err_sycl_kernel_not_function_object); KernelFunc->setInvalidDecl(); return; From 373dd699c4ea9796ce855152e8dd04ae4c534e52 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Fri, 21 Oct 2022 11:01:30 -0700 Subject: [PATCH 04/12] Add test comment --- clang/test/SemaSYCL/undefined-functor.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/clang/test/SemaSYCL/undefined-functor.cpp b/clang/test/SemaSYCL/undefined-functor.cpp index 35cfc0df9739a..40dc1ff2eb42b 100644 --- a/clang/test/SemaSYCL/undefined-functor.cpp +++ b/clang/test/SemaSYCL/undefined-functor.cpp @@ -1,4 +1,5 @@ // 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" From aa0a48be5d839347f3336693182e2cf0466448a4 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Fri, 21 Oct 2022 14:18:54 -0700 Subject: [PATCH 05/12] Refactor common code --- clang/lib/Sema/SemaSYCL.cpp | 55 ++++++++++++++++++------------------- 1 file changed, 26 insertions(+), 29 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 90fa30ac21c32..babd0d9dea562 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -976,6 +976,27 @@ 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). +CXXMethodDecl *IsCallOperatorDefined(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,32 +2420,11 @@ 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(); } -// Fetch the associated call operator of the kernel object -// (of either the lambda or the function object). -static bool IsCallOperatorDefined(const CXXRecordDecl *KernelObjType) { - const CXXMethodDecl *CallOperator = nullptr; - if (!KernelObjType) - return false; - if (KernelObjType->isLambda()) - CallOperator = KernelObjType->getLambdaCallOperator(); - else - CallOperator = getOperatorParens(KernelObjType); - return (CallOperator != nullptr); -} - class SyclKernelBodyCreator : public SyclKernelFieldHandler { SyclKernelDeclCreator &DeclCreator; llvm::SmallVector BodyStmts; @@ -2509,11 +2509,8 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { // (of either the lambda or the function object). CXXRecordDecl *KernelObj = GetSYCLKernelObjectType(KernelCallerFunc)->getAsCXXRecordDecl(); - CXXMethodDecl *WGLambdaFn = nullptr; - if (KernelObj->isLambda()) - WGLambdaFn = KernelObj->getLambdaCallOperator(); - else - WGLambdaFn = getOperatorParens(KernelObj); + CXXMethodDecl *WGLambdaFn = IsCallOperatorDefined(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 @@ -3467,11 +3464,11 @@ void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc, if (!LangOpts.SYCLIsDevice) return; - const CXXRecordDecl *KernelObj = + CXXRecordDecl *KernelObj = GetSYCLKernelObjectType(KernelFunc)->getAsCXXRecordDecl(); - bool IsKernelTypeValid = IsCallOperatorDefined(KernelObj); - if (!IsKernelTypeValid) { + CXXMethodDecl *CallOperator = IsCallOperatorDefined(KernelObj); + if (!CallOperator) { Diag(Args[0]->getExprLoc(), diag::err_sycl_kernel_not_function_object); KernelFunc->setInvalidDecl(); return; From 3eecf940a3c136d557c9fad705961912cdc82102 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Fri, 21 Oct 2022 21:18:08 -0700 Subject: [PATCH 06/12] Address review comments --- clang/lib/Sema/SemaSYCL.cpp | 8 ++++---- clang/test/SemaSYCL/undefined-functor.cpp | 13 ++++++++++++- 2 files changed, 16 insertions(+), 5 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index babd0d9dea562..fcfb0b928b9b2 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -986,7 +986,8 @@ static CXXMethodDecl *getOperatorParens(const CXXRecordDecl *Rec) { // Fetch the associated call operator of the kernel object // (of either the lambda or the function object). -CXXMethodDecl *IsCallOperatorDefined(const CXXRecordDecl *KernelObjType) { +CXXMethodDecl * +GetCallOperatorOfKernelObject(const CXXRecordDecl *KernelObjType) { CXXMethodDecl *CallOperator = nullptr; if (!KernelObjType) return CallOperator; @@ -2509,7 +2510,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { // (of either the lambda or the function object). CXXRecordDecl *KernelObj = GetSYCLKernelObjectType(KernelCallerFunc)->getAsCXXRecordDecl(); - CXXMethodDecl *WGLambdaFn = IsCallOperatorDefined(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: @@ -3467,8 +3468,7 @@ void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc, CXXRecordDecl *KernelObj = GetSYCLKernelObjectType(KernelFunc)->getAsCXXRecordDecl(); - CXXMethodDecl *CallOperator = IsCallOperatorDefined(KernelObj); - if (!CallOperator) { + 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 index 40dc1ff2eb42b..a7391f50da416 100644 --- a/clang/test/SemaSYCL/undefined-functor.cpp +++ b/clang/test/SemaSYCL/undefined-functor.cpp @@ -8,12 +8,23 @@ queue q; struct FunctorWithoutCallOperator; // expected-note {{forward declaration of 'FunctorWithoutCallOperator'}} +struct StructDefined { + int x; +}; + int main() { - // expected-error@#KernelSingleTask {{kernel parameter must be a lambda or function object}} + 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{}); + }); + } From 91dcdbb45d54deaec41d470d6667cb409c366564 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Mon, 24 Oct 2022 14:33:26 -0700 Subject: [PATCH 07/12] Add 'const' qualifier back and add more test cases. --- clang/lib/Sema/SemaSYCL.cpp | 6 +++--- clang/test/SemaSYCL/undefined-functor.cpp | 20 ++++++++++++++++++++ 2 files changed, 23 insertions(+), 3 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index fcfb0b928b9b2..3cb328211ff42 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -2508,7 +2508,7 @@ 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 = GetCallOperatorOfKernelObject(KernelObj); @@ -3042,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); } @@ -3465,7 +3465,7 @@ void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc, if (!LangOpts.SYCLIsDevice) return; - CXXRecordDecl *KernelObj = + const CXXRecordDecl *KernelObj = GetSYCLKernelObjectType(KernelFunc)->getAsCXXRecordDecl(); if (!GetCallOperatorOfKernelObject(KernelObj)) { diff --git a/clang/test/SemaSYCL/undefined-functor.cpp b/clang/test/SemaSYCL/undefined-functor.cpp index a7391f50da416..945bceec644af 100644 --- a/clang/test/SemaSYCL/undefined-functor.cpp +++ b/clang/test/SemaSYCL/undefined-functor.cpp @@ -12,6 +12,18 @@ struct StructDefined { int x; }; +class FunctorWithCallOpDeclared { + int x; + public: + void operator()() const {} +}; + +class FunctorWithCallOpDefined { + int x; + public: + void operator()() const {}; +}; + int main() { q.submit([&](sycl::handler &cgh) { @@ -27,4 +39,12 @@ int main() { cgh.single_task(StructDefined{}); }); + q.submit([&](sycl::handler &cgh) { + cgh.single_task(FunctorWithCallOpDeclared{}); + }); + + q.submit([&](sycl::handler &cgh) { + cgh.single_task(FunctorWithCallOpDefined{}); + }); + } From c7798556b624e901e297a587619ea2b5ee03d489 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Tue, 25 Oct 2022 10:22:03 -0700 Subject: [PATCH 08/12] Remove unused function parameter. --- clang/include/clang/Sema/Sema.h | 2 +- clang/lib/Sema/SemaChecking.cpp | 2 +- clang/lib/Sema/SemaSYCL.cpp | 2 +- 3 files changed, 3 insertions(+), 3 deletions(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index f6ba02c767ca6..fc1af3ac59ea3 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -13500,7 +13500,7 @@ class Sema final { bool IsMemberFunction, SourceLocation Loc, SourceRange Range, VariadicCallType CallType); - void CheckSYCLKernelCall(FunctionDecl *CallerFunc, SourceRange CallLoc, + void CheckSYCLKernelCall(FunctionDecl *CallerFunc, ArrayRef Args); bool CheckObjCString(Expr *Arg); diff --git a/clang/lib/Sema/SemaChecking.cpp b/clang/lib/Sema/SemaChecking.cpp index 945ba70a7472b..78e496f52a2cf 100644 --- a/clang/lib/Sema/SemaChecking.cpp +++ b/clang/lib/Sema/SemaChecking.cpp @@ -5967,7 +5967,7 @@ void Sema::checkCall(NamedDecl *FDecl, const FunctionProtoType *Proto, diagnoseArgDependentDiagnoseIfAttrs(FD, ThisArg, Args, Loc); if (FD && FD->hasAttr()) - CheckSYCLKernelCall(FD, Range, Args); + CheckSYCLKernelCall(FD, Args); // Diagnose variadic calls in SYCL. if (FD && FD->isVariadic() && getLangOpts().SYCLIsDevice && diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 3cb328211ff42..42642c0edb6b2 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -3451,7 +3451,7 @@ class SYCLKernelNameTypeVisitor } }; -void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc, +void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, ArrayRef Args) { QualType KernelNameType = calculateKernelNameType(getASTContext(), KernelFunc); From 7c2b3b06c1aba9d80ca7d196c1c488a290c994ac Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Tue, 25 Oct 2022 13:48:34 -0700 Subject: [PATCH 09/12] Fix function names --- clang/include/clang/Sema/Sema.h | 2 +- clang/lib/Sema/SemaChecking.cpp | 2 +- clang/lib/Sema/SemaSYCL.cpp | 2 +- clang/test/SemaSYCL/undefined-functor.cpp | 4 ++-- 4 files changed, 5 insertions(+), 5 deletions(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index fc1af3ac59ea3..f6ba02c767ca6 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -13500,7 +13500,7 @@ class Sema final { bool IsMemberFunction, SourceLocation Loc, SourceRange Range, VariadicCallType CallType); - void CheckSYCLKernelCall(FunctionDecl *CallerFunc, + void CheckSYCLKernelCall(FunctionDecl *CallerFunc, SourceRange CallLoc, ArrayRef Args); bool CheckObjCString(Expr *Arg); diff --git a/clang/lib/Sema/SemaChecking.cpp b/clang/lib/Sema/SemaChecking.cpp index 78e496f52a2cf..945ba70a7472b 100644 --- a/clang/lib/Sema/SemaChecking.cpp +++ b/clang/lib/Sema/SemaChecking.cpp @@ -5967,7 +5967,7 @@ void Sema::checkCall(NamedDecl *FDecl, const FunctionProtoType *Proto, diagnoseArgDependentDiagnoseIfAttrs(FD, ThisArg, Args, Loc); if (FD && FD->hasAttr()) - CheckSYCLKernelCall(FD, Args); + CheckSYCLKernelCall(FD, Range, Args); // Diagnose variadic calls in SYCL. if (FD && FD->isVariadic() && getLangOpts().SYCLIsDevice && diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 42642c0edb6b2..3cb328211ff42 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -3451,7 +3451,7 @@ class SYCLKernelNameTypeVisitor } }; -void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, +void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc, ArrayRef Args) { QualType KernelNameType = calculateKernelNameType(getASTContext(), KernelFunc); diff --git a/clang/test/SemaSYCL/undefined-functor.cpp b/clang/test/SemaSYCL/undefined-functor.cpp index 945bceec644af..0efe9678a6395 100644 --- a/clang/test/SemaSYCL/undefined-functor.cpp +++ b/clang/test/SemaSYCL/undefined-functor.cpp @@ -12,13 +12,13 @@ struct StructDefined { int x; }; -class FunctorWithCallOpDeclared { +class FunctorWithCallOpDefined { int x; public: void operator()() const {} }; -class FunctorWithCallOpDefined { +class FunctorWithCallOpDeclared { int x; public: void operator()() const {}; From 7d1c2269f4705effb9c77dba09c5f83765c66647 Mon Sep 17 00:00:00 2001 From: Srividya Sundaram Date: Wed, 26 Oct 2022 10:31:41 -0700 Subject: [PATCH 10/12] Update clang/lib/Sema/SemaSYCL.cpp Co-authored-by: premanandrao --- clang/lib/Sema/SemaSYCL.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 3cb328211ff42..a5aa3f03afe14 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -986,7 +986,7 @@ static CXXMethodDecl *getOperatorParens(const CXXRecordDecl *Rec) { // Fetch the associated call operator of the kernel object // (of either the lambda or the function object). -CXXMethodDecl * +static CXXMethodDecl * GetCallOperatorOfKernelObject(const CXXRecordDecl *KernelObjType) { CXXMethodDecl *CallOperator = nullptr; if (!KernelObjType) From 8e6718313783c22f65fa56a903ccdbcaaa1ec9e7 Mon Sep 17 00:00:00 2001 From: Srividya Sundaram Date: Wed, 26 Oct 2022 10:31:47 -0700 Subject: [PATCH 11/12] Update clang/test/SemaSYCL/undefined-functor.cpp Co-authored-by: premanandrao --- clang/test/SemaSYCL/undefined-functor.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/test/SemaSYCL/undefined-functor.cpp b/clang/test/SemaSYCL/undefined-functor.cpp index 0efe9678a6395..a6c16032c8b13 100644 --- a/clang/test/SemaSYCL/undefined-functor.cpp +++ b/clang/test/SemaSYCL/undefined-functor.cpp @@ -21,7 +21,7 @@ class FunctorWithCallOpDefined { class FunctorWithCallOpDeclared { int x; public: - void operator()() const {}; + void operator()() const; }; int main() { From 94f8a596070e8088a2a9fd084b133709dacc6d02 Mon Sep 17 00:00:00 2001 From: Srividya Sundaram Date: Thu, 27 Oct 2022 14:42:49 -0700 Subject: [PATCH 12/12] Update undefined-functor.cpp --- clang/test/SemaSYCL/undefined-functor.cpp | 12 +----------- 1 file changed, 1 insertion(+), 11 deletions(-) diff --git a/clang/test/SemaSYCL/undefined-functor.cpp b/clang/test/SemaSYCL/undefined-functor.cpp index a6c16032c8b13..496987a3953d0 100644 --- a/clang/test/SemaSYCL/undefined-functor.cpp +++ b/clang/test/SemaSYCL/undefined-functor.cpp @@ -18,12 +18,6 @@ class FunctorWithCallOpDefined { void operator()() const {} }; -class FunctorWithCallOpDeclared { - int x; - public: - void operator()() const; -}; - int main() { q.submit([&](sycl::handler &cgh) { @@ -38,11 +32,7 @@ int main() { // 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{}); });