feat: Initial code nav support for CUDA kernel call expressions

docs/Development.md

@@ -182,6 +182,12 @@ clang -Xclang -ast-dump=json file.c
 
 Another option is to use clang-query ([tutorial](https://devblogs.microsoft.com/cppblog/exploring-clang-tooling-part-2-examining-the-clang-ast-with-clang-query/)).
 
+**NOTE:** If running the above on CUDA code
+leads to a Clang error suggesting that CUDA could not be found,
+it's likely that the code is ill-formed. Adding flags like
+`-nocudainc` or `-nocudalib` (sometimes suggested by Clang) will
+lead to `CUDAKernelCallExpr` values not being parsed properly.
+
 ### Automated test case reduction
 
 In case of a crash, it may be possible to automatically reduce

indexer/ClangAstMacros.h

@@ -27,6 +27,7 @@
   F(TemplateTypeParm)
 
 #define FOR_EACH_EXPR_TO_BE_INDEXED(F) \
+  F(CUDAKernelCall)                    \
   F(CXXConstruct)                      \
   F(CXXDependentScopeMember)           \
   F(DeclRef)                           \

indexer/Indexer.cc

@@ -946,6 +946,21 @@ void TuIndexer::saveVarTemplateDecl(const clang::VarTemplateDecl &) {
   // Skip emitting a definition here, as we'll emit one for the inner VarDecl.
 }
 
+void TuIndexer::saveCUDAKernelCallExpr(
+    const clang::CUDAKernelCallExpr &cudaKernelCallExpr) {
+  if (auto *cudaConfig = cudaKernelCallExpr.getConfig()) {
+    if (auto *calleeDecl = cudaConfig->getCalleeDecl()) {
+      if (auto *namedDecl = llvm::dyn_cast<clang::NamedDecl>(calleeDecl)) {
+        if (auto optSymbolName =
+                this->symbolFormatter.getNamedDeclSymbol(*namedDecl)) {
+          auto symbolName = optSymbolName.value();
+          this->saveReference(symbolName, cudaConfig->getExprLoc());
+        }
+      }
+    }
+  }
+}
+
 void TuIndexer::saveCXXConstructExpr(
     const clang::CXXConstructExpr &cxxConstructExpr) {
   if (auto *cxxConstructorDecl = cxxConstructExpr.getConstructor()) {

test/BUILD

@@ -54,6 +54,8 @@ scip_clang_test_suite(
     index_data = glob([
         "index/**/*.c",
         "index/**/*.cc",
+        "index/**/*.cu",
+        "index/**/*.cuh",
         "index/**/*.h",
         "index/**/package-map.json",
     ]),

test/Snapshot.cc

@@ -116,7 +116,8 @@ bool isTuMainFilePath(std::string_view p) {
     return false;
   }
   auto ext = p.substr(dotIndex);
-  return ext == ".cc" || ext == ".cpp" || ext == ".cxx" || ext == ".c";
+  return ext == ".cc" || ext == ".cpp" || ext == ".cxx" || ext == ".c"
+         || ext == ".cu";
 }
 
 // static

test/index/cuda/cuda_stub.h

@@ -0,0 +1,57 @@
+// Based off https://sourcegraph.com/github.com/llvm/llvm-project/-/blob/clang/test/SemaCUDA/Inputs/cuda.h
+// 
+// Add common stuff for CUDA headers here.
+
+using size_t = unsigned long long;
+
+// Make this file work with nvcc, for testing compatibility.
+
+#ifndef __NVCC__
+#define __constant__ __attribute__((constant))
+#define __device__ __attribute__((device))
+#define __global__ __attribute__((global))
+#define __host__ __attribute__((host))
+#define __shared__ __attribute__((shared))
+#define __managed__ __attribute__((managed))
+#define __launch_bounds__(...) __attribute__((launch_bounds(__VA_ARGS__)))
+
+struct dim3 {
+  unsigned x, y, z;
+  __host__ __device__ dim3(unsigned x, unsigned y = 1, unsigned z = 1) : x(x), y(y), z(z) {}
+};
+
+#ifdef __HIP__
+typedef struct hipStream *hipStream_t;
+typedef enum hipError {} hipError_t;
+int hipConfigureCall(dim3 gridSize, dim3 blockSize, size_t sharedSize = 0,
+                     hipStream_t stream = 0);
+extern "C" hipError_t __hipPushCallConfiguration(dim3 gridSize, dim3 blockSize,
+                                                 size_t sharedSize = 0,
+                                                 hipStream_t stream = 0);
+extern "C" hipError_t hipLaunchKernel(const void *func, dim3 gridDim,
+                                      dim3 blockDim, void **args,
+                                      size_t sharedMem,
+                                      hipStream_t stream);
+#else
+typedef struct cudaStream *cudaStream_t;
+typedef enum cudaError {} cudaError_t;
+
+extern "C" int cudaConfigureCall(dim3 gridSize, dim3 blockSize,
+                                 size_t sharedSize = 0,
+                                 cudaStream_t stream = 0);
+extern "C" int __cudaPushCallConfiguration(dim3 gridSize, dim3 blockSize,
+                                           size_t sharedSize = 0,
+                                           cudaStream_t stream = 0);
+extern "C" cudaError_t cudaLaunchKernel(const void *func, dim3 gridDim,
+                                        dim3 blockDim, void **args,
+                                        size_t sharedMem, cudaStream_t stream);
+#endif
+
+// Host- and device-side placement new overloads.
+void *operator new(size_t, void *p) { return p; }
+void *operator new[](size_t, void *p) { return p; }
+__device__ void *operator new(size_t, void *p) { return p; }
+__device__ void *operator new[](size_t, void *p) { return p; }
+
+#endif // !__NVCC__
+

test/index/cuda/cuda_stub.snapshot.h

@@ -0,0 +1,132 @@
+  // Based off https://sourcegraph.com/github.com/llvm/llvm-project/-/blob/clang/test/SemaCUDA/Inputs/cuda.h
+//^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ definition [..] `<file>/cuda_stub.h`/
+  // 
+  // Add common stuff for CUDA headers here.
+
+  using size_t = unsigned long long;
+//      ^^^^^^ definition [..] size_t#
+
+  // Make this file work with nvcc, for testing compatibility.
+
+  #ifndef __NVCC__
+  #define __constant__ __attribute__((constant))
+//        ^^^^^^^^^^^^ definition [..] `cuda_stub.h:10:9`!
+  #define __device__ __attribute__((device))
+//        ^^^^^^^^^^ definition [..] `cuda_stub.h:11:9`!
+  #define __global__ __attribute__((global))
+//        ^^^^^^^^^^ definition [..] `cuda_stub.h:12:9`!
+  #define __host__ __attribute__((host))
+//        ^^^^^^^^ definition [..] `cuda_stub.h:13:9`!
+  #define __shared__ __attribute__((shared))
+//        ^^^^^^^^^^ definition [..] `cuda_stub.h:14:9`!
+  #define __managed__ __attribute__((managed))
+//        ^^^^^^^^^^^ definition [..] `cuda_stub.h:15:9`!
+  #define __launch_bounds__(...) __attribute__((launch_bounds(__VA_ARGS__)))
+//        ^^^^^^^^^^^^^^^^^ definition [..] `cuda_stub.h:16:9`!
+
+  struct dim3 {
+//       ^^^^ definition [..] dim3#
+    unsigned x, y, z;
+//           ^ definition [..] dim3#x.
+//              ^ definition [..] dim3#y.
+//                 ^ definition [..] dim3#z.
+    __host__ __device__ dim3(unsigned x, unsigned y = 1, unsigned z = 1) : x(x), y(y), z(z) {}
+//  ^^^^^^^^ reference [..] `cuda_stub.h:13:9`!
+//           ^^^^^^^^^^ reference [..] `cuda_stub.h:11:9`!
+//                      ^^^^ definition [..] dim3#dim3(6df00707c193238d).
+//                                    ^ definition local 0
+//                                                ^ definition local 1
+//                                                                ^ definition local 2
+//                                                                         ^ reference [..] dim3#x.
+//                                                                           ^ reference local 0
+//                                                                               ^ reference [..] dim3#y.
+//                                                                                 ^ reference local 1
+//                                                                                     ^ reference [..] dim3#z.
+//                                                                                       ^ reference local 2
+  };
+
+  #ifdef __HIP__
+  typedef struct hipStream *hipStream_t;
+  typedef enum hipError {} hipError_t;
+  int hipConfigureCall(dim3 gridSize, dim3 blockSize, size_t sharedSize = 0,
+                       hipStream_t stream = 0);
+  extern "C" hipError_t __hipPushCallConfiguration(dim3 gridSize, dim3 blockSize,
+                                                   size_t sharedSize = 0,
+                                                   hipStream_t stream = 0);
+  extern "C" hipError_t hipLaunchKernel(const void *func, dim3 gridDim,
+                                        dim3 blockDim, void **args,
+                                        size_t sharedMem,
+                                        hipStream_t stream);
+  #else
+  typedef struct cudaStream *cudaStream_t;
+//               ^^^^^^^^^^ reference [..] cudaStream#
+//               ^^^^^^^^^^ reference [..] cudaStream#
+//                           ^^^^^^^^^^^^ definition [..] cudaStream_t#
+  typedef enum cudaError {} cudaError_t;
+//             ^^^^^^^^^ definition [..] cudaError#
+//                          ^^^^^^^^^^^ definition [..] cudaError_t#
+
+  extern "C" int cudaConfigureCall(dim3 gridSize, dim3 blockSize,
+//                                 ^^^^ reference [..] dim3#
+//                                      ^^^^^^^^ definition local 3
+//                                                ^^^^ reference [..] dim3#
+//                                                     ^^^^^^^^^ definition local 4
+                                   size_t sharedSize = 0,
+//                                 ^^^^^^ reference [..] size_t#
+//                                        ^^^^^^^^^^ definition local 5
+                                   cudaStream_t stream = 0);
+//                                 ^^^^^^^^^^^^ reference [..] cudaStream_t#
+//                                              ^^^^^^ definition local 6
+  extern "C" int __cudaPushCallConfiguration(dim3 gridSize, dim3 blockSize,
+//                                           ^^^^ reference [..] dim3#
+//                                                ^^^^^^^^ definition local 7
+//                                                          ^^^^ reference [..] dim3#
+//                                                               ^^^^^^^^^ definition local 8
+                                             size_t sharedSize = 0,
+//                                           ^^^^^^ reference [..] size_t#
+//                                                  ^^^^^^^^^^ definition local 9
+                                             cudaStream_t stream = 0);
+//                                           ^^^^^^^^^^^^ reference [..] cudaStream_t#
+//                                                        ^^^^^^ definition local 10
+  extern "C" cudaError_t cudaLaunchKernel(const void *func, dim3 gridDim,
+//           ^^^^^^^^^^^ reference [..] cudaError_t#
+//                                                    ^^^^ definition local 11
+//                                                          ^^^^ reference [..] dim3#
+//                                                               ^^^^^^^ definition local 12
+                                          dim3 blockDim, void **args,
+//                                        ^^^^ reference [..] dim3#
+//                                             ^^^^^^^^ definition local 13
+//                                                              ^^^^ definition local 14
+                                          size_t sharedMem, cudaStream_t stream);
+//                                        ^^^^^^ reference [..] size_t#
+//                                               ^^^^^^^^^ definition local 15
+//                                                          ^^^^^^^^^^^^ reference [..] cudaStream_t#
+//                                                                       ^^^^^^ definition local 16
+  #endif
+
+  // Host- and device-side placement new overloads.
+  void *operator new(size_t, void *p) { return p; }
+//      ^^^^^^^^ definition [..] `operator new`(ecd71fefd6822377).
+//                   ^^^^^^ reference [..] size_t#
+//                                 ^ definition local 17
+//                                             ^ reference local 17
+  void *operator new[](size_t, void *p) { return p; }
+//      ^^^^^^^^ definition [..] `operator new[]`(ecd71fefd6822377).
+//                     ^^^^^^ reference [..] size_t#
+//                                   ^ definition local 18
+//                                               ^ reference local 18
+  __device__ void *operator new(size_t, void *p) { return p; }
+//^^^^^^^^^^ reference [..] `cuda_stub.h:11:9`!
+//                 ^^^^^^^^ definition [..] `operator new`(ecd71fefd6822377).
+//                              ^^^^^^ reference [..] size_t#
+//                                            ^ definition local 19
+//                                                        ^ reference local 19
+  __device__ void *operator new[](size_t, void *p) { return p; }
+//^^^^^^^^^^ reference [..] `cuda_stub.h:11:9`!
+//                 ^^^^^^^^ definition [..] `operator new[]`(ecd71fefd6822377).
+//                                ^^^^^^ reference [..] size_t#
+//                                              ^ definition local 20
+//                                                          ^ reference local 20
+
+  #endif // !__NVCC__
+

test/index/cuda/kernelcall.cu

@@ -0,0 +1,60 @@
+// Initially based off kernel-call.cu in the Clang tests
+// https://sourcegraph.com/github.com/llvm/llvm-project/-/blob/clang/test/SemaCUDA/kernel-call.cu
+
+#include "cuda_stub.h"
+
+__global__ void g1(int x) {}
+
+template <typename T> void t1(T arg) {
+  g1<<<arg, arg>>>(1);
+}
+
+void h1(int x) {}
+int h2(int x) { return 1; }
+
+int main(void) {
+  g1<<<1, 1>>>(42);
+  g1(42); // expected-error {{call to global function 'g1' not configured}}
+  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}}
+
+  t1(1);
+
+  h1<<<1, 1>>>(42); // expected-error {{kernel call to non-global function 'h1'}}
+
+  int (*fp)(int) = h2;
+  fp<<<1, 1>>>(42); // expected-error {{must have void return type}}
+
+  g1<<<undeclared, 1>>>(42); // expected-error {{use of undeclared identifier 'undeclared'}}
+}
+
+// Make sure we can call static member kernels.
+template <typename > struct a0 {
+  template <typename T> static __global__ void Call(T);
+};
+struct a1 {
+  template <typename T> static __global__ void Call(T);
+};
+template <typename T> struct a2 {
+  static __global__ void Call(T);
+};
+struct a3 {
+  static __global__ void Call(int);
+  static __global__ void Call(void*);
+};
+
+struct b {
+  template <typename c> void d0(c arg) {
+    a0<c>::Call<<<0, 0>>>(arg);
+    a1::Call<<<0,0>>>(arg);
+    a2<c>::Call<<<0,0>>>(arg);
+    a3::Call<<<0, 0>>>(arg);
+  }
+  void d1(void* arg) {
+    a0<void*>::Call<<<0, 0>>>(arg);
+    a1::Call<<<0,0>>>(arg);
+    a2<void*>::Call<<<0,0>>>(arg);
+    a3::Call<<<0, 0>>>(arg);
+  }
+  void e() { d0(1); }
+};