metal

Apple Metal GPU bindings for Kit - graphics and compute on macOS

Files

Dependencies

No Kit package dependencies.

System Requirements

• macOS 10.11 or later.
• Xcode Command Line Tools. Install them with xcode-select --install.
• A Metal-capable GPU or Apple GPU runtime.

Metal, QuartzCore, Foundation, and CoreGraphics are Apple system frameworks included with macOS. They are linked as native frameworks by kit.toml.

Installation

kit add gitlab.com/kit-lang/packages/kit-metal.git

Usage

import Kit.Metal as Metal

match Metal.create-default-device
  | Ok device ->
    defer Metal.release device

    println "Device Name: ${Metal.device-name device}"
    println "Unified Memory: ${Metal.has-unified-memory? device}"

    match Metal.new-command-queue device
      | Ok queue ->
        defer Metal.release queue
        println "Command queue created."
      | Err e -> println "Failed to create command queue: ${e}"

  | Err e -> println "No Metal device available: ${e}"

Compute Shader Example

import Kit.Metal as Metal

shader-source = <<~MSL
  #include <metal_stdlib>
  using namespace metal;

  kernel void add_arrays(
      device const float* a [[buffer(0)]],
      device const float* b [[buffer(1)]],
      device float* result [[buffer(2)]],
      uint id [[thread_position_in_grid]]
  ) {
      result[id] = a[id] + b[id];
  }
MSL

match Metal.create-default-device
  | Ok device ->
    defer Metal.release device

    match Metal.compile-shader device shader-source
      | Ok library ->
        defer Metal.release library

        match Metal.get-function library "add_arrays"
          | Ok kernel ->
            defer Metal.release kernel

            match Metal.new-compute-pipeline device kernel
              | Ok pipeline ->
                defer Metal.release pipeline
                println "Max threads: ${Metal.compute-pipeline-max-threads pipeline}"
              | Err e -> println "Pipeline error: ${e}"

          | Err e -> println "Function error: ${e}"

      | Err e -> println "Shader error: ${e}"

  | Err e -> println "Device error: ${e}"

API Overview

Devices

• Metal.create-default-device
• Metal.device-name device
• Metal.is-low-power? device
• Metal.has-unified-memory? device
• Metal.max-buffer-length device

Metal.create-default-device is a zero-argument API and is used as Metal.create-default-device in current examples.

Queues And Command Buffers

• Metal.new-command-queue device
• Metal.command-buffer queue
• Metal.commit command-buffer
• Metal.wait-until-completed command-buffer
• Metal.command-buffer-status command-buffer

Shaders And Pipelines

• Metal.compile-shader device source
• Metal.get-function library name
• Metal.new-compute-pipeline device function
• Metal.compute-pipeline-max-threads pipeline

Buffers And Memory Helpers

• Metal.new-buffer device length options
• Metal.buffer-contents buffer
• Metal.write-float ptr offset value
• Metal.put-float ptr offset value
• Metal.read-float ptr offset
• Metal.write-int ptr offset value
• Metal.read-int ptr offset
• Metal.float-size
• Metal.int-size

Many Metal objects are retained Objective-C objects. Release devices, queues, command buffers, buffers, libraries, functions, encoders, and pipeline states when they are no longer needed:

match Metal.new-command-queue device
  | Ok queue ->
    defer Metal.release queue
    # use queue
  | Err e -> println "Queue error: ${e}"

Development

Running Examples

kit run examples/device-info.kit
kit run examples/shader-compile.kit
kit run examples/compute.kit
kit run examples/vector-add.kit

Build an example binary:

kit build examples/device-info.kit
./device-info

Examples import ../src/metal.kit for local development. Installed package consumers should import Kit.Metal.

Running Tests

kit test

Run tests with coverage:

kit test --coverage

Running Parity

kit parity --no-spinner --failures-only

The examples are marked @check(interactive) so parity build-checks them without comparing runtime output. GPU availability, device names, and Metal runtime behavior vary by machine.

Running kit dev

kit dev

