feat(lsp): CUDA support — clangd + bundled tree-sitter grammar

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
  <<<grid, block>>> 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) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01YJ9FQ832QwftJXCD9LeFan

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"

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.

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

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, `<<<grid, block>>>` 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: `<<<grid, block>>>`
+        // 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";

tests/m4_acceptance.rs

@@ -5742,6 +5742,50 @@ fn m4_12_default_bundle_wires_commands_and_keymaps() {
     assert!(probe.get::<bool>("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::<String>("cfg_cmd").unwrap(),
+        "clangd",
+        "config.cuda targets clangd"
+    );
+    assert_eq!(probe.get::<String>("ft_cu").unwrap(), "cuda");
+    assert_eq!(probe.get::<String>("ft_cuh").unwrap(), "cuda");
+    assert_eq!(
+        probe.get::<String>("grammar_cu").unwrap(),
+        "cuda",
+        "bundled grammar resolves `.cu` to cuda"
+    );
+    assert_eq!(probe.get::<String>("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