From c3827a6494749ae407f3fb29c680b8dd469c16b8 Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Wed, 22 Mar 2023 15:57:05 +0800 Subject: [PATCH 1/6] [SYCLomatic] Emit a warning when the virtual function are called in kernel function. --- clang/lib/DPCT/ASTTraversal.cpp | 3 ++ clang/lib/DPCT/Diagnostics.inc | 2 + clang/test/dpct/device_virtual_function.cu | 60 ++++++++++++++++++++++ 3 files changed, 65 insertions(+) create mode 100644 clang/test/dpct/device_virtual_function.cu diff --git a/clang/lib/DPCT/ASTTraversal.cpp b/clang/lib/DPCT/ASTTraversal.cpp index f4ae5f8a7ead..9c2b8e989491 100644 --- a/clang/lib/DPCT/ASTTraversal.cpp +++ b/clang/lib/DPCT/ASTTraversal.cpp @@ -9326,6 +9326,9 @@ void DeviceFunctionDeclRule::runRule( if (FD->isVariadic()) { report(FD->getBeginLoc(), Warnings::DEVICE_VARIADIC_FUNCTION, false); } + if (FD->isVirtualAsWritten()) { + report(FD->getBeginLoc(), Warnings::DEVICE_VIRTUAL_FUNCTION, false); + } FuncInfo = DeviceFunctionDecl::LinkRedecls(FD); if (!FuncInfo) return; diff --git a/clang/lib/DPCT/Diagnostics.inc b/clang/lib/DPCT/Diagnostics.inc index 2bc706011460..b2d65f9d15c5 100644 --- a/clang/lib/DPCT/Diagnostics.inc +++ b/clang/lib/DPCT/Diagnostics.inc @@ -231,3 +231,5 @@ DEF_WARNING(OVERLOAD_UNSUPPORTED, 1107, "Migration for this overload of %0 is no DEF_COMMENT(OVERLOAD_UNSUPPORTED, 1107, "Migration for this overload of {0} is not supported.") DEF_WARNING(EXPERIMENTAL_FEATURE, 1108, "'%0' was migrated with the experimental feature %0 which may not be supported by all compilers or runtimes. You may need to adjust the code.") DEF_COMMENT(EXPERIMENTAL_FEATURE, 1108, "'{0}' was migrated with the experimental feature {0} which may not be supported by all compilers or runtimes. You may need to adjust the code.") +DEF_WARNING(DEVICE_VIRTUAL_FUNCTION, 1109, "Virtual functions cannot be called in a SYCL kernel or by functions called by the kernel. You may need to adjust the code.") +DEF_COMMENT(DEVICE_VIRTUAL_FUNCTION, 1109, "Virtual functions cannot be called in a SYCL kernel or by functions called by the kernel. You may need to adjust the code.") \ No newline at end of file diff --git a/clang/test/dpct/device_virtual_function.cu b/clang/test/dpct/device_virtual_function.cu new file mode 100644 index 000000000000..1af943d1d363 --- /dev/null +++ b/clang/test/dpct/device_virtual_function.cu @@ -0,0 +1,60 @@ +// RUN: dpct --format-range=none -out-root %T/device_virtual_function %s --cuda-include-path="%cuda-path/include" -- -std=c++14 -x cuda --cuda-host-only +// RUN: FileCheck %s --match-full-lines --input-file %T/device_virtual_function/device_virtual_function.dp.cpp + +#include +template +class TestVirtual { +public: + + __device__ TestVirtual() {} +// CHECK: /* +// CHECK-NEXT: DPCT1109:{{[0-9]+}}: Virtual functions cannot be called in a SYCL kernel or by functions called by the kernel. You may need to adjust the code. +// CHECK-NEXT: */ + __device__ virtual ~TestVirtual() {} +// CHECK: /* +// CHECK-NEXT: DPCT1109:{{[0-9]+}}: Virtual functions cannot be called in a SYCL kernel or by functions called by the kernel. You may need to adjust the code. +// CHECK-NEXT: */ + __device__ virtual void push(const T &&e)= 0; +}; +template +class TestSeqContainer : public TestVirtual { +public: + __device__ TestSeqContainer(int size) : index_top(-1) { m_data = new T[size]; } + + __device__ ~TestSeqContainer() { + if (m_data) delete []m_data; + } + // CHECK: /* + // CHECK-NEXT: DPCT1109:{{[0-9]+}}: Virtual functions cannot be called in a SYCL kernel or by functions called by the kernel. You may need to adjust the code. + // CHECK-NEXT: */ + __device__ virtual void push(const T &&e) { + if (m_data) { + int idx = atomicAdd(&this->index_top, 1); + m_data[idx] = e; + } + } +private: + T *m_data; + int index_top; + +}; +__global__ void func(){ + + auto seq = new TestSeqContainer(10); + seq->push(10); + delete seq; +} + +template +class Container { +public: + __device__ Container() {} + __device__ ~Container() {} +}; + +int main() { +func<<<1,1>>>(); +cudaDeviceSynchronize(); +return 0; + +} \ No newline at end of file From 3cf2cbd1bbad8cb755280a21270653b12cbb4f9d Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Thu, 23 Mar 2023 14:49:06 +0800 Subject: [PATCH 2/6] Emit a recursive warning. --- clang/lib/DPCT/ASTTraversal.cpp | 18 +++++++++++++++++ clang/lib/DPCT/Diagnostics.inc | 4 +++- clang/test/dpct/recursive_function.cu | 28 +++++++++++++++++++++++++++ 3 files changed, 49 insertions(+), 1 deletion(-) create mode 100644 clang/test/dpct/recursive_function.cu diff --git a/clang/lib/DPCT/ASTTraversal.cpp b/clang/lib/DPCT/ASTTraversal.cpp index 9c2b8e989491..d84e19249fd5 100644 --- a/clang/lib/DPCT/ASTTraversal.cpp +++ b/clang/lib/DPCT/ASTTraversal.cpp @@ -34,6 +34,8 @@ #include "llvm/Support/Casting.h" #include "llvm/Support/Path.h" #include "MemberExprRewriter.h" +#include "clang/Analysis/CallGraph.h" +#include "llvm/ADT/SCCIterator.h" #include #include @@ -9329,6 +9331,22 @@ void DeviceFunctionDeclRule::runRule( if (FD->isVirtualAsWritten()) { report(FD->getBeginLoc(), Warnings::DEVICE_VIRTUAL_FUNCTION, false); } + + // Build call graph for FunctionDecl and look for cycles in call graph. + // Emit the warning message when the recursive call exists in kernel function. + CallGraph CG; + CG.addToCallGraph(const_cast(FD)); + for (llvm::scc_iterator SCCI = llvm::scc_begin(&CG), + SCCE = llvm::scc_end(&CG); + SCCI != SCCE; ++SCCI) { + if (SCCI.hasCycle()) { + for (auto node : *SCCI) { + FunctionDecl *RecFD = node->getDefinition(); + report(RecFD->getBeginLoc(), Warnings::DEVICE_RECURSIVE_FUNCTION, false); + } + } + } + FuncInfo = DeviceFunctionDecl::LinkRedecls(FD); if (!FuncInfo) return; diff --git a/clang/lib/DPCT/Diagnostics.inc b/clang/lib/DPCT/Diagnostics.inc index b2d65f9d15c5..67f805e05687 100644 --- a/clang/lib/DPCT/Diagnostics.inc +++ b/clang/lib/DPCT/Diagnostics.inc @@ -232,4 +232,6 @@ DEF_COMMENT(OVERLOAD_UNSUPPORTED, 1107, "Migration for this overload of {0} is n DEF_WARNING(EXPERIMENTAL_FEATURE, 1108, "'%0' was migrated with the experimental feature %0 which may not be supported by all compilers or runtimes. You may need to adjust the code.") DEF_COMMENT(EXPERIMENTAL_FEATURE, 1108, "'{0}' was migrated with the experimental feature {0} which may not be supported by all compilers or runtimes. You may need to adjust the code.") DEF_WARNING(DEVICE_VIRTUAL_FUNCTION, 1109, "Virtual functions cannot be called in a SYCL kernel or by functions called by the kernel. You may need to adjust the code.") -DEF_COMMENT(DEVICE_VIRTUAL_FUNCTION, 1109, "Virtual functions cannot be called in a SYCL kernel or by functions called by the kernel. You may need to adjust the code.") \ No newline at end of file +DEF_COMMENT(DEVICE_VIRTUAL_FUNCTION, 1109, "Virtual functions cannot be called in a SYCL kernel or by functions called by the kernel. You may need to adjust the code.") +DEF_WARNING(DEVICE_RECURSIVE_FUNCTION, 1110, "Recursive functions cannot be called in a SYCL kernel or by functions called by the kernel. You may need to adjust the code.") +DEF_COMMENT(DEVICE_RECURSIVE_FUNCTION, 1110, "Recursive functions cannot be called in a SYCL kernel or by functions called by the kernel. You may need to adjust the code.") diff --git a/clang/test/dpct/recursive_function.cu b/clang/test/dpct/recursive_function.cu new file mode 100644 index 000000000000..47cf772f292a --- /dev/null +++ b/clang/test/dpct/recursive_function.cu @@ -0,0 +1,28 @@ +#include +// CHECK: /* +// CHECK-NEXT: DPCT1110:{{[0-9]+}}: Recursive functions cannot be called in a SYCL kernel or by functions called by the kernel. You may need to adjust the code. +// CHECK-NEXT: */ +__device__ int factorial(int n) { + if (n <= 1) { + return 1; + } else { + return n * factorial(n - 1); + } +} + +__global__ void test_kernel() { + factorial(10); +} + + +int factorial2(int n) { + if (n <= 1) { + return 1; + } else { + return n * factorial2(n - 1); + } +} + +int main() { + test_kernel<<<1,1>>>(); +} \ No newline at end of file From d0a9f993f392d9beb67131ef2decaab24afbcae5 Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Thu, 23 Mar 2023 14:55:20 +0800 Subject: [PATCH 3/6] update case. Signed-off-by: Chen, Sheng S --- clang/test/dpct/recursive_function.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/test/dpct/recursive_function.cu b/clang/test/dpct/recursive_function.cu index 47cf772f292a..a39f64e64447 100644 --- a/clang/test/dpct/recursive_function.cu +++ b/clang/test/dpct/recursive_function.cu @@ -1,4 +1,4 @@ -#include +#include // CHECK: /* // CHECK-NEXT: DPCT1110:{{[0-9]+}}: Recursive functions cannot be called in a SYCL kernel or by functions called by the kernel. You may need to adjust the code. // CHECK-NEXT: */ From f63464132511373bf0f0357552c2daf3ab7d86d3 Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Thu, 23 Mar 2023 16:24:50 +0800 Subject: [PATCH 4/6] Update test case. --- clang/test/dpct/recursive_function.cu | 2 ++ 1 file changed, 2 insertions(+) diff --git a/clang/test/dpct/recursive_function.cu b/clang/test/dpct/recursive_function.cu index a39f64e64447..1480237edead 100644 --- a/clang/test/dpct/recursive_function.cu +++ b/clang/test/dpct/recursive_function.cu @@ -1,3 +1,5 @@ +// RUN: dpct --format-range=none -out-root %T/recursive_function %s --cuda-include-path="%cuda-path/include" -- -std=c++14 -x cuda --cuda-host-only +// RUN: FileCheck %s --match-full-lines --input-file %T/recursive_function/recursive_function.dp.cpp #include // CHECK: /* // CHECK-NEXT: DPCT1110:{{[0-9]+}}: Recursive functions cannot be called in a SYCL kernel or by functions called by the kernel. You may need to adjust the code. From 1df18a631953525efdc3c7d01369857564ae9bd1 Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Fri, 24 Mar 2023 10:20:55 +0800 Subject: [PATCH 5/6] Update the test case to merge 2 cases. Signed-off-by: Chen, Sheng S --- clang/lib/DPCT/ASTTraversal.cpp | 6 ++++-- clang/lib/DPCT/Diagnostics.inc | 6 ++---- clang/test/dpct/recursive_function.cu | 2 +- 3 files changed, 7 insertions(+), 7 deletions(-) diff --git a/clang/lib/DPCT/ASTTraversal.cpp b/clang/lib/DPCT/ASTTraversal.cpp index 30b5c2881665..079bd59d47a4 100644 --- a/clang/lib/DPCT/ASTTraversal.cpp +++ b/clang/lib/DPCT/ASTTraversal.cpp @@ -9329,7 +9329,8 @@ void DeviceFunctionDeclRule::runRule( report(FD->getBeginLoc(), Warnings::DEVICE_VARIADIC_FUNCTION, false); } if (FD->isVirtualAsWritten()) { - report(FD->getBeginLoc(), Warnings::DEVICE_VIRTUAL_FUNCTION, false); + report(FD->getBeginLoc(), Warnings::DEVICE_UNSUPPORTED_CALL_FUNCTION, + false, "Virtual functions"); } // Build call graph for FunctionDecl and look for cycles in call graph. @@ -9342,7 +9343,8 @@ void DeviceFunctionDeclRule::runRule( if (SCCI.hasCycle()) { for (auto node : *SCCI) { FunctionDecl *RecFD = node->getDefinition(); - report(RecFD->getBeginLoc(), Warnings::DEVICE_RECURSIVE_FUNCTION, false); + report(RecFD->getBeginLoc(), Warnings::DEVICE_UNSUPPORTED_CALL_FUNCTION, + false, "Recursive functions"); } } } diff --git a/clang/lib/DPCT/Diagnostics.inc b/clang/lib/DPCT/Diagnostics.inc index 67f805e05687..d47aa273a8b0 100644 --- a/clang/lib/DPCT/Diagnostics.inc +++ b/clang/lib/DPCT/Diagnostics.inc @@ -231,7 +231,5 @@ DEF_WARNING(OVERLOAD_UNSUPPORTED, 1107, "Migration for this overload of %0 is no DEF_COMMENT(OVERLOAD_UNSUPPORTED, 1107, "Migration for this overload of {0} is not supported.") DEF_WARNING(EXPERIMENTAL_FEATURE, 1108, "'%0' was migrated with the experimental feature %0 which may not be supported by all compilers or runtimes. You may need to adjust the code.") DEF_COMMENT(EXPERIMENTAL_FEATURE, 1108, "'{0}' was migrated with the experimental feature {0} which may not be supported by all compilers or runtimes. You may need to adjust the code.") -DEF_WARNING(DEVICE_VIRTUAL_FUNCTION, 1109, "Virtual functions cannot be called in a SYCL kernel or by functions called by the kernel. You may need to adjust the code.") -DEF_COMMENT(DEVICE_VIRTUAL_FUNCTION, 1109, "Virtual functions cannot be called in a SYCL kernel or by functions called by the kernel. You may need to adjust the code.") -DEF_WARNING(DEVICE_RECURSIVE_FUNCTION, 1110, "Recursive functions cannot be called in a SYCL kernel or by functions called by the kernel. You may need to adjust the code.") -DEF_COMMENT(DEVICE_RECURSIVE_FUNCTION, 1110, "Recursive functions cannot be called in a SYCL kernel or by functions called by the kernel. You may need to adjust the code.") +DEF_WARNING(DEVICE_UNSUPPORTED_CALL_FUNCTION, 1109, "%0 cannot be called in a SYCL kernel or by functions called by the kernel. You may need to adjust the code.") +DEF_COMMENT(DEVICE_UNSUPPORTED_CALL_FUNCTION, 1109, "{0} cannot be called in a SYCL kernel or by functions called by the kernel. You may need to adjust the code.") diff --git a/clang/test/dpct/recursive_function.cu b/clang/test/dpct/recursive_function.cu index 1480237edead..be5e866955ac 100644 --- a/clang/test/dpct/recursive_function.cu +++ b/clang/test/dpct/recursive_function.cu @@ -2,7 +2,7 @@ // RUN: FileCheck %s --match-full-lines --input-file %T/recursive_function/recursive_function.dp.cpp #include // CHECK: /* -// CHECK-NEXT: DPCT1110:{{[0-9]+}}: Recursive functions cannot be called in a SYCL kernel or by functions called by the kernel. You may need to adjust the code. +// CHECK-NEXT: DPCT1109:{{[0-9]+}}: Recursive functions cannot be called in a SYCL kernel or by functions called by the kernel. You may need to adjust the code. // CHECK-NEXT: */ __device__ int factorial(int n) { if (n <= 1) { From dee547e742718b194851acd37db3fbfb7aabb0b5 Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Tue, 28 Mar 2023 15:57:15 +0800 Subject: [PATCH 6/6] Emit Warning message at callexpr. --- clang/lib/DPCT/ASTTraversal.cpp | 50 +++++++++++++++------- clang/test/dpct/device_virtual_function.cu | 3 ++ clang/test/dpct/recursive_function.cu | 3 ++ 3 files changed, 41 insertions(+), 15 deletions(-) diff --git a/clang/lib/DPCT/ASTTraversal.cpp b/clang/lib/DPCT/ASTTraversal.cpp index 079bd59d47a4..5cd27ebd350c 100644 --- a/clang/lib/DPCT/ASTTraversal.cpp +++ b/clang/lib/DPCT/ASTTraversal.cpp @@ -9250,6 +9250,28 @@ void KernelCallRule::removeTrailingSemicolon( REGISTER_RULE(KernelCallRule, PassKind::PK_Analysis) + +bool isRecursiveDeviceFuncDecl(const FunctionDecl* FD) { + // Build call graph for FunctionDecl and look for cycles in call graph. + // Emit the warning message when the recursive call exists in kernel function. + if (!FD) return false; + CallGraph CG; + CG.addToCallGraph(const_cast(FD)); + bool FDIsRecursive = false; + for (llvm::scc_iterator SCCI = llvm::scc_begin(&CG), + SCCE = llvm::scc_end(&CG); + SCCI != SCCE; ++SCCI) { + if (SCCI.hasCycle()) FDIsRecursive = true; + } + return FDIsRecursive; +} + +bool isRecursiveDeviceCallExpr(const CallExpr* CE) { + if (isRecursiveDeviceFuncDecl(CE->getDirectCallee())) + return true; + return false; +} + // __device__ function call information collection void DeviceFunctionDeclRule::registerMatcher(ast_matchers::MatchFinder &MF) { auto DeviceFunctionMatcher = @@ -9328,26 +9350,15 @@ void DeviceFunctionDeclRule::runRule( if (FD->isVariadic()) { report(FD->getBeginLoc(), Warnings::DEVICE_VARIADIC_FUNCTION, false); } + if (FD->isVirtualAsWritten()) { report(FD->getBeginLoc(), Warnings::DEVICE_UNSUPPORTED_CALL_FUNCTION, false, "Virtual functions"); } - // Build call graph for FunctionDecl and look for cycles in call graph. - // Emit the warning message when the recursive call exists in kernel function. - CallGraph CG; - CG.addToCallGraph(const_cast(FD)); - for (llvm::scc_iterator SCCI = llvm::scc_begin(&CG), - SCCE = llvm::scc_end(&CG); - SCCI != SCCE; ++SCCI) { - if (SCCI.hasCycle()) { - for (auto node : *SCCI) { - FunctionDecl *RecFD = node->getDefinition(); - report(RecFD->getBeginLoc(), Warnings::DEVICE_UNSUPPORTED_CALL_FUNCTION, - false, "Recursive functions"); - } - } - } + if(isRecursiveDeviceFuncDecl(FD)) + report(FD->getBeginLoc(), Warnings::DEVICE_UNSUPPORTED_CALL_FUNCTION, + false, "Recursive functions"); FuncInfo = DeviceFunctionDecl::LinkRedecls(FD); if (!FuncInfo) @@ -9366,6 +9377,15 @@ void DeviceFunctionDeclRule::runRule( } if (auto CE = getAssistNodeAsType(Result, "callExpr")) { + if (CE->getDirectCallee()) { + if (CE->getDirectCallee()->isVirtualAsWritten()) + report(CE->getBeginLoc(), Warnings::DEVICE_UNSUPPORTED_CALL_FUNCTION, + false, "Virtual functions"); + } + + if (isRecursiveDeviceCallExpr(CE)) + report(CE->getBeginLoc(), Warnings::DEVICE_UNSUPPORTED_CALL_FUNCTION, + false, "Recursive functions"); FuncInfo->addCallee(CE); } else if (CE = getAssistNodeAsType(Result, "PrintfExpr")) { if (FD->hasAttr()) { diff --git a/clang/test/dpct/device_virtual_function.cu b/clang/test/dpct/device_virtual_function.cu index 1af943d1d363..86fc060b7758 100644 --- a/clang/test/dpct/device_virtual_function.cu +++ b/clang/test/dpct/device_virtual_function.cu @@ -41,6 +41,9 @@ private: __global__ void func(){ auto seq = new TestSeqContainer(10); + // CHECK: /* + // CHECK-NEXT: DPCT1109:{{[0-9]+}}: Virtual functions cannot be called in a SYCL kernel or by functions called by the kernel. You may need to adjust the code. + // CHECK-NEXT: */ seq->push(10); delete seq; } diff --git a/clang/test/dpct/recursive_function.cu b/clang/test/dpct/recursive_function.cu index be5e866955ac..ec84c68fdae8 100644 --- a/clang/test/dpct/recursive_function.cu +++ b/clang/test/dpct/recursive_function.cu @@ -8,6 +8,9 @@ __device__ int factorial(int n) { if (n <= 1) { return 1; } else { + // CHECK: /* + // CHECK-NEXT: DPCT1109:{{[0-9]+}}: Recursive functions cannot be called in a SYCL kernel or by functions called by the kernel. You may need to adjust the code. + // CHECK-NEXT: */ return n * factorial(n - 1); } }