Module Ir.C_syntax

OCANNL's C-family code generator, shared by the cc, CUDA, HIP and Metal backends: renders lowered routines as C, CUDA, HIP or MSL source through a backend's C_syntax_config, and reports per-compile censuses of the renderings it chose.

type t = PPrint.document

A rendered C-family source document.

val main_logic_marker : Base.string

The comment opening every kernel's main logic: its statements, or an outlined kernel's part calls.

val warp_shuffle_stages : width:Base.Int.t -> Base.Int.t Base.list

The ordered XOR lane masks of one warp-shuffle reduction over the power-of-two width.

val integer_power_operands : Base.Float.t -> Base.int64 * Base.int64 * Base.Int.t * Base.bool

An ocannl_powi_* exponent's operands: |exponent| as (lo, hi) words, a shift, its sign.

val pool_slot_is_64 : Base.unit -> Base.bool

Whether Metal's pooled slot table holds 64-bit entries, as it does under large_models.

val pool_slot_msl_typ : Base.unit -> Base.string

The MSL element type of the pooled slot table: ulong or uint, per pool_slot_is_64.

val routine_stack_locals_bytes : Base.Int.t

Stack budget for a routine's function-scope Local arrays; heap-capable backends spill the rest.

val chunk_stack_locals_bytes : Base.Int.t

Stack budget for one pool-parallel Grid chunk's privatized tiles; past it they go to the heap.

val per_chunk_private_bytes_cap : Base.Int.t lazy_t

The resolved cc_grid_private_bytes_cap: a Grid loop needing more per chunk stays serial.

val default_grid_min_updates : Base.int

The default of cc_grid_min_updates, the work floor of a pool-parallel Grid launch.

val grid_min_updates : Base.unit -> Base.Int.t

The resolved cc_grid_min_updates: a launch estimated below it renders serially; 0 admits all.

val tile_mma_macs_per_update : Base.int

