[SYCL] Add support for templated call operator in functors (#7970)

This PR enables kernels to be defined as named function object types,
where the operator() member function is templated.

Example:
```
class FunctorWithCallOpTemplated {
  int x;
  public:
  template <int x = 0>
  void operator()() const {}
};

q.submit([&](sycl::handler &cgh) {
    cgh.single_task(FunctorWithCallOpTemplated{});
  });
```

---------

Co-authored-by: premanandrao <[email protected]>

Author: srividya-sundaram

clang/lib/Sema/SemaSYCL.cpp

@@ -991,12 +991,26 @@ 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;
+// Get the call operator function associated with the function object
+// for both templated and non-templated operator()().
+
+static CXXMethodDecl *getFunctorCallOperator(const CXXRecordDecl *RD) {
+  DeclarationName Name =
+      RD->getASTContext().DeclarationNames.getCXXOperatorName(OO_Call);
+  DeclContext::lookup_result Calls = RD->lookup(Name);
+
+  if (Calls.empty())
+    return nullptr;
+
+  NamedDecl *CallOp = Calls.front();
+
+  if (CallOp == nullptr)
+    return nullptr;
+
+  if (const auto *CallOpTmpl = dyn_cast<FunctionTemplateDecl>(CallOp))
+    return cast<CXXMethodDecl>(CallOpTmpl->getTemplatedDecl());
+
+  return cast<CXXMethodDecl>(CallOp);
 }
 
 // Fetch the associated call operator of the kernel object
@@ -1009,7 +1023,7 @@ GetCallOperatorOfKernelObject(const CXXRecordDecl *KernelObjType) {
   if (KernelObjType->isLambda())
     CallOperator = KernelObjType->getLambdaCallOperator();
   else
-    CallOperator = getOperatorParens(KernelObjType);
+    CallOperator = getFunctorCallOperator(KernelObjType);
   return CallOperator;
 }
 
@@ -2802,7 +2816,7 @@ class SyclOptReportCreator : public SyclKernelFieldHandler {
 };
 
 static bool isESIMDKernelType(const CXXRecordDecl *KernelObjType) {
-  const CXXMethodDecl *OpParens = getOperatorParens(KernelObjType);
+  const CXXMethodDecl *OpParens = getFunctorCallOperator(KernelObjType);
   return (OpParens != nullptr) && OpParens->hasAttr<SYCLSimdAttr>();
 }
 
@@ -4041,7 +4055,7 @@ void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc,
 // kernel to wrapped kernel.
 void Sema::copySYCLKernelAttrs(const CXXRecordDecl *KernelObj) {
   // Get the operator() function of the wrapper.
-  CXXMethodDecl *OpParens = getOperatorParens(KernelObj);
+  CXXMethodDecl *OpParens = getFunctorCallOperator(KernelObj);
   assert(OpParens && "invalid kernel object");
 
   typedef std::pair<FunctionDecl *, FunctionDecl *> ChildParentPair;

clang/test/SemaSYCL/undefined-functor.cpp

@@ -18,6 +18,12 @@ class FunctorWithCallOpDefined {
   void operator()() const {}
 };
 
+class FunctorWithCallOpTemplated {
+  public:
+  template <int x = 0>
+  void operator()() const {}
+};
+
 int main() {
 
   q.submit([&](sycl::handler &cgh) {
@@ -37,4 +43,8 @@ int main() {
     cgh.single_task(FunctorWithCallOpDefined{});
   });
 
+  q.submit([&](sycl::handler &cgh) {
+    cgh.single_task(FunctorWithCallOpTemplated{});
+  });
+
 }