From 0383adbf49f015a7082055ef38c9cc498af070dd Mon Sep 17 00:00:00 2001 From: rdeodhar Date: Fri, 8 Jan 2021 12:59:57 -0800 Subject: [PATCH 1/3] [SYCL] Fix for detection of free function calls. Signed-off-by: rdeodhar --- clang/lib/Sema/SemaSYCL.cpp | 11 +++-------- 1 file changed, 3 insertions(+), 8 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index c3578ab240b8f..2296738b3d44c 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; @@ -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, From 29fa487be716db60010e0f011167e56fdd88e402 Mon Sep 17 00:00:00 2001 From: rdeodhar Date: Fri, 8 Jan 2021 12:59:57 -0800 Subject: [PATCH 2/3] [SYCL] Fix for detection of free function calls. Signed-off-by: rdeodhar --- clang/lib/Sema/SemaSYCL.cpp | 11 +++------- .../CodeGenSYCL/parallel_for_this_item.cpp | 20 ++++++++++++++++++- 2 files changed, 22 insertions(+), 9 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index c3578ab240b8f..2296738b3d44c 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; @@ -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 index 422a1bad33373..bf6a4abe7412e 100755 --- 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 { @@ -72,6 +73,20 @@ // CHECK-NEXT: __SYCL_DLL_LOCAL // CHECK-NEXT: static constexpr bool callsThisItem() { 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:}; #include "sycl.hpp" @@ -108,6 +123,9 @@ int main() { // This kernel calls sycl::this_item cgh.parallel_for(range<1>(1), [=](id<1> I) { f(); }); + + // This kernel calls sycl::this_item + cgh.parallel_for(range<1>(1), [=](auto I) { this_item<1>(); }); }); return 0; From fdf0675170792b4b5752a71b2ff134fca40fadf8 Mon Sep 17 00:00:00 2001 From: rdeodhar Date: Mon, 11 Jan 2021 14:10:09 -0800 Subject: [PATCH 3/3] Added test for this_item. --- clang/test/CodeGenSYCL/parallel_for_this_item.cpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/clang/test/CodeGenSYCL/parallel_for_this_item.cpp b/clang/test/CodeGenSYCL/parallel_for_this_item.cpp index 7541dded08bdd..7c9bee0c4a69d 100755 --- a/clang/test/CodeGenSYCL/parallel_for_this_item.cpp +++ b/clang/test/CodeGenSYCL/parallel_for_this_item.cpp @@ -111,6 +111,8 @@ // 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"