Skip to content

Commit 29fa487

Browse files
committed
[SYCL] Fix for detection of free function calls.
Signed-off-by: rdeodhar <[email protected]>
1 parent 237675b commit 29fa487

File tree

2 files changed

+22
-9
lines changed

2 files changed

+22
-9
lines changed

clang/lib/Sema/SemaSYCL.cpp

Lines changed: 3 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -2705,23 +2705,18 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
27052705

27062706
// Sets a flag if the kernel is a parallel_for that calls the
27072707
// free function API "this_item".
2708-
void setThisItemIsCalled(const CXXRecordDecl *KernelObj,
2709-
FunctionDecl *KernelFunc) {
2708+
void setThisItemIsCalled(FunctionDecl *KernelFunc) {
27102709
if (getKernelInvocationKind(KernelFunc) != InvokeParallelFor)
27112710
return;
27122711

2713-
const CXXMethodDecl *WGLambdaFn = getOperatorParens(KernelObj);
2714-
if (!WGLambdaFn)
2715-
return;
2716-
27172712
// The call graph for this translation unit.
27182713
CallGraph SYCLCG;
27192714
SYCLCG.addToCallGraph(SemaRef.getASTContext().getTranslationUnitDecl());
27202715
using ChildParentPair =
27212716
std::pair<const FunctionDecl *, const FunctionDecl *>;
27222717
llvm::SmallPtrSet<const FunctionDecl *, 16> Visited;
27232718
llvm::SmallVector<ChildParentPair, 16> WorkList;
2724-
WorkList.push_back({WGLambdaFn, nullptr});
2719+
WorkList.push_back({KernelFunc, nullptr});
27252720

27262721
while (!WorkList.empty()) {
27272722
const FunctionDecl *FD = WorkList.back().first;
@@ -2759,7 +2754,7 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
27592754
bool IsSIMDKernel = isESIMDKernelType(KernelObj);
27602755
Header.startKernel(Name, NameType, StableName, KernelObj->getLocation(),
27612756
IsSIMDKernel);
2762-
setThisItemIsCalled(KernelObj, KernelFunc);
2757+
setThisItemIsCalled(KernelFunc);
27632758
}
27642759

27652760
bool handleSyclAccessorType(const CXXRecordDecl *RD,

clang/test/CodeGenSYCL/parallel_for_this_item.cpp

Lines changed: 19 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -13,7 +13,8 @@
1313
// CHECK-NEXT: "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3GNU",
1414
// CHECK-NEXT: "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3EMU",
1515
// CHECK-NEXT: "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3OWL",
16-
// CHECK-NEXT: "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3RAT"
16+
// CHECK-NEXT: "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3RAT",
17+
// CHECK-NEXT: "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3BEE"
1718
// CHECK-NEXT: };
1819

1920
// CHECK:template <> struct KernelInfo<class GNU> {
@@ -72,6 +73,20 @@
7273
// CHECK-NEXT: __SYCL_DLL_LOCAL
7374
// CHECK-NEXT: static constexpr bool callsThisItem() { return 1; }
7475
// CHECK-NEXT:};
76+
// CHECK-NEXT:template <> struct KernelInfo<class BEE> {
77+
// CHECK-NEXT: __SYCL_DLL_LOCAL
78+
// CHECK-NEXT: static constexpr const char* getName() { return "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3BEE"; }
79+
// CHECK-NEXT: __SYCL_DLL_LOCAL
80+
// CHECK-NEXT: static constexpr unsigned getNumParams() { return 0; }
81+
// CHECK-NEXT: __SYCL_DLL_LOCAL
82+
// CHECK-NEXT: static constexpr const kernel_param_desc_t& getParamDesc(unsigned i) {
83+
// CHECK-NEXT: return kernel_signatures[i+0];
84+
// CHECK-NEXT: }
85+
// CHECK-NEXT: __SYCL_DLL_LOCAL
86+
// CHECK-NEXT: static constexpr bool isESIMD() { return 0; }
87+
// CHECK-NEXT: __SYCL_DLL_LOCAL
88+
// CHECK-NEXT: static constexpr bool callsThisItem() { return 1; }
89+
// CHECK-NEXT:};
7590

7691
#include "sycl.hpp"
7792

@@ -108,6 +123,9 @@ int main() {
108123

109124
// This kernel calls sycl::this_item
110125
cgh.parallel_for<class RAT>(range<1>(1), [=](id<1> I) { f(); });
126+
127+
// This kernel calls sycl::this_item
128+
cgh.parallel_for<class BEE>(range<1>(1), [=](auto I) { this_item<1>(); });
111129
});
112130

113131
return 0;

0 commit comments

Comments
 (0)