API Reference

LavaBackend <: KA.GPU

Lava's GPU compute backend. Carries the Vulkan context and batch queues explicitly.

• dispatch_bq: where KA kernel dispatches are recorded.

• upload_bq: where CPU→GPU transfers (upload!, download!, staging) record their copies. May be the same queue as dispatch_bq (single-queue mode) or a separate async queue for true upload/compute overlap.

  LavaBackend() # default: dispatch + upload both on defaultbq LavaBackend(bq) # single queue for both LavaBackend(dispatchbq, upload_bq) # split — enables pipelining

LavaBackend() with no arguments resolves dispatch_bq / upload_bq lazily via vk_context().default_bq at every property access. Pinning would break after vk_reset_device!(): a const BACKEND = LavaBackend() created at module-load would keep a stale BatchQueue tied to the old VkDevice, and every subsequent buffer created via that backend would end up allocated on the dead device — later triggering VUID-vkCmdCopyBuffer-commonparent and a page-aligned GPUVM fault. Explicit queues passed to LavaBackend(bq) / LavaBackend(d, u) are pinned on purpose (the caller wants those exact queues, e.g. an async upload queue).

LavaArray{T,N} <: AbstractGPUArray{T,N}

GPU array backed by a Vulkan device-local buffer with BDA (Buffer Device Address).

LavaDeviceArray{T,N}

Device-side isbits array representation for GPU kernels. Contains a Ptr{T} (actually a BDA address) and dimensions.

lava_launch!(bq, f, args...; ndrange, workgroup_size=(64,1,1))

Compile and dispatch a Julia function as a Vulkan compute kernel on bq.

Arguments flow through LavaAdaptor(batch), which strips every LavaArray reached by Adapt.adapt (top level or nested inside closures / wrapper structs) to a LavaDeviceArray{T,N} AND pins its backing VkManagedBuffer into the batch. Kernels therefore see exactly the post-adapt types — the same ABI as KernelAbstractions — never Ptr{T} for what was originally a LavaArray.

Example: a = LavaArray{Float32,1}(undef, (n,)) lavalaunch!(bq, mykernel, a, b, Int32(n); ndrange=n, workgroupsize=(256,1,1)) # kernel signature: mykernel(a::LavaDeviceArray{Float32,1}, ...)

vk_flush!(bq::BatchQueue)
vk_flush!(ctx::VkContext)       # flushes ctx.default_bq

Flush a specific batch queue. Always spell the queue (or its ctx) explicitly — zero-arg convenience forms have been removed per the "Explicit arguments over implicit state" rule.

clear_kernel_cache!()

Evict this device's in-session kernel + pipeline caches so the next dispatch of each kernel recompiles from Julia source.

Use this after editing a Julia kernel under Revise — Revise invalidates the Julia method, but Lava's hash-keyed kernel cache stays populated with the old SPIR-V because hash(f, tt, workgroup_size) doesn't change when the method body changes. Unlike vk_reset_device!(), this keeps all existing LavaArrays and the Vulkan context alive.

Both caches have to go. caches.launchplans holds its own VkPipeline and is consulted before caches.linked on every dispatch, so emptying only the latter left the old pipeline running with no symptom — the function silently did nothing. That is not hypothetical: it made a SPIR-V A/B harness report "no difference" for six variants on 2026-08-02, including one that had its OpStore deleted. The Revise path happened to work anyway, because a method redefinition moves the world counter and launch_plan rejects plans from a superseded world; a caller who only clears the cache had no such luck.

Clear Lava's SPIR-V disk cache.

HWTLAS{Tri} <: Raycore.AbstractAccel

Lava-native hardware-accelerated TLAS. Concretely typed on Tri (the per-primitive triangle type, typically Raycore.Triangle{UInt32}).

Build geometry with push!(hwtlas, mesh, transform), then call Raycore.sync!(hwtlas) to upload and build the Vulkan AS. The adapted form lives in hwtlas.static_tlas as a HWAdaptedAccel{HWTLAS{Tri}}.

Mutation contract

