CUDA rides the C++ grammar via the Metal (#1121) dialect pattern:
blankCudaConstructs (offset-preserving) blanks execution-space specifiers
(__global__ family), __launch_bounds__(...), and <<<grid, block>>> launch
configs — which otherwise lex as shift operators and destroy the
host→kernel call edge entirely. Gated by .cu/.cuh extension OR by content
(looksLikeCudaSource), because much real CUDA lives in .h/.hpp headers:
cutlass launches most kernels from headers and flash-attention's launch
templates are .h. Safe by construction — no CUDA marker is valid C++
anywhere, and the launch blank is bounded + brace-balance-checked so a
stray <<< (committed merge-conflict markers) can never blank real code.
All real-world launch styles connect: plain, templated
(k<T, 256><<<...>>>), function-pointer (auto kernel = &fn<...>; with
branch reassignments each linked), dim3{...} brace-init configs, and
kernels defined through name-in-first-argument macros
(DEFINE_FLASH_FORWARD_KERNEL style — gtest TEST_F / PYBIND11_MODULE
shapes deliberately excluded by the two-lone-identifiers rule).
Two general C++ resolution wins the flow validation forced out:
- namespace blocks now prefix contained symbols' qualifiedNames
(prefix-only — no namespace nodes, avoiding #1093-style crowd-out), so
ns::fn(...) calls resolve; previously every namespace-qualified C++
call was a permanently dead edge. cutlass: +30,864 edges (~10%), node
count byte-identical.
- templated callees (fn<T, 256>(args)) strip template args at extraction
(mirroring #1043 for base classes), so they match their definitions.
Validated on llm.c (165 host→kernel launch edges, was 0),
flash-attention (run_flash_fwd → flash_fwd_kernel → compute_attn traces
in one codegraph_explore call), and NVIDIA CUTLASS; fmt as the plain-C++
control (unchanged). A/B n=2/arm: Read/Grep displacement decisive on all
three repos (flash-attention Reads 29,13 → 5,2).
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Colby McHenry
fd34f6c870