HIP for ROCm

Syntax highlighting, code navigation and completion for HIP — AMD's GPU programming language for ROCm.

The extension is self-contained: it ships its own parser and its own HIP API database, so it works immediately on any machine with no language server, no compile_commands.json, and no build configuration.

Installing, building or packaging it? See INSTALL.md.

Features

Syntax highlighting

Highlighting comes from three layers that each do what they are good at.

The base grammar. HIP is syntactically CUDA, and VS Code bundles a full CUDA fork of the C++ grammar in which __global__, __device__, __host__, __shared__, __constant__, __managed__, __restrict__ and the dimension built-ins are woven into the structural rules. Delegating to it means a __global__ void saxpy(...) line is recognised as a function definition at any nesting depth, rather than only where a bolt-on pattern happens to win.

An injection grammar adds what the base grammar has no rule for at all, so these win deterministically:

• <<<grid, block, sharedMem, stream>>> launch chevrons, which plain C++ grammars mangle into shift operators
• HIP error codes and enum constants (hipSuccess, hipErrorOutOfMemory, hipMemcpyHostToDevice, hipStreamNonBlocking, …)
• vector types (float4, ulonglong2, …)
• predefined macros (__HIP_ARCH__, __HIP_DEVICE_COMPILE__, __gfx942__, …)

Semantic tokens carry the rest. A TextMate grammar cannot reliably tell a HIP runtime call from any other function call — the C++ rules already claim those positions and which one wins depends on the surrounding syntax. Semantic tokens are resolved after tokenisation and always take priority, so this layer gives distinct, consistent colours to:

Each maps to a TextMate scope as well, so themes without semantic-highlighting support still pick up a sensible colour. Comments and string literals are never re-coloured, and a symbol of your own that happens to share a HIP name is left alone.

Code navigation

• Go to Definition / Declaration / Implementation / Type Definition for your own kernels, functions, classes, enums, typedefs, macros and variables. HIP API symbols jump into the real ROCm header at the right line.
• Find All References and document highlights.
• Outline and breadcrumbs, with kernels marked ⚡ and members nested under their class, namespace or enum.
• Go to Symbol in Workspace (Ctrl+T) across the whole indexed project.
• #include navigation — Ctrl+click a header to open it, resolved against the file's directory, the workspace, and the ROCm include tree.

Locals and parameters resolve to the declaration in scope, so a i inside one kernel does not jump to an i in another.

Completion

• The full HIP runtime API with real signatures and the doxygen documentation extracted from your installed headers.
• Device intrinsics and the device math library, offered only inside __global__/__device__ code.
• Context-sensitive lists:
  • threadIdx. → x, y, z
  • inside <<< >>> → the four launch parameters, plus in-scope dim3 and hipStream_t variables
  • an argument whose HIP parameter is an enum → that enum's constants, so hipMemcpy(..., offers hipMemcpyHostToDevice and friends first
  • #include < → headers from the ROCm include tree
• Symbols from anywhere in your workspace, with kernels inserting a launch configuration template.

Signature help

Parameter hints for HIP runtime calls (with per-parameter documentation from the headers), for device intrinsics with all their overloads, for your own functions, and for the launch configuration inside <<< >>>.

Hover

Signature, documentation, parameter list, return values and originating header for HIP API symbols; declaration, kind, enclosing scope and doc comment for your own code.

Snippets

Kernel skeletons (including a grid-stride variant), the allocate/upload/launch/ download/free sequence, HIP_CHECK, event timing, streams, shared-memory tiles, warp reductions, and device property queries.

Extras

• A hipcc problem matcher for build tasks.
• A status bar item showing the HIP version and index size.

Requirements

None to run. The extension bundles an API database generated from ROCm 7.2.

A local ROCm install is used for two optional things: jumping into headers from go-to-definition, and completing #include <hip/...> paths. It is found via hip.rocmPath, then ROCM_PATH, then /opt/rocm.

Settings

Commands

• HIP: Re-index Workspace
• HIP: Show Index Status
• HIP: Open ROCm Header for Symbol

File types

.hip, .hip.cpp and .hipcc are claimed as the hip language. .cpp/.h files keep their normal C++ identity, and HIP features switch on for them only when they include a HIP header — turn that off with hip.enableInCppFiles if you would rather the C/C++ extension handled them alone.

Coexisting with the C/C++ extension

Both can be active. Microsoft's C/C++ extension will not attach to the hip language id, so in .hip files this extension is the only provider. In .cpp files both contribute and VS Code merges the results.

Building from source

Full instructions, prerequisites and troubleshooting are in INSTALL.md. The short version:

npm install && npm run reinstall

That compiles, packages a .vsix and installs it into VS Code. Reload the window afterwards.

npm test

Five suites with no test framework — parser, API database, providers (against a mock vscode module), the TextMate grammar (tokenised with the real oniguruma engine) and the manifest plus an activate() smoke test. 119 checks in a few seconds.

How it works, and what it will not do

src/parser.ts masks comments and literals, then walks the text tracking brace depth so each { or ; can be classified from the statement header before it. That recovers declarations, their scopes and their extents without a C++ front end — fast enough to re-run on every keystroke, and it needs no build system knowledge. It resolves C++'s most vexing parse the way a reader would: inside a function body, dim3 grid(n); is a variable, not a function declaration.

The trade-off is that it is not a compiler. It does not resolve overloads by argument type, instantiate templates, expand macros, or follow #include to type an expression. Member completion after . or -> is limited to the built-in dimension variables. If you need compiler-accurate semantics, run clangd alongside it with a compile_commands.json from hipcc; this extension is designed to be useful when that is not set up.

License

MIT