update_transform! / update_transforms! write directly to the batch's GPU-resident instance_buf via a compute kernel and flag transforms_dirty. The next sync! decides between full rebuild (topology change, dirty=true) and MODE_UPDATE_KHR refit (transforms_dirty=true). No CPU-side staging.

Adapted-form invariant

sync!(hwtlas) is the sole owner of hwtlas.static_tlas. It rebuilds as efficiently as possible — in place via resize!/copyto! where the backing buffer can be reused, reallocated only when a buffer grew — and stores the result in hwtlas.static_tlas. sync! MAY reassign hwtlas.static_tlas when a buffer was reallocated.

Every consumer that hands the accel to a raytracing dispatch MUST go through hwtlas.static_tlas or Adapt.adapt(backend, hwtlas) (which reads / refreshes hwtlas.static_tlas) per dispatch. Both are cheap; sync! did the heavy lifting. NEVER cache the HWAdaptedAccel returned by adapt across mutations — consumers that cache silently see stale geometry.

Non-blocking sync!

sync!(hwtlas) does NOT call KA.synchronize(backend). Old backings are dropped via Lava.unsafe_free!, which defers destruction through hwtlas.bq's timeline (bq.deferred_as_frees / bq.deferred_frees) when prior dispatches are still in flight. Phase-B pinning of RT closure leaves (trigpu, offgpu, hwtlas, hwaccel) is what makes the timeline tracking correct on the BDA path.

