vulki

README

vulki

Vulki is a cgo-free compute library with a small owner-bound Go API, direct Vulkan execution, WGSL shader compilation, and experimental image registration.

Vulki loads Vulkan through purego, compiles WGSL to SPIR-V with gogpu/naga, and does not require cgo.

This module is experimental. APIs may change. The root compute package is usable for development, while image-registration accuracy and portability are still being validated.

When you find any errors please report them as issues.

Requirements

• Go 1.26 or newer.
• A 64-bit target.
• A Vulkan 1.1 loader and compute-capable device for GPU acceleration.
• ImageMagick's convert command only for the registration CLI self-test.

Vulkan is optional for image registration. The automatic correlator prefers the GPU and falls back to the CPU implementation when Vulkan is unavailable. The GPU path stages each packed input once, keeps every intermediate on the device, and includes both uploads and the 64-byte result readback in one Vulkan queue submission.

The Linux Vulkan path is the only GPU path with current runtime evidence. Windows, Android, Apple platforms, browser support, and other platform work are backlog only and are not currently supported.

Install

Add a package to a Go module:

go get github.com/srlehn/vulki

Install the example commands:

go install github.com/srlehn/vulki/cmd/demo@latest
go install github.com/srlehn/vulki/cmd/correlate@latest

Packages

• The root vulki package provides owner-bound devices, buffers, reusable WGSL kernels, binding sets, synchronous dispatch, and command recording without exposing Vulkan handles.
• shader compiles WGSL source to SPIR-V.
• vk contains the low-level Vulkan types, constants, loader, and function wrappers used by the direct implementation.
• registration contains an experimental FFT-based phase-correlation pipeline for estimating rotation, scale, and translation between images.

Migration from the legacy packages

• Replace compute.NewContext with vulki.Open and let each root resource remember its creating device.
• Replace Vulkan buffer flags and typed buffers with Device.NewBuffer plus explicit byte encoders matching the WGSL layout.
• Replace permanently bound compute pipelines with Device.NewKernel and reusable Kernel.NewBindings sets.
• Replace shader.Compile(source, nil) with shader.Compile(source) or the package's functional options.
• Replace the imgproc import path with registration; use its single NewCorrelator constructor and functional options.

The smallest complete compute example is in cmd/demo. Its core setup looks like this:

device, err := vulki.Open()
if err != nil {
    return err
}
defer device.Close()

input, err := device.NewBuffer(size)
if err != nil {
    return err
}
defer input.Close()

output, err := device.NewBuffer(size)
if err != nil {
    return err
}
defer output.Close()

kernel, err := device.NewKernel(vulki.KernelOptions{
    WGSL: wgslSource,
    Bindings: []vulki.BindingLayout{
        {Binding: 0, Access: vulki.BufferReadOnly},
        {Binding: 1, Access: vulki.BufferReadWrite},
    },
})
if err != nil {
    return err
}
defer kernel.Close()

Bind concrete buffers with Kernel.NewBindings, upload raw bytes through the buffer, and call Device.DispatchAndWait. cmd/demo is the complete checked example of that blocking upload, dispatch, and download flow. Use a Recorder to batch arbitrary recorded uploads and downloads, aligned inline updates, explicit compute barriers, and multiple dispatches into one queue submission. Recorder.SubmitAndWait blocks until completion. Recorder.Submit instead returns a Submission handle with non-blocking Poll and blocking Wait, so the next batch can be recorded and submitted while an earlier one executes, and Device.Submit carries several recorded batches in one fused queue submission in argument order. Recorded download destinations become valid only once completion is observed. Recorder.TimestampBegin and Recorder.TimestampEnd record up to MaxTimestampSpans flat labeled GPU timing spans per submission, and Recorder.Timestamps reports nanoseconds per label once completion is observed; devices without usable compute-queue timestamps record nothing and report ErrTimestampsUnsupported. Recorder.DispatchIndirect takes its workgroup counts from a device buffer instead of the host, so a producing kernel in the same submission can size the next stage without a readback; the counts are three uint32 values at a four-byte-aligned offset and are not validated on the host. Recorder.Fill clears or sets a four-byte-aligned range to a repeated 32-bit value and Recorder.Copy copies between two device buffers, both without a host round trip. Recorder.UploadRegions writes several disjoint destination ranges from one staging resource and one host mapping, packed tightly in argument order, so a caller with scattered payloads neither builds a carrier slice nor holds one staging resource per range; regions naming the same buffer are recorded as a single copy command. The recorder features are exercised end to end by the direct-device tests in recorder_test.go. Compatible submissions from separate goroutines may remain in flight together: disjoint buffers and shared read-only buffers can overlap, while any overlapping write remains ordered.

Command pools, command buffers, and submission fences are leased exclusively from bounded device-owned idle pools. Recorders, blocking dispatches, uploads, and downloads reuse them only after completion and required resets succeed; failed or ambiguous submission resources are never returned to idle reuse.

