From 11075914f3dac2dae9b68336032715194e22dccf Mon Sep 17 00:00:00 2001 From: Levi Neuwirth Date: Tue, 14 Jul 2026 10:58:36 +0100 Subject: [PATCH] =?UTF-8?q?feat(lsp):=20CUDA=20support=20=E2=80=94=20clang?= =?UTF-8?q?d=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