For a CPU-blocking drain use Raycore.wait_for_gpu!(hwtlas), which calls vk_flush!(hwtlas.bq) (waits on the HWTLAS's own queue specifically, not the backend-wide queue the Raycore default wait_for_gpu! uses).

HardwareAccel{TriVec <: AbstractVector}

Hardware-accelerated ray tracing context. Built from a Raycore-compatible TLAS. Parametrised on the concrete triangle-vector type (e.g. Vector{Raycore.Triangle{UInt32}}) so field accesses stay type-stable.

Fields

• tlas::LavaTLAS — Vulkan top-level acceleration structure
• triangle_data::TriVec — CPU vector of all primitives (for lookup after trace)
• blas_offsets::Vector{UInt32} — Per-BLAS offset into triangle_data
• rt_pipeline::RayTracingPipeline — Pre-compiled raygen+closesthit+miss

Usage

hw = HardwareAccel(raycore_tlas)
results = LavaArray{RTHitResult}(n_rays)
trace_closest_hits!(results, rays, hw)

RayTracingPipeline

A ray tracing pipeline defined by Julia functions for each shader stage. Shaders are compiled lazily on first trace_rays! call and cached.

Constructor

RayTracingPipeline(; raygen, closest_hit, miss, payload_type=:f32)

Example

function my_raygen(output::LavaDeviceArray{Float32,1})
    lid = Lava.lava_rt_launch_id_x()
    Lava.lava_rt_payload_store_f32(-1f0)
    Lava.lava_rt_trace_ray(...)
    t = Lava.lava_rt_payload_load_f32()
    output[lid + 1] = t
end
function my_chit()
    Lava.lava_rt_payload_store_f32(Lava.lava_rt_ray_tmax())
end
function my_miss()
    Lava.lava_rt_payload_store_f32(-1f0)
end

rt = RayTracingPipeline(raygen=my_raygen, closest_hit=my_chit, miss=my_miss)
trace_rays!(rt, tlas, output_buf; width=1920, height=1080)

trace_closest_hits!(results, rays, accel::HardwareAccel, n_rays::Integer;
                    cull_mask::UInt32 = UInt32(0xFF))

Trace n_rays rays against the hardware acceleration structure. Results are written to results buffer (one RTHitResult per ray).

results and rays can be LavaArray{RTHitResult}/LavaArray{RTRay}, or any type accepted by trace_rays!.

cull_mask is ANDed against each instance's instancemask in the TLAS. An instance is visible to a ray only when `(cullmask & instance_mask) != 0. Default0xFF` matches all instances (backward-compatible).

GraphicsPipeline

High-level graphics pipeline wrapping Julia shader functions. Compiles lazily on first use and caches the result.

Fields

• vertex, fragment: Required Julia shader functions
• geometry: Optional (func, GeometryConfig) tuple
• tess_control, tess_eval: Optional tessellation stages
• blend, cull, topology, depth: Pipeline state types

RenderWindow

A window with Vulkan surface and swapchain for presenting rendered frames. Uses GLFW for cross-platform window management.

LavaFramebuffer

Offscreen render target with color and optional depth images. Used for render-to-texture or offscreen rendering.

2D texture backed by VkImage.

Reusable sampler configuration.

present_frame!(bq::BatchQueue, win::RenderWindow)

Submit recorded draw commands and present to screen.

vk_context() -> VkContext

Get or create the global Vulkan context. Lazily initializes on first call.

vk_context(backend) -> VkContext
vk_context(a::LavaArray) -> VkContext

The device a backend or an array belongs to.

The accessor the rest of the stack should reach for instead of calling vk_context() and hoping. vk_context() stays as the convenience default for the single-device case; what has to stop is code depending on it, because a global cannot answer "which device" once there are two.

No new state. Both are derived from what these objects already carried: BatchQueue.ctx for a backend and Buffer.ctx for an array. That is worth saying because it was briefly got wrong in the other direction — a ctx field was added to LavaBackend on the belief that no path existed, which came from reading the first half of BatchQueue's field list, where ctx::Any sits sixty-odd lines down. A second copy of a fact the queue already holds can only ever disagree with it, so this derives instead.

b.dispatch_bq resolves through vk_context() when the backend is unpinned, so an unpinned backend answers "whichever device is current" — which is the correct answer for it, and the reason the queue-only constructors need nothing extra.

vk_reset_device!(; select = pick_physical_device,
                   debug = <the outgoing device's config>)

Replace the process-default Vulkan device. Destroys the old context and creates a fresh one; clears all caches (pipelines, kernels, arg buffers).

Two reasons to call it.

1. Recovery, after ERROR_DEVICE_LOST or another unrecoverable error: vk_reset_device!(). The debugging configuration carries across, so a reset in the middle of a session does not silently turn the instruments off.

2. Switching validation on or off — and this is the only way, because those settings are fixed at vkCreateInstance and cannot be applied to a device that already exists:

   Lava.vk_reset_device!(debug = DebugConfig(gpu_av = true, pool_disabled = true))
   Lava.verify_gpu_av()                     # prove the layer actually fires
   Lava.vk_reset_device!(debug = DebugConfig())   # …and back to the fast path

   See DebugConfig. There is nothing else: no enable_gpu_av, no environment variable, no post-hoc toggle. If you want a device without installing it as the default — a second device, a lavapipe reference — build it with VkContext(; select, debug) instead.

WARNING: All existing LavaArrays become INVALID after reset — their backing GPU buffers no longer exist. You must reallocate all GPU data.

dump_state(; io::IO=stdout)

Print a comprehensive summary of Lava.jl runtime state for debugging.

gpu_memory_usage() -> NamedTuple

Return current GPU memory usage statistics.

set_dispatch_logging!(enabled::Bool)

Enable or disable dispatch name logging. When enabled, each dispatch records its kernel name and parameters for crash debugging. Disabled by default for zero-alloc performance. Auto-enabled on DEVICE_LOST.

get_dispatch_log() -> Vector{String}

Return a copy of the recent dispatch log (up to 2000 entries).

@lava_printf "format" args...

Print from inside a GPU kernel via Vulkan debug printf. The format must be a string literal; arguments are matched positionally to its specifiers. Enable output with Lava.vk_reset_device!(debug = DebugConfig(printf = true)).

@kernel cpu=false function k!(out)
    i = @index(Global)
    @lava_printf "thread %u writing %f\n" UInt32(i) out[i]
end

get_printf_output() -> Vector{String}

Return captured @lava_printf output since the last clear. Drains the async callback ring first. Requires a device built with DebugConfig(printf = true).

clear_printf_output!()

Drop captured @lava_printf output (drains the ring first).