[SYCL] Fix pathological case of visiting callees of a function. (#4065)

The markdevice rewrite improved the way we were checking recursive
functions, however as an oversight didn't 'uniqify' each callee-check.

This patch ensures we only visit each callee 1x, even if it is called
multiple times.

Note that this isn't a 'perfect' fix, we could skip any function we've
ever 'seen' before in this kernel, however it results in some reduced
diagnostic quality for recursive and attribute-collection issues.

This at least reduces the 'pathological' cases that remain to just those
that are also mostly pathological for templates in general (though we
are still worse-off than template instantiations).

clang/lib/Sema/SemaSYCL.cpp

@@ -764,10 +764,8 @@ class SingleDeviceFunctionTracker {
     return FD->getMostRecentDecl();
   }
 
-  void VisitCallNode(CallGraphNode *Node,
+  void VisitCallNode(CallGraphNode *Node, FunctionDecl *CurrentDecl,
                      llvm::SmallVectorImpl<FunctionDecl *> &CallStack) {
-    FunctionDecl *CurrentDecl = GetFDFromNode(Node);
-
     // If this isn't a function, I don't think there is anything we can do here.
     if (!CurrentDecl)
       return;
@@ -842,8 +840,16 @@ class SingleDeviceFunctionTracker {
 
     // Recurse.
     CallStack.push_back(CurrentDecl);
+    llvm::SmallPtrSet<FunctionDecl *, 16> SeenCallees;
     for (CallGraphNode *CI : Node->callees()) {
-      VisitCallNode(CI, CallStack);
+      FunctionDecl *CurFD = GetFDFromNode(CI);
+
+      // Make sure we only visit each callee 1x from this function to avoid very
+      // time consuming template recursion cases.
+      if (!llvm::is_contained(SeenCallees, CurFD)) {
+        VisitCallNode(CI, CurFD, CallStack);
+        SeenCallees.insert(CurFD);
+      }
     }
     CallStack.pop_back();
   }
@@ -852,7 +858,7 @@ class SingleDeviceFunctionTracker {
   void Init() {
     CallGraphNode *KernelNode = Parent.getNodeForKernel(SYCLKernel);
     llvm::SmallVector<FunctionDecl *> CallStack;
-    VisitCallNode(KernelNode, CallStack);
+    VisitCallNode(KernelNode, GetFDFromNode(KernelNode), CallStack);
   }
 
 public:

clang/test/SemaSYCL/mark-device-finishes.cpp

@@ -0,0 +1,49 @@
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -verify -fsyntax-only -std=c++20 %s
+
+// This test validates that this actually makes it through 'MarkDevice'.  This
+// is a bit of a pathological case where we ended up visiting each call
+// individually. There is likely a similar test case that can cause us to hit
+// a pathological case in a very similar situation (where the callees aren't
+// exactly the same), but that likely causes problems with template
+// instantiations first.
+
+// expected-no-diagnostics
+
+#include "sycl.hpp"
+
+template<bool B, typename V = void>
+struct enable_if { };
+template<typename V>
+struct enable_if<true, V> {
+  using type = V;
+};
+template<bool B, typename V = void>
+using enable_if_t = typename enable_if<B, V>::type;
+
+
+template<int N, enable_if_t<N == 24, int> = 0>
+void mark_device_pathological_case() {
+  // Do nothing.
+}
+
+template<int N, enable_if_t<N < 24, int> = 0>
+void mark_device_pathological_case() {
+  // We were visiting each of these, which caused 9^24 visits.
+  mark_device_pathological_case<N + 1>();
+  mark_device_pathological_case<N + 1>();
+  mark_device_pathological_case<N + 1>();
+  mark_device_pathological_case<N + 1>();
+  mark_device_pathological_case<N + 1>();
+  mark_device_pathological_case<N + 1>();
+  mark_device_pathological_case<N + 1>();
+  mark_device_pathological_case<N + 1>();
+  mark_device_pathological_case<N + 1>();
+  mark_device_pathological_case<N + 1>();
+}
+
+int main() {
+  sycl::queue q;
+  q.submit([](sycl::handler &h) {
+    h.single_task<class kernel>([]() { mark_device_pathological_case<0>(); });
+  });
+}