[SYCLomatic] Emit a warning when the virtual function are called in k…

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 <algorithm>
 #include <iostream>
@@ -9248,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<FunctionDecl *>(FD));
+  bool FDIsRecursive = false;
+  for (llvm::scc_iterator<CallGraph *> 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 =
@@ -9326,6 +9350,16 @@ 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");
+  }
+
+  if(isRecursiveDeviceFuncDecl(FD))
+    report(FD->getBeginLoc(), Warnings::DEVICE_UNSUPPORTED_CALL_FUNCTION,
+                            false, "Recursive functions");
+
   FuncInfo = DeviceFunctionDecl::LinkRedecls(FD);
   if (!FuncInfo)
     return;
@@ -9343,6 +9377,15 @@ void DeviceFunctionDeclRule::runRule(
   }
 
   if (auto CE = getAssistNodeAsType<CallExpr>(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<CallExpr>(Result, "PrintfExpr")) {
     if (FD->hasAttr<CUDAHostAttr>()) {

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_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.")

clang/test/dpct/device_virtual_function.cu

@@ -0,0 +1,63 @@
+// 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 <cuda_runtime.h>
+template <typename T>
+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 <typename T>
+class TestSeqContainer : public TestVirtual<T> {
+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<int>(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;
+}
+
+template <class T>
+class Container {
+public:
+    __device__ Container() {}
+    __device__ ~Container() {}
+};
+
+int main() {
+func<<<1,1>>>();
+cudaDeviceSynchronize();
+return 0;
+
+}

clang/test/dpct/recursive_function.cu

@@ -0,0 +1,33 @@
+// 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 <cuda.h>
+// 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: */
+__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);
+    }
+}
+
+__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>>>();
+}