kit dev runs the package development workflow, including native build checks and tests.

Generating Documentation

kit doc

Cleaning Build Artifacts

kit task clean

Native Build Notes

This package uses the ffi-c package kind and builds c/kit_metal.m into libkit_metal.dylib.

The native build is configured in kit.toml:

• csources includes c/kit_metal.m.
• frameworks includes Metal, QuartzCore, Foundation, and CoreGraphics.
• build uses xcrun clang with -fobjc-arc and the required Apple frameworks.

Metal is a macOS framework, not a separate Unix library. Keep it in [native].frameworks; do not add Metal to a native requires list.

Local Installation

From this package directory:

kit install

This installs the package to:

~/.kit/packages/@kit/metal/

Use the package from another Kit project with:

import Kit.Metal as Metal

Troubleshooting

If installation or kit dev reports that native library Metal is not installed, verify that:

• You are running on macOS.
• Xcode Command Line Tools are installed.
• kit.toml lists Metal under [native].frameworks.
• kit.toml does not list Metal as a native requires dependency.

For compiler details, rerun with verbose output:

kit install --verbose
kit dev --verbose

License

MIT License - see LICENSE for details.

Metal is provided by Apple as part of macOS and is subject to Apple's platform terms.

Exported Functions & Types

MetalError

Metal error type for typed error handling. Variants distinguish between different failure modes.

pixel-format-invalid

Invalid pixel format.

pixel-format-a8-unorm

Single 8-bit unsigned normalized alpha channel.

pixel-format-r8-unorm

Single 8-bit unsigned normalized red channel.

pixel-format-r16-float

Single 16-bit floating point red channel.

pixel-format-rg8-unorm

Two 8-bit unsigned normalized channels (RG).

pixel-format-r32-float

Single 32-bit floating point red channel.

pixel-format-rg16-float

Two 16-bit floating point channels (RG).

pixel-format-rgba8-unorm

Four 8-bit unsigned normalized channels (RGBA).

pixel-format-rgba8-unorm-srgb

Four 8-bit unsigned normalized channels (RGBA) in sRGB color space.

pixel-format-bgra8-unorm

Four 8-bit unsigned normalized channels (BGRA).

pixel-format-bgra8-unorm-srgb

Four 8-bit unsigned normalized channels (BGRA) in sRGB color space.

pixel-format-rgb10a2-unorm

10-bit RGB with 2-bit alpha unsigned normalized.

pixel-format-rg32-float

Two 32-bit floating point channels (RG).

pixel-format-rgba16-unorm

Four 16-bit unsigned normalized channels (RGBA).

pixel-format-rgba16-float

Four 16-bit floating point channels (RGBA).

pixel-format-rgba32-float

Four 32-bit floating point channels (RGBA).

pixel-format-depth16-unorm

16-bit unsigned normalized depth.

pixel-format-depth32-float

32-bit floating point depth.

pixel-format-stencil8

8-bit unsigned integer stencil.

pixel-format-depth24-unorm-stencil8

24-bit unsigned normalized depth with 8-bit stencil.

pixel-format-depth32-float-stencil8

32-bit floating point depth with 8-bit stencil.

primitive-type-point

Render points.

primitive-type-line

Render lines.

primitive-type-line-strip

Render line strips.

primitive-type-triangle

Render triangles.

primitive-type-triangle-strip

Render triangle strips.

index-type-uint16

16-bit unsigned integer indices.

index-type-uint32

32-bit unsigned integer indices.

load-action-dont-care

Don't care about existing contents.

load-action-load

Load existing contents.

load-action-clear

Clear to a specified value.

store-action-dont-care

Don't care about storing results.

store-action-store

Store results to the attachment.

store-action-multisample-resolve

Resolve multisampled data.

storage-mode-shared

Shared storage mode (CPU and GPU can access).

storage-mode-managed

Managed storage mode (synchronized between CPU and GPU).

storage-mode-private

Private storage mode (GPU only).

storage-mode-memoryless

Memoryless storage mode (tile memory only).

