Module Sarek_codegen.Sarek_ir_intrinsic_dispatch

Shared, table-driven EIntrinsic dispatcher for the source backends (GLSL/Vulkan, WGSL/WebGPU, Metal, CUDA, OpenCL).

Before this module each backend hand-rolled a near-identical gen_intrinsic (~200-265 lines apiece): the same 25-entry thread-intrinsic list, the same "emit callee(arg0, arg1, ...)" argument loop, the same atomic add/min/max arg-count arms copy-pasted three times per backend, the same pure-registry query skeleton, and the same FFI-registry template expansion. Unifying them (audit #49) also closes audit #48: the unknown-intrinsic fall-through no longer emits a raw full_name(args) string (which the pipeline returned Ok, yielding invalid device code) - it now raises the same located unknown_intrinsic error on every backend, exactly as GLSL already did after #259.

PTX is intentionally NOT a client: it lowers intrinsics natively through Sarek_ir_ptx_expr.emit_intrinsic_native and is owned by a separate task.

type 'e spec = {
  1. framework : unit -> string;
    (*

    Framework tag passed to Sarek_pure_registry.fun_device_template. A thunk for historical reasons rather than present necessity: it was one when the tag lived in a mutable current_framework ref that a caller could rewrite between generations. Since backlog-185/200 all five implementations are fun () -> "<backend>" — a literal, the same on every call. Nothing varies it: an intermediate draft threaded a ?framework argument so a caller could override the tag, but every caller passed a backend its own name, so it was dropped. Left as a thunk because collapsing it to a string is a separate, purely cosmetic change to five call sites.

    *)
  2. gen_expr : Stdlib.Buffer.t -> 'e -> unit;
    (*

    The backend's (recursive) expression generator, for call arguments.

    *)
  3. thread_intrinsic : string -> string;
    (*

    Maps a thread-intrinsic name to its backend spelling. Only invoked for names in thread_intrinsic_names.

    *)
  4. arm : string -> (Stdlib.Buffer.t -> 'e list -> unit) option;
    (*

    Backend-specific math / atomic / cast lowering. None means "not handled here" - the dispatcher then tries post_hook and finally raises via on_unknown.

    *)
  5. pre_hook : Stdlib.Buffer.t -> full_name:string -> string list -> string -> 'e list -> bool;
    (*

    Arms tried BEFORE the pure registry. Returns true iff it emitted.

    *)
  6. post_hook : Stdlib.Buffer.t -> string list -> string -> 'e list -> bool;
    (*

    Arms tried AFTER arm, before raising. Returns true iff it emitted (the FFI-registry template path for Metal/CUDA/OpenCL).

    *)
  7. on_unknown : string -> unit;
    (*

    Raise the backend's located unknown_intrinsic error. Never returns.

    *)
  8. warp_dialect : Sarek_ir_warp_text.dialect;
    (*

    The target's spelling of the warp operations (warp_shuffle*, warp_size, lane_id). They are lowered by the shared Sarek_ir_warp_text, ahead of every other path: their IR names are fixed by Sarek_warp, so no per-backend arm is needed.

    *)
  9. invalid_arg_count : string -> int -> int -> unit;
    (*

    Raise the backend's located invalid-argument-count error, given the operation name, the expected count and the given one. Never returns. Each backend already had this as a local bad_arity for emit_atomic; putting it in the spec is what lets the shared pipeline enforce intrinsic_arity centrally instead of at ~130 identical call sites.

    *)
}
val c_signed_int : string -> string

C-dialect spelling of a thread/block/grid index as a signed int.

CUDA/HIP threadIdx.x and friends are unsigned int, OpenCL get_local_id(0) is size_t and Metal's uint3 members are unsigned, so the raw spelling turns thread_idx_x - 1l < 0l into an always-false comparison and (thread_idx_x - 5l) / 2l into a huge positive number. The Sarek type of these intrinsics is int32; the cast makes the C code agree with the interpreter and with GLSL/WGSL (which already cast).

val thread_intrinsic_names : string list

The thread/grid position intrinsics, identical across all five backends.

val intrinsic_arity : (string * int) list

Argument count of the intrinsics whose arity is the same on every backend.

Checked once, centrally, in gen_intrinsic. It has to be central: the lowering for these names is emit_call, which takes whatever argument list it is handed and writes callee(a, b, c) — so sin(x, y) used to emit verbatim on all five backends and the pipeline returned Ok. There are ~26 such names per backend and the call sites are identical, which is precisely why the check does not belong at the call sites (audit #94).

Deliberately NOT listed here:

  • the atomics, which accept both (addr, value) and (arr, idx, value) — emit_atomic checks them against the form it was given;
  • float / int_of_float / rsqrt, guarded at their arms by emit_unary, whose wrapper text differs per backend;
  • block_barrier, which ignores its arguments on every backend;
  • anything reached through pre_hook or the FFI registry, where the arity is the polyfill's or the template's business, not this table's.

A name absent from this table is unconstrained, which is the safe default: an entry here rejects code, so a wrong entry breaks a working kernel.

val full_name : string list -> string -> string

The dotted OCaml source name of an intrinsic (Float64.log10, sin, ...).

val c_int32_literal : int32 -> string

An int32 literal in C, OpenCL C or MSL source. -2147483648 is unary minus applied to 2147483648, which has type long in all three, so Int32.min_int would reach a size-checked reinterpretation (as_float, as_type<float>) as an 8-byte operand and fail to compile. (-2147483647 - 1) is an int.

val is_gpu_or_std_path : string list -> bool

The module paths under which the core warp primitives are reached.

val render_expr : gen_expr:(Stdlib.Buffer.t -> 'a -> unit) -> 'a -> string

Render one expression to a string with the backend generator.

val emit_args : gen_expr:(Stdlib.Buffer.t -> 'a -> unit) -> Stdlib.Buffer.t -> 'a list -> unit

Emit arg0, arg1, ... (comma-space separated) using the backend generator.

val emit_unary : gen_expr:(Stdlib.Buffer.t -> 'a -> unit) -> invalid_arg_count:('b -> int -> int -> unit) -> Stdlib.Buffer.t -> prefix:string -> suffix:string -> opname:'b -> 'a list -> unit

Emit a fixed-arity wrapper prefix arg suffix, failing loudly on any other argument count.

This exists because the shape it replaces did not. Four backend arms were written as

  Buffer.add_string buf "f32(" ;
  (match args with [e] -> gen_expr buf e | _ -> ()) ;
  Buffer.add_char buf ')'

— a wildcard that SUCCEEDS on the wrong argument count, emitting f32() with no argument at all and returning normally, so the pipeline yields Ok and the defect surfaces as a shader-compiler error with no connection to the kernel that caused it. That is the same failure mode as the raw-name fall-through audit #48 closed for unknown intrinsics; it survived inside the arms that WERE handled (audit #94).

val is_stdlib_int32_path : string list -> bool

The module paths under which a kernel reaches the stdlib Int32 module: Int32.x as written, the library-qualified form, and the FFI-free metadata twin used by the transpiler.

val emit_int32_bitcast : gen_expr:(Stdlib.Buffer.t -> 'a -> unit) -> invalid_arg_count:(string -> int -> int -> unit) -> to_float:string -> to_int:string -> Stdlib.Buffer.t -> string list -> string -> 'a list -> bool

Int32.float_of_bits / Int32.bits_of_float on a backend that lowers them natively rather than through the FFI registry (GLSL, WGSL). to_float and to_int are the backend's spellings of the two reinterpretations. The match is scoped to the stdlib Int32 path so another module's intrinsic of the same bare name is left to its own lowering. Returns true iff it emitted.

val emit_call : gen_expr:(Stdlib.Buffer.t -> 'a -> unit) -> Stdlib.Buffer.t -> string -> 'a list -> unit

Emit a plain call callee(arg0, arg1, ...).

val emit_atomic : gen_expr:(Stdlib.Buffer.t -> 'a -> unit) -> invalid_arg_count:('b -> 'c -> int -> unit) -> Stdlib.Buffer.t -> callee:string -> prefix:string -> suffix:string -> opname:'b -> expected:'c -> allow_array:bool -> 'a list -> unit

Emit a binary atomic, factoring the add/min/max arg-count arms that were copy-pasted three times inside every backend's gen_intrinsic. Output is callee(<prefix>addr, value)<suffix> for the two-argument form, or callee(<prefix>arr[idx], value)<suffix> for the optional three-argument array form.

val count_placeholders : string -> int
val emit_registry_template : gen_expr:(Stdlib.Buffer.t -> 'a -> unit) -> framework:string -> invalid_arg_count:(string -> int -> int -> string) -> Stdlib.Buffer.t -> string list -> string -> 'a list -> bool

Expand a Metal/CUDA/OpenCL FFI-registry (Sarek_registry) device template for path.name. Returns false when the registry has no template (caller then raises the unknown-intrinsic error); true once emitted.

val gen_intrinsic : 'e spec -> Stdlib.Buffer.t -> string list -> string -> 'e list -> unit

The shared dispatch pipeline every backend routes through. Order matches the former hand-rolled bodies: pre_hook, pure registry (path-qualified), shared thread list, arm, post_hook, else raise via on_unknown (audit #48).