How many Tile_mma multiply-accumulates count as one scalar update (Cost_model's constant).

val grid_update_count : ?cap:Base.int -> Low_level.t -> Base.int

Estimated scalar updates of one launch of a Grid body, saturating at cap.

type mma_rendering =
  1. | Mma_intrinsics
  2. | Mma_intrinsics_ldmatrix
    (*

    The intrinsic arms fed by warp-cooperative ldmatrix loads over a Swizzle_b128 staged tile (gh-ocannl-481 item 3) instead of per-lane gathers. Recorded distinctly because the gh-476 sweep must be able to pin which of the two load paths a timing measured.

    *)
  3. | Mma_register_tiled
  4. | Mma_scalar_fallback

How one Tile_mma statement rendered: intrinsics, a register tile, or the scalar fallback.

val sexp_of_mma_rendering : mma_rendering -> Sexplib0.Sexp.t
val equal_mma_rendering : mma_rendering -> mma_rendering -> Base.bool
val mma_census : Base.unit -> (Base.string * mma_rendering) Base.list Base.ref

This domain's Tile_mma census: (kernel name, rendering) entries, most recent first.

val mma_scope_census : Base.unit -> Base.string Base.list Base.ref

This domain's census of persistent-fragment scopes the fragment hook rendered, one kernel name each.

val is_tensorized_rendering : mma_rendering -> Base.bool

Whether a rendering is a real tensor-core or SIMD-register-tile emission, not the scalar fallback.

type tensorization =
  1. | Not_requested
    (*

    Codegen emitted no Tile_mma statement at all: nothing about this routine claims tensor cores, and a timing of it is not a tensorized timing.

    *)
  2. | Tensorized
    (*

    At least one Tile_mma statement rendered to a real tensor-core or SIMD-register-tile emission (Mma_intrinsics, Mma_intrinsics_ldmatrix, Mma_register_tiled). The C backends' register-tiled rendering counts: it is the SIMD-tile emission those backends have, and the point of the label is fallback-vs-not, not which ISA.

    *)
  3. | Scalar_fallback
    (*

    Tile_mma statements were emitted and every one of them declined to the lane-0 scalar fallback. A timing labeled tensorized in this state measured scalar code.

    *)

Whether a routine's Tile_mma statements tensorized: none emitted, any honoured, or all fell back.

val sexp_of_tensorization : tensorization -> Sexplib0.Sexp.t
val equal_tensorization : tensorization -> tensorization -> Base.bool
val tensorization_name : tensorization -> Base.string

The label a census line prints for a tensorization state.

type mma_summary = {
  1. renderings : (Base.string * mma_rendering) Base.list;
    (*

    The census entries of one compile, in emission order (kernel name, rendering). Fissioned segments of one routine contribute their kernels to the same summary.

    *)
  2. tensorization : tensorization;
  3. statements : Base.int;
    (*

    Tile_mma statements emitted: List.length renderings.

    *)
  4. scalar_fallbacks : Base.int;
    (*

    Of those, how many rendered as the lane-0 scalar fallback.

    *)
  5. fragment_scopes : Base.int;
    (*

    gh-ocannl-1707: persistent-fragment regions (Low_level.fragment_region) the backend's fragment hook rendered, the accumulator register-resident across the whole reduction. Not counted: a region whose hook declined (its Tile_mmas alias the target per statement, or fall back), and a pad-masked contraction's fragment, which is never marked as one.

    *)
}

What a compile's mma_census says about the routine it produced.

val sexp_of_mma_summary : mma_summary -> Sexplib0.Sexp.t
val summarize_census : ?fragment_scopes:Base.int -> (Base.string * mma_rendering) Base.List.t -> mma_summary

Summarize census entries and the fragment_scopes count into an mma_summary.

val empty_mma_summary : mma_summary

The summary of a compile that emitted no Tile_mma statement.

val merge_mma_summaries : mma_summary Base.List.t -> mma_summary

One summary of several compiles' census entries and fragment scopes together.

val mma_summary_string : mma_summary -> Base.string

The one-line census a timing report prints: the tensorization label, then per-rendering counts.

val with_census : (Base.unit -> 'a) -> 'a * mma_summary

Run f collecting the Tile_mma census, nesting additively; return its result and summary.

type peel_skip = Low_level.accum_skip =
  1. | Skip_debug_logging
  2. | Skip_dead_level
  3. | Skip_accum_pinned

Why a site's width decision declined a base the peel reached: Low_level.accum_skip re-exported.

val sexp_of_peel_skip : peel_skip -> Sexplib0.Sexp.t
val equal_peel_skip : peel_skip -> peel_skip -> Base.bool
type peel_verdict = {
  1. levels : Base.int;
    (*

    Loop levels the localized scope spans: the level being rendered, plus those Low_level.peel_accum_nest took below it. One for a reduction localized at its own level; two where the accumulated cell let the peel swallow the enclosing level too.

    *)
  2. guards : Low_level.peel_guard_verdict Base.list;
    (*

    The verdict of each If among them, outermost first.

    *)
}

What a localizing site decided: the loop levels its scope spans and its guards' verdicts.

val sexp_of_peel_verdict : peel_verdict -> Sexplib0.Sexp.t
val equal_peel_verdict : peel_verdict -> peel_verdict -> Base.bool
type peel_site =
  1. | Peel_localized of peel_verdict
    (*

    The site localized, deciding this.

    *)
  2. | Peel_ceded of peel_verdict
    (*

    The width decision reached a base the serial rendering would have localized, deciding this — and the site rendered a form that holds the accumulator at that same residency somewhere other than a minted scope: the warp-shuffle tree or the SIMD accumulator grid (gh-ocannl-754). Not Peel_localized, since no scope was minted; not a skip, since the width question was asked and answered wide. A test can therefore pin that the shuffle or the grid took its width from the one shared decision rather than from a recognizer of its own.

    *)
  3. | Peel_refused of Low_level.peel_refusal
    (*

    Low_level.peel_accum_nest refused, with why.

    *)
  4. | Peel_not_attempted of peel_skip
    (*

    The peel reached a base, and localization was not attempted, for this reason. Every rendering of the level then keeps the accumulator at storage width (gh-ocannl-754).

    *)

One reduction site's peel outcome: localized, ceded to a wide rendering, refused, or not attempted.

val sexp_of_peel_site : peel_site -> Sexplib0.Sexp.t
val equal_peel_site : peel_site -> peel_site -> Base.bool
val peel_census_enabled : Base.unit -> Base.bool Base.ref

This domain's switch for collecting the peel_census.

val peel_census : Base.unit -> (Base.string * peel_site) Base.list Base.ref

This domain's reduction-peel census: (kernel name, site) entries, most recent first.

val is_localized_peel : peel_site -> Base.bool

Whether a site localized its accumulator into a minted scope.

type peel_summary = {
  1. sites : (Base.string * peel_site) Base.list;
    (*

    The census entries of one compile, in emission order (kernel name, site). Fissioned segments of one routine contribute their kernels to the same summary.

    *)
  2. localized : Base.int;
  3. ceded : Base.int;
    (*

    Sites whose width decision was wide and whose rendering held the accumulator outside a scope (gh-ocannl-754).

    *)
  4. declined : Base.int;
    (*

    Sites that did not localize or cede, refused and not-attempted together.

    *)
}

What a compile's peel_census says about the routine it produced.

val sexp_of_peel_summary : peel_summary -> Sexplib0.Sexp.t
type accum_width =
  1. | Accum_not_a_nest of Low_level.peel_refusal
    (*

    Low_level.peel_accum_nest reached no accumulation base under the level.

    *)
  2. | Accum_base of {
    1. tn : Tnode.t;
    2. idcs : Indexing.axis_index Base.array;
    3. base : [ `Update of Low_level.scalar_t | `Scope of Low_level.scope_id * Low_level.t Base.list ];
    4. debug : Base.string;
    5. rebuild : Low_level.t -> Low_level.t;
    6. verdict : peel_verdict;
      (*

      What the peel decided, the level being rendered counted in (levels >= 1).

      *)
    7. pinned : peel_skip Base.option;
      (*

      Some why: the serial rendering keeps the accumulator in its storage cell, narrowing every step, so no rendering may hold it wider. None: the serial rendering localizes it at the residency, and so may any other.

      *)
    }

How one reduction level's accumulator is held, decided once for every rendering of that level.

type statement_accum = {
  1. sa_tn : Tnode.t;
  2. sa_idcs : Indexing.axis_index Base.array;
  3. sa_op : Ops.binop;
  4. sa_contrib : Low_level.scalar_t;
  5. sa_verdict : peel_verdict;
  6. sa_pinned : peel_skip Base.option;
}

An accum_width base that is a single accumulation statement at the level itself.

val peel_summary_string : peel_summary -> Base.string

The one-line peel census: how many sites localized, ceded or declined, and the distinct verdicts.

val with_peel_census : (Base.unit -> 'a) -> 'a * peel_summary

Run f collecting the peel_census, nesting additively; return its result and summary.

type volatility_site =
  1. | Volatile_accumulation_reads of Base.string
    (*

    A reduction-shaped scope local for which the backend selected expression-level volatile pointer casts on any device reads in its accumulating loop, by the local's emitted identifier: the localized accumulation form (gh-ocannl-731, gh-ocannl-820).

    *)
  2. | Plain_accumulator of Base.string
    (*

    A reduction-shaped scope local whose reads stay plain — either this backend does not request the workaround, or no materialized device read was emitted in its accumulating loop. The accumulator itself stays register-resident in both cases.

    *)
  3. | Volatile_rmw_reads of Base.string
    (*

    A device-memory read-modify-write at an address invariant across an enclosing serial loop, whose device reads use expression-level volatile pointer casts.

    *)

One serial accumulation's volatility-census entry, by the emitted identifier.

val sexp_of_volatility_site : volatility_site -> Sexplib0.Sexp.t
val equal_volatility_site : volatility_site -> volatility_site -> Base.bool
type volatility_summary = {
  1. entries : (Base.string * volatility_site) Base.list;
    (*

    The census entries of one compile, in emission order (kernel name, site). Fissioned segments of one routine contribute their kernels to the same summary.

    *)
  2. requested : Base.bool;
    (*

    Whether the backend this compile ran on asks for the workaround (C_syntax_config.volatile_serial_accumulation). false makes every localized site a Plain_accumulator and suppresses the device-memory RMW arm entirely.

    *)
  3. volatile_accumulations : Base.int;
  4. plain_accumulations : Base.int;
  5. volatile_rmw_reads : Base.int;
}

What a compile's volatility census says about the routine it produced.

val sexp_of_volatility_summary : volatility_summary -> Sexplib0.Sexp.t
val volatility_census_enabled : Base.unit -> Base.bool Base.ref

This domain's switch for collecting the volatility_census.

val volatility_census : Base.unit -> (Base.string * volatility_site) Base.list Base.ref

This domain's volatility census: (kernel name, site) entries, most recent first.

val volatility_requested : Base.unit -> Base.bool Base.ref

Whether the backend compiling on this domain requests the volatile-read workaround.

val volatility_summary_string : volatility_summary -> Base.string

The one-line volatility census: accumulations, and how many use the volatile-read workaround.

val with_volatility_census : (Base.unit -> 'a) -> 'a * volatility_summary

Run f collecting the volatility census, nesting additively; return its result and summary.

type grid_serial_reason =
  1. | Grid_opaque_body
    (*

    An opaque statement or a barrier sits under the loop body.

    *)
  2. | Grid_scope_local_escapes
    (*

    A scope local is set under the loop but declared outside it, so its function-scope storage would race across chunks.

    *)
  3. | Grid_local_accessed_outside of Base.string
    (*

    This local is written grid-variantly (it fails the shared rule) and is also accessed outside the loop body (it fails the privatization rule).

    *)
  4. | Grid_local_not_covering of Base.string
    (*

    This local fails the shared rule, and its first access per iteration is not a standalone covering write (it fails the privatization rule).

    *)
  5. | Grid_private_over_cap of {
    1. private_bytes : Base.int;
    2. cap : Base.int;
    }
    (*

    Every local the loop needs privatized could be, but their combined per-chunk footprint private_bytes exceeds cap, the resolved cc_grid_private_bytes_cap (per_chunk_private_bytes_cap).

    *)
  6. | Grid_below_min_updates of {
    1. updates : Base.int;
    2. floor : Base.int;
    }
    (*

    The launch's estimated scalar updates (grid_update_count, saturating at floor) fall below floor, the resolved cc_grid_min_updates.

    *)
  7. | Grid_offset_range of {
    1. from_ : Base.int;
    2. to_ : Base.int;
    }
    (*

    The loop does not start at 0, which the chunked rendering assumes.

    *)

Why an outermost Grid loop rendered serially on a pool-backed backend.

val sexp_of_grid_serial_reason : grid_serial_reason -> Sexplib0.Sexp.t
val equal_grid_serial_reason : grid_serial_reason -> grid_serial_reason -> Base.bool
type grid_site =
  1. | Grid_pool of {
    1. private_bytes : Base.int;
    }
    (*

    Rendered on the pool, each chunk privatizing private_bytes of tiles (0 when none).

    *)
  2. | Grid_serial of grid_serial_reason

One outermost Grid loop's decision: rendered on the pool, or serial and why.

val sexp_of_grid_site : grid_site -> Sexplib0.Sexp.t
val equal_grid_site : grid_site -> grid_site -> Base.bool
type grid_entry = {
  1. kernel : Base.string;
  2. loop : Base.string;
  3. site : grid_site;
}

One outermost Grid loop: the kernel it is in, its loop index's identifier, the decision.

val sexp_of_grid_entry : grid_entry -> Sexplib0.Sexp.t
type grid_summary = {
  1. loops : grid_entry Base.list;
    (*

    The census entries of one compile, in emission order. Fissioned segments of one routine contribute their kernels to the same summary.

    *)
  2. pool : Base.bool;
    (*

    Whether any kernel of this compile was rendered with pool-parallel Grid loops available at all: a pool-backed backend (cc_parallel_grid not none) with more than one chunk and no debug_log_from_routines. false means no entry was recorded, and on a GPU backend that a Grid loop is a hardware grid, not a serial loop.

    *)
  3. pool_rendered : Base.int;
  4. serial : Base.int;
}

What a compile's Grid census says about the routine it produced.

val sexp_of_grid_summary : grid_summary -> Sexplib0.Sexp.t
val grid_summary_string : grid_summary -> Base.string

The one-line Grid census: Grid loops rendered on the pool, and why each other one is serial.

val with_grid_census : (Base.unit -> 'a) -> 'a * grid_summary

Run f collecting the Grid census, nesting additively; return its result and summary.

type mma_space = [
  1. | `Device
  2. | `Shared
  3. | `Thread
  4. | `Fragment of Base.string
]

The address space of a tile-MMA operand as the emission hooks see it.

type mma_layout = [
  1. | `Plain
  2. | `Swizzled_b128
]

The physical layout of a tile-MMA operand's storage, which the emission hooks trust as given.

type mma_operand = PPrint.document * Base.int * mma_space * mma_layout

A tile-MMA operand: tile-base pointer, leading-dimension stride, address space, layout.

type mma_source = Base.int * mma_space * mma_layout

A tile-MMA input operand described but not addressed: stride, address space, layout.

type mma_emission = a_ptr:PPrint.document -> b_ptr:PPrint.document -> PPrint.document

The emitting half of an accepted C_syntax_config.mma_syntax call, awaiting the a/b addresses.

type async_copy_syntax = {
  1. ac_copy : dst:PPrint.document -> src:PPrint.document -> bytes:Base.int -> PPrint.document;
    (*

    One element-sized asynchronous global→workgroup-shared copy statement; dst and src are element addresses (&ident[offset]) and bytes the element size — 4 or 8 today. The hardware also copies 16, but a 16-byte copy requires a 16-byte-aligned destination, and plain workgroup-shared declarations align to the element type only (4 for uint4x32_t) — so 16 stays out of the per-element eligibility until a rendering guarantees the alignment (Codex P2 on PR #317). The copy is byte-for-byte: eligibility (same storage precision on both sides, no value transformation) is the caller's check.

    *)
  2. ac_wait_all : Base.string;
    (*

    Statement completing every asynchronous copy issued so far by the calling thread, committed or not (CUDA cp.async.wait_all). Cross-thread visibility still needs the workgroup barrier the caller emits right after.

    *)
}

Asynchronous global-to-workgroup-shared copies for staging software-pipelined tiles.

type select_blend =
  1. | Blend_truth of PPrint.document -> PPrint.document -> PPrint.document -> PPrint.document
  2. | Blend_over_locals of Base.string
    (*

    How a backend renders a select evaluating both arms: a truth blend, or arms bound to locals.

    *)
type param_guard_dialect =
  1. | C_guards
  2. | Cxx_guards
    (*

    How a dialect spells a builtin's parameter guard: GNU C static assertion, or C++ braced init.

    *)
val blend_local_a : Base.string

The local a Blend_over_locals blend binds its first arm to; tests find blends by it.

module type C_syntax_config = sig ... end

What a backend supplies to the C-family renderer: its dialect's spellings, capabilities and hooks.

val codegen_capabilities : (module C_syntax_config) -> Backend_intf.codegen_capabilities

The codegen capabilities a backend's config publishes, derived from its own hooks.

val default_binop_syntax : Ops.prec -> Ops.binop -> PPrint.document -> PPrint.document -> PPrint.document

The C-family layout of a binop's Ops.binop_c_syntax spelling, shared by every backend.

val rng_binop_syntax : backend:Base.string -> call:(Base.String.t -> 'a) -> Ops.prec -> Ops.binop -> 'a

Render an RNG binop (the Threefry variants, the uniform lane conversion) as its builtin's call.

val all_precs : Ops.prec Base.List.t

Every storage-bearing precision: an alias of Ops.storage_precs.

val operand_conditionality_violations : ternop_syntax: (Ops.prec -> Ops.ternop -> PPrint.document -> PPrint.document -> PPrint.document -> PPrint.ToBuffer.document) -> binop_syntax: (Ops.prec -> Ops.binop -> PPrint.document -> PPrint.document -> PPrint.ToBuffer.document) -> select_blend_syntax:(Ops.prec -> select_blend Base.option) -> Base.string Base.List.t

Where a backend's operator renderings break the operand-evaluation contract; [] when nowhere.

val builtin_idents : ('a * 'b * 'c) Base.List.t -> 'a Base.List.t

The names a builtins table defines (its entries' keys), for a backend's ident_blacklist.

val builtin_definitions : (Base.String.t * 'a * Base.String.t Base.List.t) Base.List.t -> keys:Base.String.t Base.List.t -> 'a Base.List.t

The definitions keys need from a builtins table, dependencies included, in table order.

val vec_widen_macro : store_prec:Ops.prec -> prec:Ops.prec -> lanes:Base.int -> Base.String.t Base.option

The cc builtin widening a 16-bit-storage vector to lanes lanes of prec; None otherwise.

val vec_narrow_macro : store_prec:Ops.prec -> prec:Ops.prec -> lanes:Base.int -> Base.String.t Base.option

The cc builtin narrowing lanes lanes of prec to 16-bit storage; None otherwise.

val partial_staging_idents : Base.string * Base.string

The (bits, conv) locals a register tile's partial bf16/fp16 column is staged through.

val derived_name : (Base.String.t * Base.String.t) -> Base.String.t -> Base.String.t

The identifier prefix ^ base ^ suffix a derivation template spells from base.

val mem_alias_derivation : Base.string * Base.string

Derivation template of a local array's alias pointer.

val heap_local_derivation : Base.string * Base.string

Derivation template of the allocation a heap-declared local array points into.

val grid_chunk_derivations : (Base.string * Base.string) * (Base.string * Base.string) * (Base.string * Base.string)

Derivation templates of a pool-chunked Grid loop's chunk index, lower and upper bounds.

val lane_reduce_derivations : (Base.string * Base.string) * (Base.string * Base.string)

Derivation templates of a lane all-reduce's value and partials.

val warp_reduce_derivations : (Base.string * Base.string) * (Base.string * Base.string)

Derivation templates of a warp-shuffle reduction's value and partials.

val mma_fragment_derivation : Base.string * Base.string

Derivation template of an MMA fragment register, applied to its node's uid.

val minted_derivations : (Base.string * Base.string) Base.list

Every derivation template the emitter mints names with, for the kernel code cache's renaming.

val vec_widen_rows_macros : store_prec:Ops.prec -> prec:Ops.prec -> lanes:Base.int -> (Base.string * Base.string) Base.option

The (pack, row) cc builtins a register tile's A column widens through; None: per-row widening.

val cpp_keywords : Base.string Base.list

The words C++ reserves beyond C's, for the CUDA, HIP and MSL backends' blacklists.

val implementation_reserved : Base.String.t -> Base.bool

Whether name lies in the implementation-reserved namespace: __x or _X.

val system_macro_idents : Base.string Base.list

Predefined compiler macros outside the reserved namespace (linux, unix, ...).

type routine_name_probe = {
  1. prelude : Base.string;
    (*

    The source every kernel follows: the unconditional includes, with every conditional include ahead of them, so a name any kernel's headers declare is measured.

    *)
  2. kernel_definition : Base.string -> Base.string;
    (*

    A minimal kernel DEFINITION of the given name, one line, as the backend declares one.

    *)
  3. compile : Base.string -> (Base.unit, Base.string) Base.Result.t;
    (*

    The backend's runtime compiler on a whole source; Error carries its report.

    *)
  4. kernel_ident : Base.string -> Base.string;
    (*

    The backend's routine-name mangling (C_syntax.kernel_ident).

    *)
  5. routine_floor_idents : Base.string Base.list;
    (*

    What routine names avoid besides vendor_routine_idents: the node-name blacklist and the C library's names. The reserved namespace (implementation_reserved) is a rule, not a list.

    *)
  6. vendor_routine_idents : Base.string Base.list;
    (*

    The backend's harvested table.

    *)
}

What tools/vendor_idents.exe needs from a backend to harvest its toolchain's file-scope names.

val vector_type_names : prefixes:Base.String.t Base.List.t -> scalars:Base.String.t Base.List.t -> widths:Base.Int.t Base.List.t -> Base.String.t Base.List.t

Every <prefix><scalar><width> spelling: GPU dialects' vector-type names, for name blacklists.

val diagnostic_scopes : Base.String.t -> ([> `Error | `Other | `Unscoped ] * Base.String.t) Base.List.t

Each line of a compiler report paired with its diagnostic's severity, or `Unscoped.

val routine_name_in_diagnostics : routine_names:Base.String.t Base.List.t -> Base.String.t -> Base.String.t Base.option

The first of routine_names whose declaration a compiler's error report blames, if any.

val codegen_bug_claim_prefix : Base.string

The opening of the sentence a compile-failure report blames OCANNL's codegen with.

val codegen_bug_claim : artifact:Base.String.t -> Base.String.t

The one-line codegen-bug sentence, asking for artifact.

val missing_include : Base.String.t -> Base.bool

Whether a failed compile's report says a header could not be opened: a toolchain fault.

val attribute_routine_name_failure : ?report:Base.string -> routine_names:(Base.String.t * Base.String.t) Base.List.t -> Base.exn -> Stdlib.Printexc.raw_backtrace -> 'a

Re-raise a failed compile, as a Utils.User_error naming the routine when its name caused it.

val printf_support_unless_uniform : Base.unit -> Base.bool

full_printf_support for a %g-capable dialect: true unless prefer_backend_uniformity is set.

val source_mentions : marker:Base.string -> Base.string -> Base.bool

Whether source's code tokens (not comments or literals) contain marker's tokens in sequence.

val prepend_conditional_includes : conditional_includes:(Base.string * Base.string) Base.List.t -> Base.string -> Base.String.t

Prepend each header whose marker source mentions and that source does not already include.

val with_compiler_options : compiler:Base.String.t -> options:Base.string Base.list -> enrich:(Base.exn -> suffix:Base.String.t -> Base.exn Base.Option.t) -> (Base.unit -> 'a) -> 'a

Run compile, attaching the compiler's options to an exception enrich recognizes.

module Over_locals : sig ... end

The Blend_over_locals renderings, shared by the emitter and operand_conditionality_violations.

module Param_guards : sig ... end

Builtin parameter guards: a test seam checking builtin calls pass arguments without narrowing.

module Source_tokens : sig ... end

C source tokens for builtin and header selection: comments and literals are not uses.

module Pure_C_config (_ : sig ... end) : sig ... end

The plain-C defaults of C_syntax_config, the cc backend's, which other backends override.

module C_syntax (_ : C_syntax_config) : sig ... end

The C-family renderer for one backend's config: routine naming, prelude assembly, kernel emission.

module Compile_driver (_ : functor (_ : sig ... end) -> C_syntax_config) : sig ... end

Renders and compiles routines for a backend, attributing name-caused compile failures to the name.