resource-storage-mode-shared

Resource with shared storage mode.

resource-storage-mode-managed

Resource with managed storage mode.

resource-storage-mode-private

Resource with private storage mode.

resource-cpu-cache-mode-default

Default CPU cache mode.

resource-cpu-cache-mode-write-combined

Write-combined CPU cache mode.

texture-type-1d

1D texture.

texture-type-1d-array

1D texture array.

texture-type-2d

2D texture.

texture-type-2d-array

2D texture array.

texture-type-2d-multisample

2D multisampled texture.

texture-type-cube

Cube texture.

texture-type-cube-array

Cube texture array.

texture-type-3d

3D texture.

texture-usage-unknown

Unknown usage.

texture-usage-shader-read

Shader read usage.

texture-usage-shader-write

Shader write usage.

texture-usage-render-target

Render target usage.

texture-usage-pixel-format-view

Pixel format view usage.

command-buffer-status-not-enqueued

Not yet enqueued.

command-buffer-status-enqueued

Enqueued.

command-buffer-status-committed

Committed.

command-buffer-status-scheduled

Scheduled.

command-buffer-status-completed

Completed successfully.

command-buffer-status-error

Completed with error.

blend-factor-zero

Zero blend factor.

blend-factor-one

One blend factor.

blend-factor-source-color

Source color blend factor.

blend-factor-one-minus-source-color

One minus source color blend factor.

blend-factor-source-alpha

Source alpha blend factor.

blend-factor-one-minus-source-alpha

One minus source alpha blend factor.

blend-factor-dest-color

Destination color blend factor.

blend-factor-one-minus-dest-color

One minus destination color blend factor.

blend-factor-dest-alpha

Destination alpha blend factor.

blend-factor-one-minus-dest-alpha

One minus destination alpha blend factor.

blend-factor-source-alpha-saturated

Source alpha saturated blend factor.

blend-operation-add

Add blend operation.

blend-operation-subtract

Subtract blend operation.

blend-operation-reverse-subtract

Reverse subtract blend operation.

blend-operation-min

Min blend operation.

blend-operation-max

Max blend operation.

compare-function-never

Never pass.

compare-function-less

Pass if less.

compare-function-equal

Pass if equal.

compare-function-less-equal

Pass if less or equal.

compare-function-greater

Pass if greater.

compare-function-not-equal

Pass if not equal.

compare-function-greater-equal

Pass if greater or equal.

compare-function-always

Always pass.

cull-mode-none

No culling.

cull-mode-front

Cull front faces.

cull-mode-back

Cull back faces.

winding-clockwise

Clockwise winding.

winding-counter-clockwise

Counter-clockwise winding.

triangle-fill-mode-fill

Fill triangles.

triangle-fill-mode-lines

Draw triangle edges as lines.

sampler-filter-nearest

Nearest neighbor filtering.

sampler-filter-linear

Linear filtering.

sampler-mip-filter-not-mipmapped

No mipmapping.

sampler-mip-filter-nearest

Nearest mipmap level.

sampler-mip-filter-linear

Linear mipmap interpolation.

sampler-address-mode-clamp-to-edge

Clamp to edge.

sampler-address-mode-mirror-clamp-to-edge

Mirror clamp to edge.

sampler-address-mode-repeat

Repeat.

sampler-address-mode-mirror-repeat

Mirror repeat.

sampler-address-mode-clamp-to-zero

Clamp to zero.

sampler-address-mode-clamp-to-border

Clamp to border color.

vertex-format-invalid

Invalid vertex format.

vertex-format-float

Float vertex format.

vertex-format-float2

Float2 vertex format.

vertex-format-float3

Float3 vertex format.

vertex-format-float4

Float4 vertex format.

vertex-format-int

Int vertex format.

vertex-format-int2

Int2 vertex format.

vertex-format-int3

Int3 vertex format.

vertex-format-int4

Int4 vertex format.

vertex-step-function-constant

Constant vertex step.

vertex-step-function-per-vertex

Per-vertex step.

