fix: Try to emit refs for unresolved call expressions

indexer/ClangAstMacros.h

@@ -32,6 +32,7 @@
   F(CXXDependentScopeMember)           \
   F(DeclRef)                           \
   F(Member)                            \
+  F(UnresolvedLookup)                  \
   F(UnresolvedMember)
 
 #define FOR_EACH_TYPE_TO_BE_INDEXED(F) \

indexer/Indexer.cc

@@ -1022,6 +1022,32 @@ void TuIndexer::saveMemberExpr(const clang::MemberExpr &memberExpr) {
   this->saveReference(optSymbol.value(), memberExpr.getMemberLoc(), namedDecl);
 }
 
+void TuIndexer::saveUnresolvedLookupExpr(
+    const clang::UnresolvedLookupExpr &unresolvedLookupExpr) {
+  // The number 32 has been chosen somewhat arbitrarily.
+  // C++ has 20 fundamental types (see
+  // https://stackoverflow.com/a/51777680/2682729) so 32 should cover situations
+  // where overloads for all fundamental types are present, and then some. Set a
+  // limit here to reduce risk of blowing up index sizes and indexing time in
+  // case there are a LOT of unresolved lookup expressions.
+  size_t MAX_UNRESOLVED_DECL_LIMIT = 16;
+  size_t count = 0;
+  for (auto *namedDecl : unresolvedLookupExpr.decls()) {
+    if (!namedDecl) {
+      continue;
+    }
+    auto optSymbol = this->symbolFormatter.getNamedDeclSymbol(*namedDecl);
+    if (!optSymbol) {
+      continue;
+    }
+    this->saveReference(*optSymbol, unresolvedLookupExpr.getNameLoc());
+    ++count;
+    if (count == MAX_UNRESOLVED_DECL_LIMIT) {
+      break;
+    }
+  }
+}
+
 void TuIndexer::saveUnresolvedMemberExpr(
     const clang::UnresolvedMemberExpr &unresolvedMemberExpr) {
   this->trySaveMemberReferenceViaLookup(

test/index/cuda/kernelcall.cu

@@ -58,3 +58,15 @@ struct b {
   }
   void e() { d0(1); }
 };
+
+namespace x {
+  namespace y {
+    template <typename DType, int layout>
+    __global__ void mykernel(const int nthreads, const DType *in_data, DType *out_data) {}
+  }
+}
+
+template <typename DType, int layout>
+void call_mykernel2() {
+  x::y::mykernel<DType, layout><<<0, 0>>>(0, nullptr, nullptr);
+}

test/index/cuda/kernelcall.snapshot.cu

@@ -35,6 +35,7 @@
 //       ^ reference [..] dim3#dim3(6df00707c193238d).
 //          ^ reference [..] dim3#dim3(6df00707c193238d).
     g1(42); // expected-error {{call to global function 'g1' not configured}}
+//  ^^ reference [..] g1(d4f767463ce0a6b3).
     g1<<<1>>>(42); // expected-error {{too few execution configuration arguments to kernel function call}}
     g1<<<1, 1, 0, 0, 0>>>(42); // expected-error {{too many execution configuration arguments to kernel function call}}
 
@@ -103,6 +104,7 @@
 //                          ^^^ reference local 10
       a1::Call<<<0,0>>>(arg);
 //    ^^ reference [..] a1#
+//        ^^^^ reference [..] a1#Call(9b289cee16747614).
 //               ^ reference [..] dim3#dim3(6df00707c193238d).
 //                 ^ reference [..] dim3#dim3(6df00707c193238d).
 //                      ^^^ reference local 10
@@ -113,6 +115,8 @@
 //                         ^^^ reference local 10
       a3::Call<<<0, 0>>>(arg);
 //    ^^ reference [..] a3#
+//        ^^^^ reference [..] a3#Call(5d22bdacc48458e8).
+//        ^^^^ reference [..] a3#Call(d4f767463ce0a6b3).
 //               ^ reference [..] dim3#dim3(6df00707c193238d).
 //                  ^ reference [..] dim3#dim3(6df00707c193238d).
 //                       ^^^ reference local 10
@@ -149,3 +153,36 @@
 //       ^ definition [..] b#e(49f6e7a06ebc5aa8).
 //             ^^ reference [..] b#d0(d4f767463ce0a6b3).
   };
+
+  namespace x {
+//          ^ definition [..] x/
+    namespace y {
+//            ^ definition [..] x/y/
+      template <typename DType, int layout>
+//                       ^^^^^ definition local 12
+//                                  ^^^^^^ definition local 13
+      __global__ void mykernel(const int nthreads, const DType *in_data, DType *out_data) {}
+//    ^^^^^^^^^^ reference [..] `cuda_stub.h:12:9`!
+//                    ^^^^^^^^ definition [..] x/y/mykernel(36fc24b3817d5bcc).
+//                                       ^^^^^^^^ definition local 14
+//                                                       ^^^^^ reference local 12
+//                                                              ^^^^^^^ definition local 15
+//                                                                       ^^^^^ reference local 12
+//                                                                              ^^^^^^^^ definition local 16
+    }
+  }
+
+  template <typename DType, int layout>
+//                   ^^^^^ definition local 17
+//                              ^^^^^^ definition local 18
+  void call_mykernel2() {
+//     ^^^^^^^^^^^^^^ definition [..] call_mykernel2(49f6e7a06ebc5aa8).
+    x::y::mykernel<DType, layout><<<0, 0>>>(0, nullptr, nullptr);
+//  ^ reference [..] x/
+//     ^ reference [..] x/y/
+//        ^^^^^^^^ reference [..] x/y/mykernel(36fc24b3817d5bcc).
+//                 ^^^^^ reference local 17
+//                        ^^^^^^ reference local 18
+//                                  ^ reference [..] dim3#dim3(6df00707c193238d).
+//                                     ^ reference [..] dim3#dim3(6df00707c193238d).
+  }

test/index/functions/templates.snapshot.cc

@@ -55,6 +55,7 @@
 //     ^^ definition [..] h1(9b289cee16747614).
 //        ^ reference local 7
 //          ^ definition local 8
+//               ^^ reference [..] h0(9b289cee16747614).
 //                  ^ reference local 7
 //                     ^ reference local 8