Host-visible staging buffers come from a second device-owned pool that both Buffer transfers and recorded transfers draw from. Device.ReserveTransfer pre-creates one for a known size and direction, so a caller that knows its geometry during setup can move buffer creation, memory allocation, and the first touch of the mapping off the first transfer. Reserving again for a size an idle reservation already covers does nothing, reserving while that reservation is in flight creates another, and a reservation is retained until the device closes instead of being evicted by the ordinary idle limits. There is a bounded number of live reservations, and a request beyond it returns an error rather than appearing to succeed. Recorded staging returns to the pool only once its submission completes, so overlapping submissions cannot share one reservation and a caller that keeps several of a size in flight needs that many reservations of it.

KernelOptions.RequireFullSubgroups asks for fully populated subgroups, which makes deriving a subgroup index from the local invocation index valid. It needs a Vulkan 1.3 device supporting computeFullSubgroups, fails with an error matching ErrFullSubgroupsUnsupported when the device cannot guarantee them so a caller can fall back, and requires the kernel's x workgroup size to be a multiple of Limits.SubgroupSize. KernelOptions.RequiredSubgroupSize pins the subgroup size instead of deriving it in the shader, so a kernel can divide its workgroup into a known number of subgroups. It takes a power of two within Limits.MinSubgroupSize and Limits.MaxSubgroupSize, and reports ErrRequiredSubgroupSizeUnsupported on a device that cannot pin one, which Limits.RequiredSubgroupSizeSupported predicts.

Pipeline-overridable WGSL constants are not currently supported. The vendored Naga 0.18.0 accepts override declarations and can resolve their values in its intermediate representation, but its SPIR-V backend does not yet preserve those resolved constants or an overridden workgroup size. Constants that must vary between kernels therefore need to be baked into separate WGSL sources.

Use BufferUniform for a fixed-size var<uniform> parameter block. Uniform bindings are read-only, and the bound buffer must not exceed Limits.MaxUniformBufferSize; runtime-sized arrays remain storage-buffer bindings. Limits.MinUniformBufferOffsetAlignment reports the device's native uniform-buffer offset alignment.

Device.Info reports the selected adapter and its compute limits, including device-local memory capacity, the subgroup size, its supported range, the supported subgroup operation classes, and whether the device can guarantee full subgroups, so a kernel can pick a workgroup size or choose between a subgroup path and a fallback without building a kernel to find out. Subgroup values are zero when the implementation does not report them, which means unknown rather than unsupported.

Failures keep their Vulkan cause inspectable with errors.Is: ErrOutOfDeviceMemory marks recoverable device memory exhaustion, ErrDeviceLost marks a lost device, and ErrDeviceUnavailable marks a device that refuses further submissions after a failed fence wait. Device.Err reports that unavailable state and its cause.

Pipeline cache

Open creates one application-managed Vulkan pipeline cache per device and NewKernel persists it after each successful pipeline creation. This is enabled by default so compiled pipelines can be reused by later processes and by differently named executables. The default file is os.UserCacheDir()/vulki/pipeline-<pipelineCacheUUID>.bin.

Set VULKI_PIPELINE_CACHE=off to disable both the in-memory cache and disk persistence. Set VULKI_PIPELINE_CACHE_PATH to an exact file path to override the default location. Cache files are device- and driver-specific opaque data; Vulki validates their standard header before use. Missing, corrupt, mismatched, unreadable, or unwritable files are ignored, so cache failures do not change the result of Open or NewKernel.

Commands

Run the WGSL compute demo, which doubles 256 float32 values and verifies the readback:

go run ./cmd/demo

Estimate the transform between two PNG images:

go run ./cmd/correlate image-a.png image-b.png

The command uses -backend auto by default. Use -backend vulkan to require Vulkan or -backend cpu to bypass Vulkan explicitly:

go run ./cmd/correlate -backend cpu image-a.png image-b.png

The two input images must have the same pixel dimensions. Inputs with different dimensions are rejected until the higher-resolution-image semantics from the reference algorithm are implemented explicitly.

Phase-correlation peaks are normalized to the theoretical [0, 1] range. Matches at or below the paper's 0.03 validity threshold return registration.ErrLowConfidence instead of a transform.

Run the randomized registration self-test for one PNG. Add -save to write a stacked comparison image:

go run ./cmd/correlate -save image.png

The registration command is a research tool. Its reported transform should be checked against known inputs before relying on it.

Library users get the same GPU-first behavior from registration.NewCorrelator(maxW, maxH). Pass registration.WithBackend(registration.BackendVulkan) or registration.WithBackend(registration.BackendCPU) to require a backend, or use registration.WithDevice(device) to borrow an existing root device. Correlator.Backend reports the implementation actually in use, and FallbackReason explains an automatic CPU selection. PhaseCorrelate accepts any image.Image implementation and converts non-RGBA inputs internally.

Development

go test ./...
go vet ./...

CPU registration tests do not establish support for untested platforms. GPU tests skip when no suitable Vulkan loader, driver, or device is available; a skip is not GPU evidence. When spirv-val is installed, the image-processing shader test also validates every generated module against the Vulkan 1.1 SPIR-V rules.

Dependencies are vendored, so the test and build commands use the checked-in dependency source by default.

License

Vulki is available under the MIT License.

Documentation

Overview

Package vulki provides explicit, cgo-free GPU compute resources.

