From a6b390ab16969d98eea9fc2dc95892e5a62aa4da Mon Sep 17 00:00:00 2001 From: Varun Gandhi Date: Mon, 2 Oct 2023 10:17:18 -0700 Subject: [PATCH 1/3] test: Add test for kernel call in template --- test/index/cuda/kernelcall.cu | 12 ++++++++++ test/index/cuda/kernelcall.snapshot.cu | 32 ++++++++++++++++++++++++++ 2 files changed, 44 insertions(+) diff --git a/test/index/cuda/kernelcall.cu b/test/index/cuda/kernelcall.cu index a54e078e..2676e633 100644 --- a/test/index/cuda/kernelcall.cu +++ b/test/index/cuda/kernelcall.cu @@ -58,3 +58,15 @@ struct b { } void e() { d0(1); } }; + +namespace x { + namespace y { + template + __global__ void mykernel(const int nthreads, const DType *in_data, DType *out_data) {} + } +} + +template +void call_mykernel2() { + x::y::mykernel<<<0, 0>>>(0, nullptr, nullptr); +} diff --git a/test/index/cuda/kernelcall.snapshot.cu b/test/index/cuda/kernelcall.snapshot.cu index 69aee45c..fa5d3344 100644 --- a/test/index/cuda/kernelcall.snapshot.cu +++ b/test/index/cuda/kernelcall.snapshot.cu @@ -149,3 +149,35 @@ // ^ definition [..] b#e(49f6e7a06ebc5aa8). // ^^ reference [..] b#d0(d4f767463ce0a6b3). }; + + namespace x { +// ^ definition [..] x/ + namespace y { +// ^ definition [..] x/y/ + template +// ^^^^^ 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 +// ^^^^^ definition local 17 +// ^^^^^^ definition local 18 + void call_mykernel2() { +// ^^^^^^^^^^^^^^ definition [..] call_mykernel2(49f6e7a06ebc5aa8). + x::y::mykernel<<<0, 0>>>(0, nullptr, nullptr); +// ^ reference [..] x/ +// ^ reference [..] x/y/ +// ^^^^^ reference local 17 +// ^^^^^^ reference local 18 +// ^ reference [..] dim3#dim3(6df00707c193238d). +// ^ reference [..] dim3#dim3(6df00707c193238d). + } From 8b0f1c1a1a6cfe9db5981360980d0db805bfd213 Mon Sep 17 00:00:00 2001 From: Varun Gandhi Date: Mon, 2 Oct 2023 15:06:11 -0700 Subject: [PATCH 2/3] fix: Try to emit refs for unresolved call exprs --- indexer/ClangAstMacros.h | 1 + indexer/Indexer.cc | 26 ++++++++++++++++++++++ test/index/cuda/kernelcall.snapshot.cu | 5 +++++ test/index/functions/templates.snapshot.cc | 1 + 4 files changed, 33 insertions(+) diff --git a/indexer/ClangAstMacros.h b/indexer/ClangAstMacros.h index ee2f23f6..a60d6d30 100644 --- a/indexer/ClangAstMacros.h +++ b/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) \ diff --git a/indexer/Indexer.cc b/indexer/Indexer.cc index f857bfff..0a669239 100644 --- a/indexer/Indexer.cc +++ b/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( diff --git a/test/index/cuda/kernelcall.snapshot.cu b/test/index/cuda/kernelcall.snapshot.cu index fa5d3344..cb05b1a1 100644 --- a/test/index/cuda/kernelcall.snapshot.cu +++ b/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 @@ -176,6 +180,7 @@ x::y::mykernel<<<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). diff --git a/test/index/functions/templates.snapshot.cc b/test/index/functions/templates.snapshot.cc index 9711c5f4..c6963717 100644 --- a/test/index/functions/templates.snapshot.cc +++ b/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 From 2af09420d35be174349c1b41abbdef75afec2b5e Mon Sep 17 00:00:00 2001 From: Varun Gandhi Date: Mon, 2 Oct 2023 15:17:27 -0700 Subject: [PATCH 3/3] style: Reformat --- indexer/Indexer.cc | 16 ++++++++-------- 1 file changed, 8 insertions(+), 8 deletions(-) diff --git a/indexer/Indexer.cc b/indexer/Indexer.cc index 0a669239..57aeebb3 100644 --- a/indexer/Indexer.cc +++ b/indexer/Indexer.cc @@ -1022,17 +1022,17 @@ void TuIndexer::saveMemberExpr(const clang::MemberExpr &memberExpr) { this->saveReference(optSymbol.value(), memberExpr.getMemberLoc(), namedDecl); } -void TuIndexer::saveUnresolvedLookupExpr(const clang::UnresolvedLookupExpr &unresolvedLookupExpr) { +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. + // 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()) { + for (auto *namedDecl : unresolvedLookupExpr.decls()) { if (!namedDecl) { continue; }