# [`std::gpu`](/silk/docs/std/gpu/) [`std::gpu`](/silk/docs/std/gpu/) is the public host-side API for pure-Silk GPU discovery, device memory, streams, kernel launch, and synchronization. GPU-v1 mixed Linux x86_64 executables use the provider recorded in their bundle: AMDHSA through HIP or NVIDIA PTX through the CUDA Driver API. Importing the module adds no link-time HIP or CUDA dependency. ## Exported API - `is_supported() -> bool` returns whether the bundled provider runtime can be opened and reports at least one device. It does not pre-load every embedded target artifact; target/device compatibility is verified when a named kernel is loaded for launch. - `device_count() -> int` returns the number of provider devices, or `0` when the provider is unavailable. - `current_device() -> int` returns the selected zero-based device index, or `-1` when no device is selected. - `select_device(index: int) -> int` selects one enumerated device for new allocations, default-stream launches, and device-wide synchronization. - `last_error() -> string` returns the runtime's borrowed diagnostic text for the most recent failed operation on the current process. A later GPU operation may overwrite the borrowed text. - `launch(kernel_name: string, grid_size: u64, workgroup_size: u64, ...args: u64) -> int` packs zero to 32 explicit `u64` kernel arguments, loads the named embedded kernel, validates the argument byte count recorded in the executable bundle, and enqueues a one-dimensional dispatch. `grid_size` is the total number of work items; `grid_size / workgroup_size` provider blocks are used. Both values must be nonzero, `grid_size` must be divisible by `workgroup_size`, and `workgroup_size` must not exceed 1024. `grid_size` must not exceed `4294967295`, and the resulting block count must not exceed `2147483647`. The divisibility contract prevents unrequested trailing work items from escaping the requested buffer bounds; the work-item and block limits give every advertised provider the same exact one-dimensional geometry range. - `synchronize() -> int` waits for all previously submitted work on the selected device. - `DispatchResult { launch_status: int, synchronize_status: int }` records the two independent statuses from one automatically completed dispatch. - `is_ok() -> bool` is true only when both statuses are zero. - `launch_failed() -> bool` is true when `launch_status` is nonzero. - `synchronize_failed() -> bool` is true when `synchronize_status` is nonzero. - `launch_and_synchronize(kernel_name: string, grid_size: u64, workgroup_size: u64, ...args: u64) -> DispatchResult` calls `launch`, then calls `synchronize` exactly once even if launch failed, and returns both statuses. It is the public replaceable-stdlib operation behind checked `gpu` forms. - `Buffer { address: u64, byte_len: u64, device_index: int }` owns a provider device allocation and records its creation device. - `Buffer.alloc(byte_len) -> Buffer` returns an invalid zero-address buffer on failure. - `is_valid() -> bool` reports whether the allocation succeeded. - `upload(host_address, byte_count) -> int` copies host memory to the start of the device allocation. - `download(host_address, byte_count) -> int` copies from the allocation to host memory. - `close() -> int` releases the allocation and invalidates the handle. - `drop()` atomically releases a still-live allocation on its creation device during scope cleanup without changing the caller's selected device. - `Stream { handle: u64, device_index: int }` owns one provider stream on the device selected when it was created. - `Stream.create() -> Stream` returns an invalid stream on failure. - `is_valid() -> bool` reports whether the stream owns a runtime handle. - `synchronize() -> int` waits only for work in that stream. - `close() -> int` releases the stream and invalidates it. - `drop()` atomically releases a still-live stream on its creation device during scope cleanup without changing the caller's selected device. - `launch_on(stream, kernel_name, grid_size, workgroup_size, ...args) -> int` enqueues through an owning stream without changing the packed kernel ABI. All status-returning operations use `0` for success and a nonzero failure status. Call `last_error()` for diagnostic text; programs must not parse that text for control flow. The language-level `gpu (grid=..., workspace=...) { kernel(args...); }` form resolves and checks `kernel` at compile time, then dispatches through this module's `launch_and_synchronize` implementation. Statement position discards its `DispatchResult`; value position returns that result so launch and completion failure can independently affect control flow without repeating a kernel-name string. Separate manual `launch` and `synchronize` calls remain supported for overlap and phase-local `last_error()` access. The checked form does not replace buffer ownership, discovery, or `last_error`. `Buffer.upload_on(stream, host_address, byte_count)` and `Buffer.download_on(stream, host_address, byte_count)` enqueue asynchronous copies on a stream owned by the same device. The caller synchronizes that stream before reusing host storage or observing downloaded bytes. Every explicit-stream operation requires a live nonzero `Stream`; an invalid stream returns a nonzero status and never falls back to the default stream. `host_address` is an explicit raw address. A caller may obtain it with the normal Silk raw conversion for an array, slice, or struct whose lifetime covers the copy. A device buffer's `address` may be passed as a `u64` kernel argument. The standard library packs arguments in declaration order, so applications do not construct provider launch-argument memory or depend on a HIP/CUDA ABI. ## Launch lifecycle 1. Build the executable with `--gpu-target`. 2. Optionally call `is_supported()` to handle machines without the bundled provider runtime or any reported device. A target mismatch remains an observable launch failure rather than an availability skip. 3. Allocate/copy any device buffers. 4. Use a checked `gpu` value or call `launch_and_synchronize` when automatic completion and both statuses are wanted. For explicit overlap, call `launch` with the lexical kernel name and one trailing `u64` value per parameter. 5. After a manual launch, call `synchronize` before observing results or releasing resources used by the dispatch. The runtime caches loaded modules/functions for the process lifetime. Buffer ownership remains explicit and independent of that cache. `DispatchResult` values are ordinary per-call values and remain stable when other tasks use the GPU API. The initial bundled runtime state and `last_error()` storage are process-global. The shipped runtime serializes each public GPU operation with a process-wide mutex, so launch blocks and manual calls may execute from multiple task-worker threads without racing runtime initialization, allocation tracking, kernel caching, or provider calls. Because launch and synchronization are distinct operations, another task may submit work between them; synchronization waits for all prior work on the selected device. The borrowed `last_error()` text may still be overwritten after an operation returns, so programs must read it before another task performs a GPU operation. Because automatic synchronization is the later phase, `last_error()` after a checked block is not a durable copy of an earlier launch diagnostic. Use `launch_status` for control flow; use the separate manual calls when diagnostic text must be read between phases. The shipped runtime accepts up to 64 reported provider devices and tracks up to 256 simultaneous live `Buffer` allocations, 128 streams, and 128 cached kernel functions. Exceeding a bound returns a nonzero failure status and records a diagnostic in `last_error()`. ## Errors and portability The shipped runtime rejects non-Linux hosts, a missing or incomplete selected HIP/CUDA provider library, no reported provider device, an executable without a valid Silk GPU bundle, an unknown kernel name, a target artifact that the selected device cannot load, invalid grid/work-group sizes, copies larger than the owning buffer, cross-device resource use, a launch argument byte count that does not match the kernel declaration, and errors returned by provider operations. An alternative `std::` package may implement this API using another runtime without changing application imports. The compiler-owned bundle format and `attr(device=gpu)` semantics remain the same. See [gpu execution](/silk/docs/language/gpu-execution/) for placement rules and [gpu device](/silk/docs/std/gpu-device/) for the semantic device-operation surface.