Open acquires a compute device. Resources created from a Device remember their owner and may be closed explicitly; closing the Device closes any remaining children in reverse creation order. The current implementation uses Vulkan directly through the low-level vk package.

Buffer uploads and downloads, DispatchAndWait, and Recorder.SubmitAndWait block until the requested queue work completes. Recorder.Submit and Device.Submit instead return a Submission whose Wait or Poll observes completion later, and Device.Submit carries several recorded batches in one queue submission in argument order. Recorder uploads copy their input immediately, while recorded download destinations become valid only once completion is observed. Recorder.DispatchIndirect reads its workgroup counts from a device buffer, and Recorder.Fill and Recorder.Copy write and move device memory as recorded commands, so a pipeline can stay on the device without host round trips between stages. Recorder timestamp spans measure labeled device-side time; Timestamps reports them once completion is observed, or an error matching ErrTimestampsUnsupported when the compute queue cannot write usable timestamps. Device serializes calls that access its Vulkan queue, while submissions using disjoint buffers or only shared read-only buffers may remain in flight concurrently. Overlapping writes retain submission order. Individual child resources and recorders must not be used concurrently with Close; Recorder also rejects use after submission or abort. If a submitted fence cannot establish completion, later submissions fail and Device.Close retains the uncertain resources through its device-idle cleanup. Refused submissions match ErrDeviceUnavailable and Device.Err reports the sticky cause; ErrOutOfDeviceMemory and ErrDeviceLost classify the underlying Vulkan failures for errors.Is.

The public API is experimental. The direct Vulkan path is cgo-free and has runtime evidence on Linux. All other platforms are unsupported backlog work.

Index

• Constants
• Variables
• type BindingLayout
• type BindingSet
  • func (set *BindingSet) Close() error
• type Buffer
  • func (b *Buffer) Bytes() ([]byte, error)
  • func (b *Buffer) Close() error
  • func (b *Buffer) Download(destination []byte) error
  • func (b *Buffer) DownloadAt(offset uint64, destination []byte) error
  • func (b *Buffer) Size() uint64
  • func (b *Buffer) Upload(data []byte) error
  • func (b *Buffer) UploadAt(offset uint64, data []byte) error
• type BufferAccess
• type BufferBinding
  • func BindBuffer(binding uint32, buffer *Buffer) BufferBinding
• type Device
  • func Open() (*Device, error)
  • func (d *Device) Close() error
  • func (d *Device) Closed() bool
  • func (d *Device) DispatchAndWait(kernel *Kernel, bindings *BindingSet, groups Workgroups) error
  • func (d *Device) Err() error
  • func (d *Device) Info() DeviceInfo
  • func (d *Device) NewBuffer(size uint64) (*Buffer, error)
  • func (d *Device) NewKernel(options KernelOptions) (*Kernel, error)
  • func (d *Device) NewRecorder() (*Recorder, error)
  • func (d *Device) ReserveTransfer(direction TransferDirection, size uint64) error
  • func (d *Device) Submit(recorders ...*Recorder) (*Submission, error)
• type DeviceInfo
• type DeviceType
• type Kernel
  • func (k *Kernel) Close() error
  • func (k *Kernel) NewBindings(bindings ...BufferBinding) (*BindingSet, error)
• type KernelOptions
• type Limits
• type Recorder
  • func (r *Recorder) Abort() error
  • func (r *Recorder) Barrier(buffers ...*Buffer) error
  • func (r *Recorder) Close() error
  • func (r *Recorder) Copy(destination *Buffer, destinationOffset uint64, source *Buffer, ...) error
  • func (r *Recorder) Dispatch(kernel *Kernel, bindings *BindingSet, groups Workgroups) error
  • func (r *Recorder) DispatchIndirect(kernel *Kernel, bindings *BindingSet, arguments *Buffer, offset uint64) error
  • func (r *Recorder) Download(buffer *Buffer, offset uint64, destination []byte) error
  • func (r *Recorder) Fill(buffer *Buffer, offset, size uint64, value uint32) error
  • func (r *Recorder) Submit() (*Submission, error)
  • func (r *Recorder) SubmitAndWait() error
  • func (r *Recorder) TimestampBegin(label string) error
  • func (r *Recorder) TimestampEnd(label string) error
  • func (r *Recorder) Timestamps() ([]TimestampSpan, error)
  • func (r *Recorder) Update(buffer *Buffer, offset uint64, data []byte) error
  • func (r *Recorder) Upload(buffer *Buffer, offset uint64, data []byte) error
  • func (r *Recorder) UploadRegions(regions ...UploadRegion) error
• type SubgroupFeatures
• type Submission
  • func (s *Submission) Poll() (bool, error)
  • func (s *Submission) Wait() error
• type TimestampSpan
• type TransferDirection
  • func (direction TransferDirection) String() string
• type UploadRegion
• type Workgroups

Examples

• Device.DispatchAndWait
• Recorder

Constants

const DispatchIndirectCommandSize uint64 = uint64(unsafe.Sizeof(vk.DispatchIndirectCommand{}))

DispatchIndirectCommandSize is the byte size of the dispatch arguments read by DispatchIndirect: three consecutive uint32 workgroup counts for the x, y, and z dimensions.

