diff --git a/.claude/skills/agent-eval/corpus.json b/.claude/skills/agent-eval/corpus.json index 8b4331f..def7c3e 100644 --- a/.claude/skills/agent-eval/corpus.json +++ b/.claude/skills/agent-eval/corpus.json @@ -515,5 +515,28 @@ "files": "~400", "question": "How does an OpenZeppelin AccessControl-protected function check the caller's role? Trace from the onlyRole modifier through hasRole to the role storage." } + ], + "CUDA": [ + { + "name": "llm.c", + "repo": "https://github.com/karpathy/llm.c", + "size": "Small", + "files": "~76", + "question": "How does the attention forward pass reach the GPU in the CUDA training path? Trace from attention_forward to the kernels it launches, and explain where softmax happens." + }, + { + "name": "flash-attention", + "repo": "https://github.com/Dao-AILab/flash-attention", + "size": "Medium", + "files": "~900", + "question": "How does a Python call to flash_attn_func reach the CUDA kernel that computes forward attention? Trace the dispatch path from the Python API through the C++ binding to the kernel launch." + }, + { + "name": "cutlass", + "repo": "https://github.com/NVIDIA/cutlass", + "size": "Large", + "files": "~2700", + "question": "When a cutlass device-level GEMM (cutlass::gemm::device::Gemm) is invoked, how does it reach the GPU kernel? Trace from the operator() call to the kernel entry point and its launch site." + } ] } diff --git a/CHANGELOG.md b/CHANGELOG.md index 86cab81..2a7dc37 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -11,6 +11,9 @@ and adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0.html). ### New Features +- CodeGraph now indexes **CUDA** (`.cu`, `.cuh`) — kernels, device/host functions, structs, and classes become symbols, and the host→kernel call edge survives the `<<>>` launch syntax, so questions like "how does this call reach the GPU kernel?" trace across the CPU/GPU boundary instead of going dark at the launch site. Real-world launch styles all connect: templated launches (`my_kernel<<>>(args)`), launches through a local function pointer (`auto kernel = &my_kernel<...>; ... kernel<<>>(args)` — each branch-assigned target linked), brace-initialized launch configs (`<<>>`), and kernels defined through a name-in-first-argument macro (flash-attention's `DEFINE_FLASH_FORWARD_KERNEL(kernel_name, ...) { ... }` style), which now index under their real kernel names. CUDA that lives in plain `.h`/`.hpp` headers — where much real-world device code sits, launch-template headers included — is recognized by content and indexed the same way. Validated on llm.c, flash-attention, and NVIDIA CUTLASS. (#387, #648) +- C++ symbols defined inside `namespace` blocks now carry the namespace in their qualified name (`flash::compute_attn`, C++17 `namespace a::b {` included), and namespace-qualified calls (`ns::fn(...)`) resolve to their definitions — previously such calls never linked at all, which hid much of the call graph in namespace-heavy C++ codebases from callers and impact analysis. +- C++ calls that spell out template arguments (`fn(args)`) now link to the function they instantiate, the same normalization templated base classes already had. - CodeGraph now indexes **Solidity** (`.sol`) — contracts, libraries, interfaces, structs, enums, modifiers, events, errors, and state variables become first-class symbols, with call edges that follow `emit`, `revert`, modifier guards (`onlyOwner`-style, including base-constructor chains like `constructor() ERC20(...)`), and library/method calls (including `using` directives). `import` directives resolve to the imported file, so cross-contract questions like "trace `transferFrom` through allowance and balance updates" or "how does `onlyRole` reach the role storage?" work out of the box on real Solidity codebases (validated on solmate, solady, and OpenZeppelin Contracts). Thanks @naiba. (#374, #648) - Erlang behaviour dispatch is now followed through the graph: a framework call through a variable module — cowboy's `Handler:init`/`Middleware:execute` folds, a plugin manager's `Mod:callback(...)` — links to the repo's implementations of the behaviour that declares that callback, so flow traces and impact cross the OTP callback boundary instead of stopping at it. The links are precision-gated: the callback arity must match, exactly one behaviour may own that callback shape (a collision stays unlinked rather than guessed), the implementer must actually export the callback, and the fan-out is bounded — a behaviour with hundreds of implementers stays a visibly dynamic boundary. Every bridged hop is labeled as dynamic dispatch with its wiring site, never shown as a plain static call. - CodeGraph now indexes **Erlang** (`.erl`, `.hrl`) — functions, with clauses and arities of the same name grouped as one symbol spanning all of them, plus records with their fields, `-type`/`-opaque` aliases, `-define` macros, and `-spec` signatures attached to every function. Cross-module `mod:fn(...)` calls resolve to the target module's function, `fun name/arity` values are captured as references (so callback registrations like `lists:foreach(fun submit/1, ...)` link up), `-include`/`-include_lib` connect to the header files they pull in, `-behaviour` declarations link a callback module to its behaviour (and only ever to a module — a same-named macro or function elsewhere in the repo is never mistaken for one), and `-export` lists (plus `-compile(export_all)`) drive each function's public/private flag. OTP's indirection idioms are followed where the target is static: `spawn`/`apply`/`proc_lib`/`timer`/`rpc` calls that name their target as `(Module, Function, Args)` arguments produce call edges, and `gen_server:call`/`cast` connects to the target module's `handle_call`/`handle_cast` — its own when targeting `?MODULE` (including the `-define(SERVER, ?MODULE)` idiom), and the named module when a registered name follows OTP's name-the-server-after-its-module convention (`gen_server:call(other_mod, ...)`, directly or through a `-define(STORE, other_mod)` macro); a registered name that matches no module stays unlinked. Macros participate in the graph too: a `-define` body's calls belong to the macro, each `?MACRO(...)` use site links into the call chain (and bare `?CONSTANT` reads are tracked as references), so a call path hidden behind a macro — `set_password → ?SQL_UPSERT_T → sql_query_t` — traces end-to-end and "where is this macro used" is answerable. escripts index like any module (the shebang line is understood), and OTP application resource files (`.app.src`, `.app`) join the graph: `{mod, ...}` links an app to its callback module and `{applications, [...]}` connects umbrella sibling apps — resolving only ever to modules, so an OTP app name like `ssl` is never mistaken for a same-named function. Truly dynamic dispatch (`Mod:handle(...)`, message sends, var-module spawns) is deliberately left unlinked rather than guessed. `codegraph_explore` also understands Erlang-native symbol spelling in queries — `mod:fn/3` and `init/2` find the symbols they name. (#635, #648) diff --git a/README.md b/README.md index ae1078d..e26ed81 100644 --- a/README.md +++ b/README.md @@ -244,7 +244,7 @@ The reliable, universal payoff is **surgical context and speed**: CodeGraph coll | **Full-Text Search** | Find code by name instantly across your entire codebase, powered by FTS5 | | **Impact Analysis** | Trace callers, callees, and the full impact radius of any symbol before making changes | | **Always Fresh** | File watcher uses native OS events (FSEvents/inotify/ReadDirectoryChangesW) with debounced auto-sync — the graph stays current as you code, zero config | -| **20+ Languages** | TypeScript, JavaScript, Python, Go, Rust, Java, C#, VB.NET, PHP, Ruby, C, C++, Objective-C, Metal, Swift, Kotlin, Scala, Dart, Lua, Luau, R, Erlang, CFML, COBOL, Solidity, Svelte, Vue, Astro, Liquid, Pascal/Delphi | +| **20+ Languages** | TypeScript, JavaScript, Python, Go, Rust, Java, C#, VB.NET, PHP, Ruby, C, C++, CUDA, Objective-C, Metal, Swift, Kotlin, Scala, Dart, Lua, Luau, R, Erlang, CFML, COBOL, Solidity, Svelte, Vue, Astro, Liquid, Pascal/Delphi | | **Framework-aware Routes** | Recognizes web-framework routing files and links URL patterns to their handlers across 17 frameworks | | **Mixed iOS / React Native / Expo** | Closes cross-language flows that static parsing misses: Swift ↔ ObjC bridging, React Native legacy bridge + TurboModules + Fabric view components, native → JS event emitters, Expo Modules | | **100% Local** | No data leaves your machine. No API keys. No external services. SQLite database only | @@ -703,6 +703,7 @@ is written): | C++ | `.cpp`, `.hpp`, `.cc` | Full support | | Objective-C | `.m`, `.mm`, `.h` | Partial support (classes, protocols, methods, `@property`, `#import`, message sends; `.mm` ObjC++ may parse incompletely) | | Metal | `.metal` | Full support (vertex/fragment/kernel functions, structs, type aliases, call edges — MSL parses as C++, with `[[attribute]]` annotations handled) | +| CUDA | `.cu`, `.cuh` | Full support (kernels and device/host functions, structs, classes, host→kernel call edges through `<<>>` launch syntax — templated launches, function-pointer launches (`auto kernel = &fn<...>`), `dim3{...}` configs, and macro-defined kernels included; `__global__`/`__device__`/`__launch_bounds__` specifiers handled; CUDA in plain `.h`/`.hpp` headers recognized by content) | | Swift | `.swift` | Full support | | Kotlin | `.kt`, `.kts` | Full support | | Scala | `.scala`, `.sc` | Full support (classes, traits, methods, type aliases, Scala 3 enums) | diff --git a/__tests__/extraction.test.ts b/__tests__/extraction.test.ts index f68dfc5..eb51c86 100644 --- a/__tests__/extraction.test.ts +++ b/__tests__/extraction.test.ts @@ -11,7 +11,7 @@ import * as os from 'os'; import { CodeGraph } from '../src'; import { extractFromSource, scanDirectory, buildDefaultIgnore, discoverEmbeddedRepoRoots, buildScopeIgnore } from '../src/extraction'; import { detectLanguage, isLanguageSupported, getSupportedLanguages, initGrammars, loadAllGrammars, isSourceFile } from '../src/extraction/grammars'; -import { stripCppTemplateArgs, blankCppExportMacros, blankCppInlineMacros, blankMetalAttributes, recoverMangledCppName } from '../src/extraction/languages/c-cpp'; +import { stripCppTemplateArgs, blankCppExportMacros, blankCppInlineMacros, blankMetalAttributes, blankCudaConstructs, recoverMangledCppName } from '../src/extraction/languages/c-cpp'; import { normalizePath } from '../src/utils'; beforeAll(async () => { @@ -107,6 +107,13 @@ describe('Language Detection', () => { expect(isSourceFile('Renderer/Shaders.metal')).toBe(true); }); + it('should detect CUDA files as C++ (#387)', () => { + expect(detectLanguage('kernels/scan.cu')).toBe('cpp'); + expect(detectLanguage('include/reduce.cuh')).toBe('cpp'); + expect(isSourceFile('csrc/flash_attn/softmax.cu')).toBe(true); + expect(isSourceFile('include/block_reduce.cuh')).toBe(true); + }); + it('should detect Erlang files', () => { expect(detectLanguage('src/my_server.erl')).toBe('erlang'); expect(detectLanguage('include/records.hrl')).toBe('erlang'); @@ -2946,6 +2953,333 @@ kernel void computeBlur(texture2d inTexture [[texture(0)]], }); }); + describe('CUDA extraction (#387)', () => { + // CUDA parses with the C++ grammar. Three CUDA-only shapes misparse: + // execution-space specifiers (`__global__ void f(…)`) shunt the real return + // type into an ERROR node, `__shared__ float tile[256]` mangles the declared + // name to `float`, and — the critical one — `k<<>>(args)` lexes + // as shift operators around an empty-named template so NO call_expression + // (and therefore no host→kernel call edge) exists. blankCudaConstructs + // (preParse, `.cu`/`.cuh`-gated) blanks all three so extraction matches + // plain C++. + const CUDA = `#include +#include "kernels.cuh" + +__constant__ float d_scale[16]; + +__device__ __forceinline__ float warp_reduce_sum(float val) { + for (int offset = 16; offset > 0; offset /= 2) { + val += __shfl_down_sync(0xffffffff, val, offset); + } + return val; +} + +__global__ void scale_kernel(float* out, const float* __restrict__ in, int n) { + int i = blockIdx.x * blockDim.x + threadIdx.x; + __shared__ float tile[256]; + if (i < n) { + tile[threadIdx.x] = in[i]; + __syncthreads(); + out[i] = warp_reduce_sum(tile[threadIdx.x]) * d_scale[0]; + } +} + +__global__ void __launch_bounds__(256, 4) bounded_kernel(float* data, int n) { + if (blockIdx.x * blockDim.x + threadIdx.x < n) data[0] *= 2.0f; +} + +template +__global__ void templated_kernel(T* data, int n) { + if (blockIdx.x * BLOCK + threadIdx.x < n) data[0] += T(1); +} + +class GpuBuffer { +public: + explicit GpuBuffer(size_t n) { cudaMalloc(&ptr_, n * sizeof(float)); } + ~GpuBuffer() { cudaFree(ptr_); } +private: + float* ptr_ = nullptr; +}; + +void launch_scale(float* out, const float* in, int n, cudaStream_t stream) { + dim3 block(256); + dim3 grid((n + block.x - 1) / block.x); + scale_kernel<<>>(out, in, n); + bounded_kernel<<>>(out, n); + templated_kernel<<>>(out, n); +} +`; + + it('extracts kernels, device functions, and host→kernel launch calls from a .cu file', () => { + const result = extractFromSource('kernels/scan.cu', CUDA); + expect(result.errors).toHaveLength(0); + + const functions = result.nodes.filter((n) => n.kind === 'function').map((n) => n.name); + expect(functions).toEqual( + expect.arrayContaining([ + 'warp_reduce_sum', + 'scale_kernel', + 'bounded_kernel', + 'templated_kernel', + 'launch_scale', + ]) + ); + expect(result.nodes.filter((n) => n.kind === 'class').map((n) => n.name)).toContain('GpuBuffer'); + expect(result.nodes.find((n) => n.kind === 'import')?.name).toBe('cuda_runtime.h'); + + // No misparse artifacts: pre-blank, `__shared__ float tile[256]` parsed + // with `float` as the declared name. (Top-level C++ variables aren't + // extracted as nodes — matching plain-C++ behavior is the target.) + expect(result.nodes.map((n) => n.name)).not.toContain('float'); + + // Blanking is offset-preserving, so positions stay exact. + expect(result.nodes.find((n) => n.name === 'scale_kernel')!.startLine).toBe(13); + + // THE point of CUDA support: every `<<<…>>>` launch form — plain, + // launch-bounds, multi-line config, and templated — emits a `calls` + // reference, so the host→kernel edge exists in the graph. Pre-blank, + // the chevrons lexed as shifts and none of these existed. + const calls = result.unresolvedReferences + .filter((r) => r.referenceKind === 'calls') + .map((r) => r.referenceName); + // The templated launch is normalized to the bare kernel name (template + // args stripped at extraction, like base-class extends refs — #1043), so + // it resolves to the kernel the template was defined as. + expect(calls).toEqual( + expect.arrayContaining([ + 'scale_kernel', + 'bounded_kernel', + 'templated_kernel', + 'warp_reduce_sum', + ]) + ); + }); + + it('blankCudaConstructs blanks every CUDA form, offset- and newline-preserving', () => { + const inp = [ + '__global__ void __launch_bounds__(256, 4) step(float* p) {', + ' __shared__ float tile[32];', + '}', + '__host__ __device__ int both() { return 0; }', + 'void run(float* p, int n) {', + ' step<<>>(p);', + '}', + ].join('\n'); + const out = blankCudaConstructs(inp); + expect(out.length).toBe(inp.length); // every byte offset preserved + expect(out.split('\n').length).toBe(inp.split('\n').length); // newlines survive the multi-line launch config + expect(out).not.toMatch(/__global__|__launch_bounds__|__shared__|__host__|__device__|<<<|>>>/); + // Collapsing blank runs gives plain C++ back. + expect(out.split('\n').map((l) => l.replace(/ +/g, ' ').trimEnd())).toEqual([ + ' void step(float* p) {', + ' float tile[32];', + '}', + ' int both() { return 0; }', + 'void run(float* p, int n) {', + ' step', + ' (p);', + '}', + ]); + }); + + it('blankCudaConstructs never touches non-CUDA chevrons or identifiers', () => { + for (const c of [ + 'std::cout << "a" << b << c;', // shift chains — never three consecutive < + 'auto x = f(a >> 3, b >> 3);', // right shifts + 'std::vector>> deep;', // template >>> closer with no <<< opener + 'printf("<<>> anywhere + 'int __restrict__like = 1;', // dunder-ish identifier not in the specifier list + 'int z = 1;', // nothing CUDA at all — early-return path + ]) { + expect(blankCudaConstructs(c)).toBe(c); + } + // A stray `<<<` (committed merge-conflict marker) must not blank the code + // between markers. Two independent guards: statements between markers + // carry `;` (excluded from the span)… + const conflict = [ + '<<<<<<< HEAD', + 'int a = compute(1);', + '=======', + 'int a = compute(2);', + '>>>>>>> feature-branch', + ].join('\n'); + expect(blankCudaConstructs(conflict)).toBe(conflict); + // …and a `;`-free region still fails the brace-balance check (the `{`s + // opened between the markers never close before the `>>>`). + const semicolonFree = [ + '<<<<<<< HEAD', + 'void foo() {', + '=======', + 'void bar() {', + '>>>>>>> feature-branch', + ].join('\n'); + expect(blankCudaConstructs(semicolonFree)).toBe(semicolonFree); + }); + + it('blanks brace-initialized launch configs (`dim3{…}`), balanced-only', () => { + const inp = 'run_it<<>>(data, n);'; + const out = blankCudaConstructs(inp); + expect(out.length).toBe(inp.length); + expect(out.replace(/ +/g, ' ')).toBe('run_it (data, n);'); + }); + + it('recovers the real kernel name from a macro-definition idiom, gtest/pybind untouched', () => { + const code = `#define DEFINE_MY_FWD_KERNEL(kernelName, ...) \\ +template \\ +__global__ void kernelName(const Params params) + +DEFINE_MY_FWD_KERNEL(fwd_kernel, bool Is_causal, int kBlockM) { + do_work(params); +} + +TEST_F(MyFixture, HandlesEmptyInput) { + check(1); +} +`; + const result = extractFromSource('kernels/impl.cu', code); + const functions = result.nodes.filter((n) => n.kind === 'function').map((n) => n.name); + // The macro invocation's first argument is the defined name. + expect(functions).toContain('fwd_kernel'); + expect(functions).not.toContain('DEFINE_MY_FWD_KERNEL'); + // gtest's TEST_F(Fixture, Name) has TWO lone identifiers — ambiguous, so + // it keeps the macro name rather than guessing. + expect(functions).toContain('TEST_F'); + expect(functions).not.toContain('MyFixture'); + }); + + it('links launches through a local function pointer to the real kernel(s)', () => { + // The flash-attention launch-template shape end-to-end: a macro-defined + // kernel + `auto kernel = &fn<…>` + branch reassignment + launch through + // the local. The call refs must name the real kernels, not `kernel`. + const code = `template +__global__ void fwd_kernel(T* data, int n) { + if (blockIdx.x * blockDim.x + threadIdx.x < n) data[0] += T(1); +} + +template +__global__ void fwd_splitkv_kernel(T* data, int n) { + if (blockIdx.x * blockDim.x + threadIdx.x < n) data[0] += T(2); +} + +template +void run_fwd(T* data, int n, cudaStream_t stream) { + auto kernel = &fwd_kernel; + if (n % 2 == 0) { + kernel = &fwd_kernel; + } else if (n % 3 == 0) { + kernel = &fwd_splitkv_kernel; + } + kernel<<<(n + 255) / 256, 256, 0, stream>>>(data, n); +} +`; + const result = extractFromSource('kernels/launch.cu', code); + expect(result.errors).toHaveLength(0); + const calls = result.unresolvedReferences + .filter((r) => r.referenceKind === 'calls') + .map((r) => r.referenceName); + // Every DISTINCT branch target recorded once — the two fwd_kernel<…> + // instantiations strip to one target, the splitkv branch adds a second. + // The local's name never leaks as a callee. + expect(calls.filter((c) => c === 'fwd_kernel')).toHaveLength(1); + expect(calls.filter((c) => c === 'fwd_splitkv_kernel')).toHaveLength(1); + expect(calls).not.toContain('kernel'); + }); + + it('CUDA blanking is gated by extension or content — plain C++ shift/template chevrons are untouched', () => { + const cpp = `#include +int shift_it(int a, int b) { return a << b << 1; } +std::vector>> matrix() { return {}; } +`; + const result = extractFromSource('math.cpp', cpp); + expect(result.errors).toHaveLength(0); + const functions = result.nodes.filter((n) => n.kind === 'function').map((n) => n.name); + expect(functions).toEqual(expect.arrayContaining(['shift_it', 'matrix'])); + }); + + it('CUDA in extension-less headers is caught by content: launch templates in .h connect host→kernel', () => { + // Much real CUDA lives in .h: cutlass launches most of its kernels from + // headers, flash-attention's launch templates are .h, llm.c keeps device + // helpers in C-detected .h. `looksLikeCudaSource` content-gates the same + // blank there — no CUDA marker is valid C++ anywhere, so this can't + // affect a genuinely-plain C++ header. + const header = `#pragma once +#include + +template +__global__ void fill_kernel(T* out, T value, int n) { + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) out[i] = value; +} + +template +void launch_fill(T* out, T value, int n, cudaStream_t stream) { + fill_kernel<<<(n + 255) / 256, 256, 0, stream>>>(out, value, n); +} +`; + const result = extractFromSource('include/fill_launch_template.h', header); + expect(result.errors).toHaveLength(0); + const functions = result.nodes.filter((n) => n.kind === 'function').map((n) => n.name); + expect(functions).toEqual(expect.arrayContaining(['fill_kernel', 'launch_fill'])); + const calls = result.unresolvedReferences + .filter((r) => r.referenceKind === 'calls') + .map((r) => r.referenceName); + expect(calls).toContain('fill_kernel'); + }); + }); + + describe('C++ namespace qualifiedName prefixing', () => { + // C++ namespaces previously left no trace in qualifiedNames, so a + // namespace-qualified call (`flash::compute(...)`) could never match its + // definition — every `ns::fn()` call site was a permanently dead edge. + // The namespace name now prefixes contained symbols' qualifiedNames + // (prefix-only: no namespace node is minted — `namespace cutlass {` opens + // in thousands of files and a node per block would crowd search, #1093). + it('prefixes contained symbols and handles nesting; anonymous stays bare', () => { + const code = `namespace flash { +namespace detail { +void helper() {} +} +void compute_attn(int x) { detail::helper(); } +class Softmax { +public: + void rescale() {} +}; +} +namespace { +void file_local() {} +} +void global_fn() { flash::compute_attn(1); } +`; + const result = extractFromSource('dispatch.cpp', code); + expect(result.errors).toHaveLength(0); + const byName = new Map(result.nodes.map((n) => [n.name, n])); + expect(byName.get('compute_attn')?.qualifiedName).toBe('flash::compute_attn'); + expect(byName.get('helper')?.qualifiedName).toBe('flash::detail::helper'); + expect(byName.get('Softmax')?.qualifiedName).toBe('flash::Softmax'); + // Class scope still stacks under the namespace prefix. + expect(byName.get('rescale')?.qualifiedName).toBe('flash::Softmax::rescale'); + // Anonymous namespace contents and true globals stay bare. + expect(byName.get('file_local')?.qualifiedName).toBe('file_local'); + expect(byName.get('global_fn')?.qualifiedName).toBe('global_fn'); + // The qualified call refs are emitted as spelled. + const calls = result.unresolvedReferences + .filter((r) => r.referenceKind === 'calls') + .map((r) => r.referenceName); + expect(calls).toEqual(expect.arrayContaining(['flash::compute_attn', 'detail::helper'])); + }); + + it('C++17 nested namespace form prefixes as written', () => { + const code = `namespace a::b { +int f() { return 1; } +} +`; + const result = extractFromSource('nested.cpp', code); + expect(result.nodes.find((n) => n.name === 'f')?.qualifiedName).toBe('a::b::f'); + }); + }); + describe('C++ forward declarations do not mint phantom class nodes (#1093)', () => { // `class Foo;` parses as a bodiless class_specifier. Repeated across headers, // each forward decl minted a phantom bodiless `class` node that crowded out — diff --git a/src/extraction/grammars.ts b/src/extraction/grammars.ts index 7dd014d..973f902 100644 --- a/src/extraction/grammars.ts +++ b/src/extraction/grammars.ts @@ -125,6 +125,13 @@ export const EXTENSION_MAP: Record = { // structs, and calls. MSL-specific `[[attribute]]` annotations are blanked // pre-parse for `.metal` files (see blankMetalAttributes in c-cpp.ts). (#1121) '.metal': 'cpp', + // CUDA ≈ C++ plus execution-space specifiers (`__global__` …) and + // `<<>>` kernel-launch syntax: the C++ grammar extracts its + // functions/structs/classes/calls once blankCudaConstructs (pre-parse; gated + // by these extensions OR by content for CUDA living in `.h`/`.hpp` headers — + // see c-cpp.ts) blanks the CUDA-only tokens. (#387) + '.cu': 'cpp', + '.cuh': 'cpp', // XML: file-level tracking; the MyBatis extractor matches `` // shape and emits SQL-statement nodes (other XML returns empty). '.xml': 'xml', diff --git a/src/extraction/languages/c-cpp.ts b/src/extraction/languages/c-cpp.ts index 983d1a5..f7111a7 100644 --- a/src/extraction/languages/c-cpp.ts +++ b/src/extraction/languages/c-cpp.ts @@ -27,7 +27,54 @@ function findDeclaratorQualifiedId(declarator: SyntaxNode): SyntaxNode | undefin return undefined; } +/** + * Recover the real function name from the macro-definition idiom + * `MACRO_NAME(real_name, typed args…) { body }` — flash-attention's + * `DEFINE_FLASH_FORWARD_KERNEL(flash_fwd_kernel, bool Is_dropout, …) { … }` + * being the motivating case: tree-sitter parses the invocation as a + * function_definition NAMED after the macro, so every such kernel shared one + * name (`DEFINE_FLASH_FORWARD_KERNEL`) and the launch sites' calls to the real + * names (`flash_fwd_kernel<…><<<…>>>`) could never resolve. + * + * Deliberately narrow so name-in-first-arg is unambiguous — ALL of: + * - the parsed name is macro-shaped: ALL-CAPS with at least one underscore + * (`TEST` never matches; K&R C definitions have lowercase names); + * - the first "parameter" is a LONE identifier (no type, no declarator) + * containing a lowercase letter — the name being defined; + * - at least one more parameter follows and NONE of them is another lone + * identifier — a second bare arg means the first isn't the name (gtest's + * `TEST_F(Fixture, Name)`, `PYBIND11_MODULE(ext, m)`, + * google-benchmark's `BENCHMARK_DEFINE_F(Fix, name)` all bail here). + */ +function recoverCppMacroDefinedName(node: SyntaxNode, source: string): string | undefined { + if (node.type !== 'function_definition') return undefined; + const declarator = getChildByField(node, 'declarator'); + if (declarator?.type !== 'function_declarator') return undefined; + const inner = getChildByField(declarator, 'declarator'); + if (inner?.type !== 'identifier') return undefined; + const macroName = getNodeText(inner, source); + if (!/^[A-Z][A-Z0-9]*(?:_[A-Z0-9]+)+$/.test(macroName)) return undefined; + const params = getChildByField(declarator, 'parameters'); + if (!params || params.namedChildCount < 2) return undefined; + const loneIdentText = (p: SyntaxNode): string | null => + p.type === 'parameter_declaration' && + p.namedChildCount === 1 && + p.namedChild(0)?.type === 'type_identifier' + ? getNodeText(p.namedChild(0)!, source) + : null; + const first = params.namedChild(0); + const name = first ? loneIdentText(first) : null; + if (!name || !/[a-z]/.test(name)) return undefined; + for (let i = 1; i < params.namedChildCount; i++) { + const p = params.namedChild(i); + if (p && loneIdentText(p) !== null) return undefined; + } + return name; +} + function extractCppQualifiedMethodName(node: SyntaxNode, source: string): string | undefined { + const macroDefined = recoverCppMacroDefinedName(node, source); + if (macroDefined) return macroDefined; const declarator = getChildByField(node, 'declarator'); if (!declarator) return undefined; const qid = findDeclaratorQualifiedId(declarator); @@ -123,6 +170,8 @@ function extractCppReturnType(node: SyntaxNode, source: string): string | undefi } export const cExtractor: LanguageExtractor = { + // CUDA in C-detected headers (content-gated blank; see preParseCSource). + preParse: preParseCSource, // Universal net: recover a real name from any macro-mangled function name. recoverMangledName: recoverMangledCppName, functionTypes: ['function_definition'], @@ -384,14 +433,108 @@ export function blankMetalAttributes(source: string): string { return source.replace(METAL_ATTRIBUTE_RE, (m) => ' '.repeat(m.length)); } +/** + * Blank CUDA-specific constructs before parsing `.cu`/`.cuh` files (parsed with + * the C++ grammar). Three shapes tree-sitter-cpp can't reconcile, each replaced + * with equal-length whitespace so every byte offset survives (#387): + * + * 1. Execution-space / storage specifiers: in `__global__ void step(…)` or + * `__shared__ float tile[256]` the specifier parses as the declaration's + * TYPE and shunts the real return/value type into an ERROR node — mangling + * signatures and, for `__shared__` arrays, the declared name itself. Blanked + * unconditionally (no following-token lookahead) so extended lambdas + * (`[=] __device__ (int i) { … }`) recover too. `__restrict__` is deliberately + * absent: the grammar already parses it natively as a type_qualifier. + * 2. `__launch_bounds__(…)` between specifier and declarator — same misparse. + * The parenthesized form is blanked first; a bare leftover token is caught + * by the specifier list. + * 3. Kernel-launch configs `step<<>>(args)`: the + * chevrons lex as shift operators around an empty-named template, so no + * call_expression exists and the host→kernel call edge — the main reason to + * index CUDA at all — is lost. Blanking the `<<<…>>>` span leaves + * `step (args)`, a plain call the grammar + * parses natively (templated launches `k<<<…>>>(…)` included). + * + * The launch-config match is deliberately bounded — statement/brace characters + * excluded, span capped, newlines preserved by the replacer — so a stray `<<<` + * (a committed merge-conflict marker, a string literal) can never blank a run + * of real code: an unmatched launch degrades to the status quo for that call + * site (no call edge), never to corruption. Applied to `.cu`/`.cuh` files and — + * because much real CUDA lives in extension-less headers (cutlass launches the + * majority of its kernels from `.h`; flash-attention's launch templates are + * `.h`; llm.c keeps device helpers in C-detected `.h`) — to any C/C++-family + * file whose CONTENT carries a strong CUDA marker (`looksLikeCudaSource`). + * Unlike Metal's `[[attribute]]` (legal C++ syntax elsewhere, hence Metal's + * strict extension gate), no CUDA marker is valid C++ anywhere: `<<<` isn't + * legal syntax and the dunder specifiers are implementation-reserved names no + * real codebase defines — so a content-triggered blank on a non-CUDA file can + * only ever whitespace tokens inside comments or strings, which parse the same. + */ +const CUDA_LAUNCH_BOUNDS_RE = /\b__launch_bounds__\s*\([^()\n]*\)/g; +const CUDA_SPECIFIER_RE = + /\b__(?:global|device|host|constant|shared|managed|grid_constant|forceinline|noinline|launch_bounds)__\b/g; +// `;` stays excluded (launch configs are expressions; a stray `<<<` spanning +// real statements always crosses one) and the span is capped. Braces are +// allowed through the regex — `k<<>>(…)` is a real +// launch shape — but the replacer only blanks a BALANCED match: a merge +// conflict's `<<<<<<< … >>>>>>>` region that dodges every `;` still opens +// braces it never closes, so it fails the balance check and stays untouched. +const CUDA_LAUNCH_CONFIG_RE = /<<<[^;]{0,400}?>>>/g; +export function blankCudaConstructs(source: string): string { + let out = source; + if (out.indexOf('__') !== -1) { + out = out + .replace(CUDA_LAUNCH_BOUNDS_RE, (m) => ' '.repeat(m.length)) + .replace(CUDA_SPECIFIER_RE, (m) => ' '.repeat(m.length)); + } + if (out.indexOf('<<<') !== -1) { + out = out.replace(CUDA_LAUNCH_CONFIG_RE, (m) => { + let depth = 0; + for (let i = 0; i < m.length; i++) { + const ch = m.charCodeAt(i); + if (ch === 0x7b /* { */) depth++; + else if (ch === 0x7d /* } */ && --depth < 0) return m; + } + return depth === 0 ? m.replace(/[^\n]/g, ' ') : m; + }); + } + return out; +} + +/** Strong content markers for CUDA source in files without a CUDA extension + * (headers). The dunders are execution-space specifiers that only nvcc defines; + * `cudaStream_t` is the runtime's stream handle, pervasive in launcher headers + * that themselves declare no kernel. Deliberately excludes weak markers (`dim3`, + * `<<<`) that could plausibly appear in non-CUDA text. */ +function looksLikeCudaSource(source: string): boolean { + return ( + source.indexOf('__global__') !== -1 || + source.indexOf('__device__') !== -1 || + source.indexOf('__constant__') !== -1 || + source.indexOf('cudaStream_t') !== -1 + ); +} + /** C/C++ source pre-processing before tree-sitter: recover both macro-annotated - * class definitions and macro-prefixed function definitions — plus, for `.metal` - * shaders (parsed with the C++ grammar), MSL attribute annotations. Offset-preserving. */ + * class definitions and macro-prefixed function definitions — plus the non-C++ + * surface of the dialects parsed with the C++ grammar: `.metal` MSL attribute + * annotations, and CUDA specifiers + launch syntax (by `.cu`/`.cuh` extension + * or by content, for CUDA living in `.h`/`.hpp` headers). Offset-preserving. */ function preParseCppSource(source: string, filePath?: string): string { const blanked = blankCppInlineMacros(blankCppExportMacros(source)); - return filePath && filePath.toLowerCase().endsWith('.metal') - ? blankMetalAttributes(blanked) - : blanked; + const lower = filePath ? filePath.toLowerCase() : ''; + if (lower.endsWith('.metal')) return blankMetalAttributes(blanked); + if (lower.endsWith('.cu') || lower.endsWith('.cuh') || looksLikeCudaSource(source)) { + return blankCudaConstructs(blanked); + } + return blanked; +} + +/** C source pre-processing: C-detected headers in CUDA projects (llm.c keeps + * `__device__` helpers and kernel prototypes in plain `.h`) get the same + * content-gated CUDA blank as C++. */ +function preParseCSource(source: string): string { + return looksLikeCudaSource(source) ? blankCudaConstructs(source) : source; } export const cppExtractor: LanguageExtractor = { diff --git a/src/extraction/tree-sitter.ts b/src/extraction/tree-sitter.ts index 62e66b8..ce72d13 100644 --- a/src/extraction/tree-sitter.ts +++ b/src/extraction/tree-sitter.ts @@ -382,6 +382,20 @@ export class TreeSitterExtractor { private errors: ExtractionError[] = []; private extractor: LanguageExtractor | null = null; private nodeStack: string[] = []; // Stack of parent node IDs + // C/C++ enclosing `namespace ns { … }` names, prepended to every contained + // symbol's qualifiedName (see visitNode). Prefix-only by design — no + // namespace NODE is created: `namespace cutlass {` opens in thousands of + // files, and a node per block would flood search with same-named symbols + // (the #1093 crowd-out failure mode). Always empty outside C/C++. + private namespacePrefix: string[] = []; + // C++ local function-pointer bindings, per enclosing symbol: + // `auto kernel = &flash_fwd_kernel<…>;` recorded as callerId → kernel → + // {flash_fwd_kernel}, so a later `kernel<<>>(params)` (or plain + // `kernel(args)`) in the same body emits calls refs to the real target(s) + // instead of an unresolvable local name. Branch reassignments accumulate — + // each assigned target is a genuine possible callee. Same-body locality is + // the precision guard (the #932 table-dispatch philosophy scoped to locals). + private cppLocalFnPtrs = new Map>>(); private methodIndex: Map | null = null; // lookup key → node ID for Pascal defProc lookup // Function-as-value capture (#756): per-language spec + candidates collected // during the walk, gated & flushed into unresolvedReferences at end-of-file @@ -909,6 +923,31 @@ export class TreeSitterExtractor { if (skipChildren) return; } + // C++ namespace blocks: carry the namespace name as a qualifiedName prefix + // while walking the body, so `namespace flash { void compute_attn(); }` + // indexes compute_attn with qualifiedName `flash::compute_attn` and a + // namespace-qualified call (`flash::compute_attn(...)`) resolves by exact + // qualified match instead of never resolving — C++ namespaces previously + // left no trace in qualifiedNames at all, so every `ns::fn()` call site + // was a permanently dead edge (surfaced by #387 flow validation on + // flash-attention/cutlass, whose kernel dispatch is namespace-qualified). + // C++17 nested forms (`namespace a::b {`) prefix as written. An anonymous + // namespace falls through to the generic walk — its contents stay bare, + // matching how call sites spell them. + if (this.language === 'cpp' && nodeType === 'namespace_definition') { + const nameNode = getChildByField(node, 'name'); + const nsName = nameNode ? getNodeText(nameNode, this.source) : ''; + if (nsName) { + this.namespacePrefix.push(nsName); + for (let i = 0; i < node.namedChildCount; i++) { + const child = node.namedChild(i); + if (child) this.visitNode(child); + } + this.namespacePrefix.pop(); + return; + } + } + // Function-as-value capture (#756) — independent of the dispatch ladder // below (the captured container types have no other handler there), so it // can never shadow or be shadowed by an extraction branch. @@ -1345,7 +1384,8 @@ export class TreeSitterExtractor { private buildQualifiedName(name: string): string { // Build a qualified name from the semantic hierarchy only (no file path). // The file path is stored separately in filePath and pollutes FTS if included here. - const parts: string[] = []; + // C/C++ enclosing namespaces prefix first (empty for every other language). + const parts: string[] = [...this.namespacePrefix]; for (const nodeId of this.nodeStack) { const node = this.nodes.find((n) => n.id === nodeId); if (node && node.kind !== 'file') { @@ -4183,6 +4223,44 @@ export class TreeSitterExtractor { if (conv && conv[1]) calleeName = conv[1]; } + // C/C++ templated callees — a direct templated call (`fn(args)`, + // the shape every CUDA kernel-launch site takes once its `<<<…>>>` config + // is blanked) or a qualified one (`ns::fn(args)`) — carry template + // arguments in the callee text, which can never match the bare name the + // function was DEFINED as, so the call edge silently never resolves. Strip + // them: the same normalization base-class `extends` refs already get + // (#1043). `operator<`/`operator<<` callees are excluded — their `<` is the + // operator itself, not a template-argument list. + if ( + calleeName && + calleeName.includes('<') && + (this.language === 'cpp' || this.language === 'c') && + !calleeName.includes('operator') + ) { + calleeName = stripCppTemplateArgs(calleeName); + } + + // C++ call/launch through a local function pointer: `auto kernel = + // &flash_fwd_kernel<…>; … kernel<<>>(params);` — the callee + // is an unresolvable local name. When the same enclosing symbol bound the + // local from `&fn` (each branch assignment counts), emit the call against + // every recorded target instead of the local. + if (calleeName && this.language === 'cpp' && /^[A-Za-z_]\w*$/.test(calleeName)) { + const targets = this.cppLocalFnPtrs.get(callerId)?.get(calleeName); + if (targets && targets.size > 0) { + for (const target of targets) { + this.unresolvedReferences.push({ + fromNodeId: callerId, + referenceName: target, + referenceKind: 'calls', + line: node.startPosition.row + 1, + column: node.startPosition.column, + }); + } + return; + } + } + if (calleeName) { this.unresolvedReferences.push({ fromNodeId: callerId, @@ -4704,6 +4782,42 @@ export class TreeSitterExtractor { flush(); } + /** + * Record a C++ local function-pointer binding (`local = &fn` / `&fn<…>` / + * `&ns::fn<…>`) for the CURRENT enclosing symbol, so calls through the local + * resolve to the real target (see cppLocalFnPtrs). Only the address-of shape + * is accepted — a bare-identifier RHS (`auto x = y;`) is any value copy, and + * linking through it would guess. + */ + private recordCppFnPtrBinding(localName: string, value: SyntaxNode | null): void { + if (!value || value.type !== 'pointer_expression') return; + if (value.child(0)?.type !== '&') return; // `*p` dereference, not address-of + const arg = getChildByField(value, 'argument') ?? value.namedChild(0); + if ( + !arg || + (arg.type !== 'identifier' && + arg.type !== 'template_function' && + arg.type !== 'qualified_identifier') + ) { + return; + } + const callerId = this.nodeStack[this.nodeStack.length - 1]; + if (!callerId) return; + const target = stripCppTemplateArgs(getNodeText(arg, this.source)); + if (!target || target === localName) return; + let locals = this.cppLocalFnPtrs.get(callerId); + if (!locals) { + locals = new Map(); + this.cppLocalFnPtrs.set(callerId, locals); + } + let targets = locals.get(localName); + if (!targets) { + targets = new Set(); + locals.set(localName, targets); + } + targets.add(target); + } + private visitFunctionBody(body: SyntaxNode, _functionId: string): void { if (!this.extractor) return; @@ -4763,6 +4877,35 @@ export class TreeSitterExtractor { this.extractInstantiation(node); } + // C++ local function-pointer bindings (see cppLocalFnPtrs): record + // `auto kernel = &fn<…>;` declarations and `kernel = &other_fn<…>;` + // branch reassignments so a call/launch through the local links to the + // real target(s). The body walker sees these in source order, and C++ + // requires declaration-before-use, so the map is always populated before + // the call that consumes it. + if (this.language === 'cpp' && this.nodeStack.length > 0) { + if (nodeType === 'declaration') { + for (let i = 0; i < node.namedChildCount; i++) { + const child = node.namedChild(i); + if (child?.type !== 'init_declarator') continue; + const decl = getChildByField(child, 'declarator'); + if (decl?.type !== 'identifier') continue; + this.recordCppFnPtrBinding( + getNodeText(decl, this.source), + getChildByField(child, 'value') + ); + } + } else if (nodeType === 'assignment_expression') { + const left = getChildByField(node, 'left'); + if (left?.type === 'identifier') { + this.recordCppFnPtrBinding( + getNodeText(left, this.source), + getChildByField(node, 'right') + ); + } + } + } + // Static-member / value-read: `Enum.value`, `Type.CONST`, `Foo::BAR`. this.extractStaticMemberRef(node);