[SYCL] Fix detection of free function calls (#3003)

This change simplifies the code to detect calls to "free functions", i.e., functions such as this_item().
The existing code does not work in all cases to detect the Kernel function,
while the corrected code uses an already known Kernel function.

Signed-off-by: rdeodhar <[email protected]>

Author: rdeodhar

clang/lib/Sema/SemaSYCL.cpp

@@ -2705,23 +2705,18 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
 
   // Sets a flag if the kernel is a parallel_for that calls the
   // free function API "this_item".
-  void setThisItemIsCalled(const CXXRecordDecl *KernelObj,
-                           FunctionDecl *KernelFunc) {
+  void setThisItemIsCalled(FunctionDecl *KernelFunc) {
     if (getKernelInvocationKind(KernelFunc) != InvokeParallelFor)
       return;
 
-    const CXXMethodDecl *WGLambdaFn = getOperatorParens(KernelObj);
-    if (!WGLambdaFn)
-      return;
-
     // The call graph for this translation unit.
     CallGraph SYCLCG;
     SYCLCG.addToCallGraph(SemaRef.getASTContext().getTranslationUnitDecl());
     using ChildParentPair =
         std::pair<const FunctionDecl *, const FunctionDecl *>;
     llvm::SmallPtrSet<const FunctionDecl *, 16> Visited;
     llvm::SmallVector<ChildParentPair, 16> WorkList;
-    WorkList.push_back({WGLambdaFn, nullptr});
+    WorkList.push_back({KernelFunc, nullptr});
 
     while (!WorkList.empty()) {
       const FunctionDecl *FD = WorkList.back().first;
@@ -2772,7 +2767,7 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
     bool IsSIMDKernel = isESIMDKernelType(KernelObj);
     Header.startKernel(Name, NameType, StableName, KernelObj->getLocation(),
                        IsSIMDKernel);
-    setThisItemIsCalled(KernelObj, KernelFunc);
+    setThisItemIsCalled(KernelFunc);
   }
 
   bool handleSyclAccessorType(const CXXRecordDecl *RD,

clang/test/CodeGenSYCL/parallel_for_this_item.cpp

@@ -14,7 +14,8 @@
 // CHECK-NEXT:   "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3EMU",
 // CHECK-NEXT:   "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3OWL",
 // CHECK-NEXT:   "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3RAT",
-// CHECK-NEXT:   "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3FOX"
+// CHECK-NEXT:   "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3FOX",
+// CHECK-NEXT:   "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3BEE"
 // CHECK-NEXT: };
 
 // CHECK:template <> struct KernelInfo<class GNU> {
@@ -97,6 +98,22 @@
 // CHECK-NEXT: __SYCL_DLL_LOCAL
 // CHECK-NEXT:    static constexpr bool callsAnyThisFreeFunction() { return 1; }
 // CHECK-NEXT:};
+// CHECK-NEXT:template <> struct KernelInfo<class BEE> {
+// CHECK-NEXT:  __SYCL_DLL_LOCAL
+// CHECK-NEXT:    static constexpr const char* getName() { return "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3BEE"; }
+// CHECK-NEXT:  __SYCL_DLL_LOCAL
+// CHECK-NEXT:    static constexpr unsigned getNumParams() { return 0; }
+// CHECK-NEXT:  __SYCL_DLL_LOCAL
+// CHECK-NEXT:    static constexpr const kernel_param_desc_t& getParamDesc(unsigned i) {
+// CHECK-NEXT:    return kernel_signatures[i+0];
+// CHECK-NEXT:  }
+// CHECK-NEXT:  __SYCL_DLL_LOCAL
+// CHECK-NEXT:    static constexpr bool isESIMD() { return 0; }
+// CHECK-NEXT:  __SYCL_DLL_LOCAL
+// CHECK-NEXT:    static constexpr bool callsThisItem() { return 1; }
+// CHECK-NEXT: __SYCL_DLL_LOCAL
+// CHECK-NEXT:    static constexpr bool callsAnyThisFreeFunction() { return 1; }
+// CHECK-NEXT:};
 
 #include "sycl.hpp"
 
@@ -135,8 +152,10 @@ int main() {
     cgh.parallel_for<class RAT>(range<1>(1), [=](id<1> I) { f(); });
 
     // This kernel does not call sycl::this_item, but does call this_id
-    cgh.parallel_for<class FOX>(range<1>(1),
-                                [=](id<1> I) { this_id<1>(); });
+    cgh.parallel_for<class FOX>(range<1>(1), [=](id<1> I) { this_id<1>(); });
+
+    // This kernel calls sycl::this_item
+    cgh.parallel_for<class BEE>(range<1>(1), [=](auto I) { this_item<1>(); });
   });
 
   return 0;