const MaxTimestampSpans = 64

MaxTimestampSpans is the maximum number of flat timestamp spans one Recorder submission can contain. Each span consumes two timestamp queries.

Variables

var ErrDeviceLost = errors.New("vulki: device lost")

ErrDeviceLost matches, via errors.Is, operations that failed with Vulkan's VK_ERROR_DEVICE_LOST. A lost device does not recover; the caller should close it and open a new Device.

var ErrDeviceUnavailable = errors.New("vulki: cannot submit while earlier queue completion is unknown")

ErrDeviceUnavailable matches, via errors.Is, submissions refused because an earlier submitted fence wait failed and queue completion is unknown. The device stays unavailable until it is closed. Device.Err reports the cause.

var ErrFullSubgroupsUnsupported = errors.New("vulki: device cannot guarantee full subgroups")

ErrFullSubgroupsUnsupported matches, via errors.Is, a kernel that requested full subgroups from a device that cannot guarantee them. It separates the unsupported device from a shader or pipeline failure, so a caller can fall back to a kernel that does not depend on full subgroups.

var ErrOutOfDeviceMemory = errors.New("vulki: out of device memory")

ErrOutOfDeviceMemory matches, via errors.Is, operations that failed with Vulkan's VK_ERROR_OUT_OF_DEVICE_MEMORY. It signals recoverable device memory pressure: releasing device resources and retrying can succeed.

var ErrRequiredSubgroupSizeUnsupported = errors.New("vulki: device cannot pin a compute subgroup size")

ErrRequiredSubgroupSizeUnsupported matches, via errors.Is, a kernel that pinned its subgroup size on a device that cannot accept one for compute. It separates the unsupported device from a size the device rejects as invalid, which is reported as a plain validation error.

var ErrTimestampsUnsupported = errors.New("vulki: device timestamps are unsupported")

ErrTimestampsUnsupported matches, via errors.Is, timestamp results requested from a device whose compute queue cannot write usable timestamp queries. Recording timestamp spans on such a device is a harmless no-op.

Types

type BindingLayout

type BindingLayout struct {
	Binding uint32
	Access  BufferAccess
}

BindingLayout declares one buffer binding in descriptor set zero.

type BindingSet

type BindingSet struct {
	// contains filtered or unexported fields
}

BindingSet is a concrete set of buffers bound to one Kernel.

func (*BindingSet) Close

func (set *BindingSet) Close() error

Close releases the binding set. Repeated calls return nil.

type Buffer

type Buffer struct {
	// contains filtered or unexported fields
}

Buffer is a fixed-size storage buffer owned by its creating Device.

func (*Buffer) Bytes

func (b *Buffer) Bytes() ([]byte, error)

Bytes allocates and downloads the complete buffer contents.

func (*Buffer) Close

func (b *Buffer) Close() error

Close releases the buffer and any lazily allocated transfer resources. Repeated calls return nil.

func (*Buffer) Download

func (b *Buffer) Download(destination []byte) error

Download copies bytes from the beginning of the buffer into destination and blocks until the copy completes.

func (*Buffer) DownloadAt

func (b *Buffer) DownloadAt(offset uint64, destination []byte) error

DownloadAt copies bytes from byte offset into destination and blocks until the copy completes.

func (*Buffer) Size

func (b *Buffer) Size() uint64

Size returns the fixed buffer size in bytes. It returns zero for a nil Buffer.

func (*Buffer) Upload

func (b *Buffer) Upload(data []byte) error

Upload copies data to the beginning of the buffer and blocks until the copy completes.

func (*Buffer) UploadAt

func (b *Buffer) UploadAt(offset uint64, data []byte) error

UploadAt copies data to byte offset and blocks until the copy completes.

type BufferAccess

type BufferAccess uint8

BufferAccess describes the descriptor and access class of a buffer binding.

const (
	// BufferReadOnly declares a read-only storage-buffer binding.
	BufferReadOnly BufferAccess = iota + 1
	// BufferReadWrite declares a read-write storage-buffer binding.
	BufferReadWrite
	// BufferUniform declares a read-only uniform-buffer binding for a fixed-size
	// parameter block. WGSL runtime-sized arrays require a storage binding.
	BufferUniform
)

type BufferBinding

type BufferBinding struct {
	// contains filtered or unexported fields
}

BufferBinding associates a declared binding number with a Buffer.

func BindBuffer

func BindBuffer(binding uint32, buffer *Buffer) BufferBinding

BindBuffer creates one buffer binding for Kernel.NewBindings.

type Device

type Device struct {
	// contains filtered or unexported fields
}

Device owns a logical compute device and the resources created from it. Its zero value is closed and safe to close again.

func Open

func Open() (*Device, error)

Open acquires the first Vulkan physical device with a compute queue and creates one logical compute device. The caller must close the returned Device.

func (*Device) Close

func (d *Device) Close() error

Close waits for active queue work, closes remaining child resources in reverse creation order, and releases the native device. Cleanup continues after a wait or child error. Repeated calls after cleanup return nil.

func (*Device) Closed

func (d *Device) Closed() bool

Closed reports whether the Device is nil, closing, or closed.