vertex-step-function-per-instance

Per-instance step.

create-default-device

Create the default system GPU device. Returns Err if no Metal-capable GPU is available.

Result Ptr MetalError

device-name

Get the device name.

Ptr -> String

is-low-power?

Check if device is low-power (integrated GPU).

Ptr -> Bool

is-headless?

Check if device is headless (no display).

Ptr -> Bool

is-removable?

Check if device is removable (eGPU).

Ptr -> Bool

has-unified-memory?

Check if device has unified memory.

Ptr -> Bool

max-buffer-length

Get maximum buffer size.

Ptr -> Long

max-threadgroup-memory-length

Get maximum threadgroup memory length.

Ptr -> Long

recommended-max-working-set-size

Get recommended maximum working set size.

Ptr -> Long

new-command-queue

Create a new command queue.

Ptr -> Result Ptr MetalError

new-command-queue-with-max-count

Create a new command queue with maximum buffer count.

Ptr -> Int -> Result Ptr MetalError

set-command-queue-label

Set a label for the command queue (for debugging).

Ptr -> String -> Unit

command-buffer

Create a new command buffer from a queue.

Ptr -> Result Ptr MetalError

command-buffer-unretained

Create a command buffer with unretained references (performance optimization).

Ptr -> Result Ptr MetalError

set-command-buffer-label

Set a label for the command buffer (for debugging).

Ptr -> String -> Unit

enqueue

Enqueue the command buffer for scheduling.

Ptr -> Unit

commit

Commit a command buffer for execution.

Ptr -> Unit

wait-until-scheduled

Wait for command buffer to be scheduled.

Ptr -> Unit

wait-until-completed

Wait for command buffer to complete.

Ptr -> Unit

present-drawable

Present a drawable when the command buffer completes.

Ptr -> Ptr -> Unit

command-buffer-status

Get command buffer status.

Ptr -> Int

command-buffer-error

Get command buffer error message (if any).

Ptr -> Option String

new-buffer

Create a buffer with specified length.

Ptr -> Int -> Int -> Result Ptr MetalError

new-buffer-shared

Create a buffer with shared storage mode.

Ptr -> Int -> Result Ptr MetalError

new-buffer-private

Create a buffer with private storage mode.

Ptr -> Int -> Result Ptr MetalError

buffer-contents

Get buffer contents pointer.

Ptr -> Ptr

buffer-length

Get buffer length.

Ptr -> Long

buffer-did-modify-range

Notify that a range of the buffer was modified.

Ptr -> Int -> Int -> Unit

set-buffer-label

Set a label for the buffer (for debugging).

Ptr -> String -> Unit

new-texture-descriptor

Create a texture descriptor.

Ptr

configure-texture-2d

Configure a 2D texture descriptor.

Ptr -> PositiveInt -> PositiveInt -> Int -> Int -> Ptr

new-texture

Create a texture from a descriptor.

Ptr -> Ptr -> Result Ptr MetalError

texture-width

Get texture width.

Ptr -> Int

texture-height

Get texture height.

Ptr -> Int

set-texture-label

Set a label for the texture (for debugging).

Ptr -> String -> Unit

compile-shader

Compile Metal shader source code.

Ptr -> String -> Result Ptr MetalError

load-library

Load a pre-compiled shader library from a file.

Ptr -> NonEmptyString -> Result Ptr MetalError

get-function

Get a function from a shader library.

Ptr -> NonEmptyString -> Result Ptr MetalError

new-render-pipeline-simple

Create a simple render pipeline state.

Ptr -> Ptr -> Ptr -> Int -> Result Ptr MetalError

new-render-pass-descriptor

Create a render pass descriptor.

Ptr

set-color-attachment

Set color attachment for render pass.

Ptr -> NonNegativeInt -> Ptr -> Int -> Int -> Float -> Float -> Float -> Float -> Unit

set-color-attachment-clear

Set color attachment with clear color.

Ptr -> Ptr -> Float -> Float -> Float -> Float -> Unit

