diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 907af88b74a58..1dde494bc1bad 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -2705,15 +2705,10 @@ 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()); @@ -2721,7 +2716,7 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { std::pair; llvm::SmallPtrSet Visited; llvm::SmallVector 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, diff --git a/clang/test/CodeGenSYCL/parallel_for_this_item.cpp b/clang/test/CodeGenSYCL/parallel_for_this_item.cpp index dc484b33bfbb4..7c9bee0c4a69d 100755 --- a/clang/test/CodeGenSYCL/parallel_for_this_item.cpp +++ b/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 { @@ -97,6 +98,22 @@ // CHECK-NEXT: __SYCL_DLL_LOCAL // CHECK-NEXT: static constexpr bool callsAnyThisFreeFunction() { return 1; } // CHECK-NEXT:}; +// CHECK-NEXT:template <> struct KernelInfo { +// 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(range<1>(1), [=](id<1> I) { f(); }); // This kernel does not call sycl::this_item, but does call this_id - cgh.parallel_for(range<1>(1), - [=](id<1> I) { this_id<1>(); }); + cgh.parallel_for(range<1>(1), [=](id<1> I) { this_id<1>(); }); + + // This kernel calls sycl::this_item + cgh.parallel_for(range<1>(1), [=](auto I) { this_item<1>(); }); }); return 0;