func (*Device) DispatchAndWait

func (d *Device) DispatchAndWait(kernel *Kernel, bindings *BindingSet, groups Workgroups) error

DispatchAndWait records one compute dispatch and blocks until it completes.

Example

package main

import (
	"encoding/binary"

	"github.com/srlehn/vulki"
)

const doubleWGSL = `
@group(0) @binding(0) var<storage, read_write> values: array<u32>;

@compute @workgroup_size(1)
fn main(@builtin(global_invocation_id) id: vec3<u32>) {
    values[id.x] = values[id.x] * 2u;
}
`

func main() {
	device, err := vulki.Open()
	if err != nil {
		return
	}
	defer device.Close()

	buffer, err := device.NewBuffer(4)
	if err != nil {
		return
	}
	defer buffer.Close()
	input := make([]byte, 4)
	binary.LittleEndian.PutUint32(input, 21)
	if err := buffer.Upload(input); err != nil {
		return
	}

	kernel, err := device.NewKernel(vulki.KernelOptions{
		WGSL: doubleWGSL,
		Bindings: []vulki.BindingLayout{
			{Binding: 0, Access: vulki.BufferReadWrite},
		},
	})
	if err != nil {
		return
	}
	defer kernel.Close()
	bindings, err := kernel.NewBindings(vulki.BindBuffer(0, buffer))
	if err != nil {
		return
	}
	defer bindings.Close()

	_ = device.DispatchAndWait(kernel, bindings, vulki.Workgroups{X: 1, Y: 1, Z: 1})
}

func (*Device) Err

func (d *Device) Err() error

Err reports why the Device no longer accepts queue submissions. It returns nil while submissions are available. After a submitted fence wait fails, Err returns an error matching ErrDeviceUnavailable that wraps the original causes, so errors.Is can detect a wrapped ErrDeviceLost. Err does not report lifecycle state; use Closed for that.

func (*Device) Info

func (d *Device) Info() DeviceInfo

Info returns a copy of the immutable device information captured by Open.

func (*Device) NewBuffer

func (d *Device) NewBuffer(size uint64) (*Buffer, error)

NewBuffer creates a fixed-size storage buffer. Host-transfer resources are allocated lazily on the first upload or download.

func (*Device) NewKernel

func (d *Device) NewKernel(options KernelOptions) (*Kernel, error)

NewKernel compiles WGSL and creates a reusable compute kernel. EntryPoint defaults to main.

func (*Device) NewRecorder

func (d *Device) NewRecorder() (*Recorder, error)

NewRecorder begins a command recording owned by d.

func (*Device) ReserveTransfer

func (d *Device) ReserveTransfer(direction TransferDirection, size uint64) error

ReserveTransfer pre-creates a pooled staging resource for transfers of up to size bytes in the given direction, so the first transfer of that size pays only the copy instead of also paying buffer creation, memory allocation, and the first touch of the host-visible mapping.

A reservation is idempotent: it does nothing when the pool already holds an unused reservation of sufficient capacity for that direction. A reservation that a transfer is currently using is not unused, so reserving again while the first reservation is in flight creates a second resource, up to a bounded number of live reservations.

Reserved resources are exempt from the ordinary idle retention limits, so a reservation is never silently discarded and never evicts ordinary pooling. They are destroyed when the device closes. Reserving beyond the reservation limit returns an error instead of appearing to succeed.

Both Buffer transfers and recorded transfers draw from this pool, so a reservation serves whichever of them asks for a fitting size first.

A reservation serves one transfer at a time. Recorded staging returns to the pool only when its submission completes, so submissions that overlap cannot share one reservation: while the first holds it, the second takes another idle resource of sufficient capacity or creates one, and a created resource above the ordinary idle bounds is destroyed after use instead of pooled. A caller that keeps several submissions of a size in flight needs that many reservations of it.

func (*Device) Submit

func (d *Device) Submit(recorders ...*Recorder) (*Submission, error)

Submit ends recording on every recorder and submits all recorded batches in argument order as one queue submission, without waiting for completion. Batches execute in submission order on the shared queue, observing the barriers they recorded. On success the recorders accept no further commands and the returned Submission owns completion tracking. On error every recorder is aborted and its resources are released.

Each recorder must belong to this device, appear only once, and not be used concurrently by another goroutine, as documented on Recorder.

type DeviceInfo

type DeviceInfo struct {
	// Implementation names the active native compute implementation.
	Implementation string
	// AdapterName is the native physical-device name.
	AdapterName string
	// DeviceType identifies the broad class of the physical device.
	DeviceType DeviceType
	// APIVersion is the Vulkan API version reported by the adapter.
	APIVersion uint32
	// DriverVersion is the implementation-defined native driver version.
	DriverVersion uint32
	// VendorID is the PCI vendor identifier when the implementation reports one.
	VendorID uint32
	// DeviceID is the implementation-defined physical-device identifier.
	DeviceID uint32
	// DeviceLocalMemoryBytes is the total size in bytes of the physical
	// device's device-local memory heaps. It reports capacity, not current
	// availability.
	DeviceLocalMemoryBytes uint64
	// Limits contains the portable compute limits used by Vulki.
	Limits Limits
}

