[SYCL] Fix for detection of free function calls.

Signed-off-by: rdeodhar <[email protected]>
intel · Jan 11, 2021 · 29fa487 · 29fa487
1 parent 237675b
commit 29fa487
Show file tree

Hide file tree

Showing 2 changed files with 22 additions and 9 deletions.
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/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;
@@ -2759,7 +2754,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,

diff --git a/clang/test/CodeGenSYCL/parallel_for_this_item.cpp b/clang/test/CodeGenSYCL/parallel_for_this_item.cpp
@@ -13,7 +13,8 @@
 // CHECK-NEXT:   "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3GNU",
 // CHECK-NEXT:   "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3EMU",
 // CHECK-NEXT:   "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3OWL",
-// CHECK-NEXT:   "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3RAT"
+// CHECK-NEXT:   "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3RAT",
+// CHECK-NEXT:   "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3BEE"
 // CHECK-NEXT: };
 
 // CHECK:template <> struct KernelInfo<class GNU> {
@@ -72,6 +73,20 @@
 // CHECK-NEXT: __SYCL_DLL_LOCAL
 // CHECK-NEXT:    static constexpr bool callsThisItem() { 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:};
 
 #include "sycl.hpp"
 
@@ -108,6 +123,9 @@ int main() {
 
     // This kernel calls sycl::this_item
     cgh.parallel_for<class RAT>(range<1>(1), [=](id<1> I) { f(); });
+
+    // This kernel calls sycl::this_item
+    cgh.parallel_for<class BEE>(range<1>(1), [=](auto I) { this_item<1>(); });
   });
 
   return 0;