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 6a1b21e..a4b6516 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..16ec76a 100644 --- a/builtin/runtime/lsp.lua +++ b/builtin/runtime/lsp.lua @@ -82,6 +82,29 @@ 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 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 -- pulls its configuration via `workspace/configuration` (now -- answered, #13) under the `gopls` section; an empty section means @@ -169,6 +192,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..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,43 @@ 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__` + // 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. + // + // 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_c::HIGHLIGHT_QUERY, + tree_sitter_cpp::HIGHLIGHT_QUERY, + tree_sitter_cuda::HIGHLIGHTS_QUERY, + ], }, ]; @@ -601,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(); @@ -799,6 +843,116 @@ 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, + // 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") + .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 + .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() + ); + } + + #[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 eb05196..50b3162 100644 --- a/tests/m4_acceptance.rs +++ b/tests/m4_acceptance.rs @@ -5747,6 +5747,63 @@ 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 = {} + 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. + 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::("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!( + 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