DeviceInfo is an immutable snapshot of the selected compute device.

type DeviceType

type DeviceType uint32

DeviceType identifies the broad class of a physical compute device. Values match Vulkan's VkPhysicalDeviceType constants.

const (
	// DeviceTypeOther identifies a device that does not fit another class.
	DeviceTypeOther DeviceType = 0
	// DeviceTypeIntegratedGPU identifies an integrated GPU.
	DeviceTypeIntegratedGPU DeviceType = 1
	// DeviceTypeDiscreteGPU identifies a discrete GPU.
	DeviceTypeDiscreteGPU DeviceType = 2
	// DeviceTypeVirtualGPU identifies a virtual GPU.
	DeviceTypeVirtualGPU DeviceType = 3
	// DeviceTypeCPU identifies a CPU Vulkan implementation.
	DeviceTypeCPU DeviceType = 4
)

type Kernel

type Kernel struct {
	// contains filtered or unexported fields
}

Kernel is a reusable compute pipeline and binding layout owned by a Device.

func (*Kernel) Close

func (k *Kernel) Close() error

Close releases the kernel. It returns an error while binding sets still reference it. Repeated calls return nil.

func (*Kernel) NewBindings

func (k *Kernel) NewBindings(bindings ...BufferBinding) (*BindingSet, error)

NewBindings binds one Buffer to every binding declared by the Kernel.

type KernelOptions

type KernelOptions struct {
	WGSL       string
	EntryPoint string
	Bindings   []BindingLayout
	// RequireFullSubgroups asks the implementation to launch the kernel with
	// fully populated subgroups, which makes deriving a subgroup index from the
	// local invocation index valid. Vulkan requires the workgroup size in x to
	// be a multiple of Limits.SubgroupSize under this option, and the size
	// lives in the WGSL rather than in host code. NewKernel fails with an error
	// matching ErrFullSubgroupsUnsupported when the device cannot guarantee
	// full subgroups. Limits.FullSubgroupsSupported reports that in advance.
	RequireFullSubgroups bool
	// RequiredSubgroupSize pins the subgroup size the kernel runs with, so a
	// shader can divide its workgroup into a known number of subgroups instead
	// of deriving one from the subgroup_size builtin and hoping it divides
	// evenly. Zero leaves the choice to the implementation. A non-zero value
	// must be a power of two within Limits.MinSubgroupSize and
	// Limits.MaxSubgroupSize. NewKernel fails with an error matching
	// ErrRequiredSubgroupSizeUnsupported when the device cannot pin a size,
	// which Limits.RequiredSubgroupSizeSupported reports in advance.
	RequiredSubgroupSize uint32
}

KernelOptions describes a reusable WGSL compute kernel.

type Limits

type Limits struct {
	// MaxStorageBufferSize is the maximum storage-buffer range in bytes.
	MaxStorageBufferSize uint64
	// MaxUniformBufferSize is the maximum uniform-buffer range in bytes.
	MaxUniformBufferSize uint64
	// MinUniformBufferOffsetAlignment is the required byte alignment for a
	// uniform-buffer descriptor offset.
	MinUniformBufferOffsetAlignment uint64
	// MaxComputeWorkGroupCount is the maximum dispatch count per dimension.
	MaxComputeWorkGroupCount [3]uint32
	// MaxComputeWorkGroupInvocations is the maximum invocations in one workgroup.
	MaxComputeWorkGroupInvocations uint32
	// MaxComputeWorkGroupSize is the maximum workgroup size per dimension.
	MaxComputeWorkGroupSize [3]uint32
	// SubgroupSize is the number of invocations in one subgroup. It is zero
	// when the implementation does not report subgroup properties, which means
	// unknown rather than unsupported.
	SubgroupSize uint32
	// MinSubgroupSize and MaxSubgroupSize bound the subgroup sizes a compute
	// pipeline may run with. Implementations without subgroup size control
	// report SubgroupSize for both. They are zero exactly when SubgroupSize is.
	MinSubgroupSize uint32
	MaxSubgroupSize uint32
	// SubgroupOperations reports the subgroup operation classes the device
	// supports. It is zero when subgroup properties are unknown.
	SubgroupOperations SubgroupFeatures
	// FullSubgroupsSupported reports whether KernelOptions.RequireFullSubgroups
	// can succeed on this device, so a host can probe the capability without
	// building a kernel. It is false below Vulkan 1.3 and whenever the feature
	// query is unavailable.
	FullSubgroupsSupported bool
	// RequiredSubgroupSizeSupported reports whether
	// KernelOptions.RequiredSubgroupSize can succeed on this device. It is
	// false below Vulkan 1.3, whenever the feature query is unavailable, and
	// when the device does not accept a required size for compute.
	RequiredSubgroupSizeSupported bool
}

Limits contains the portable compute limits used by Vulki.

type Recorder

type Recorder struct {
	// contains filtered or unexported fields
}

Recorder batches buffer updates, barriers, and dispatches into one queue submission. A Recorder is not safe for concurrent method calls. Separate recorders may submit concurrently when their buffer access is compatible.

Example

package main

import (
	"encoding/binary"

	"github.com/srlehn/vulki"
)

