From 11075914f3dac2dae9b68336032715194e22dccf Mon Sep 17 00:00:00 2001 From: Levi Neuwirth Date: Tue, 14 Jul 2026 10:58:36 +0100 Subject: [PATCH 1/2] =?UTF-8?q?feat(lsp):=20CUDA=20support=20=E2=80=94=20c?= =?UTF-8?q?langd=20+=20bundled=20tree-sitter=20grammar?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Opening a .cu/.cuh file previously resolved to no language, so no server attached and there was no highlighting. Wire CUDA end to end, mirroring the existing C/C++ path: - Bundle tree-sitter-cuda (0.21) as a new BUILTIN_LANGUAGES entry claiming .cu/.cuh, with its own HIGHLIGHTS_QUERY. A dedicated grammar rather than reusing cpp: the C++ grammar errors on the <<>> kernel-launch syntax. The crate rides tree-sitter-language 0.1 (its tree-sitter dep is dev-only), so it shares the ABI crate with the other grammars — no second tree-sitter in the graph. - pmacs.lsp.config.cuda targets clangd (the same binary that serves C/C++; language_id "cuda" so clangd enters its CUDA parse mode), and .cu/.cuh filetype fallbacks map to "cuda" to keep the LSP id stable if the grammar is ever dropped. LspStyleView layers clangd's CUDA semantic tokens on top, exactly as for C/C++. Bite-verified acceptance: - cuda_grammar_loads_and_parses_kernel_launch — proves the 0.21 grammar's ABI is accepted by the 0.26 core (set_language succeeds at runtime, which the compile step cannot confirm) and that the entry wired the CUDA grammar, not a cpp fallback: the <<<...>>> launch parses without error, whereas the cpp grammar reports an error on the same source (verified out of band). - builtin_languages_include_cuda / language_for_path_resolves_cuda_ extensions — entry presence and .cu/.cuh detection. - m4_12_default_bundle_wires_cuda — config.cuda targets clangd and the filetype + grammar detection resolve to "cuda" through the loaded runtime. Co-Authored-By: Claude Opus 4.8 (1M context) Claude-Session: https://claude.ai/code/session_01YJ9FQ832QwftJXCD9LeFan --- Cargo.lock | 11 +++++ Cargo.toml | 10 +++++ builtin/runtime/lsp.lua | 22 ++++++++++ src/syntax.rs | 89 +++++++++++++++++++++++++++++++++++++++++ tests/m4_acceptance.rs | 44 ++++++++++++++++++++ 5 files changed, 176 insertions(+) diff --git a/Cargo.lock b/Cargo.lock index 7faecf7..c5f9f6f 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -2552,6 +2552,7 @@ dependencies = [ "tree-sitter", "tree-sitter-c", "tree-sitter-cpp", + "tree-sitter-cuda", "tree-sitter-lua", "tree-sitter-md", "tree-sitter-rust", @@ -3735,6 +3736,16 @@ dependencies = [ "tree-sitter-language", ] +[[package]] +name = "tree-sitter-cuda" +version = "0.21.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "715eecfee69b15991de5b9f78009c6d4cb34e18d20d028304a75d38528cddb45" +dependencies = [ + "cc", + "tree-sitter-language", +] + [[package]] name = "tree-sitter-language" version = "0.1.7" diff --git a/Cargo.toml b/Cargo.toml index 51b8b56..f469664 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -144,6 +144,16 @@ tree-sitter-lua = "0.5" # matching the LSP filetype map in `builtin/runtime/lsp.lua`. tree-sitter-c = "0.24" tree-sitter-cpp = "0.23" +# CUDA (`.cu`/`.cuh`). CUDA is C++ with device extensions, but its own +# grammar recognizes `__global__`/`__device__` kernels, `<<<...>>>` +# launch syntax, and CUDA builtins that the C++ grammar would misparse. +# Loading is lazy like the others; `LspStyleView` layers clangd's CUDA +# semantic tokens on top (clangd serves `.cu`/`.cuh` off the same +# binary as C/C++). The crate exports `LANGUAGE`/`HIGHLIGHTS_QUERY` +# (plural, the rust/lua idiom) and rides `tree-sitter-language 0.1`, so +# it shares the ABI crate with the other grammars — no second +# `tree-sitter` in the graph (its own `tree-sitter` dep is dev-only). +tree-sitter-cuda = "0.21" # T M9.7: markdown grammar so prompt result buffers with # `_meta.format = "markdown"` get structured highlighting through the # same M4 path as rust/lua — no special-case painter in Lua. diff --git a/builtin/runtime/lsp.lua b/builtin/runtime/lsp.lua index 928ee37..ec58e40 100644 --- a/builtin/runtime/lsp.lua +++ b/builtin/runtime/lsp.lua @@ -82,6 +82,21 @@ pmacs.lsp.config.cpp = pmacs.lsp.config.cpp or { args = { "--background-index" }, } +-- CUDA (`.cu`/`.cuh`) via clangd — the same binary serves it; `config.cuda` +-- is a separate entry only so the `language_id` sent in `didOpen` is +-- `cuda` (clangd keys its CUDA parse mode off both the id and the `.cu` +-- extension). Like C/C++, clangd takes the project model from +-- `compile_commands.json` / `compile_flags.txt`, so no `settings` here. +-- For real analysis clangd must also locate a CUDA toolkit: it probes +-- common install roots (e.g. `/usr/local/cuda`), and a project can pin +-- `--cuda-path=` / the GPU arch through its compile flags; absent those, +-- navigation and hover still work but diagnostics may be noisy. Users +-- override from init.lua before a CUDA file opens. +pmacs.lsp.config.cuda = pmacs.lsp.config.cuda or { + command = "clangd", + args = { "--background-index" }, +} + -- Go via gopls. `gopls` with no args serves LSP over stdio. gopls -- pulls its configuration via `workspace/configuration` (now -- answered, #13) under the `gopls` section; an empty section means @@ -169,6 +184,13 @@ pmacs.lsp.filetypes.h = pmacs.lsp.filetypes.h or "c" for _, ext in ipairs({ "cpp", "cc", "cxx", "hpp", "hh", "hxx", "ipp", "inl", "cppm" }) do pmacs.lsp.filetypes[ext] = pmacs.lsp.filetypes[ext] or "cpp" end +-- CUDA. pmacs bundles a CUDA grammar, so `language_for_path` already +-- resolves `.cu`/`.cuh` to `cuda` and this map is never consulted for +-- them in practice; the entries are the LSP-only fallback that keeps +-- the language id stable if that grammar is ever dropped (same role as +-- the `lua` entry below). +pmacs.lsp.filetypes.cu = pmacs.lsp.filetypes.cu or "cuda" +pmacs.lsp.filetypes.cuh = pmacs.lsp.filetypes.cuh or "cuda" -- Go. pmacs.lsp.filetypes.go = pmacs.lsp.filetypes.go or "go" -- Tier 1 single-binary servers. TypeScript / JavaScript distinguish diff --git a/src/syntax.rs b/src/syntax.rs index 70106c5..2470c02 100644 --- a/src/syntax.rs +++ b/src/syntax.rs @@ -404,6 +404,22 @@ pub const BUILTIN_LANGUAGES: &[LanguageEntry] = &[ loader: || tree_sitter_cpp::LANGUAGE.into(), highlights_query: tree_sitter_cpp::HIGHLIGHT_QUERY, }, + // CUDA (`.cu` source, `.cuh` header). A dedicated grammar rather + // than reusing `cpp`: CUDA extends C++ with `__global__`/`__device__` + // qualifiers, `<<>>` kernel-launch syntax, and builtin + // types the C++ grammar misparses. Neither extension collides with + // an entry above, so ordering is irrelevant here. `LspStyleView` + // layers clangd's CUDA semantic tokens on top, exactly as for C/C++. + // + // Note the const name: `tree-sitter-cuda` exposes `HIGHLIGHTS_QUERY` + // (plural, the `tree-sitter-rust`/`tree-sitter-lua` idiom), NOT the + // singular `HIGHLIGHT_QUERY` that `tree-sitter-c`/`-cpp`/`-md` use. + LanguageEntry { + name: "cuda", + extensions: &["cu", "cuh"], + loader: || tree_sitter_cuda::LANGUAGE.into(), + highlights_query: tree_sitter_cuda::HIGHLIGHTS_QUERY, + }, ]; /// Registry that the Lua surface ([`crate::lua_bindings::install_parse`]) @@ -799,6 +815,79 @@ mod tests { ); } + #[test] + fn builtin_languages_include_cuda() { + // Regression guard mirroring `builtin_languages_include_c_and_cpp`: + // the CUDA entry must keep claiming its canonical extensions and + // shipping a highlights query, so `.cu`/`.cuh` get lexical + // styling (with clangd's semantic tokens layered on top) instead + // of falling back to no grammar at all. + let cuda = BUILTIN_LANGUAGES + .iter() + .find(|l| l.name == "cuda") + .expect("`cuda` language entry must be present"); + assert!(cuda.extensions.contains(&"cu"), "`cuda` claims `.cu`"); + assert!(cuda.extensions.contains(&"cuh"), "`cuda` claims `.cuh`"); + assert!( + !cuda.highlights_query.is_empty(), + "`cuda` ships a non-empty highlights query" + ); + } + + #[test] + fn cuda_grammar_loads_and_parses_kernel_launch() { + // ABI acceptance: a `tree-sitter-cuda` 0.21 grammar must be + // accepted by our `tree-sitter` 0.26 core — `set_language` + // succeeds and a tree is produced. This is the runtime check the + // compile step cannot give us (a too-old grammar ABI fails only + // here, at parse time). Grammar identity: `<<>>` + // kernel-launch syntax is CUDA-specific; the C++ grammar parses + // it as chained comparison/shift operators and flags an error, so + // an error-free parse proves the entry wired the CUDA grammar, + // not a C++ fallback. + let reg = SyntaxRegistry::new(); + let language = reg + .language("cuda") + .expect("`cuda` language loads from BUILTIN_LANGUAGES"); + let mut buf = fresh_buffer("kernel.cu"); + buf.apply_edit(EditOp::Insert { + pos: 0, + bytes: b"__global__ void add(int *c) { c[threadIdx.x] = 1; }\n\ + int main() { add<<<1, 256>>>(0); return 0; }\n", + }) + .unwrap(); + let view = ParseView::new(&buf, language, "cuda".to_owned()); + let handle = view.handle(); + let _vid = buf.attach_view(Box::new(view)); + let bundle = parse_synchronously(&handle); + assert_eq!( + bundle.tree.root_node().kind(), + "translation_unit", + "CUDA grammar (C-derived) roots at translation_unit" + ); + assert!( + !bundle.tree.root_node().has_error(), + "CUDA grammar parses the `<<<...>>>` kernel launch without error" + ); + } + + #[test] + fn language_for_path_resolves_cuda_extensions() { + // `.cu`/`.cuh` resolve to the CUDA grammar through the same + // extension-detection path as every other bundled language, so + // the LSP filetype fallback in `lsp.lua` is never consulted for + // them in practice. + let reg = SyntaxRegistry::new(); + assert_eq!( + reg.language_name_for_path("kernel.cu").as_deref(), + Some("cuda") + ); + assert_eq!( + reg.language_name_for_path("device.cuh").as_deref(), + Some("cuda") + ); + } + #[test] fn byte_to_point_handles_first_line() { let src = b"hello world"; diff --git a/tests/m4_acceptance.rs b/tests/m4_acceptance.rs index 70d674a..5990ca6 100644 --- a/tests/m4_acceptance.rs +++ b/tests/m4_acceptance.rs @@ -5742,6 +5742,50 @@ fn m4_12_default_bundle_wires_commands_and_keymaps() { assert!(probe.get::("cmd_sig").unwrap()); } +/// The default LSP bundle wires CUDA: `pmacs.lsp.config.cuda` targets +/// clangd (the same binary that serves C/C++), and the `.cu`/`.cuh` +/// filetype fallbacks map to `cuda`. Because pmacs also bundles a CUDA +/// tree-sitter grammar, `pmacs.parse.language_for_path` resolves those +/// extensions to `cuda` directly — so the fallback map is +/// belt-and-suspenders, but is asserted here to keep the LSP language +/// id stable if the grammar is ever dropped. +#[test] +fn m4_12_default_bundle_wires_cuda() { + use pmacs::editor::EditorState; + + let s = EditorState::new(); + let probe: mlua::Table = s + .lua_host + .lua() + .load( + r" + local out = {} + out.cfg_cmd = pmacs.lsp.config.cuda and pmacs.lsp.config.cuda.command + out.ft_cu = pmacs.lsp.filetypes.cu + out.ft_cuh = pmacs.lsp.filetypes.cuh + -- Grammar-backed detection (bundled CUDA grammar) wins first. + out.grammar_cu = pmacs.parse.language_for_path('kernel.cu') + out.grammar_cuh = pmacs.parse.language_for_path('device.cuh') + return out + ", + ) + .eval() + .expect("probe cuda wiring"); + assert_eq!( + probe.get::("cfg_cmd").unwrap(), + "clangd", + "config.cuda targets clangd" + ); + assert_eq!(probe.get::("ft_cu").unwrap(), "cuda"); + assert_eq!(probe.get::("ft_cuh").unwrap(), "cuda"); + assert_eq!( + probe.get::("grammar_cu").unwrap(), + "cuda", + "bundled grammar resolves `.cu` to cuda" + ); + assert_eq!(probe.get::("grammar_cuh").unwrap(), "cuda"); +} + /// Typing-perf: the default bundle coalesces full-document /// `didChange` notifications instead of sending one per keystroke /// (each send copies the whole buffer several times and writes From ea3641bba2a9d69ee7b3d2c9f6559d5d482c3e01 Mon Sep 17 00:00:00 2001 From: Levi Neuwirth Date: Tue, 14 Jul 2026 11:55:58 +0100 Subject: [PATCH 2/2] =?UTF-8?q?fix(lsp):=20PR=20#114=20round=201=20?= =?UTF-8?q?=E2=80=94=20.cuh=20AST=20via=20fallbackFlags,=20real=20C/C++=20?= =?UTF-8?q?highlights?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Two functional gaps from review: 1. Standalone .cuh files got no clangd AST. clangd selects the compiler language from the file extension, not the LSP languageId: it knows .cu (-> -x cuda) but not .cuh, so a header with no compile command fails with fe_expected_compiler_job. config.cuda now sets init_options.fallbackFlags = { "-xcuda" }, which supplies -x cuda for any file this server opens that lacks a compile_commands.json entry (a real compile command still wins). This CUDA server only ever serves .cu/.cuh, so the fallback cannot mis-flag C/C++. 2. The CUDA highlights query was only a delta. tree-sitter-cuda's HIGHLIGHTS_QUERY opens with `; inherits: cpp` and defines only the CUDA-specific captures (launch brackets, __global__/__device__) — two capture classes. pmacs does not resolve `inherits:`, so ordinary C/C++ syntax went unhighlighted. LanguageEntry.highlights_query is now &[&str] (fragments joined base-first); the cuda entry carries [c, cpp, cuda], compiling to ~16 capture classes. Fragments are newline-joined, never bare-concatenated — a fragment can end mid `; comment`, and abutting the next fragment's first token would corrupt the query. Existing single-query grammars become one-element slices (byte-identical effective query; no behavior change). Tests: - cuda_highlights_resolve_c_and_cpp_captures — asserts the COMPILED cuda query carries the C base `@variable` capture and >= 8 capture classes, not merely a non-empty query (the CUDA delta alone has 2 and no `variable`, so this fails without the base prepend). - builtin_languages_include_cuda — now asserts the entry composes the c + cpp + cuda fragments. - m4_12_default_bundle_wires_cuda — now asserts config.cuda.init_options.fallbackFlags[1] == "-xcuda". Gates green: fmt; clippy -D warnings; test --lib (1515); --features crdt (1689); m4_acceptance --skip basedpyright (101); GPU (59); full workspace sweep; git diff --check. Co-Authored-By: Claude Opus 4.8 (1M context) Claude-Session: https://claude.ai/code/session_01YJ9FQ832QwftJXCD9LeFan --- builtin/runtime/lsp.lua | 26 +++++---- src/syntax.rs | 113 +++++++++++++++++++++++++++++++--------- tests/m4_acceptance.rs | 15 +++++- 3 files changed, 120 insertions(+), 34 deletions(-) diff --git a/builtin/runtime/lsp.lua b/builtin/runtime/lsp.lua index ec58e40..16ec76a 100644 --- a/builtin/runtime/lsp.lua +++ b/builtin/runtime/lsp.lua @@ -83,18 +83,26 @@ pmacs.lsp.config.cpp = pmacs.lsp.config.cpp or { } -- CUDA (`.cu`/`.cuh`) via clangd — the same binary serves it; `config.cuda` --- is a separate entry only so the `language_id` sent in `didOpen` is --- `cuda` (clangd keys its CUDA parse mode off both the id and the `.cu` --- extension). Like C/C++, clangd takes the project model from --- `compile_commands.json` / `compile_flags.txt`, so no `settings` here. --- For real analysis clangd must also locate a CUDA toolkit: it probes --- common install roots (e.g. `/usr/local/cuda`), and a project can pin --- `--cuda-path=` / the GPU arch through its compile flags; absent those, --- navigation and hover still work but diagnostics may be noisy. Users --- override from init.lua before a CUDA file opens. +-- is a separate entry so the `language_id` sent in `didOpen` is `cuda`. +-- clangd, however, picks the compiler *language* from the file extension, +-- NOT that `language_id`: it recognizes `.cu` (→ `-x cuda`) but not `.cuh`, +-- so a standalone header with no compile command otherwise fails to build +-- an AST (`fe_expected_compiler_job`). `initializationOptions.fallbackFlags +-- = { "-xcuda" }` supplies `-x cuda` for any file this server opens that +-- lacks a `compile_commands.json` entry, which fixes `.cuh` (and bare `.cu`) +-- headers; a real compile command still wins where present. This server +-- only ever serves `.cu`/`.cuh`, so the fallback can't mis-flag C/C++. +-- Like C/C++, clangd takes the project model from `compile_commands.json` +-- / `compile_flags.txt`, so no `settings` here. For real analysis clangd +-- must also locate a CUDA toolkit: it probes common install roots (e.g. +-- `/usr/local/cuda`), and a project can pin `--cuda-path=` / the GPU arch +-- through its compile flags; absent those, navigation and hover still work +-- but diagnostics may be noisy. Users override from init.lua before a CUDA +-- file opens. pmacs.lsp.config.cuda = pmacs.lsp.config.cuda or { command = "clangd", args = { "--background-index" }, + init_options = { fallbackFlags = { "-xcuda" } }, } -- Go via gopls. `gopls` with no args serves LSP over stdio. gopls diff --git a/src/syntax.rs b/src/syntax.rs index 2470c02..bd7e3b1 100644 --- a/src/syntax.rs +++ b/src/syntax.rs @@ -312,9 +312,9 @@ impl ParseViewHandle { /// [`tree_sitter::Language`] so the C-side grammar object isn't /// touched until the first buffer of that language is opened --- /// "load grammar lazily" per the M4.2 acceptance criterion. The -/// [`Self::highlights_query`] string is `include_str!`'d at compile -/// time so it ships in the binary; T M4.3 compiles it into a -/// [`tree_sitter::Query`] on first highlight attach. +/// [`Self::highlights_query`] fragments ship as `&'static str` +/// constants in the binary; T M4.3 concatenates and compiles them +/// into a [`tree_sitter::Query`] on first highlight attach. pub struct LanguageEntry { /// Canonical language name. Used by [`SyntaxRegistry::language`] /// lookups, surfaced through Lua as the grammar label. @@ -328,10 +328,17 @@ pub struct LanguageEntry { /// once per registry lifetime --- the result is cached under /// `name` after the first invocation. pub loader: fn() -> tree_sitter::Language, - /// Source of the bundled `highlights.scm` query (T M4.3). Empty - /// string means the grammar has no highlight query (rare; in - /// that case the highlight view runs but emits nothing). - pub highlights_query: &'static str, + /// Bundled `highlights.scm` query fragments (T M4.3), concatenated + /// in order (base grammar first) to form the effective query. Most + /// grammars ship one self-contained fragment. A grammar whose + /// bundled query is a tree-sitter `; inherits: ` delta lists + /// the inherited base queries ahead of its own, because pmacs does + /// not resolve `inherits:` directives — CUDA, for instance, ships a + /// two-capture delta over C++ and must carry the C and C++ queries + /// explicitly or ordinary C/C++ syntax goes unhighlighted. An empty + /// slice (or all-empty fragments) means no highlights: the view + /// runs but emits nothing. + pub highlights_query: &'static [&'static str], } /// Bundled grammars (T M4.2 + M4.3). The order is significant only @@ -340,8 +347,10 @@ pub struct LanguageEntry { /// /// Adding a grammar: /// 1. Add `tree-sitter-foo = "X.Y"` to `Cargo.toml`. -/// 2. Add one [`LanguageEntry`] here, including -/// `tree_sitter_foo::HIGHLIGHTS_QUERY`. +/// 2. Add one [`LanguageEntry`] here, with +/// `highlights_query: &[tree_sitter_foo::HIGHLIGHTS_QUERY]` (or the +/// inherited base queries ahead of it, if `foo`'s bundled query is +/// a `; inherits:` delta — see the `cuda` entry). /// 3. (Done.) The Lua side picks up the new grammar through the /// `buffer.after-load` hook automatically and the highlight /// overlay attaches in the same step. @@ -350,13 +359,13 @@ pub const BUILTIN_LANGUAGES: &[LanguageEntry] = &[ name: "rust", extensions: &["rs"], loader: || tree_sitter_rust::LANGUAGE.into(), - highlights_query: tree_sitter_rust::HIGHLIGHTS_QUERY, + highlights_query: &[tree_sitter_rust::HIGHLIGHTS_QUERY], }, LanguageEntry { name: "lua", extensions: &["lua"], loader: || tree_sitter_lua::LANGUAGE.into(), - highlights_query: tree_sitter_lua::HIGHLIGHTS_QUERY, + highlights_query: &[tree_sitter_lua::HIGHLIGHTS_QUERY], }, // T M9.7: markdown grammar for prompt-result buffers with // `_meta.format = "markdown"`. Uses only the block grammar @@ -375,7 +384,7 @@ pub const BUILTIN_LANGUAGES: &[LanguageEntry] = &[ name: "markdown", extensions: &["md", "markdown"], loader: || tree_sitter_md::LANGUAGE.into(), - highlights_query: tree_sitter_md::HIGHLIGHT_QUERY_BLOCK, + highlights_query: &[tree_sitter_md::HIGHLIGHT_QUERY_BLOCK], }, // T M_B3 — C / C++. Lexical highlighting (keywords / strings / // operators) so the grid TUI shows code-shaped C++ on first open. @@ -396,13 +405,13 @@ pub const BUILTIN_LANGUAGES: &[LanguageEntry] = &[ name: "c", extensions: &["c", "h"], loader: || tree_sitter_c::LANGUAGE.into(), - highlights_query: tree_sitter_c::HIGHLIGHT_QUERY, + highlights_query: &[tree_sitter_c::HIGHLIGHT_QUERY], }, LanguageEntry { name: "cpp", extensions: &["cpp", "cc", "cxx", "hpp", "hh", "hxx", "ipp", "inl", "cppm"], loader: || tree_sitter_cpp::LANGUAGE.into(), - highlights_query: tree_sitter_cpp::HIGHLIGHT_QUERY, + highlights_query: &[tree_sitter_cpp::HIGHLIGHT_QUERY], }, // CUDA (`.cu` source, `.cuh` header). A dedicated grammar rather // than reusing `cpp`: CUDA extends C++ with `__global__`/`__device__` @@ -414,11 +423,25 @@ pub const BUILTIN_LANGUAGES: &[LanguageEntry] = &[ // Note the const name: `tree-sitter-cuda` exposes `HIGHLIGHTS_QUERY` // (plural, the `tree-sitter-rust`/`tree-sitter-lua` idiom), NOT the // singular `HIGHLIGHT_QUERY` that `tree-sitter-c`/`-cpp`/`-md` use. + // + // The CUDA `highlights.scm` opens with `; inherits: cpp` and defines + // only the CUDA-specific captures (`<<<...>>>` launch brackets, the + // `__global__`/`__device__` modifiers) — two capture classes on its + // own. pmacs does not resolve `inherits:`, so the C and C++ base + // queries are prepended explicitly; the three compile together into + // ~16 capture classes against the CUDA grammar (which is a superset + // of C++). Order is base-first (C, then C++, then CUDA) so later + // fragments refine earlier ones. Without this, ordinary C/C++ syntax + // in a `.cu` file would go almost entirely unhighlighted. LanguageEntry { name: "cuda", extensions: &["cu", "cuh"], loader: || tree_sitter_cuda::LANGUAGE.into(), - highlights_query: tree_sitter_cuda::HIGHLIGHTS_QUERY, + highlights_query: &[ + tree_sitter_c::HIGHLIGHT_QUERY, + tree_sitter_cpp::HIGHLIGHT_QUERY, + tree_sitter_cuda::HIGHLIGHTS_QUERY, + ], }, ]; @@ -617,14 +640,19 @@ impl SyntaxRegistry { } let language = self.language(lang_name)?; let entry = BUILTIN_LANGUAGES.iter().find(|e| e.name == lang_name); - let source = entry.map_or("", |e| e.highlights_query); - if source.is_empty() { + // Fragments are joined with a newline, never bare-concatenated: a + // fragment can end mid-`; comment` or without a trailing newline, + // and abutting it against the next fragment's first token would + // corrupt the query (e.g. `@variable; Functions` swallows the + // next line into a comment). + let source = entry.map_or_else(String::new, |e| e.highlights_query.join("\n")); + if source.trim().is_empty() { self.queries .borrow_mut() .insert(lang_name.to_owned(), Err("no highlights query".to_owned())); return None; } - let compiled = tree_sitter::Query::new(&language, source) + let compiled = tree_sitter::Query::new(&language, &source) .map(Arc::new) .map_err(|e| format!("compile {lang_name} highlights: {e:?}")); let result = compiled.as_ref().ok().cloned(); @@ -818,10 +846,10 @@ mod tests { #[test] fn builtin_languages_include_cuda() { // Regression guard mirroring `builtin_languages_include_c_and_cpp`: - // the CUDA entry must keep claiming its canonical extensions and - // shipping a highlights query, so `.cu`/`.cuh` get lexical - // styling (with clangd's semantic tokens layered on top) instead - // of falling back to no grammar at all. + // the CUDA entry must keep claiming its canonical extensions and, + // because its bundled query is a `; inherits: cpp` delta, prepend + // the C and C++ base queries explicitly (see the entry comment and + // `cuda_highlights_resolve_c_and_cpp_captures`). let cuda = BUILTIN_LANGUAGES .iter() .find(|l| l.name == "cuda") @@ -829,8 +857,45 @@ mod tests { assert!(cuda.extensions.contains(&"cu"), "`cuda` claims `.cu`"); assert!(cuda.extensions.contains(&"cuh"), "`cuda` claims `.cuh`"); assert!( - !cuda.highlights_query.is_empty(), - "`cuda` ships a non-empty highlights query" + cuda.highlights_query + .contains(&tree_sitter_c::HIGHLIGHT_QUERY), + "`cuda` prepends the C base highlights (it does not resolve `inherits:`)" + ); + assert!( + cuda.highlights_query + .contains(&tree_sitter_cpp::HIGHLIGHT_QUERY), + "`cuda` prepends the C++ base highlights" + ); + assert!( + cuda.highlights_query + .contains(&tree_sitter_cuda::HIGHLIGHTS_QUERY), + "`cuda` carries its own CUDA-specific highlights delta" + ); + } + + #[test] + fn cuda_highlights_resolve_c_and_cpp_captures() { + // Finding-2 regression: the bundled CUDA `highlights.scm` is only + // a two-capture `; inherits: cpp` delta (launch brackets + CUDA + // modifiers). pmacs does not resolve `inherits:`, so the entry + // prepends the C and C++ base queries; assert the COMPILED query + // actually carries ordinary C/C++ captures (the C base's + // `@variable`) and far more than the delta's two capture classes — + // not merely that some query is non-empty. + let reg = SyntaxRegistry::new(); + let query = reg + .highlights_query("cuda") + .expect("cuda highlights compile"); + let names = query.capture_names(); + assert!( + names.contains(&"variable"), + "combined query carries the C base `@variable` capture; got {names:?}" + ); + assert!( + names.len() >= 8, + "combined C+C+++CUDA query resolves many capture classes, not the \ + CUDA delta's two; got {} ({names:?})", + names.len() ); } diff --git a/tests/m4_acceptance.rs b/tests/m4_acceptance.rs index 5990ca6..2ba2fd9 100644 --- a/tests/m4_acceptance.rs +++ b/tests/m4_acceptance.rs @@ -5760,7 +5760,15 @@ fn m4_12_default_bundle_wires_cuda() { .load( r" local out = {} - out.cfg_cmd = pmacs.lsp.config.cuda and pmacs.lsp.config.cuda.command + local cfg = pmacs.lsp.config.cuda + out.cfg_cmd = cfg and cfg.command + -- `.cuh` (and bare `.cu`) headers are not recognized as CUDA + -- by clangd's extension-based language selection, so the + -- server must pass `-x cuda` via fallbackFlags for files with + -- no compile command. + out.fallback = cfg and cfg.init_options + and cfg.init_options.fallbackFlags + and cfg.init_options.fallbackFlags[1] out.ft_cu = pmacs.lsp.filetypes.cu out.ft_cuh = pmacs.lsp.filetypes.cuh -- Grammar-backed detection (bundled CUDA grammar) wins first. @@ -5776,6 +5784,11 @@ fn m4_12_default_bundle_wires_cuda() { "clangd", "config.cuda targets clangd" ); + assert_eq!( + probe.get::("fallback").unwrap(), + "-xcuda", + "config.cuda forces `-x cuda` so standalone `.cuh`/`.cu` headers get an AST" + ); assert_eq!(probe.get::("ft_cu").unwrap(), "cuda"); assert_eq!(probe.get::("ft_cuh").unwrap(), "cuda"); assert_eq!(