new-render-encoder

Create a render command encoder.

Ptr -> Ptr -> Result Ptr MetalError

set-render-pipeline

Set the render pipeline state.

Ptr -> Ptr -> Unit

set-vertex-buffer

Set a vertex buffer.

Ptr -> Ptr -> Int -> NonNegativeInt -> Unit

set-viewport

Set the viewport.

Ptr -> Float -> Float -> Float -> Float -> Float -> Float -> Unit

set-scissor-rect

Set the scissor rect.

Ptr -> Int -> Int -> PositiveInt -> PositiveInt -> Unit

set-cull-mode

Set cull mode.

Ptr -> Int -> Unit

set-front-facing-winding

Set front facing winding.

Ptr -> Int -> Unit

draw-primitives

Draw primitives.

Ptr -> Int -> Int -> Int -> Unit

draw-primitives-instanced

Draw primitives with instancing.

Ptr -> Int -> Int -> Int -> Int -> Unit

draw-indexed-primitives

Draw indexed primitives.

Ptr -> Int -> Int -> Int -> Ptr -> Int -> Unit

end-render-encoding

End render encoding.

Ptr -> Unit

new-compute-pipeline

Create a compute pipeline state.

Ptr -> Ptr -> Result Ptr MetalError

compute-pipeline-max-threads

Get maximum threads per threadgroup.

Ptr -> Int

compute-pipeline-execution-width

Get thread execution width.

Ptr -> Int

new-compute-encoder

Create a compute command encoder.

Ptr -> Result Ptr MetalError

set-compute-pipeline

Set the compute pipeline state.

Ptr -> Ptr -> Unit

set-compute-buffer

Set a buffer for compute.

Ptr -> Ptr -> Int -> NonNegativeInt -> Unit

dispatch-threads

Dispatch threads.

Ptr -> Int -> Int -> Int -> Int -> Int -> Int -> Unit

dispatch-threadgroups

Dispatch threadgroups.

Ptr -> Int -> Int -> Int -> Int -> Int -> Int -> Unit

end-compute-encoding

End compute encoding.

Ptr -> Unit

new-blit-encoder

Create a blit command encoder.

Ptr -> Result Ptr MetalError

copy-buffer-to-buffer

Copy buffer to buffer.

Ptr -> Ptr -> Int -> Ptr -> Int -> PositiveInt -> Unit

generate-mipmaps

Generate mipmaps for a texture.

Ptr -> Ptr -> Unit

fill-buffer

Fill buffer with a value.

Ptr -> Ptr -> Int -> Int -> Int -> Unit

synchronize-resource

Synchronize a managed resource.

Ptr -> Ptr -> Unit

end-blit-encoding

End blit encoding.

Ptr -> Unit

new-metal-layer

Create a new CAMetalLayer.

Ptr

set-layer-device

Set the device for a metal layer.

Ptr -> Ptr -> Unit

set-layer-pixel-format

Set the pixel format for a metal layer.

Ptr -> Int -> Unit

set-layer-drawable-size

Set the drawable size for a metal layer.

Ptr -> Float -> Float -> Unit

set-layer-framebuffer-only

Set framebuffer-only mode.

Ptr -> Bool -> Unit

next-drawable

Get the next drawable from a metal layer.

Ptr -> Result Ptr MetalError

drawable-texture

Get the texture from a drawable.

Ptr -> Ptr

release

Release a Metal object.

Ptr -> Unit

retain

Retain a Metal object.

Ptr -> Unit

write-float

Write a float to memory at byte offset.

Ptr -> Int -> Float -> Unit

put-float

Write a float to memory at byte offset.

Ptr -> Int -> Float -> Unit

read-float

Read a float from memory at byte offset.

Ptr -> Int -> Float

write-int

Write an int to memory at byte offset.

Ptr -> Int -> Int -> Unit

read-int

Read an int from memory at byte offset.

Ptr -> Int -> Int

float-size

Size of a float in bytes (for offset calculations).

int-size

Size of an int in bytes (for offset calculations).