const doubleWGSL = `
@group(0) @binding(0) var<storage, read_write> values: array<u32>;

@compute @workgroup_size(1)
fn main(@builtin(global_invocation_id) id: vec3<u32>) {
    values[id.x] = values[id.x] * 2u;
}
`

func main() {
	device, err := vulki.Open()
	if err != nil {
		return
	}
	defer device.Close()
	buffer, err := device.NewBuffer(4)
	if err != nil {
		return
	}
	defer buffer.Close()
	kernel, err := device.NewKernel(vulki.KernelOptions{
		WGSL: doubleWGSL,
		Bindings: []vulki.BindingLayout{
			{Binding: 0, Access: vulki.BufferReadWrite},
		},
	})
	if err != nil {
		return
	}
	defer kernel.Close()
	bindings, err := kernel.NewBindings(vulki.BindBuffer(0, buffer))
	if err != nil {
		return
	}
	defer bindings.Close()
	recorder, err := device.NewRecorder()
	if err != nil {
		return
	}
	defer recorder.Close()

	input := make([]byte, 4)
	binary.LittleEndian.PutUint32(input, 21)
	if err := recorder.Update(buffer, 0, input); err != nil {
		return
	}
	if err := recorder.Dispatch(kernel, bindings, vulki.Workgroups{X: 1, Y: 1, Z: 1}); err != nil {
		return
	}
	if err := recorder.Barrier(buffer); err != nil {
		return
	}
	if err := recorder.Dispatch(kernel, bindings, vulki.Workgroups{X: 1, Y: 1, Z: 1}); err != nil {
		return
	}
	_ = recorder.SubmitAndWait()
}

func (*Recorder) Abort

func (r *Recorder) Abort() error

Abort discards an unsubmitted recording and releases its resources. Repeated calls return nil.

func (*Recorder) Barrier

func (r *Recorder) Barrier(buffers ...*Buffer) error

Barrier records a compute-to-compute memory dependency for buffers.

func (*Recorder) Close

func (r *Recorder) Close() error

Close is equivalent to Abort for a recording and is otherwise a no-op.

func (*Recorder) Copy

func (r *Recorder) Copy(
	destination *Buffer, destinationOffset uint64,
	source *Buffer, sourceOffset uint64,
	size uint64,
) error

Copy records a device-side copy of size bytes between two buffers owned by this device, without a host round trip. Source and destination may be the same buffer, but Vulkan forbids overlapping source and destination ranges within one buffer, so Copy rejects them. A zero size records nothing.

func (*Recorder) Dispatch

func (r *Recorder) Dispatch(kernel *Kernel, bindings *BindingSet, groups Workgroups) error

Dispatch records one dispatch using kernel and bindings.

func (*Recorder) DispatchIndirect

func (r *Recorder) DispatchIndirect(kernel *Kernel, bindings *BindingSet, arguments *Buffer, offset uint64) error

DispatchIndirect records one dispatch whose workgroup counts are read from arguments at byte offset instead of from the host. The offset must be divisible by four and the DispatchIndirectCommandSize bytes at that offset must lie inside the buffer. An earlier dispatch or recorded write in the same recording may produce the counts; DispatchIndirect orders that write before the indirect read. Vulkan reads the counts after this call returns, so they are not validated here and the producing kernel must keep them within Limits.MaxComputeWorkGroupCount.

func (*Recorder) Download

func (r *Recorder) Download(buffer *Buffer, offset uint64, destination []byte) error

Download records a readback from buffer at byte offset. The destination is filled after GPU completion, when SubmitAndWait returns or, after Submit, when Submission.Wait or Submission.Poll observes completion. The caller must not access or modify destination until then.

func (*Recorder) Fill

func (r *Recorder) Fill(buffer *Buffer, offset, size uint64, value uint32) error

Fill records a device-side fill of size bytes at byte offset with a repeated 32-bit value, which clears a counter or a region without a host upload. Offset and size must be divisible by four. A zero size records nothing.

func (*Recorder) Submit

func (r *Recorder) Submit() (*Submission, error)

Submit ends recording, submits once without waiting, and hands completion tracking to the returned Submission. It may be called only while recording. After Submit succeeds the Recorder accepts no further commands; the caller must observe completion through Submission.Wait or Submission.Poll before reusing recorded download destinations. On error the recording is aborted and its resources are released.

func (*Recorder) SubmitAndWait

func (r *Recorder) SubmitAndWait() error

SubmitAndWait ends recording, submits once, waits for completion, and closes the Recorder. It may be called only while recording. If submission succeeds but completion cannot be established, referenced resources remain retained until the Device is closed.

func (*Recorder) TimestampBegin

func (r *Recorder) TimestampBegin(label string) error

TimestampBegin records the start of a labeled GPU timing span. Spans are flat and sequential: the previous span must be ended first, and at most MaxTimestampSpans spans may be recorded per submission. On a device without usable compute timestamps the span records nothing and Timestamps later reports ErrTimestampsUnsupported; recording never fails for that reason.

func (*Recorder) TimestampEnd

func (r *Recorder) TimestampEnd(label string) error

TimestampEnd records the end of the currently open GPU timing span. The label must match the span's TimestampBegin label.

