README
¶
metal
Compute kernels on a Mac's GPU, from pure Go, with no cgo.
A Mac has a large general-purpose processor sitting idle in it. A Go program that grinds through pixels on the CPU is not merely slower than it could be — it is taking cores away from everything else the machine is doing, which on a laptop that is also syncing files and drawing a browser is the part a person actually notices.
Measured on an M4 Max, one 4K frame of image work (a depth pass, three separable blurs, and two synthesised views):
| per frame | processor time per frame | |
|---|---|---|
| CPU, sixteen cores | 65 ms | 82 ms |
| GPU, this package | 4 ms | 0.16 ms |
Sixteen times faster, and five hundred times cheaper in CPU. The second column is the one that matters: the work stops competing with the rest of the machine almost entirely.
Kernels are compiled at run time
There is no build step and no toolchain to install. A kernel is Metal Shading
Language in a Go string, handed to the system compiler when the program starts,
and a kernel that will not compile answers with the compiler's own message —
program_source:4:9: error: use of undeclared identifier, not "it did not
work".
const source = `
#include <metal_stdlib>
using namespace metal;
struct P { uint n; uchar add; };
kernel void bump(device uchar *d [[buffer(0)]], constant P &p [[buffer(1)]],
uint i [[thread_position_in_grid]]) {
if (i >= p.n) return;
d[i] = d[i] + p.add;
}`
dev, err := metal.Default()
defer dev.Close()
lib, err := dev.Compile(source)
pipe, err := lib.Pipeline("bump")
buf, err := dev.NewBuffer(4096)
copy(buf.Bytes(), input)
p := struct {
N uint32
Add uint8
_ [3]uint8
}{N: 4096, Add: 7}
err = dev.Run(func(e *metal.Encoder) {
e.Use(pipe)
e.Buffer(0, buf)
metal.Constant(e, 1, &p)
e.Dispatch(4096)
})
fmt.Println(buf.Bytes()[0]) // input[0] + 7
Nothing is uploaded or downloaded
On a chip with unified memory — every Apple Silicon Mac — Buffer.Bytes()
hands back a Go slice over the very bytes the GPU reads. Writing into it is
the upload; reading out of it after Run is the download. The copy that
usually eats a GPU's advantage on small work never happens at all.
The scope is the API
A Metal command buffer and its encoder are autoreleased objects with an order
they must be used in — encode, end encoding, commit, wait — and every way of
getting it wrong is a leak or a hang. Run owns all of it, including the
autorelease pool, and there is no way to keep either object past the call.
Whatever the closure gets wrong — dispatching before choosing a kernel, binding
a closed buffer, asking for four axes — is reported as an error from Run, and
the first complaint is the one you get, because the later ones are usually its
consequences.
Kernels must check their own bounds
A grid is dispatched in whole threadgroups, so the last group along an axis runs past the end of the work whenever the group size does not divide the grid. Every kernel begins by comparing its thread position against the real size and returning early. This is how Metal is written anyway; it is stated here because the alternative is memory corruption that only appears at unusual image sizes.
The threadgroup shape is not guessed. It comes from the compiled pipeline's own
threadExecutionWidth and maxTotalThreadsPerThreadgroup, so a kernel that
uses many registers — and is therefore given a smaller budget — is dispatched
accordingly.
Everywhere else
On any platform that is not macOS, every constructor returns ErrUnsupported
and the rest are no-ops. A program that offers a GPU path and a portable one
cross-compiles without build tags of its own.
ErrNoDevice is separate, and real: some virtual machines run macOS with no
Metal device at all.
What this is not
Compute only. There is no rendering, no textures, no command-buffer pipelining across frames — those are worth adding when something here needs them, and not before.
Install
go get github.com/go-macos/metal
CGO_ENABLED=0. The only dependencies are purego and go-macos/objc.
Documentation
¶
Overview ¶
Package metal runs compute kernels on a Mac's GPU from pure Go, with no cgo.
A Mac has a large, idle, general-purpose processor in it, and a program that grinds through pixels on the CPU is not merely slower — it is taking cores away from everything else the machine is doing. That is the case this package answers. The shape of it, on an M4 Max, for a 4K frame of image work:
CPU, sixteen cores 65 ms per frame, 82 ms of processor time GPU, this package 4 ms per frame, 0.16 ms of processor time
Sixteen times faster, and five hundred times cheaper in CPU. The second number is the one that matters on a machine that is also syncing files and drawing a browser.
Kernels are Metal Shading Language, compiled at RUN TIME by the system compiler — there is no build step, no toolchain to install, and nothing shipped but the source string.
dev, err := metal.Default()
lib, err := dev.Compile(source)
pipe, err := lib.Pipeline("addOne")
buf, err := dev.NewBuffer(4096)
err = dev.Run(func(e *metal.Encoder) {
e.Use(pipe)
e.Buffer(0, buf)
e.Dispatch(4096)
})
fmt.Println(buf.Bytes()[0])
On a chip with unified memory — every Apple Silicon Mac — a buffer's bytes are the same bytes the GPU reads. Buffer.Bytes hands back a Go slice over them, so there is no upload and no download, and the copy that usually eats a GPU's advantage never happens.
Kernels must check their own bounds ¶
A grid is dispatched in whole threadgroups, so the last group can run past the end of the work. Every kernel here must begin by comparing its thread position against the real size and returning early, exactly as it would if it were written against Metal directly.
Everywhere else ¶
On a platform that is not macOS every constructor returns ErrUnsupported and the rest are no-ops, so a program that offers a GPU path and a portable one cross-compiles without build tags of its own.
Index ¶
Constants ¶
This section is empty.
Variables ¶
var ErrNoDevice = errors.New("metal: no GPU on this machine")
ErrNoDevice is returned when macOS reports no usable GPU. It happens: some virtual machines present no Metal device at all.
var ErrUnsupported = errors.New("metal: only macOS has Metal")
ErrUnsupported is returned by every constructor on a platform that has no Metal, so that a caller can fall back rather than fail to build.
Functions ¶
func Constant ¶
Constant passes a small value straight into the kernel — the sizes, strides and coefficients that change every frame and would be silly to allocate a buffer for. It is a function rather than a method because Go has no generic methods, and the alternative is unsafe.Sizeof at every call site.
The Go type must match the kernel struct field for field: Go and Metal agree on the layout of fixed-width scalars, and agree on nothing else.
Types ¶
type Buffer ¶
type Buffer struct {
// contains filtered or unexported fields
}
Buffer is memory both processors can read.
type Device ¶
type Device struct {
// contains filtered or unexported fields
}
Device is one GPU and the queue of work sent to it.
func (*Device) UnifiedMemory ¶
UnifiedMemory reports whether buffer memory is shared with the CPU rather than copied across a bus. It is true on every Apple Silicon Mac.
type Encoder ¶
type Encoder struct {
// contains filtered or unexported fields
}
Encoder records the work of a single submission.
type Library ¶
type Library struct {
// contains filtered or unexported fields
}
Library is a set of kernels compiled together from one source string.
type Pipeline ¶
type Pipeline struct {
// contains filtered or unexported fields
}
Pipeline is one kernel, compiled and ready to dispatch. It carries the two numbers the hardware reports about itself — the width a group of threads executes in and the most threads a group may hold — which is what lets Dispatch choose a threadgroup shape without the caller guessing.
Source Files