Metal_bindings (sarek-metal.Sarek_metal.Metal

Library Loading

val metal_lib : Dl.library option Stdlib.Lazy.t

Load Metal framework dynamically (lazy)

val is_available : unit -> bool

Check if Metal library is available

val get_metal_lib : unit -> Dl.library

Get Metal library, raising if not available

Objective-C Runtime Helpers

val objc_lib : Dl.library option Stdlib.Lazy.t

We need libobjc for calling Objective-C methods

val get_objc_lib : unit -> Dl.library

val objc_msgSend_lazy : 
  (unit Ctypes_static.ptr ->
    unit Ctypes_static.ptr ->
    unit Ctypes_static.ptr)
    lazy_t

objc_msgSend - the core Objective-C message dispatch

val objc_msgSend : 
  unit Ctypes_static.ptr ->
  unit Ctypes_static.ptr ->
  unit Ctypes_static.ptr

val sel_registerName_lazy : (string -> unit Ctypes_static.ptr) lazy_t

sel_registerName - register a selector

val sel_registerName : string -> unit Ctypes_static.ptr

val objc_getClass_lazy : (string -> unit Ctypes_static.ptr) lazy_t

objc_getClass - get a class by name

val objc_getClass : string -> unit Ctypes_static.ptr

Foundation Helpers

val foundation_lib : Dl.library option Stdlib.Lazy.t

Load Foundation framework for NSString, NSArray, etc

val get_foundation_lib : unit -> Dl.library

val nsstring_from_cstring : string -> unit Ctypes_static.ptr

Helper: Create NSString from C string

val cstring_from_nsstring : unit Ctypes_static.ptr -> string

Helper: Get C string from NSString

val nserror_description : unit Ctypes.ptr -> string

Helper: Get NSError description

Metal Device API

val mtl_create_system_default_device_lazy : 
  (unit -> Metal_types.mtl_device) lazy_t

MTLCreateSystemDefaultDevice - get default GPU

val mtl_create_system_default_device : unit -> Metal_types.mtl_device

val mtl_copy_all_devices_lazy : (unit -> unit Ctypes_static.ptr) lazy_t

MTLCopyAllDevices - get all Metal devices

val mtl_copy_all_devices : unit -> unit Ctypes_static.ptr

val mtl_device_name : unit Ctypes_static.ptr -> string

Device property getters via objc_msgSend

val mtl_device_max_threads_per_threadgroup : 
  unit Ctypes_static.ptr ->
  Metal_types.mtl_size Ctypes.structure

val mtl_device_max_threadgroup_memory_length : unit Ctypes_static.ptr -> int

Command Queue API

val mtl_device_new_command_queue : 
  unit Ctypes_static.ptr ->
  unit Ctypes_static.ptr

val mtl_command_queue_command_buffer : 
  unit Ctypes_static.ptr ->
  unit Ctypes_static.ptr

Command Buffer API

val mtl_command_buffer_compute_command_encoder : 
  unit Ctypes_static.ptr ->
  unit Ctypes_static.ptr

val mtl_command_buffer_commit : unit Ctypes_static.ptr -> unit

val mtl_command_buffer_wait_until_completed : unit Ctypes_static.ptr -> unit

Buffer API

val mtl_device_new_buffer_with_length : 
  unit Ctypes_static.ptr ->
  int ->
  Unsigned.uint64 ->
  Metal_types.mtl_buffer

val mtl_buffer_contents : unit Ctypes_static.ptr -> unit Ctypes_static.ptr

val mtl_buffer_length : unit Ctypes_static.ptr -> int

Library API

MTLMathMode / MTLMathFloatingPointFunctions

READ FROM THE SDK, not guessed: MTLLibrary.h:241-246 and 258-262 on macOS 15.6.1 (Command Line Tools SDK). The previous revision of this file declined to use setMathMode: precisely because these values could not be checked; they can now.

val mtl_math_mode_safe : int64

val mtl_math_floating_point_functions_precise : int64

val mtl_compile_options_conformant : 
  unit ->
  Metal_types.mtl_compile_options option

Build an MTLCompileOptions configured for Sarek's float semantics, or None if that cannot be done on this host.

WHY (backlog #125). Metal's defaults trade accuracy for speed, and until this change mtl_device_new_library_with_source took an _options argument and IGNORED it while Metal_api passed null anyway. So every Sarek Metal kernel took those defaults and there was no route to change them — not "no policy", but an UNSETTABLE wrong default.

MEASURED on Apple M4 / macOS 15.6.1 (24G90) / Apple clang 17.0.0. A freshly constructed MTLCompileOptions reports mathMode = 2 (MTLMathModeFast) and mathFloatingPointFunctions = 0 (...Fast). BOTH defaults are the fast one, which is why both are set here — setting mathMode alone still leaves single-precision math functions resolving to metal::fast.

THE OPTIONS ARE HONOURED, and that had to be established rather than assumed: a compile that succeeds proves plumbing, not semantics. Over 65536 inputs on sqrt(a) + 1/a, against the default:

mathMode=Safe alone changes 16017 results;
mathMode=Safe + fpFunctions=Precise changes 22135.

So Metal is NOT in the class of rusticl's OpenCL FP options, which are accepted and discarded (docs/fp-contraction-policy.md §9.2).

WHAT THESE OPTIONS DO NOT BUY: CONTRACTION. Measured on the same device, a*b+c is contracted into an fma on all 8773 observable elements under EVERY setting tried, including mathMode=Safe. Contraction is defeated in the generated source instead, by #pragma METAL fp contract(off) (Sarek_ir_metal.metal_fp_contract_pragma, measured 0/8773). Do not read this function as a contraction defence; it is not one.

WHY BOTH SPELLINGS. fastMathEnabled is deprecated since macOS 15.0 in favour of mathMode, but mathMode does not exist before macOS 15.0 / iOS 18.0, so the deprecated property is the only route on older systems. The modern pair is preferred when present and the boolean is the fallback. They are EQUIVALENT, measured: fastMathEnabled = NO and mathMode=Safe + fpFunctions=Precise are BIT-IDENTICAL over 65536 elements of sqrt + reciprocal + sin + log + exp. The fallback is therefore not a degraded path, and that is measured rather than believed.

FAIL-SOFT BY CONSTRUCTION. Every failure path returns None, and None makes the caller pass null, which is EXACTLY the behaviour that shipped before this change: a missing libobjc, a missing class, an OS responding to neither selector, or a null allocation all degrade to the old behaviour rather than propagating.

val mtl_device_new_library_with_source : 
  unit Ctypes_static.ptr ->
  string ->
  Metal_types.mtl_compile_options option ->
  (Metal_types.mtl_library, string) Stdlib.result

Compile MSL source into an MTLLibrary.

options is now HONOURED. Before backlog #125 this function's parameter was named _options and dropped on the floor, so the only thing it could ever pass was null — see mtl_compile_options_conformant for what that meant.

val mtl_library_new_function_with_name : 
  unit Ctypes_static.ptr ->
  string ->
  Metal_types.mtl_function

Compute Pipeline API

val mtl_device_new_compute_pipeline_state : 
  unit Ctypes_static.ptr ->
  Metal_types.mtl_function ->
  (Metal_types.mtl_compute_pipeline_state, string) Stdlib.result

val mtl_compute_pipeline_state_max_total_threads_per_threadgroup : 
  unit Ctypes_static.ptr ->
  int

val mtl_compute_pipeline_state_threadgroup_memory_length : 
  unit Ctypes_static.ptr ->
  int

Compute Command Encoder API

val mtl_compute_command_encoder_set_compute_pipeline_state : 
  unit Ctypes_static.ptr ->
  Metal_types.mtl_compute_pipeline_state ->
  unit

val mtl_compute_command_encoder_set_buffer : 
  unit Ctypes_static.ptr ->
  Metal_types.mtl_buffer ->
  int ->
  int ->
  unit

val mtl_compute_command_encoder_set_bytes : 
  unit Ctypes_static.ptr ->
  unit Ctypes_static.ptr ->
  int ->
  int ->
  unit

val mtl_compute_command_encoder_dispatch_threads : 
  unit Ctypes_static.ptr ->
  Metal_types.mtl_size Ctypes.structure ->
  Metal_types.mtl_size Ctypes.structure ->
  unit

val mtl_compute_command_encoder_end_encoding : unit Ctypes_static.ptr -> unit

Memory Management

val release : unit Ctypes.ptr -> unit

Release any NSObject/Metal object

val retain : unit Ctypes_static.ptr -> unit Ctypes_static.ptr

Retain any NSObject/Metal object