func (*Recorder) Timestamps

func (r *Recorder) Timestamps() ([]TimestampSpan, error)

Timestamps returns the labeled GPU timing spans measured by this recorder, in recording order. It is valid once submission completion has been observed through SubmitAndWait, Submission.Wait, or Submission.Poll. On a device without usable compute timestamps it returns an error matching ErrTimestampsUnsupported when spans were recorded.

func (*Recorder) Update

func (r *Recorder) Update(buffer *Buffer, offset uint64, data []byte) error

Update records an inline buffer update. Offset and data length must be divisible by four, and data must not exceed 65536 bytes.

func (*Recorder) Upload

func (r *Recorder) Upload(buffer *Buffer, offset uint64, data []byte) error

Upload records an arbitrary-size host upload to buffer at byte offset. The input is copied before Upload returns and may be reused immediately. It is the single-region case of UploadRegions.

func (*Recorder) UploadRegions

func (r *Recorder) UploadRegions(regions ...UploadRegion) error

UploadRegions records one host upload that writes several disjoint ranges, using one staging resource and one host-visible mapping for the whole batch instead of one per range. Regions naming the same destination buffer are recorded as a single copy command; each further destination buffer adds one more copy command. All inputs are copied before UploadRegions returns and may be reused immediately.

Staging holds the region data packed tightly in argument order, so the batch acquires exactly the sum of the region lengths and a matching Device.ReserveTransfer reservation serves it without allocating.

Destination ranges must not overlap, because Vulkan forbids writing one location through more than one region of a copy. Empty regions are validated and then ignored, and a batch with no non-empty region records nothing and returns nil.

type SubgroupFeatures

type SubgroupFeatures uint32

SubgroupFeatures is a set of subgroup operation classes. Its values match Vulkan's VkSubgroupFeatureFlagBits, so a device may also report bits for extension classes that have no constant here.

const (
	// SubgroupBasic covers subgroup barriers, elect, and the built-in subgroup
	// size and invocation identifiers.
	SubgroupBasic SubgroupFeatures = 0x00000001
	// SubgroupVote covers all, any, and all-equal.
	SubgroupVote SubgroupFeatures = 0x00000002
	// SubgroupArithmetic covers reductions and inclusive or exclusive scans.
	SubgroupArithmetic SubgroupFeatures = 0x00000004
	// SubgroupBallot covers ballots and broadcasts.
	SubgroupBallot SubgroupFeatures = 0x00000008
	// SubgroupShuffle covers shuffle and shuffle-xor.
	SubgroupShuffle SubgroupFeatures = 0x00000010
	// SubgroupShuffleRelative covers shuffle-up and shuffle-down.
	SubgroupShuffleRelative SubgroupFeatures = 0x00000020
	// SubgroupClustered covers clustered reductions.
	SubgroupClustered SubgroupFeatures = 0x00000040
	// SubgroupQuad covers quad broadcast and swap.
	SubgroupQuad SubgroupFeatures = 0x00000080
)

type Submission

type Submission struct {
	// contains filtered or unexported fields
}

Submission tracks one in-flight queue submission created by Recorder.Submit or Device.Submit. Completion must be observed through Wait or Poll before recorded download destinations are valid. Wait and Poll are safe for concurrent use; once completion is observed the result is final and repeated calls return it unchanged. A Submission that is never observed keeps its resources retained until the Device is closed.

func (*Submission) Poll

func (s *Submission) Poll() (bool, error)

Poll reports whether the submission has completed, without blocking. When it returns true the submission result is final and recorded download destinations are filled.

func (*Submission) Wait

func (s *Submission) Wait() error

Wait blocks until the submission completes, fills recorded download destinations, releases the submitted recorders, and returns the submission result. If completion cannot be established, the affected resources remain retained until the Device is closed and later submissions are refused.

type TimestampSpan

type TimestampSpan struct {
	// Label names the span as passed to TimestampBegin.
	Label string
	// Duration is the device-side time between the span's begin and end
	// timestamps.
	Duration time.Duration
}

TimestampSpan reports one labeled GPU timing span measured by a Recorder.

type TransferDirection

type TransferDirection uint8

TransferDirection selects which side of a host-visible staging transfer a resource serves. Upload and download staging use different buffer usage and different memory-type preferences, so they are pooled separately.

const (
	// TransferUpload stages host memory for copies into a device buffer.
	TransferUpload TransferDirection = iota
	// TransferDownload stages device memory for copies back to the host.
	TransferDownload
)

func (TransferDirection) String

func (direction TransferDirection) String() string

String names the direction for error messages.

type UploadRegion

type UploadRegion struct {
	// Buffer receives Data and must belong to the recording device.
	Buffer *Buffer
	// Offset is the byte offset in Buffer at which Data lands.
	Offset uint64
	// Data is copied into staging before UploadRegions returns and may be
	// reused immediately. An empty Data records nothing.
	Data []byte
}

UploadRegion is one destination range of a batched host upload recorded by Recorder.UploadRegions.

type Workgroups

type Workgroups struct {
	X uint32
	Y uint32
	Z uint32
}

Workgroups is a three-dimensional compute dispatch size.