Module Sarek_metal.Metal_bindings

Load Metal framework dynamically (lazy)

Check if Metal library is available

Get Metal library, raising if not available

We need libobjc for calling Objective-C methods

objc_msgSend - the core Objective-C message dispatch

sel_registerName - register a selector

objc_getClass - get a class by name

Selectors and typed objc_msgSend stubs on the launch path are built once. foreign creates a libffi call description on every evaluation and sel_registerName takes a runtime lock: done per call they dominated the host cost of a Metal launch.

objc_msgSend with the OCaml runtime lock released for the call, for the messages that block until the GPU is done (waitUntilCompleted). While the lock is held, every other domain stops at its next minor collection until the call returns. The general safety argument is the note on Hip_bindings.foreign_hip_blocking_lazy. For this binding: both arguments are C pointers (the receiver and a selector from sel_registerName), no OCaml heap pointer is passed, and the receiver must stay retained by the caller until the call returns (every caller waits on a command buffer it owns a retain on). Use it only for blocking messages: a release/reacquire costs more than a short message.

Send a no-argument, void-returning message through a cached selector.

send_void for a message that blocks on the device.

Autorelease pools: objects Metal returns autoreleased (command buffers, encoders) are freed when the enclosing pool is popped; OCaml threads have no pool of their own, so without one they leak.

Load Foundation framework for NSString, NSArray, etc

Helper: Create NSString from C string

Helper: Get C string from NSString

Helper: Get NSError description

MTLCreateSystemDefaultDevice - get default GPU

MTLCopyAllDevices - get all Metal devices

Device property getters via objc_msgSend

maxBufferLength: the largest MTLBuffer this device can allocate, in bytes.

recommendedMaxWorkingSetSize: how many bytes of resources the device can use without degrading performance (on Apple silicon, the share of unified memory the GPU may keep resident).

MTLCommandBufferStatusError: the command buffer did not complete.

status of a command buffer (MTLCommandBufferStatus, an NSUInteger).

error of a command buffer (an NSError, or null), as its description.

GPUEndTime of a completed command buffer, in seconds on the CLOCK_UPTIME_RAW (mach absolute time) clock; 0 when not available.

Seconds on CLOCK_UPTIME_RAW, the clock GPUEndTime is expressed in.

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.

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.

Releases the OCaml runtime lock for the compile (see msg_void_blocking_lazy for the safety argument): dev, the NSString and the options are Objective-C objects, the error cell is allocated C memory kept reachable until the call returns.

Releases the OCaml runtime lock for the GPU-code compile; all arguments are Objective-C objects or an allocated error cell (see msg_void_blocking_lazy).

Release any NSObject/Metal object

Retain any NSObject/Metal object

Force every lazily built stub and selector above. Called once when the backend registers (one domain, before user code runs): in OCaml 5 a lazy forced concurrently from two domains, or re-entrantly from a GC finaliser releasing a buffer, raises Lazy.Undefined.