minc Language Reference
minc is a minimal C replacement for building native software. It compiles directly to native executables for x64 Windows (PE), x64/ARM64 Linux (ELF), ARM64 macOS (Mach-O), ARM64 iOS, ARM64 Android (.so), and WebAssembly.
Types
Primitive types
| Type | Size | Description |
|---|---|---|
i8 |
1 byte | Signed 8-bit integer |
i16 |
2 bytes | Signed 16-bit integer |
i32 |
4 bytes | Signed 32-bit integer |
i64 |
8 bytes | Signed 64-bit integer |
u8 |
1 byte | Unsigned 8-bit integer |
u16 |
2 bytes | Unsigned 16-bit integer |
u32 |
4 bytes | Unsigned 32-bit integer |
u64 |
8 bytes | Unsigned 64-bit integer |
f32 |
4 bytes | 32-bit float (IEEE 754) |
f64 |
8 bytes | 64-bit double (IEEE 754) |
bool |
1 byte | Boolean (true/false) |
void |
0 | No value (return type) |
Vector and matrix types
Built-in SIMD-friendly types for graphics and math. All elements are 4 bytes. float4 and int4/uint4 operations compile to packed 128-bit SIMD instructions.
| Type | Size | Description |
|---|---|---|
float2 |
8 bytes | 2x f32 vector |
float3 |
12 bytes | 3x f32 vector |
float4 |
16 bytes | 4x f32 vector (SIMD) |
int2 |
8 bytes | 2x i32 vector |
int3 |
12 bytes | 3x i32 vector |
int4 |
16 bytes | 4x i32 vector (SIMD) |
uint2 |
8 bytes | 2x u32 vector |
uint3 |
12 bytes | 3x u32 vector |
uint4 |
16 bytes | 4x u32 vector (SIMD) |
f64x2 |
16 bytes | 2x f64 vector (SIMD) |
i64x2 |
16 bytes | 2x i64 vector (SIMD) |
u64x2 |
16 bytes | 2x u64 vector (SIMD) |
i8x16 |
16 bytes | 16x i8 vector (SIMD) |
float4x4 |
64 bytes | 4x4 f32 matrix (SIMD, column-major) |
float3 pos = float3{1.0f, 2.0f, 3.0f};
float4 color = float4{1.0f, 0.0f, 0.0f, 1.0f};
int4 indices = int4{0, 1, 2, 3};
A vector literal takes one value or all of its components. One value splats across every lane; a vector-typed value contributes all of its own. Literals concatenate:
float4 grey = float4{0.5f}; // 0.5 in all four lanes
float4 opaque = float4{color.xyz, 1.0f}; // 3 components + 1
float4 pair = float4{uv, uv}; // 2 + 2
float4 zero = float4{}; // all zeros
A count that lands in between is an error; unlike a struct or array literal, a vector does not zero-fill a short list:
float4 bad = float4{1.0f, 2.0f}; // error: takes one value
// or all of its components
Matrices require all of their values. These rules are the same inside
@shader functions. A vector literal is identical on the CPU
and the GPU.
The 256-bit wide types map to native AVX2 instructions under
--target windows-avx2 / linux-avx2 / uefi-x64-avx2. On other x64
targets they lower to two 128-bit halves; on ARM64 and wasm to scalar
code. The operations on them are the f32x8_* / i32x8_* / i8x32
SIMD intrinsics; f32x8 carries no operators of its own, i8x32 has
& | ^ ~ and i32x8 has + - & | ^ ~.
| Type | Size | Description |
|---|---|---|
f32x8 |
32 bytes | 8x f32 vector |
i32x8 |
32 bytes | 8x i32 vector |
i8x32 |
32 bytes | 32x i8 vector |
Component access
Single components via .x, .y, .z, .w (or .r, .g, .b, .a):
f32 x = pos.x; // scalar access
pos.y = 5.0f; // scalar write
Swizzle
Multi-component swizzle returns a new vector. Supports .xyzw and .rgba naming
(no mixing). Reordering and duplication allowed:
float2 xy = pos.xy; // first two components
float3 bgr = color.bgr; // reversed color channels
float2 xx = pos.xx; // duplicated component
float4 rev = color.wzyx; // full reverse
Arithmetic
Component-wise +, -, *, / and unary -. Scalar broadcast supported:
float4 a = float4{1.0f, 2.0f, 3.0f, 4.0f};
float4 b = float4{10.0f, 20.0f, 30.0f, 40.0f};
float4 sum = a + b; // {11, 22, 33, 44}
float4 scaled = a * 3.0f; // {3, 6, 9, 12}
float4 neg = -a; // {-1, -2, -3, -4}
int4/uint4 support the integer operators + - * & | ^ << >> ~
and unary -, all compiling to packed 128-bit SIMD. The shift
count is a uniform scalar; per-lane vector counts are not allowed.
<< is sign-agnostic. >> is arithmetic for int4, logical for
uint4.
int4 v = int4{1, 2, 3, 4};
int4 w = (v & 6) | (v << 2); // pand / pslld / por
int4 n = ~v; // bitwise NOT
Integer-vector division scalarizes (SSE2/NEON have no packed integer
divide): one divide per lane, following scalar integer rules. For an
unsigned power-of-two divisor use >>.
SIMD Vectors (f64x2, i64x2, u64x2)
The 2-wide 64-bit SIMD vectors. f64x2 supports component-wise
+, -, *, /, and the same operators against a scalar f64
(broadcast, either order). Components are .x and .y.
i64x2 / u64x2 support the operators that map to packed SIMD:
+ - & | ^ ~, unary -, and << >> by a uniform scalar count
(>> is logical for u64x2, arithmetic for i64x2). A vector shift
count past the lane width follows the scalar rule in each lane: 0 for
<< and logical >>, the sign for arithmetic >>. * / % are
rejected (no 64-bit packed multiply or divide); scalarize those
explicitly. i8x16 has the bitwise & | ^ ~ only; cast() to
u64x2 or int4 reaches lane arithmetic.
i64x2 v = i64x2{5000000000, -3000000000};
i64x2 d = (v + v) - i64x2{1, 1}; // packed 2x64
i64x2 s = v >> 2; // arithmetic, sign-preserving
u64x2 m = u64x2{0xFF00, 3} & u64x2{0x0FF0, 1};
For more info see SIMD intrinsics in built-in functions.
Literals
42 // integer (coerces to any int type)
0xFF // hexadecimal
0b1010 // binary
1_000_000 // digit separators
3.14 // f64 float
3. // f64 float (trailing dot, = 3.0)
.5 // f64 float (leading dot, = 0.5)
0.5f // f32 float (suffix)
3.f // f32 float (suffix after trailing dot)
2f // f32 float (= 2.0f, integer-with-suffix)
1e6 // f64 scientific (= 1000000.0)
1.5e-3 // f64 scientific with negative exponent
2.5E+10 // f64 scientific (uppercase E, explicit '+')
1.5e3f // f32 scientific
"hello" // string (type str)
'A' // ASCII char literal (Unicode codepoint)
'Γ©' // UTF-8 char literal (codepoint 0xE9)
'π' // non-BMP char literal (codepoint 0x1F389)
'\u{1F389}' // Unicode escape in char literal
true, false // bool
null // null pointer
__line__ // line number of the token (integer literal)
__file__ // basename of the current file (string literal)
UTF-8 in source
Source files are byte-transparent UTF-8. Raw multi-byte characters
pass through unchanged in string literals, comments, and character
literals. Identifiers are restricted to ASCII
([A-Za-z_][A-Za-z_0-9]*), by design.
- String escapes:
\n \t \r \0 \\ \' \" \xNN \u{HEX}where\u{HEX}is 1-6 hex digits naming a Unicode codepoint; the string literal emits the UTF-8 encoding of that codepoint (1-4 bytes). Surrogates (U+D800-U+DFFF) and codepoints aboveU+10FFFFare rejected. - Char literals carry a Unicode codepoint in an integer-literal
node.
u32 c = 'π';is a clean assign. Narrowing to a type that can't hold the codepoint is a compile error:u8 c = 'π';rejects withcharacter literal value 127881 does not fit in u8. - A UTF-8 BOM (
EF BB BF) at the start of a source or included file is silently skipped. Mid-file BOMs are not skipped.
Pointers
i32* p = &x; // pointer to i32
*p = 42; // dereference
p.field // auto-dereference (no -> needed)
p + n // pointer arithmetic (advances by n * sizeof(*p))
q - p // element distance, an i64; both point at the same element type
*p++ = expr; // write to *p, then p++ (postfix yields old)
*p-- = expr; // write to *p, then p-- (postfix yields old)
Arrays
i32[10] arr; // fixed-size array
i32[4] arr = {1, 2, 3, 4}; // array with initializer
i32[8] part = {1, 2}; // partial initializer; the rest zeroes
arr[0] // indexing (bounds-checked by default)
arr = {5, 6, 7, 8}; // whole-array assignment
i32[4] b = arr; // value copy
b = arr; // whole-array copy assignment (lengths must match)
An array converts to a pointer to its first element where a T* is
expected. A ternary with two array arms does the same in a pointer
context, so the arms may differ in length.
u8[4] small = {1, 2, 3, 4};
u8[8] large = {5, 6, 7, 8, 9, 10, 11, 12};
u8* tab = wide ? large : small; // u8*, aliases one of the two
var copy = flag ? small : small2; // u8[4], a copy of the chosen one
Slices
[]i32 s; // slice: pointer + length
s.ptr // data pointer
s.len // element count
s[i] // indexed access
s[1..4] // subslice [start, end)
A fixed array has .len as well: i32[4] a; a.len is the constant 4.
An array, slice or pointer can be written through a slice target. This writes the elements given, starting at the offset, and leaves everything else alone:
i32[6] a = {1, 2, 3, 4, 5, 6};
a[2..] = {7, 8}; // a is now {1, 2, 7, 8, 5, 6}
a[1..4] = {7, 8, 9}; // a[1], a[2], a[3]; the rest untouched
An open range takes its length from the initializer. A range with
both bounds must be filled exactly. Non-literal bounds are left to
the runtime bounds check. The elements may be of any type, structs,
unions and vectors included (ps[0..2] = { P{1, 2}, {3, 4} }). The
right-hand side is always a brace list; copy from another array with a
loop.
That is different from whole-array assignment, which replaces the array and zeroes anything the initializer did not cover:
a = {9}; // a is now {9, 0, 0, 0, 0, 0}
Multi-dimensional arrays
A shape after a colon declares a multi-dimensional array. The leftmost dimension is the outermost, the same order the subscripts read:
i32:[2][3] m = { {1, 2, 3}, {4, 5, 6} }; // 2 rows of 3
i32 v = m[1][2]; // row 1, column 2 (C order)
u8:[16][32][32] vox; // 3-D
type Grid = f32:[64][64]; // shapes work anywhere a type does
i32[3] row = m[1]; // partial index: copy a row out
m[0] = row; // copy a row in
The layout is one contiguous row-major block: the last dimension
varies fastest and sizeof is the product of all dimensions. Each
dimension is bounds-checked independently. Nested brace init,
whole-array copy, and row copy work per level.
A shape needs at least two dimensions. f32:[4] is an error; a
one-dimensional array is f32[4]. The C-style f32[2][3] is also an
error.
A trailing * makes a pointer to an array or shape type:
f32 corner(f32:[2][3]* a) { return a[0][1][2]; }
corner(&m); // pointer to the whole shape
void scale(f32[3]* row) { } // pointer to one row
scale(m); // a shape value decays to its first row
scale(&m[1]);
float4x4 is a distinct type, not f32:[4][4]: it is column-major
and SIMD-backed, with no implicit conversion between the two.
Strings
minc has two string types: str (borrowed view) and string (owned, heap-allocated).
str β borrowed view
str s = "hello"; // UTF-8 string view: { u8* data; i32 len }
s.len // byte length (5)
s.data // raw u8 pointer
str is a built-in struct { u8* data; i32 len }, a non-owning view.
String literals are type str. Not null-terminated by default.
String literals are read-only. Identical literals share one copy, and on
native targets the bytes live in the executable's code section. A
store through a literal's pointer faults. Copy into a string or a
buffer to modify text.
Copying a str copies the pointer and length (16 bytes), not the underlying data.
Functions that only read strings should take str parameters.
string β owned string
string s = string("hello"); // heap-allocates and copies
defer free(s); // freed at scope exit
print("{}\n", s); // implicit string β str conversion
string has the same layout as str ({ u8* data; i32 len }) but owns its data.
The compiler tracks ownership and enforces cleanup:
string s = string("hello");
// error: owned string 's' must be freed or moved before scope exit
Ownership rules:
- string locals must be freed or moved before scope exit
- free(s) β frees the string's data
- defer free(s) β idiomatic cleanup (freed at scope exit, usable until then).
A string with a pending defer free cannot be moved or returned.
- using string s = ...; β freed at scope exit unless moved out with
move(s) (see Using)
- move(s) β transfers ownership, invalidates source
- return s β implicit move (transfers to caller)
- string β str β implicit conversion (safe borrow for function calls)
Tracking follows control flow. A free on a branch that returns does not clear other paths. Reassigning a freed variable makes it live again. The new value needs its own free.
string base = format("base");
for i32 i = 0; i < count; i++ {
if match(i) {
free(base);
return i; // freed on this path
}
}
free(base); // freed when the loop finishes
string s = format("a");
free(s);
s = format("b"); // live again, needs its own free
free(s);
A free that can reach a later use is an error. A local declared inside a loop body dies with each iteration:
for i32 i = 0; i < count; i++ {
string cand = make(i);
if skip(cand) {
free(cand);
continue; // ok: cand dies with the iteration
}
consume(cand);
free(cand);
}
// Construction
string s = string("literal"); // from string literal
string s2 = string(some_str); // from str view (copies)
string s3 = str_concat(a, b); // functions that allocate return string
Using a string after free(s) or move(s) is a compile error
("use of freed string" / "use of moved string").
String utilities (lib/str.mc)
#include "lib/str.mc"
// Concatenation (returns owned string)
string full = str_concat("hello ", "world");
defer free(full);
// String builder
str_buf sb;
str_buf_init(&sb);
str_buf_add(&sb, "hello ");
str_buf_add(&sb, "world");
str result = str_buf_to_str(&sb); // borrows sb's buffer
// ... use result ...
str_buf_free(&sb);
// String formatting (built-in, returns owned string)
string msg = format("hello {}", name); // single placeholder
string line = format("{} + {} = {}", a, b, a+b); // multiple
defer free(msg);
defer free(line);
// format() supports any printable type: integers of every width, enums, floats, bools,
// str, string, pointers. A struct, union or array argument is an error.
// Use {} as placeholder β arguments are matched left-to-right.
// The same {} syntax works with print() and eprint() (which write to stdout/stderr):
print("x = {}\n", x); // prints to stdout
eprint("error: {}\n", msg); // prints to stderr
// Float formatting note: `{}` for f32/f64 is a *display* format, not
// a round-trip format: up to 6 fractional digits, trailing zeros
// trimmed (at least one digit kept: "3.0"), "nan"/"inf"/"-inf" for
// specials, scientific notation outside the six-digit window.
// For exact, shortest round-trip output, import the Ryu formatters:
// import format_f64; string s = format_f64(0.1); // "0.1"
// import format_f32; string s = format_f32(0.1f); // "0.1"
// Each returns the shortest decimal string that parses back to the
// same bits. f32_to_str / f64_to_str write into a caller buffer
// instead (decimal for 1e-4 <= |x| < 1e21, scientific outside).
// Conversion
str view = str_from_cstr(c_string); // u8* β str (no copy)
u8* cstr = str_to_cstr(s); // str β null-terminated u8* (allocates)
Declarations
Variables
i32 x = 42; // explicit type
var y = 42; // type inference, integer literals default to i32,
// i64 above the i32 range and u64 above the i64
// range (18446744073709551615, 0x8000000000000000)
const i32 MAX = 100; // compile-time constant
i32 g_count = 0; // global variable
threadlocal i64 t_count = 0; // one copy per thread
A name is declared once per scope; a nested block may shadow it.
threadlocal gives a file-scope variable one copy per thread,
initialised from its constant initialiser; &x is this thread's copy.
On the uefi targets the runtime provides the blocks
(doc/DESIGN_runtime_contract.md).
Functions
i32 add(i32 a, i32 b) {
return a + b;
}
void greet() {
print("hello\n");
}
A function with a return type must return on every path. Control reaching the end of the body is a compile error:
i32 pick(i32 x) {
if x > 0 { return 1; }
} // error: missing return in function 'pick'
Function overloading
Functions with the same name can be overloaded if their parameter types differ. Resolution is by exact match (no implicit conversions for disambiguation):
f32 dot(float2 a, float2 b) { return a.x * b.x + a.y * b.y; }
f32 dot(float3 a, float3 b) { return a.x * b.x + a.y * b.y + a.z * b.z; }
f32 dot(float4 a, float4 b) { return a.x * b.x + a.y * b.y + a.z * b.z + a.w * b.w; }
float3 v = float3{1.0f, 0.0f, 0.0f};
f32 d = dot(v, v); // calls float3 overload
An overloaded function used as a value takes the overload its
function-pointer type names, in a declaration, an assignment or an
argument: fn(float3, float3): f32 f = dot;. Without a matching
function-pointer type it is a compile error.
A generic function may share the name of an overload set, as when two
imported modules each declare show. The overloads are tried first;
when none matches, the call instantiates the template.
An enum member argument is its enum where a candidate expects that
enum at its position: with f(i32) and f(Mode), f(Fast) calls
f(Mode).
Structs
struct Point {
i32 x;
i32 y;
}
Point p = Point{3, 4}; // positional literal
Point p2 = Point{.y = 10, .x = 5}; // named fields
p = {7, 8}; // assignment RHS: type inferred from target
var (x, y) = make_point(1, 2); // var destructuring
A field that is itself an array or struct takes a nested {...}
initializer; the inner brace-init has no type prefix and resolves
against the field type. Nesting is arbitrary (struct in array in
struct β¦). An array initializer may supply fewer elements than the
array holds. The remaining elements are set to zero.
struct Row { i32[3] v; i32 n; }
struct Grid { Row[2] rows; }
Row r = Row{ {1, 2, 3}, 0 }; // array field
Grid g = Grid{ {Row{{1,2,3},0}, {{4,5,6},1}} }; // mix of explicit and bare braces
Comma-separated field declarations share the parsed type. Mirrors
local + global multi-decl: i32 a, b, c; declares three fields all
of type i32. Useful for packed records like RGBA color components:
struct InstanceData {
f32 pos_x, pos_y, pos_z;
u8 r, g, b, a;
}
Structs and arrays are zero-initialized by default when declared without an initializer.
Use noinit to skip zero-initialization for performance-critical code:
Point p; // p.x = 0, p.y = 0 (zero-init)
noinit Point q; // uninitialized (faster, use with care)
i32[1024] buf; // all zeros
noinit u8[4096] scratch; // uninitialized buffer
Use unused to suppress unused-variable warnings (e.g., platform-specific variables):
unused bool has_feature = false; // no warning if unused on this platform
unused i32 hlsl_flags = 0; // used on Windows, not Linux
Variables prefixed with _ also suppress the warning: i32 _reserved = 0;
Enums
enum Color { RED, GREEN, BLUE } // auto-incremented (0, 1, 2)
enum Flags { A = 1, B = 2, C = 4 } // explicit values
// Values are integer constant expressions: literals (incl. negatives
// and char literals), parens, unary `! ~ -`, and the full C integer
// operator set (`+ - * / % << >> & | ^ == != < <= > >= && ||`).
// Earlier members of the same enum are in scope. Two enums are
// distinct types: a member of one is not a value of the other, and
// cast() converts through the integer.
enum Errno {
OK = 0,
EBADF = -9,
EAGAIN = -11,
LAST = EAGAIN - 1 // -12; auto-increment continues from here
}
enum Bits {
MASK = (1 << 8) - 1, // 255
HIGH = 1 << 31
}
Enum values are i32 constants.
An enum name and an enum-typed variable convert by different rules:
enum Feature { NONE, FACE, WIDE = 300 }
u8 a = FACE; // OK: a member is an integer literal
u8 b = WIDE; // ERROR: 300 doesn't fit u8, same as any literal
Feature f = WIDE;
i32 c = f; // OK: same width
i64 d = f; // OK: widens like i32
f64 e = f; // OK: widens like i32 (lossless)
u8 g = f; // ERROR: narrowing, needs cast(u8, f)
f32 h = f; // ERROR: i32 -> f32 loses precision, needs a cast
u32 limit = 400;
if f < limit { } // ERROR: mixed signed/unsigned, same as i32 vs u32
if FACE < limit { } // OK: a member is a literal, and literals coerce
i32 raw = 1;
Feature back = raw; // OK: same width
An enum member and a union variant may share a name; the name is the one its context expects (an enum declaration, case label or comparison operand takes the member, a union context the variant).
A member is an i32 constant literal. It coerces wherever a literal
does, and where its own enum is expected it is that enum: Box<E> b =
wrap(B); infers T = E, while wrap(B) on its own is a Box<i32>.
An enum-typed value is an i32 value: it widens where i32 widens
and needs a cast() to narrow.
Tagged unions
union Option<T> {
Some(T),
None,
}
union Result<T, E> {
Ok(T),
Err(E),
}
union Token {
Number(i32),
Ident(u8*, i32),
Plus,
Eof,
}
Construction:
Option<i32> x = Some(42);
Option<i32> y = None;
Result<i32, i32> r = Ok(100);
Token t = Number(123);
Token t2 = Plus;
A generic union's variant takes its type from where it lands: a
declaration, an assignment, a return, a call argument, or the payload
of an outer variant (Option<Option<i32>> n = Some(Some(1));). Each
payload must match its field type; a bare brace payload takes the
field's type (Some({ 1, 2 }) for a struct payload).
A generic struct's literal names its type arguments: Box<i32>{ 4 },
Pair2<i32, f64>{ .a = 1, .b = 2.5 }. A bare brace takes the declared
type, at file scope as in a function: Pair g = { 1, 2 };.
A union may be initialized at file scope like a struct, with constant
payloads: Option<i32> g = Some(3);, Token t = Plus;, Shape[2] s =
{ Rect(2, 3), Dot };, threadlocal Option<i32> tl = Some(1);.
Switch with pattern matching (every variant must be covered):
switch x {
case Some(val): { print("{}\n", val); }
case None: { print("none\n"); }
}
Example: error handling with Result
union Result<T, E> {
Ok(T),
Err(E),
}
Result<i32, str> parse_int(str input) {
// ... parsing logic ...
if valid { return Ok(value); }
return Err("invalid number");
}
i32 main() {
Result<i32, str> r = parse_int("42");
switch r {
case Ok(val): { print("parsed: {}\n", val); }
case Err(msg): { print("error: {}\n", msg); }
}
return 0;
}
Unions are ordinary values: they can be passed, returned, stored in arrays and structs, and used with generics.
Unsafe unions
unsafe_union is a C-style union. All members share offset 0,
sizeof is the largest member, there is no tag and no active-member
tracking. Reading one member after writing another reinterprets the
bytes.
unsafe_union FloatBits { u32 i; f32 f; }
FloatBits b;
b.i = 0x40490FDB;
f32 x = b.f; // β 3.14159 β bytes reinterpreted as f32
unsafe_union Mix { u8 lo; u32 w; u64 q; } // sizeof(Mix) == 8
Usable as a variable, parameter, return value, struct field, array
element, or behind a pointer, with .member access.
Initializer literals use struct-literal syntax. At most one field may be initialized; the remaining bytes come from the implicit zero-fill:
unsafe_union FB { u32 u; f32 f; }
const FB almostone = FB{ .u = 0x3f7fffff }; // named (recommended)
FB minval = FB{ (127 - 13) << 23 }; // positional β first field
FB zero = FB{}; // empty β all-zero storage
FB declared; // implicit zero (no init)
noinit FB raw; // skip zero-fill
Multi-field initializers like FB{ .u = 1, .f = 2.0f } or FB{ 1, 2.0f }
are rejected. noinit FB x = FB{...}; is rejected (same rule as for
other types).
Anonymous nested struct/union members
A struct member's type can be an inline anonymous struct { β¦ } or
unsafe_union { β¦ }. Two shapes:
Named field, anonymous type β nested-field access via the field name:
struct Packet {
i32 header;
struct { i32 flag; i32 seq; } stuff;
i32 payload_len;
}
Packet p;
p.stuff.flag = 1;
Packet q = Packet{ .header = 1, .stuff = { .flag = 2, .seq = 3 } };
Transparent (no field name) β the inner aggregate's fields are
accessible directly on the enclosing type (C11 anonymous-member
form). Combined with unsafe_union:
struct Color {
unsafe_union {
struct { u8 r; u8 g; u8 b; u8 a; }
u32 rgba;
}
}
Color c;
c.r = 0xff;
c.rgba = 0x01020304;
Color d = Color{ .r = 1, .g = 2, .b = 3, .a = 4 };
Color e = Color{ .rgba = 0xAABBCCDD };
&c.r and &c.rgba address overlapping storage, the same rule as
a plain unsafe_union.
Rules: the inline aggregate must be anonymous, and only struct and
unsafe_union can be transparent members. Promoted names may not
collide with each other or with the parent's direct fields. A type
with transparent members takes named-field initializers only.
Function pointers
fn(i32, i32): i32 op = add;
i32 result = op(3, 4);
Function pointers can also be used as struct fields:
struct Handler {
fn(i32, i32): i32 op;
i32 id;
}
Handler h;
h.op = &add;
i32 result = h.op(3, 4); // call through struct field
Type aliases
type Size = i64;
type Byte = u8;
type IntPtr = i32*;
type BinOp = fn(i32, i32): i32;
Size x = 42; // same as i64 x = 42
BinOp op = &add; // alias for function pointer type
struct Calc { BinOp f; } // alias as struct field type
The alias name becomes interchangeable with the underlying type.
Aliases may be declared at file scope or inside a function body. A body-level alias is scoped to its enclosing block:
i32 main() {
type Row = i32[3]; // local to main
Row[2] grid = { {1,2,3}, {4,5,6} };
{
type Row = f32[2]; // shadows the outer Row in this block
Row v = {0.5f, 1.5f};
}
// outer Row is in effect again here
return grid[1][2];
}
Two differences from a file-scope alias:
- No forward references. A local alias is visible only to statements after it, like a local variable. A file-scope alias can be used anywhere in the file, including above its declaration.
- Shadowing is allowed β of a file-scope type, or of an alias from an enclosing block. Redeclaring the same name twice in one block is an error.
Local aliases are a naming convenience only; they declare no storage and generate no code.
Statements
Control flow
if condition { ... }
else if condition { ... }
else { ... }
while condition { ... }
for i32 i = 0; i < 10; i++ { ... }
for ; i < 10; i++ { ... } // no init (use existing variable)
for i32 i in 0..10 { ... } // range-based (0 to 9 inclusive)
for i32 x in items { ... } // each element of a slice or fixed array (a copy)
switch value {
case 1, 2, 3: { ... } // multi-value case
case 4: { ...; fallthrough; } // explicit fall-through (last stmt)
case 5: { ... }
default: { ... }
}
break case; // exits innermost switch case
fallthrough; // last stmt of case; falls to next
break; // exits nearest enclosing loop
continue;
return expr;
A range loop counts an integer of the declared type. A narrow counter
(u8, i16, ...) steps at its own width, and a literal end past its
range is an error, since the counter could never reach it.
No implicit fall-through; opt in per case with fallthrough;. Any
statement after break / break case / continue / return /
fallthrough at the same block level is a compile error
(unreachable statement). Braces required on every case body. Two
cases naming the same value are an error.
Defer
i64 fd = open("file.txt", 0);
defer close(fd);
// close(fd) executes automatically at block exit, LIFO order
Multi-statement cleanup via the block form:
defer {
close(handle);
counter = counter - 1;
log_exit("done");
}
Statements inside defer { ... } run in source order at scope exit.
The block as a whole follows the same LIFO order as single-statement
defers relative to its siblings.
return, break, and continue are rejected inside defer { }:
they would skip later defers. Loops nested inside the deferred body
can still break / continue against their own loop.
Leaving blocks early runs their defers. return evaluates its value,
then runs every pending defer in the function. The value is read
before the defers run, whatever its type: a defer that writes the
returned variable, or a field of it, does not change what the caller
receives. break and continue run the defers of each
block they leave, innermost first, out to the body of the loop they
target (a labeled one included). break case and fallthrough do the
same out to the case body.
defer free(x) ownership tracking applies only to the single-
statement form. defer { free(a); free(b); } does not mark a or
b as defer-freed. Use one defer free(x); per resource for the
tracking.
Using
using string path = path_join(dir, name);
using var buf = alloc<u8>(n);
using T x = e; declares x and frees it with free at the exit of
the enclosing block, in the same LIFO order as a defer free(x); on
the next line. The type is string or a pointer, and the initializer
is required. Other cleanups use defer.
The value leaves only through move():
string load(str path) {
using string s = file_read_str(path);
if s.len == 0 { return string("empty"); } // s is freed
return move(s); // s moves to the caller
}
using Node* n = new(Node);
if keep { g_cache = move(n); } // n is zeroed; the exit frees null
move(x) returns the value and zeroes x, so the free at block exit
does nothing; reading x after the move is an error. move() accepts
any using variable, pointers included.
return x, free(x), defer free(x), assignment and ++ / -- on a
using variable are errors. using takes a single declaration.
@strict_float
@strict_float on a function pins its floating-point evaluation to
exactly what the source says: the optimizer performs no FMA
contraction and no float-loop vectorization or reduction reordering
on its operations. Calls to it may still be inlined β the spliced
operations keep their strict semantics at every call site. The result
bits are identical on every target and at every optimization level.
Use it for reproducible numeric kernels.
@strict_float
f64 kernel(f64 a, f64 b, f64 c) {
return a * b + c; // always mul then add, on every target
}
Without @strict_float, x64 and arm64 still agree with each other:
they contract the same operations and split reductions the same way,
so plain float arithmetic gives the same bits on both at a given
optimization level. Those bits differ from the @strict_float result
and from wasm. The estimate builtins (rsqrt4_fast, etc.) are
target-defined.
The annotation covers the function's own operations only. A kernel is
reproducible when every function it calls is @strict_float too.
@must_use and ignore
@must_use on a function declaration emits a warning when the
caller drops its return value:
@must_use
i32 try_parse(u8* s) { ... }
i32 main() {
try_parse(input); // warning: result of 'try_parse' is unused
i32 r = try_parse(input); // OK β assigned
if try_parse(input) == 0 { } // OK β used in a condition
return try_parse(input); // OK β returned
ignore try_parse(input); // OK β explicitly dropped
}
ignore <expr>; evaluates the expression for its side effects and
discards the value. It silences @must_use at one call site. The
warning is also suppressed inside defer { ... } (the deferred call
has no caller to assign to).
@inline
@inline on a function puts its body into every direct call site,
whatever its size and whatever it calls itself, like C's
always_inline:
@inline
Match find_match(u8* ip, u8* end, Params p) { ... }
Calls through a function pointer still reach an out-of-line copy.
These are compile errors: @inline on a function that calls itself,
on @inline functions that call each other in a cycle, on a variadic
function, on an @interrupt handler, and on a function that switches
fibers. Without the annotation the compiler decides by size and
shape.
Bare blocks
Bare { } blocks create a new scope. Variables declared inside are not visible
outside, and defer statements fire at block exit:
i32 x = 1;
{
i32 y = 2;
defer print("leaving\n");
// y is visible here
}
// y is out of scope; defer has fired
Compile-time conditionals
when os(windows) { ... }
else when os(linux) { ... }
else when os(macos) { ... }
else { ... }
when arch(x64) { ... }
when arch(arm64) { ... }
when arch(wasm32) { ... }
when defined(DEBUG) { ... }
// Shader backend (d3d11, metal, opengl, opengles, webgpu)
when gpu(opengl) { ... }
// Combine with || and && and !
when os(linux) || os(macos) { ... }
when os(windows) && arch(x64) { ... }
when !os(wasm) { ... }
// Integer expressions over -D / @define values: a C-#if-style
// grammar of integer literals, config names, the predicates above,
// and ! ~ - (unary) then (lowβhigh) || && | ^ & == != < <= > >=
// << >> + - * / % with parentheses. Non-zero is truthy.
@define "API_VERSION" 3
@define "MAX_STRIDE" 16
when API_VERSION >= 3 { ... } else { ... }
when defined(USE_SSE) && SSE_LEVEL >= 2 { ... }
when (NCHANNELS * 4) > MAX_STRIDE { ... }
&& / || short-circuit: the dead operand is parsed but not
evaluated. when defined(X) && X >= 3 is legal even when X is
undefined. A bare config name that was never @defined / -D'd,
outside a dead branch, is an error; use defined(NAME) to test
presence. Unlike C, an undefined identifier is not implicitly zero.
Division / modulo by zero is an error.
@define "NAME" with no value is @define "NAME" 1. @define "NAME"
42 sets an integer value (negative and _-separated literals
allowed). On the command line, -DNAME=42 sets a value; -DNAME
sets 1.
Dead branches are skipped at parse time (no runtime overhead).
Available os values: windows, linux, macos, wasm, ios, android.
Available arch values: x64, arm64, wasm32.
Compiler version
@minc_min_version "0.9.18" // refuse to compile with anything older
when MINC_VERSION >= 9018 { ... } // gate on the running compiler
@minc_min_version : the oldest compiler that can build the file.
MINC_VERSION encoded as major*1000000 + minor*1000 + patch.
0.9.18 is 9018, 1.0.0 would be 1000000.
Expressions
Operators
| Category | Operators |
|---|---|
| Arithmetic | + - * / % |
| Inc/Dec | ++x x++ --x x-- (statement + expression position; see below) |
| Comparison | == != < > <= >= |
| Logical | && \|\| ! |
| Bitwise | & \| ^ ~ << >> |
| Assignment | = += -= *= /= %= &= \|= ^= <<= >>= |
| Ternary | condition ? true_expr : false_expr |
| Cast | cast(Type, expr) |
| Sizeof | sizeof(Type) or sizeof(expr) (returns size in bytes) |
| Address-of | &expr |
| Deref | *expr |
| Member | expr.field |
| Index | expr[index] |
Evaluation order
Operands, call arguments, struct literal fields and array initializer
elements are evaluated left to right. An assignment evaluates its
target first, then the value: a[next()] = next() indexes with the
first call, and a[j] = j++ stores into the old a[j]. && and ||
stop at the first operand that decides the result.
Increment / decrement
Prefix and postfix ++ / -- work in both statement and expression
position. The operand must be an lvalue of integer (non-bool) or
pointer type. Prefix yields the new value; postfix yields the old
value, then increments / decrements.
i32 i = 5;
i32 a = i++; // a = 5, i = 6 (postfix yields old)
i32 b = --i; // b = 5, i = 5 (prefix yields new)
array[idx++] = x; // write to array[idx], then idx++
while --n > 0 { } // decrement first, then compare
u8 c = *--end; // step back one byte, read it
Pointer increments step by sizeof(*T). p++ on an i32*
advances by 4 bytes.
The operand's address is evaluated once. Calls in it run once:
next_slot()[i]++ and counts[hash(key)]++ call each function a
single time, as += does. A field of a struct returned by value from
a call has no address. make_point().x++ is an error.
Named arguments
add(a: 3, b: 4) // named
add(b: 4, a: 3) // reordered
add(3, b: 4) // mixed positional + named
Type conversions
| Conversion | Rule |
|---|---|
Widening within a signedness: i8βi16βi32βi64, u8βu16βu32βu64 |
Implicit |
Unsigned β wider signed: u8βi16/i32/i64, u16βi32/i64, u32βi64 |
Implicit |
i32/u32 β f64, f32 β f64 |
Implicit |
Same-width reinterpret: i32βu32, i64βu64, β¦ |
Implicit on assignment only; bits unchanged (-1 β 4294967295) |
Signed β wider unsigned: i8/i16/i32 β a wider uβ¦ |
Implicit on assignment only; sign-extends, bit-identical to cast() |
Anything lossy: narrowing, i64βf64, floatβint, ptrβint |
Explicit cast() required |
"On assignment" means places with a target type: init, return,
argument passing, store. In an expression there is no target width
to convert toward. Mixed signed/unsigned still errors:
i32 + u64 is rejected; cast one side.
i64 x = 42; // i32 -> i64 (implicit)
i32 y = cast(i32, x); // i64 -> i32 (explicit, narrowing)
f64 f = y; // i32 -> f64 (implicit, i32 fits in f64)
i32 n = cast(i32, 3.14); // f64 -> i32 (explicit, truncates)
u32 r = y; // i32 -> u32 (implicit, same-width reinterpret)
u64 sz = y; // i32 -> u64 (implicit, sign-extends)
u64 bad = x + sz; // ERROR: i64 + u64 mixed sign in an expression
Integer promotion
(Runtime value behavior of the promoted types is under "Numerics" below.)
Binops on two narrow operands (u8, u16, i8, i16):
| Operators | Result type |
|---|---|
+ - * / % << >> |
i32 / u32 (the result can exceed the operand width); an opposite-sign pair (u16 + i16) is i32, since both values fit |
& \| ^ |
operand width (u8 ^ u8 is u8) |
unary ~ |
operand width (~a is a u8 for u8 a; 254 when a is 1) |
unary - |
i32 / u32, like the arithmetic binops (-i16 is i32) |
compound assignment (+=, <<=, β¦) |
lvalue width; no cast needed, the store wraps |
Byte-assembly works without per-byte casts:
u32 v = b[0] | (b[1] << 8) | (b[2] << 16) | (b[3] << 24);
u32 sum = u8_a + u8_b; // 255 + 16 = 271
Storing a promoted result into a narrower location is narrowing and
requires cast(); compound assignment is exempt:
u8 bits = 0x80;
bits = bits << 1; // error: u32 -> u8 needs a cast
bits = cast(u8, bits << 1); // wraps to 0 at u8 width
bits <<= 1; // OK: wraps at u8 width, no cast
Promotion stops at 32 bits: u32 op u32 stays u32 and wraps, and
assigning the result to a u64 cannot recover the lost bits. Widen
before the operation to keep them:
u64 hi = u32_val << 32; // 0 β shifted at u32 width
u64 hi = cast(u64, u32_val) << 32; // operates at u64 width
Mixed signed/unsigned
(Per-operator runtime details are under "Numerics" below.)
Mixing signed and unsigned integer operands:
| Operators | Mixed-sign operands |
|---|---|
== != < > <= >= |
Error |
+ - * / % |
Error |
>> |
Error (sign-sensitive: arithmetic vs logical shift) |
& \| ^ <<, same width |
OK; the result takes the left operand's type |
| Any operator vs an integer literal or enum member | OK; the literal coerces |
A literal coerces to the other operand's type, and a fold of two literals takes the wider operand's type, the unsigned one on a tie. A decimal literal above the i64 range is a u64 and does not fit a signed target; a hex or binary pattern fills either 64-bit type.
i32 a = -1;
u32 b = 100;
if a < b { } // ERROR: mixed signed/unsigned comparison
i32 c = a + b; // ERROR: mixed signed/unsigned arithmetic
i32 d = a >> b; // ERROR: >> is sign-sensitive
if b == 0 { } // OK: literal coerces
if b < SG_INVALID_ID { } // OK: enum member coerces
u32 out = b | a; // OK: same-width bitwise β u32
i32 e = a >> 5; // OK: literal coerces to a's type
Use explicit cast() to resolve mixed-sign errors:
if cast(u32, a) < b { } // OK: explicit cast
i32 d = cast(i32, b) >> 4; // OK: arithmetic shift chosen explicitly
Floating-point semantics
(IEEE 754 runtime behavior is under "Numerics" below.)
minc applies these floating-point optimizations by default:
-
FMA contraction:
a * b + cmay fuse into one FMA instruction. FMA rounds once instead of twice. Results can differ by at most 1 ULP from the unfused form. This is the same contraction license as MSVC/fp:contractand clang's default-ffp-contract=on. -
Float strength reduction:
x * 2.0may be replaced withx + x. -
SIMD vectorization: scalar float loops may be auto-vectorized to packed SIMD instructions, which can change the order of floating-point operations.
Programs that depend on exact IEEE 754 two-rounding semantics (e.g., Kahan compensated summation, Dekker exact multiplication) should be aware of these transformations. There is no flag to disable FP contraction.
Generics
Monomorphized generics with <T> syntax. Zero runtime overhead: each
instantiation generates specialized code at compile time.
Generic functions
T identity<T>(T x) { return x; }
void swap<T>(T* a, T* b) { T tmp = *a; *a = *b; *b = tmp; }
// Type inference from arguments
swap(&a, &b); // infers T from pointer type
// Explicit type argument (turbofish syntax)
identity<i64>(42); // forces T = i64
identity<i64>(1 << 33); // the argument is typed by the i64 parameter
fn(i64): i64 f = identity<i64>; // the instantiation as a value
// Bidirectional inference from return type
Pair<i32> p = make_default_pair(10); // infers T = i32 from expected type
Generic structs
struct Pair<T> { T first; T second; }
Pair<i32> p;
Pair<i64> q;
// Multi-parameter
struct Triple<A, B, C> { A x; B y; C z; }
// Nested generics
Pair<Pair<i32>> nested;
// Generic fields may reference other generic types, including the
// struct's own parameters β and the struct itself
struct Slot<V> { V val; i32 state; }
struct Map<V> { Slot<V>* slots; i32 cap; }
struct Node<T> { T val; Node<T>* next; }
Type constraints
Constraints restrict which types can be used with a generic parameter:
T add<T: Numeric>(T a, T b) { return a + b; }
T neg<T: Signed>(T x) { return 0 - x; }
T half<T: Float>(T x) { return x / cast(T, 2); }
struct NumBox<T: Numeric> { T val; }
| Constraint | Types |
|---|---|
Numeric |
i8, i16, i32, i64, u8, u16, u32, u64, f32, f64 |
Integer |
i8, i16, i32, i64, u8, u16, u32, u64 |
Signed |
i8, i16, i32, i64, f32, f64 |
Unsigned |
u8, u16, u32, u64 |
Float |
f32, f64 |
For the narrow integer types the arithmetic promotes to 32 bits like
everywhere else, so a body that returns T writes cast(T, a + b).
Calling a constrained function with the wrong type is a compile error:
add<str>("a", "b"); // error: str does not satisfy Numeric
Modules
Import
// all public symbols directly in scope
import helpers;
import "lib/helpers.mc";
// Selective: only named symbols
import { Vec, vec_push } from "lib/vec.mc"; // from a path (relative to importer)
import { sinf, cosf } from math; // from a library name (lib-search)
// Qualified: access via prefix
import math = "lib/math_helper.mc";
math.add(2, 3);
A qualified import binds the module's public functions, overload
sets, generic functions, globals, consts, types and enum members under
the prefix (math.add(2, 3), math.max<i64>(a, b), math.counter +=
1, geo.Point p = geo.Point{ 1, 2 }, geo.Box<i32> b, geo.Fast,
case geo.Line(n):).
The prefix is the
only way to reach them: in every other file the plain name is the
importer's own declaration, or undefined. Two modules imported this
way may declare the same name, and either may share a name with the
importing file; each module's types are its own namespace. Inside the
module its names stay plain. A module loaded flat anywhere keeps plain
names everywhere. One alias names one module per file.
Quoted import "file.mc"; resolves relative to the importing file;
an absolute path is used as written.
Bare import helpers; searches in order (first hit wins):
| # | Location | |
|---|---|---|
| 1 | helpers.mc next to the importing file |
sibling |
| 2 | lib/helpers.mc from cwd, walking up ancestors |
project lib |
| 3 | compiler's bundled lib/ |
stdlib |
The selective from clause takes either form: from "path.mc" is
path-relative like a quoted import; from name uses the bare-import lib-search.
The listed names are visible only in the importing file; the rest of the
module stays hidden but fully compiled. Listed functions may call
unlisted ones. A full import of the same module, from any file and in
any order, makes all of it visible everywhere. Listing a name the module
doesn't define, or a private one, is an error; so is defining a name
that collides with a hidden name of a selectively-imported module.
Types and enum values are not filtered.
Private
// Single declaration
private i32 helper() { return 42; }
// Block
private {
i32 internal_a() { ... }
struct ScratchState { i32 pos; }
type OpaqueHandle = void;
}
Private declarations are visible only within their declaring file.
This covers functions, globals, and types (struct, enum,
union, unsafe_union, type aliases), including the values of a
private enum and the constructors of a private union. Using a
private type from another file is a compile error ("type 'X' is
private to
A shared library (--shared) exports only export-marked functions
(see Export). A private function is never visible to dlsym /
GetProcAddress.
A private declaration never collides with a same-name declaration in
another file: two files may each declare a private i32 g_state or
a private helper, and each resolves to its own, with its own storage.
Inside its declaring file a private definition shadows an imported
public one of the same name; other files still reach the public name.
This holds uniformly for types, functions, globals, consts, enum
values, and generic templates (instantiable only from their file,
with separate instantiations per file).
extern DLL imports follow the same rule, and a private extern's
signature may differ from another file's declaration of the same
import; both bind the same symbol through one import-table entry.
Public externs for one name must match: same library, same symbol,
and the same signature up to integer signedness and void* versus
typed pointers; a mismatch is a compile error at the second
declaration. Dll-less extern { ... } forward declarations are
exempt; they pair with a definition by name and stay visible
program-wide.
Type redefinition
Declaring one type name twice with different definitions is a
compile error at the later declaration: "conflicting redefinition of
type 'X' (previously defined in
- Identical re-declarations. Two files may declare the same
struct Vec2 { f32 x; f32 y; }. privateshadowing. A private type never conflicts with a same-name type in another file. Mark a local typeprivateto coexist with a library's public name.- The C
typedef enumidiom. Anenum X { ... }plus an integer aliastype X = i32;is one C declaration split in two, common in transpiled headers. The alias resolves the name; the enum carries the value constants. - Forward declarations.
struct X;is compatible with any later definition. Until it is defined only pointers toXcan be used; a value,sizeof(X)ornew(X)of a struct that is never defined is an error.
Export
// Single declaration
export i32 frame_tick(i32 dt_ms) { ... }
// Block
export {
i32 sokol_main_call() { ... }
void on_event(i32 ev_ptr) { ... }
}
Exported functions are the module's ABI: they are surfaced at the platform module boundary so an external host can call them, and they are roots for dead-code elimination. An exported function is kept even when nothing inside the module calls it.
On WASM, each export-marked function appears in the module's
exports section and is callable as instance.exports.frame_tick(...)
from JavaScript. A main() function, if present, is auto-exported. A
WASM module may also omit main entirely and surface its API through
export'd functions or a _start().
On native targets, a shared library (--shared: PE DLL, ELF .so,
Mach-O dylib) exports exactly the export-marked functions, resolved
by GetProcAddress / dlsym. Everything else is internal and subject
to dead-code elimination like in an executable. A shared library with
no exported function is a compile error. In an executable the keyword
only pins the function as a DCE root.
Only functions are exported. export on a global is a compile error:
a library shares state through exported accessor functions.
Include
#include "path/to/file.mc" // textual inclusion (include-once)
#include makes the included file's declarations visible (C-style
usage). private declarations stay scoped to their declaring file;
inclusion does not lift privacy. The path resolves relative to the
including file; an absolute path is used as written.
API Version Tag
Library headers can declare a version, and #include can require a specific version.
This catches mismatches when a program is compiled against the wrong version of a library.
// In the library header (lib/vec.mc):
@api-version 3
// In the program that includes it:
#include "lib/vec.mc" @api-version 3
If the included file declares a different @api-version than required, or has no
@api-version at all, the compiler emits an error. If no version is required on the
#include line, the check is skipped.
Built-in functions
Memory
void* alloc(i64 size) // raw heap allocate (uninitialized, untyped)
T* alloc<T>(count) // typed uninit: count * sizeof(T), returns T*
T* new(T) // single zero-initialized T, returns T*
T* new(T[count]) // array of count zero-initialized T, returns T*
void free(void* ptr) // heap free (also accepts string directly)
void memcpy(u8* dst, u8* src, i64 n)
void memset(u8* dst, u8 val, i64 n)
string string(str s) // allocate + copy β owned string
string move(string s) // transfer ownership, invalidate source
alloc, new and realloc return memory aligned to at least 16
bytes on every target, which covers any scalar, any struct of them, and
a 128-bit SIMD load. A stronger alignment needs an explicit aligned
allocator.
On Windows and macOS these call the C runtime allocator. On wasm,
Linux and Android the heap is minc's own, lib/mem_heap.mc: a
two-level segregated-fit allocator with quick lists for small blocks,
O(1) alloc and free, blocks split on alloc and merged on free. The
compiler compiles it into the program; it is not imported by hand.
A program can supply its own allocator on any target by defining all
three entry points. Every allocation the compiler emits, including
new, string, format and dynamic arrays, then goes through them:
void* __minc_alloc(i64 n); // 16-aligned, uninitialized; null on failure
void __minc_free(void* p); // p from __minc_alloc, or null
void* __minc_realloc(void* p, i64 n); // C semantics; null on failure, p intact
Defining one or two of the three is an error. They must be thread-safe
if the program uses threads, and they must not call alloc or anything
that allocates. import mem_heap; on Windows or macOS opts into minc's
allocator there. An allocator gets fresh memory from the environment
through void* __heap_grow(i64 n): at least n bytes, 16-aligned, or
null. On uefi-x64 the builtin calls __minc_heap_grow(i64 n).
alloc<T>(count) and new(T[count]) fold the sizeof(T) multiplication
and the pointer cast into the builtin:
i32* a = cast(i32*, alloc(n * 4)); // untyped β explicit casts
i32* b = alloc<i32>(n); // typed, uninit (no zero fill)
i32* c = new(i32[n]); // typed, zero-init
I/O
i64 stdout()
i64 stderr()
i64 stdin() // console or piped input
i64 open(u8* path, i32 mode) // 0=read, 1=write/create
i32 read(i64 handle, u8* buf, i32 n) // bytes read; -1 on error
i32 write(i64 handle, u8* buf, i32 n) // bytes written; -1 on error
void close(i64 handle)
i32 remove(u8* path) // delete a file; 0=success, -1=error
bool file_exists(u8* path)
read/write take an opaque handle from stdin()/stdout()/
stderr()/open(), not a POSIX fd number; the handle's numeric
value is platform-specific. An integer literal as the handle argument
(write(1, ...)) is a compile error.
print("hello\n")
print("{} + {} = {}\n", 2, 3, 5) // compile-time format expansion
eprint("error: {}\n", code) // to stderr
Supported format types: i32, i64, u32, u64, f64, bool, str, string, pointers.
Program
void exit(i32 code)
void abort() // terminate abnormally, no status
i32 get_argc()
u8* get_arg(i32 index) // null if out of range; owned by the runtime, never freed
exit is a normal exit with return value code.
abort is a failure state, see below.
| target | what happens | what the parent sees |
|---|---|---|
| linux, android | kill(getpid(), SIGABRT) |
signal death, core dump |
| macos, ios | abort() in libSystem |
signal death, crash report |
| windows | __fastfail(FAST_FAIL_FATAL_APP_EXIT) |
status 0xC0000409 |
| wasm | the unreachable instruction |
module trap |
| uefi | ud2 |
#UD to the trap handler |
A debugger attached to the process breaks at the abort rather than
watching it exit. Nothing is flushed and no cleanup runs: defer
blocks in progress do not fire.
Both are noreturn. A function whose only exit path is abort()
needs no return statement:
i32 must_parse(u8* s) {
i32 v = 0;
if !parse(s, &v) { abort(); }
return v;
}
A failed bounds check ends the program the same way as abort, after
writing bounds check failed to stderr.
Bit manipulation
i32 popcount(x) // set bits, at the operand's width
i32 clz(x) // leading zeros, at the operand's width
i32 ctz(x) // trailing zeros, at the operand's width
T bswap(T x) // byte swap of a 16-, 32- or 64-bit integer
u64 mulhi(u64 a, u64 b) // bits 64..127 of the 128-bit product a * b
i64 mulhi(i64 a, i64 b) // the same for the signed product
void mac128(u64* acc, u64 a, u64 b) // acc[0..2] += a * b, carry into acc[2]
x is an integer of any width and the count is taken at that width:
clz(cast(u8, 1)) is 7, and clz(0) and ctz(0) are the width. A
literal is i32, or i64 when it needs 64 bits. bswap keeps the
operand's type; a one-byte value has no byte order and is rejected.
One instruction for the 32- and 64-bit forms on x64 and arm64; the
narrow forms mask and use the 32-bit instruction.
mulhi is the high half of the 128-bit product; the low half is
a * b. Two unsigned operands give the unsigned high half as u64,
two signed operands the signed one as i64 (mulhi(-1, 2) is -1,
not 2^64 - 1); mixing the two is an error, and a literal follows the
other operand. Narrower operands extend to 64 bits with their own
signedness. One instruction on x64 (mul / imul) and arm64 (umulh
/ smulh), 32-bit products on wasm. mac128 is unsigned; a signed
argument is reinterpreted as the same-width assignment rule says.
mac128 adds the 128-bit product into the three words at acc with
the carry chained through the third word (mul; add; adc; adc on x64,
mul; umulh; adds; adcs; adc on arm64, compares on wasm). The third
word takes only the carry out of the second: it is the column
accumulator of a multi-word multiply, not a full third limb.
Hardware hints
void cpu_pause() // spin-loop relax: x64 pause, arm64 yield, wasm no-op
void prefetch(void* addr) // software prefetch, T0 locality: x64 prefetcht0,
// arm64 PRFM PLDL1KEEP, wasm no-op
prefetch hints the CPU to pull the cache line at addr into all
cache levels ahead of use. It never faults (an invalid address is
simply ignored by the hardware) and it changes no program state.
Use it a few iterations ahead when walking index lists into large
records (the classic pattern: prefetch(&records[indices[i + 8]])).
Hardware entropy
bool cpu_has_random() // the CPU has a hardware entropy instruction
bool cpu_random(u64* out) // one 64-bit word from it; false if it could not
On x64 (Windows, Linux, UEFI) these are RDRAND: cpu_has_random is the
CPUID bit, and cpu_random stores one word through out and returns
true, or returns false with *out untouched when the CPU has no source.
On arm64 and wasm both current return false. This is intended for kernel code.
AES and carry-less multiply probes
bool cpu_has_aes() // the AES round instructions (AES-NI, FEAT_AES)
bool cpu_has_clmul() // the 64x64 carry-less multiply (PCLMULQDQ, FEAT_PMULL)
bool cpu_has_vaes() // the AES rounds on 256 bits (VAES; x64 only)
bool cpu_has_vpclmulqdq() // the carry-less multiply on 256 bits (VPCLMULQDQ; x64 only)
The two 256-bit probes read CPUID leaf 7's ECX bits 9 and 10 on x64
and fold to false on arm64 and wasm. They say when the i8x32 forms of
the crypto rows run two blocks per instruction; those forms are correct
wherever the 128-bit probe says yes, since without VAES they lower to
the 128-bit instruction on each half.
On x64 these read CPUID.01H:ECX bits 25 and 1. On macOS and iOS they
ask sysctlbyname for hw.optional.arm.FEAT_AES / FEAT_PMULL; on
Linux, Android and UEFI arm64 they read ID_AA64ISAR0_EL1, which the
kernel exposes to user space. On wasm both fold to false. They gate the
aesenc / clmul family in the SIMD intrinsics section: a call to one
of those builtins on a part whose probe says no executes an instruction
the part does not have.
SHA-256
bool cpu_has_sha256() // SHA-NI / FEAT_SHA256
void sha256_compress(u32* state, u8* blocks, i64 n) // n blocks of 64 bytes
sha256_compress runs the SHA-256 compression function over n
consecutive 64-byte blocks and updates the eight state words in place
(state[0] is a in FIPS 180-4's names); an n of zero or less
leaves the state as it was. It is the block function a SHA-256
implementation calls once per block; padding and the length word stay
with the caller. It is a block-level builtin rather than one
per instruction because x64 and arm64 split the state differently
(ABEF/CDGH against ABCD/EFGH) and a per-instruction family would carry
that into every program. On x64 the probe reads CPUID leaf 7's EBX bit
29, after leaf 0 has confirmed the part has a leaf 7, and the routine
is SHA-NI; on arm64 the probe asks sysctlbyname
for hw.optional.arm.FEAT_SHA256 (macOS, iOS) or reads
ID_AA64ISAR0_EL1's SHA2 field (Linux, Android, UEFI), and the routine
is the SHA256H family. A call on a part whose probe says no executes
instructions the part does not have. On wasm the probe is false and the
call is a compile error; keep the portable block function in the else
arm.
Time-stamp counter and CPUID (x64)
u64 cpu_rdtsc() // RDTSC: the time-stamp counter
void cpu_cpuid(u32 leaf, u32 subleaf, u32* out) // CPUID: out[0..3] = EAX, EBX, ECX, EDX
These are the x86 instructions, on Windows, Linux and UEFI x64. A call
on arm64 or wasm is a compile error; gate it with when arch(x64).
cpu_rdtsc counts cycles at a rate the CPU defines and does not
serialise. Use qpc() / qpf() for elapsed time; cpu_rdtsc is for
cycle counts, and for building a clock where there is no OS, with
cpu_cpuid to find the invariant-TSC bit (leaf 80000007H, EDX bit 8)
and the frequency (leaf 15H, 16H or a hypervisor leaf).
Math (builtins)
f64 sqrt(f64 x) // hardware sqrtsd
f64 fabs(f64 x) // hardware sign-bit clear
f32 sqrtf(f32 x) // hardware sqrtss
f32 fabsf(f32 x) // hardware sign-bit clear (f32)
f64 floor(f64 x) // hardware round toward -inf
f64 ceil(f64 x) // hardware round toward +inf
f64 trunc(f64 x) // hardware round toward zero
f32 floorf(f32 x) // f32 forms of the same
f32 ceilf(f32 x)
f32 truncf(f32 x)
All of these are IEEE exactly-defined and bit-identical on every target.
SIMD intrinsics
Streaming load/store, lane reductions, and dot-products for the
packed vec types: the 128-bit int4 / uint4 / f64x2 / i64x2 /
u64x2 / i8x16, the float4 host intrinsics, and the 256-bit
f32x8 / i32x8 / i8x32. Each call is one SIMD instruction, or a
small fixed sequence; codegen is the same shape on x64 and ARM64. The
256-bit types use native AVX2 only under --target windows-avx2 /
linux-avx2 / uefi-x64-avx2 (see Wide SIMD below).
Streaming load/store
One 16-byte SIMD access over the named type's lanes. The function name encodes the lane type; the pointer must match.
| Function | Returns | Notes |
|---|---|---|
int4_load(i32*) |
int4 |
4Γi32, one movdqu |
uint4_load(u32*) |
uint4 |
4Γu32 |
f64x2_load(f64*) |
f64x2 |
2Γf64 |
i64x2_load(i64*) |
i64x2 |
2Γi64 |
u64x2_load(u64*) |
u64x2 |
2Γu64 |
int4_store(i32*, int4) |
void |
matching store |
uint4_store(u32*, uint4) |
void |
|
f64x2_store(f64*, f64x2) |
void |
|
i64x2_store(i64*, i64x2) |
void |
|
u64x2_store(u64*, u64x2) |
void |
f64x2 a = f64x2_load(p);
f64x2 b = f64x2_load(q);
f64x2 r = a * 2.0 + b; // component-wise, scalar broadcast
f64x2_store(p, r);
uint4 v = uint4_load(&a[i]); // one 16-byte load
v = v * 1103515245 + 12345;
uint4_store(&a[i], v);
Sum and reduce (int4 / uint4)
sum4 adds the four lanes with 32-bit wrap: int4 gives i32,
uint4 gives u32. sum4_wide widens each lane first, then adds
in 64 bits: int4 sign-extends to i64, uint4 zero-extends to
u64. Use sum4_wide when the lane total can exceed 32 bits.
accum4 carries a deferred 2x64 sum across a loop; reduce4
collapses the carrier to a scalar once at the end. The carrier is
a real u64x2/i64x2 SIMD vreg. The fold is one packed add
per iteration with no horizontal work. Reach for them whenever a
per-iteration sum4/sum4_wide would do horizontal work that can
wait until after the loop.
| Function | Returns | Notes |
|---|---|---|
sum4(int4) |
i32 |
4-lane horizontal add, 32-bit wrap |
sum4(uint4) |
u32 |
4-lane horizontal add, 32-bit wrap |
sum4_wide(int4) |
i64 |
sign-extend each lane to i64 then sum |
sum4_wide(uint4) |
u64 |
zero-extend each lane to u64 then sum |
accum4(u64x2 acc, uint4 v) |
u64x2 |
deferred sum; zero-extend lanes |
accum4(i64x2 acc, int4 v) |
i64x2 |
deferred sum; sign-extend lanes |
reduce4(u64x2) |
u64 |
final reduce of accum4 carrier |
reduce4(i64x2) |
i64 |
final reduce of accum4 carrier |
int4_to_float4(int4) |
float4 |
per-lane signed intβfloat |
float4_to_int4(float4) |
int4 |
per-lane floatβint, round half-to-even, saturating, NaNβ0 |
int4 v = int4{1, 2, 3, 4};
i32 s = sum4(v); // 10, 32-bit wrap
i64 w = sum4_wide(v); // 10, widened
u64x2 acc;
for i32 i = 0; i < n; i = i + 4 {
uint4 v = uint4_load(&a[i]);
v = v * 1103515245 + 12345; // map
uint4_store(&a[i], v); // write the mapped lanes back
acc = accum4(acc, v); // fold into the carried total
}
u64 total = reduce4(acc); // one reduction, after the loop
float4 intrinsics
Host intrinsics on float4. They work without
#include "linear.mc"; the float2/float3 forms of dot still
resolve to the library overloads.
| Function | Returns | Notes |
|---|---|---|
dot(float4, float4) |
f32 |
sum of the four lane products |
normalize(float4) |
float4 |
unit-length vector |
cross(float4, float4) |
float4 |
3D cross product, w is 0 |
min4(a, b) / max4(a, b) |
float4 |
per lane; picks the SECOND operand on NaN, equal, and Β±0 ties (x64 vminps/vmaxps semantics, identical on every target) |
sqrt4(v) |
float4 |
per-lane square root, IEEE rounded |
rsqrt4(v) |
float4 |
exact 1.0f / sqrt per lane, bit-identical on every target |
rsqrt4_fast(v) |
float4 |
native reciprocal-sqrt approximation; values differ per target (x64 vrsqrtps, ARM64 FRSQRTE, wasm falls back to rsqrt4) |
cmpeq4/cmpgt4/cmpge4/cmplt4/cmple4(a, b) |
float4 |
per-lane mask: all-ones when true, all-zeros when false; quiet ordered (NaN compares false) |
and4/or4/xor4(a, b) |
float4 |
bitwise on all 128 bits |
andnot4(a, b) |
float4 |
(~a) & b (SSE andnot operand order) |
select4(mask, a, b) |
float4 |
per-BIT select: (mask & a) \| (~mask & b) |
movemask4(v) |
i32 |
bit i = lane i's top (sign) bit |
splat4(s) |
float4 |
f32 broadcast to all four lanes |
int8 dot-accumulate (128-bit)
i8x16 is a 16-lane signed-byte vector. dot_acc_i8 is a widening
dot-accumulate: each of the four i32 lanes of acc gets the sum of
four i8Γi8 products added in, accumulating into the first argument.
It is exact over i8 inputs (no saturation).
| Function | Returns | Notes |
|---|---|---|
i8x16_load(i8*) |
i8x16 |
16Γi8, one 16-byte load |
i8x16_store(i8*, i8x16) |
void |
matching store |
dot_acc_i8(int4 acc, i8x16, i8x16) |
int4 |
RMW; += four i8Γi8 products per lane |
int4 acc;
for i32 i = 0; i < n; i = i + 16 {
acc = dot_acc_i8(acc, i8x16_load(&a[i]), i8x16_load(&b[i]));
}
i64 total = sum4_wide(acc);
Byte lanes, AES and carry-less multiply (128-bit)
i8x16 also has the bitwise operators & | ^ ~, and cast() between
any two 128-bit vector types (i8x16, u64x2, i64x2, int4,
uint4, float4, f64x2) is a reinterpret of the same 16 bytes that
emits no instruction. Lane 0 is the first byte in memory on every
target; a u64x2 has its low 64 bits in lane 0.
The byte-lane rows compile on every target, x64, arm64 and wasm:
| Function | Returns | Notes |
|---|---|---|
byte_shuffle(i8x16 v, i8x16 idx) |
i8x16 |
lane i = v[idx[i] & 15], or 0 when bit 7 of idx[i] is set; bits 4..6 are ignored (vpshufb semantics everywhere) |
byte_shuffle2(i8x16 lo, i8x16 hi, i8x16 idx) |
i8x16 |
lane i = byte idx[i] of the 32-byte table lo:hi, or 0 when idx[i] is 32 or more (TBL semantics everywhere: one tbl on arm64, two swizzles on wasm, two vpshufb plus a blend and a mask on x64) |
byte_shl(i8x16 v, n) |
i8x16 |
the vector moved n bytes toward higher lanes, zero fill |
byte_shr(i8x16 v, n) |
i8x16 |
toward lower lanes, zero fill |
byte_align(i8x16 hi, i8x16 lo, n) |
i8x16 |
bytes n..n+15 of the 32-byte pair lo:hi |
n is a compile-time integer constant from 0 to 16 (an immediate in
the encoding); anything else is a compile error.
Lane arithmetic on the same container, one instruction on every target. Lane 0 is byte 0, or bytes 0-1 for the 16-bit rows; every lane wraps, except the pack, which saturates.
| Function | Returns | Notes |
|---|---|---|
i8x16_add(i8x16 a, i8x16 b) |
i8x16 |
byte lanes added |
i8x16_sub(i8x16 a, i8x16 b) |
i8x16 |
byte lanes subtracted |
i8x16_blend(i8x16 a, i8x16 b, i8x16 mask) |
i8x16 |
lane i = b[i] where bit 7 of mask[i] is set, else a[i] (vpblendvb semantics) |
i8x16_eq(i8x16 a, i8x16 b) |
i8x16 |
all-ones byte where the lanes are equal, zero elsewhere |
i8x16_movemask(i8x16 v) |
i32 |
bit i = bit 7 of lane i (vpmovmskb) |
i16x8_add(i8x16 a, i8x16 b) |
i8x16 |
16-bit lanes added |
i16x8_sub(i8x16 a, i8x16 b) |
i8x16 |
16-bit lanes subtracted |
i16x8_mul(i8x16 a, i8x16 b) |
i8x16 |
the low 16 bits of each 16-bit lane product |
i16x8_shl(i8x16 v, n) |
i8x16 |
each 16-bit lane left by n bits |
u16x8_shr(i8x16 v, n) |
i8x16 |
each 16-bit lane right by n bits, zero fill |
i16x8_shr(i8x16 v, n) |
i8x16 |
each 16-bit lane right by n bits, sign fill |
i16x8_pack_u8(i8x16 a, i8x16 b) |
i8x16 |
the sixteen signed 16-bit lanes of a then b, saturated to bytes 0..255 (vpackuswb) |
n is a constant from 0 to 16, or any integer expression read at run
time. A count of 16 or more gives the zero vector for the logical
shifts and the sign fill for i16x8_shr, on every target; a negative
count is undefined. A constant shifts by an immediate; a run-time count
goes through a register (vpsllw xmm, xmm, xmm on x64, ushl /
sshl on arm64, i16x8.shl on wasm).
The same family on i8x32, 32 byte lanes and sixteen 16-bit lanes:
i8x32_add, i8x32_sub, i8x32_eq, i8x32_blend, i8x32_movemask
(32 bits), i16x16_add, i16x16_sub, i16x16_mul, i16x16_shl,
u16x16_shr, i16x16_shr and i16x16_pack_u8. Each lane means what
the i8x16 row means; the pack saturates each 128-bit half on its
own, as vpackuswb ymm does. Under the -avx2 targets each is one
ymm instruction; elsewhere the i8x16 row on each half.
The crypto rows need the probe above to have said yes, and are a
compile error on --target wasm. A library that builds everywhere puts
its hardware path under when arch(x64) || arch(arm64) and chooses on
the probe inside it, with the software path in the else arm.
| Function | Returns | Notes |
|---|---|---|
aesenc(i8x16 block, i8x16 round_key) |
i8x16 |
one middle round: SubBytes, ShiftRows, MixColumns, XOR the key |
aesenclast(i8x16 block, i8x16 round_key) |
i8x16 |
the final round, no MixColumns |
aesdec(i8x16 block, i8x16 round_key) |
i8x16 |
one inverse middle round (equivalent inverse cipher: middle keys through aesimc) |
aesdeclast(i8x16 block, i8x16 round_key) |
i8x16 |
the final inverse round |
aesimc(i8x16 round_key) |
i8x16 |
InvMixColumns, for the decryption schedule |
aeskeygenassist(i8x16 v, rcon) |
i8x16 |
the key-expansion helper; rcon a constant 0..255 |
clmul(u64x2 a, u64x2 b, which) |
u64x2 |
64x64 -> 128 carry-less product of one half of each; which is 0x00, 0x01, 0x10 or 0x11 (bit 0 picks a's high half, bit 4 b's) |
The rounds carry the x86 meaning on arm64 too: the same source gives the same bytes on both. A ten-round block written as ten calls stays in one register from the first round to the last.
i8x16 encrypt_block(i8x16 s, i8x16* rk) { // AES-128, rk[0..10]
s = s ^ rk[0];
for i32 r = 1; r <= 9; r++ { s = aesenc(s, rk[r]); }
return aesenclast(s, rk[10]);
}
byte_shuffle, the four AES rounds and clmul also take i8x32
operands and return i8x32, each 128-bit lane on its own: under the
-avx2 targets one VPSHUFB / VAESENC / VPCLMULQDQ on a ymm
register, two blocks per instruction where the part has VAES and
VPCLMULQDQ (cpu_has_vaes, cpu_has_vpclmulqdq); elsewhere the
128-bit row on each half. clmul on i8x32 takes the lanes' 64-bit
halves as bytes 0-7, 8-15, 16-23 and 24-31, with the same which; the
i8x16 to u64x2 casts and i8x32_pack / i8x32_lo / i8x32_hi
move a u64x2 pair in and out. The bytes are the same through both
lowerings.
Wide SIMD (windows-avx2 / linux-avx2 / uefi-x64-avx2)
*** [EXPERIMENTAL] ***
The 256-bit f32x8 / i32x8 / i8x32 intrinsics. Under the
-avx2 targets each maps to one AVX2 instruction or a small fixed
sequence; on other x64 targets they lower to two 128-bit ops, and to
scalar code on ARM64 and wasm. f32x8_fma accumulates into its first
argument. An -avx2 build checks CPUID once at the top of main (or
sokol_main) and exits with a message on a part without AVX2; on
uefi-x64-avx2 the check sits in main and reports through the
runtime contract's __minc_write and __minc_exit. An image that
enters at efi_main alone carries no check.
| Function | Returns | Notes |
|---|---|---|
f32x8_load(f32*) |
f32x8 |
8Γf32, one 32-byte load |
f32x8_store(f32*, f32x8) |
void |
matching store |
i32x8_load(i32*) |
i32x8 |
8Γi32 load |
i32x8_store(i32*, i32x8) |
void |
matching store |
i8x32_load(i8*) |
i8x32 |
32Γi8 load |
i8x32_store(i8*, i8x32) |
void |
matching store |
i8x32_pack(i8x16 lo, i8x16 hi) |
i8x32 |
lo in bytes 0-15, hi in 16-31 |
i8x32_lo(i8x32) / i8x32_hi(i8x32) |
i8x16 |
bytes 0-15 / 16-31 |
byte_shuffle(i8x32 v, i8x32 idx) |
i8x32 |
within each 128-bit lane |
i8x32_add/sub/eq, i8x32_blend, i8x32_movemask, i16x16_* |
the lane family, see the byte-lane rows above | |
aesenc(i8x32, i8x32) and the other rounds, clmul(i8x32, i8x32, which) |
i8x32 |
see the crypto rows above |
i8x32 & | ^ ~, i32x8 + - & | ^ ~ |
one ymm op; cast() between i8x32 and i32x8 is free |
|
f32x8_add/sub/mul(f32x8, f32x8) |
f32x8 |
component-wise |
f32x8_min/max(f32x8, f32x8) |
f32x8 |
per-lane min / max |
f32x8_div(f32x8, f32x8) |
f32x8 |
per-lane divide |
f32x8_rcp(f32x8) |
f32x8 |
12-bit reciprocal approximation |
f32x8_rsqrt(f32x8) |
f32x8 |
12-bit 1/sqrt approximation |
f32x8_sqrt(f32x8) |
f32x8 |
full-precision per-lane sqrt |
f32x8_exp(f32x8) |
f32x8 |
per-lane exp |
f32x8_splat(f32) |
f32x8 |
broadcast scalar to 8 lanes |
f32x8_pack4(f32, f32, f32, f32) |
f32x8 |
pack 4 scalars, duplicated across both halves |
f32x8_lane4_K(f32x8) |
f32x8 |
broadcast lane K (K=0..3) of each 128-bit half |
i32x8_to_f32x8(i32x8) |
f32x8 |
per-lane signed intβfloat |
f32x8_to_i32x8(f32x8) |
i32x8 |
per-lane floatβint, round half-to-even |
f32x8_fma(f32x8 acc, f32x8, f32x8) |
f32x8 |
RMW; acc + a*b per lane |
sum8(f32x8) |
f32 |
horizontal sum of 8 lanes |
sum8(i32x8) |
i32 |
horizontal sum of 8 lanes |
dot_i8x32(i8x32, i8x32) |
i32x8 |
widening dot; operands in [-127,127] |
dot_acc_i8(i32x8 acc, i8x32, i8x32) |
i32x8 |
RMW; widening dot over 32 lanes |
f32x8 acc = f32x8_splat(0.0f);
for i32 i = 0; i < n; i = i + 8 {
acc = f32x8_fma(acc, f32x8_load(&a[i]), f32x8_load(&b[i]));
}
f32 total = sum8(acc);
The eight lanes of dot_i8x32 and dot_acc_i8 hold partial sums
whose grouping of the 32 products is target-defined, and the two
builtins group differently on x64. Reduce with sum8; a lane on its
own is not portable, and lanes of the two builtins do not add up
lane-wise. The same holds for the four lanes of the i8x16 form.
Threading (builtins)
void thread_create(Thread* t, fn(void*): void entry, void* arg)
void thread_join(Thread* t)
void thread_sleep(i32 ms)
void mutex_init(void* m)
void mutex_lock(void* m)
void mutex_unlock(void* m)
void mutex_destroy(void* m)
Standard library
Include with #include or import:
| Library | File | Description |
|---|---|---|
| Vec | lib/vec.mc |
Generic dynamic array Vec<T> |
| String | lib/str.mc |
String operations (find, slice, trim, compare, builder) |
| Math | lib/math.mc |
sin, cos, tan, exp, log, pow, floor, ceil, trunc, round + helpers |
| Float fmt | lib/format_f64.mc / lib/format_f32.mc |
Ryu shortest-round-trip float-to-string (format_f64, format_f32) |
| File | lib/file.mc |
File read/write (whole file), file_exists |
| Memory | lib/mem.mc |
Arena allocator and pool allocator |
| Thread | lib/thread.mc |
OS threads and mutexes |
| Thread pool | lib/thread_pool.mc |
Persistent worker pool: parallel_for and parallel_for_chunk; idle workers block |
| Atomic | lib/atomic.mc |
Atomic load/store/CAS/RMW with MemOrder (relaxed β seq_cst) |
| Fiber | lib/fiber.mc |
Cooperative coroutines (fiber_create, fiber_switch, fiber_yield) |
| Linear | lib/linear.mc |
Vector/matrix/quaternion math (dot, cross, normalize, perspective, look_at, quaternions) |
| Inflate | lib/inflate.mc |
DEFLATE decompressor (RFC 1951) |
| Deflate | lib/deflate.mc |
DEFLATE compressor (RFC 1951; fixed-Huffman + LZ77) |
| Zlib | lib/zlib.mc |
zlib + gzip wrappers around inflate/deflate (CRC32, Adler-32) |
| PNG | lib/png.mc |
PNG decoder (every color type and bit depth, Adam7, tRNS β RGBA8) + encoder (RGBA8) |
| JPEG | lib/jpeg.mc |
JPEG decoder (baseline, extended, progressive; YCbCr, gray, Adobe RGB/CMYK/YCCK; EXIF orientation) + encoder (4:2:0, quality 1-100, Annex K tables) |
| Sokol | lib/sokol_all.mc |
Cross-platform windowing + GPU (sokol_app / sokol_gfx) |
| Obj-C | lib/objc_runtime.mc |
Objective-C runtime bindings for Cocoa/UIKit/Metal (macOS / iOS only) |
Vec
#include "lib/vec.mc"
Vec<i32> v;
vec_init<i32>(&v, 16); // initial capacity 16
vec_push<i32>(&v, 42); // append
i32 x = vec_get<i32>(&v, 0); // read
vec_set<i32>(&v, 0, 99); // write
i32 top = vec_pop<i32>(&v); // pop last
vec_free<i32>(&v); // free
String library
#include "lib/str.mc"
str s = "hello world";
bool eq = str_equal(s, "hello world");
i32 idx = str_find(s, "world"); // 6
str sub = str_slice(s, 0, 5); // "hello"
bool sw = str_starts_with(s, "hello");
bool ew = str_ends_with(s, "world");
bool ct = str_contains(s, "lo w");
str trimmed = str_trim(" hi ");
// String builder
str_buf sb;
str_buf_init(&sb);
str_buf_add(&sb, "hello ");
str_buf_add(&sb, "world");
str result = str_buf_to_str(&sb);
str_buf_free(&sb);
Math library
#include "lib/math.mc"
f64 r = sqrt(4.0); // 2.0 (builtin, hardware)
f64 a = fabs(0.0 - 3.14); // 3.14 (builtin, hardware)
f64 s = sin(PI / 2.0); // 1.0 (platform library)
f64 c = cos(0.0); // 1.0
f64 e = exp(1.0); // 2.718... (bit-identical on every target)
f64 l = log(E); // 1.0
f64 p = pow(2.0, 10.0); // 1024.0
f64 f = floor(3.7); // 3.0 (builtin, hardware)
f64 m = fmod(7.5, 2.0); // 1.5 (exact on every target)
i32 a = abs_i32(0 - 5); // 5
f64 c = clamp_f64(15.0, 0.0, 10.0); // 10.0
f64 l = lerp(0.0, 10.0, 0.5); // 5.0
Linear math library
#include "lib/linear.mc"
// Vector operations (overloaded for float2/3/4)
f32 d = dot(a, b); // dot product
f32 l = length(v); // magnitude
float3 n = normalize(v); // unit vector
float3 c = cross(a, b); // cross product (float3 only)
// Matrix operations (float4x4, column-major)
float4x4 id = identity();
float4x4 proj = perspective(fovy, aspect, near, far);
float4x4 view = look_at(eye, target, up);
float4x4 mvp = mul(proj, view);
float4x4 rx = rotate_x(angle);
// Quaternion operations (float4: x, y, z, w)
float4 q = quat_identity();
float4 q = quat_axis_angle(axis, angle);
float4 q = quat_mul(a, b);
float3 v = quat_rotate(q, point);
float4x4 m = quat_to_mat4(q);
float4 q = quat_slerp(a, b, t);
File I/O
#include "lib/file.mc"
FileData fd = file_read("input.txt");
if fd.data != null {
// process fd.data[0..fd.len]
free(fd.data);
}
string contents = file_read_str("input.txt"); // returns owned string
defer free(contents);
file_write("output.txt", fd);
file_write_str("output.txt", "hello");
i64 n = file_size("input.txt"); // bytes, -1 if unknown
FileData.len is an i64. file_read handles files of any size.
A string carries an i32 length. file_read_str returns empty
for a file past 2 GB rather than a truncated string.
Memory allocators
#include "lib/mem.mc"
// Arena (bump allocator)
Arena a;
arena_create(&a, 4096);
void* p = arena_alloc_mem(&a, 64);
arena_reset(&a); // reuse without freeing
arena_destroy(&a);
// Pool (fixed-size blocks)
Pool p;
pool_create(&p, sizeof(Node), 100);
Node* n = cast(Node*, pool_alloc(&p));
pool_free(&p, n);
pool_destroy(&p);
Threading
#include "lib/thread.mc"
void worker(void* arg) {
i32 id = cast(i32, cast(i64, arg));
print("worker {}\n", id);
}
i32 main() {
Thread t1;
Thread t2;
thread_create(&t1, worker, cast(void*, cast(i64, 1)));
thread_create(&t2, worker, cast(void*, cast(i64, 2)));
thread_join(&t1);
thread_join(&t2);
return 0;
}
Fibers
import fiber;
void my_fiber(void* arg) {
print("step 1\n");
fiber_yield();
print("step 2\n");
fiber_yield();
print("step 3\n");
}
i32 main() {
Fiber* f = fiber_create(my_fiber, null);
fiber_switch(f); // "step 1"
fiber_switch(f); // "step 2"
fiber_switch(f); // "step 3"
if fiber_done(f) { print("done\n"); }
fiber_destroy(f);
return 0;
}
fiber_switch resumes a fiber until it yields or returns, and does
nothing once it is done. fiber_done is true after the entry function
has returned. fiber_destroy frees the fiber's stack.
A fiber can switch to another fiber. fiber_yield and the return of
the entry function hand control back to whoever switched the fiber in,
main or another fiber.
The stack is 64 KiB, or the size given as a third argument,
fiber_create(entry, arg, 1 << 20), rounded up to whole pages. A guard
page below it makes an overflow fault at the guard instead of writing
into the heap. On uefi the stack is a heap block without a guard.
C interop
DLL imports (Windows)
extern "kernel32.dll" void Sleep(i32 ms);
extern "ucrtbase.dll" f64 sin(f64 x);
Shared library imports (Linux/Android)
extern "libc.so.6" i32 getpid();
extern "libm.so.6" f64 sin(f64 x);
// Block syntax for multiple imports from the same library
extern "libm.so.6" {
f64 sin(f64 x);
f64 cos(f64 x);
f64 sqrt(f64 x);
f64 pow(f64 base, f64 exp);
}
Framework imports (macOS/iOS)
extern "libSystem.B.dylib" f64 sin(f64 x);
extern "libSystem.B.dylib" i64 clock_gettime_nsec_np(i32 clock_id);
Data symbol imports
A declarator with no parameter list imports a data symbol, bound by the dynamic linker at load time:
extern "libSystem.B.dylib" void* environ;
extern "libc.so.6" u8** environ;
extern "Foundation" {
void* NSDefaultRunLoopMode;
void* NSRunLoopCommonModes;
}
extern "QuartzCore" void* kCAFilterNearest;
Reading the name yields the value stored at the symbol; &name
yields the symbol's address. The named library is linked
automatically. Windows DLLs are not covered: the PE writer imports
functions only, so a DLL's data symbol is read through an accessor
function. A bare framework name (or "Foundation.framework")
expands to its full system path, same rule as @link "Foundation".
Path-shaped values and explicit suffixes (.dylib, .so, .o)
keep their spelling. On macOS a relative library is recorded relative
to the loading image: libfoo.dylib and foo.dll become
@loader_path/libfoo.dylib, sub/libfoo.dylib becomes
@loader_path/sub/libfoo.dylib; absolute and @rpath/... paths are
written as given.
A from "symbol" clause binds the data symbol under a different
minc name, same as for function externs. This is the way to import
a data symbol whose real name is a reserved built-in:
extern "libc.so.6" void* c_stdout from "stdout";
extern "libSystem.B.dylib" void* c_stdout from "__stdoutp";
Supported on macOS, iOS, Linux, and Android (executables and
--shared libraries alike). Windows and --target wasm reject the
declaration with a compile error naming the symbol and the target.
The same shorthand works for function externs:
extern "Foundation" void* objc_getClass(u8* name);
Renaming an imported symbol (from)
An extern declaration can bind a foreign dynamic symbol with a
different minc call-name with a from "symbol" clause. The name
in the declarator is what you call; the from operand is the
imported symbol:
extern "libc.so.6" void libc_free(void* ptr) from "free";
extern "libstdc++.so.6" void cpp_delete(void* p) from "_ZdlPv";
The operand is a string. It covers symbols that aren't valid minc identifiers (mangled C++ names, decorated names).
This also exists to import foreign symbols that collide with a
built-in. Built-in names (alloc, free, realloc, write, β¦)
are reserved: an extern that declares one of them is a compile
error. Use from syntax to rename the extern.
// error: extern 'free' shadows the builtin
// note: calling libc free on minc alloc:ed memory is invalid
extern "libc.so.6" void free(void* ptr);
// ok: libc's free, callable as libc_free; the builtin free() is untouched
extern "libc.so.6" void libc_free(void* ptr) from "free";
The same reservation applies to every other declaration in the value
namespace: a function (generic templates included), a global
variable, an enum value, a union variant, or a data extern
(extern "dll" T name;) named after a built-in is a compile error.
The rule holds regardless of visibility β private declarations are
rejected the same way, in any file of the program. Data externs take
the same from rename as function externs. A foreign data symbol
with a reserved name is imported under a distinct minc name.
Linked object imports
Functions defined in linked .o files use bare extern (no library name):
extern void sokol_gfx_setup();
extern void my_c_function(i32 x);
Static linking
# Windows: compile C to .obj, link into minc binary
zig cc -c helper.c -o helper.obj
minc app.mc --link helper.obj -o app.exe
# Linux: compile C to .o
gcc -c helper.c -o helper.o -fno-pie
minc app.mc --link helper.o -o app
# macOS: compile C to .o or link .dylib
clang -c helper.c -o helper.o
minc app.mc --link helper.o -o app
# macOS: dynamic library
minc app.mc --link libhelper.dylib -o app
Exposing minc functions to C
A minc function registered as a callback in a C library, installed as an Objective-C method, or exported from a shared library needs no annotation. Struct arguments and return values follow the platform's C ABI in every minc function on every native target:
- ARM64 (macOS, iOS, Linux, Android): AAPCS64. An HFA struct of 1-4
same-type floats (
NSPoint,NSSize,NSRect,CGSize, β¦) passes in V0-V7 and returns in V0-V3; a 9-16 byte non-HFA struct (NSRange, β¦) passes in an X-register pair. - Linux x64: the System V eightbyte classification. A struct up to 8 bytes passes in one GPR or XMM by its class, 9-16 bytes in a register pair, and anything larger by value on the stack.
- Windows x64: a 1, 2, 4 or 8-byte struct passes by value in one register; larger structs pass by pointer.
Sender and receiver use that convention for direct, indirect, and cross-image calls alike. This example needs nothing declared on it:
void my_method_impl(NSRect rect, u64 flags) {
// receives the NSRect per AAPCS64, in V0-V3
}
It is the convention rather than an opt-in because a function pointer
that reaches minc across an image boundary, through a symbol lookup or a
vtable built at run time, involves no source-level &fn for a compiler to
notice. Both sides have to agree without being told. Both key on the
target rather than on a marking.
Vector builtins need a little care: float4, int4, f32x8 and friends
are internal SIMD types with no portable C calling convention. An
extern declaration carrying one is rejected rather than silently
miscompiled. Pass a matching struct instead, or assert the signature
yourself with a fn-pointer cast, which is how the jit tests hand float4
and int4 across an image boundary. Passing them between minc functions
is unaffected.
Shaders
GPU shaders written in minc syntax. The compiler generates HLSL (Windows/D3D11), GLSL (Linux/WebGL2), or Metal MSL (macOS/iOS) for the target platform, and WGSL for WebGPU.
Shader functions
Use @shader annotations to define vertex, fragment, or compute shader functions:
struct VsOut {
float4 pos;
float4 color;
}
@shader vertex
VsOut cube_vs(
@attr(0) float4 position,
@attr(1) float4 color,
@uniform float4x4 mvp
) {
VsOut outp;
outp.pos = mul(mvp, position);
outp.color = color;
return outp;
}
@shader fragment
float4 cube_fs(VsOut input) {
return input.color;
}
@shader compute(64, 1, 1)
void my_compute(@storage(rgba8) RWTexture2D img, @uniform f32 scale) {
uint2 pixel = thread_id().xy;
img[pixel] = float4{scale, scale, scale, 1.0f};
}
Shader functions don't generate native code; they compile to GPU shader source
text. The compiler auto-generates a funcname_shader global (of type ShaderMeta)
for each @shader function. The first field of the VS output struct is treated as
the clip-space position (gl_Position / SV_Position).
Texture2D, Texture3D, Texture2DArray, TextureCube, Texture2DMS,
RWTexture2D, and
Sampler are built-in handle types only inside @shader function bodies.
Outside, those names can be used as ordinary identifiers.
ShaderMeta.uniforms points at a flat array of ShaderUniformDesc (one entry per
field across all uniform blocks); ShaderMeta.uniforms_count is its length.
Plain @uniform float4x4 mvp produces one descriptor; struct
@uniform(N) PerFrame frame { mat4 mvp; vec4 light_dir; } produces one per field
(with each field's name + offset). Each descriptor's type_kind is one of:
| Value | Kind | Source-language type |
|---|---|---|
| 1 | FLOAT | f32 |
| 2 | FLOAT2 | float2 |
| 3 | FLOAT3 | float3 |
| 4 | FLOAT4 | float4 |
| 5 | INT | i32 |
| 6 | INT2 | int2 |
| 7 | INT3 | int3 |
| 8 | INT4 | int4 |
| 9 | MAT4 | float4x4 |
| 10 | UINT | u32 |
| 11 | UINT2 | uint2 |
| 12 | UINT3 | uint3 |
| 13 | UINT4 | uint4 |
The values are runtime-agnostic; adapters map them to their own per-uniform
type system (sokol's sg_uniform_type uses the same numbers). import shader;
exposes a ShaderUType enum with these constants for code that walks the
descriptors.
An array field carries its element type in type_kind and its length in
array_count, which is 1 for every other field. std140 pads array elements to
16 bytes. An array field must be float4, int4, uint4, or float4x4.
Limitation: nested structs in @uniform structs are rejected; flatten to
scalar, vector, matrix, or array fields.
ShaderMeta.bindings points at a flat array of ShaderBinding (one entry per
@texture / @sampler / @storage / @buffer / @rwbuffer parameter, in
declaration order); bindings_count is its length. Each entry carries the
binding's name, kind, slot, image type, storage-image format string, access
mode, structured-buffer element size, stage, sample kind, and sampler kind.
Runtime adapters (e.g. sokol_make_shader in lib/sokol_all.mc) walk the
array to build their backend-specific descriptors. import shader; exposes
ShaderBindingKind, ShaderImageType, ShaderBindingAccess,
ShaderSampleKind, and ShaderSamplerKind enums.
Parameter annotations
| Annotation | Stage | Description |
|---|---|---|
@attr(N) |
Vertex | Vertex attribute at location N |
@uniform |
VS/FS | Uniform buffer variable (each gets its own block) |
@storage(fmt[, slot][, rw]) |
Compute | Storage image (e.g., rgba8); writeonly by default, readwrite with the rw keyword |
@texture(N[, unfilterable]) |
Fragment | Read-only texture at slot N; unfilterable when the bound pixel format can't be filtered |
@sampler(N[, nonfiltering]) |
Fragment | Texture sampler at slot N; nonfiltering pairs with an unfilterable texture |
@flat |
Varying | Disable interpolation (for integer varyings) |
@buffer(N) |
VS/FS/Compute | Read-only structured buffer |
@rwbuffer(N) |
Compute | Read-write structured buffer |
@shared |
Compute | Group shared memory (local variable) |
@depth |
Fragment output | Fragment depth (f32 field on the return struct) |
@point_size |
Vertex output | Rasterized point size (f32 field on the return struct); GL / GLES / Metal only |
@blend_src |
Fragment output | Second blend source for the SRC1 blend factors (float4); not available on GLES |
Backend-specific shader features
@point_size has no D3D11 or WebGPU equivalent. Writing it
unguarded is a compile error on those backends rather than a silent
no-op. when gpu(...) names the shader backend so one source can opt
in explicitly:
struct VSOut {
float4 pos;
when gpu(opengl) || gpu(opengles) || gpu(metal) { @point_size f32 psize; }
}
The backend defaults from the build target and @gpu overrides it.
Values: d3d11, metal, opengl, opengles, webgpu.
Depth and unfilterable sampling
A runtime adapter needs to know whether a texture is sampled for depth
comparison and whether its sampler compares. The compiler derives both.
A texture read through sample_cmp is a depth texture, and the sampler
in that call is a comparison sampler. The result is reported in
ShaderBinding.sample_kind and .sampler_kind.
Filterability follows from the pixel format bound at run time, so declare it on the parameter. A depth image read as plain data, or a 32-bit float target, is unfilterable on backends that filter an ordinary color texture:
@shader fragment
float4 debug_view(VOut inp,
@texture(0, unfilterable) Texture2D depth_tex,
@sampler(0, nonfiltering) Sampler depth_smp) {
return sample(depth_tex, depth_smp, inp.uv);
}
Texture2D<u32> and Texture2D<i32> carry this in the type already.
Read them with tex[coord], which the checker enforces.
Each binding takes its kind from one source. Marking a texture
unfilterable while sampling it with sample_cmp, or tagging an
integer texture, is a compile error.
Structs shared with a shader (@gpu_layout)
A storage buffer is shared memory. The CPU writes elements and the shader reads them at the offsets its own language mandates, and the two disagree on vector alignment:
| Field type | minc | GPU (std430 / MSL / WGSL) |
|---|---|---|
scalars (f32, i32, u32) |
4 | 4 |
float2 / int2 / uint2 |
4 | 8 |
float3 / int3 / uint3 |
4 | 16 |
float4 / int4 / uint4 |
4 | 16 |
minc packs vectors to 4 since its SIMD memory operations are unaligned; the shader languages mandate the wider alignments.
@gpu_layout applies the GPU rules on the CPU side. A single memcpy
into a buffer is correct:
@gpu_layout
struct sb_vertex {
float3 pos; // offset 0
float4 color; // offset 16, or 12 without the annotation
} // 32 bytes, or 28 without
A struct used as a @buffer or @rwbuffer element must match the GPU
layout. A struct with every field on a 16-byte boundary, or an
all-scalar struct, already matches; use the annotation for the rest.
Where the layouts differ the compiler reports the field and the offset
it expects. Resolve it with the annotation or with hand-written padding.
The same rule covers @uniform blocks. Every shader dialect reads a
uniform block at these offsets; HLSL's cbuffer packing and WGSL's
uniform rules mandate them. A struct like {f32 mode; float4x4 mvp;}
needs the annotation: it puts mvp at 16 and sizeof at 80, and
uploading sizeof(p) bytes of the struct is then correct on every
backend. A struct @uniform must be the only parameter in its block.
Bare @uniform params sharing an explicit slot follow the same offset
rule. Since a parameter list can't be annotated, a divergent bare
block is a compile error.
An array member strides by its element alignment. float3[2]
occupies 32 bytes and the field after it starts at 32. A pointer into
such an array steps by the padded stride. Use var p = &arr[0]; to keep
that stride, or index the array directly. A bare float3* means a
12-byte step and is a compile error.
Every code shape that would mix the two strides is a compile error:
bare float3* pointers into the array, passing it to a packed-stride
parameter, and whole-array assignment or memcpy between a padded
array and a packed one. Copy element-wise instead, or copy between
values of the same layout.
Padding costs memory. For large arrays of 3-component vectors prefer
separate scalars: f32 x; f32 y; f32 z; packs to 12 bytes under both
the GPU rules and minc's.
Multisample textures
A multisample texture holds one value per sample, and is read a sample
at a time with load_sample(tex, coord, sample). Declare the parameter
Texture2DMS:
@shader fragment
float4 resolve(VOut inp, @texture(0) Texture2DMS tex) {
int2 uv = int2{cast(i32, inp.pos.x), cast(i32, inp.pos.y)};
float4 avg = (load_sample(tex, uv, 0) + load_sample(tex, uv, 1) +
load_sample(tex, uv, 2) + load_sample(tex, uv, 3)) * 0.25f;
return avg;
}
sample_mask() returns the rasterizer's coverage mask for the current
fragment, one bit per sample. A custom resolve can treat partly
covered pixels differently. It is a fragment builtin.
Reads of a multisample texture go through load_sample;
texture_size(tex) takes no LOD, since it has one level. The
sample() family rejects a multisample texture at compile time
(filtering needs a single value per coordinate).
ShaderBinding.multisampled is 1 for such a binding, which is how a
runtime adapter describes the bound image.
GLSL exposes multisample samplers from ES 3.1 and the coverage mask
from ES 3.2. A shader using either declares that version. WebGL2 is
ES 3.0 and has neither. --target wasm reports a shader that uses
them as a compile error. WebGPU has multisample textures and takes the
same shader unchanged; select it with @gpu "webgpu".
Helper functions and out-params
A shader helper is an ordinary function called from @shader code. A
pointer parameter is how a helper returns more than one value. Each
dialect spells it its own way: inout T in HLSL and GLSL, thread T&
in MSL, and ptr<function, T> in WGSL.
void probe(float3 p, f32* dist, float3* normal) {
*dist = length(p);
*normal = normalize(p);
}
@shader fragment
float4 shade(VOut inp) {
f32 d = 0.0f;
float3 n = float3{0.0f, 0.0f, 0.0f};
probe(inp.color.xyz, &d, &n);
return float4{n.x, n.y, n.z, d};
}
An out-param stands for one value. Write through it with *p = v, read
it with *p, and pass it on as &variable or by forwarding an
out-param the caller already holds. Pointer arithmetic, indexing, and
pointer locals are compile errors, as is a pointer parameter on a
@shader entry point.
Shader builtins
These functions are available inside @shader functions. Use GPU-standard
names (not C-style sinf / cosf).
Math:
sin, cos, tan, asin, acos, atan, atan2, sqrt, rsqrt, abs,
exp, log, log2, pow, floor, ceil, round, frac, sign, saturate
Interpolation:
lerp / mix, step, smoothstep, clamp, min, max
Vector:
dot, cross, normalize, length, distance, reflect, refract
Conversion: radians, degrees
Matrix: mul (matrix Γ vector, matrix Γ matrix)
Derivatives (fragment only): ddx, ddy, fwidth
Texture:
sample(tex, smp, uv), sample_level(tex, smp, uv, lod),
sample_cmp(tex, smp, uv, cmp), sample_offset(tex, smp, uv, offset),
gather(tex, smp, uv), texture_size(tex[, lod]),
load_sample(tex, coord, sample)
texture_size returns int3 for Texture3D and Texture2DArray,
where the third component is the depth or the layer count, and int2
for the other kinds. The optional LOD selects a mip level. A storage
image has one level and takes no LOD.
Stage builtins:
thread_id(), group_id(), local_id() (compute),
vertex_id(), instance_id() (vertex),
frag_coord(), front_facing(), sample_mask(),
discard; (fragment; aborts the current fragment),
group_barrier(), memory_barrier() (compute)
Subgroup (compute): the hardware SIMD group, 32 lanes on Apple GPUs.
| Builtin | Returns | Meaning |
|---|---|---|
subgroup_lane() |
u32 |
Index of this lane in the group |
subgroup_size() |
u32 |
Lanes per group |
subgroup_sum(x), subgroup_max(x), subgroup_min(x) |
type of x |
Reduce a scalar f32/i32/u32 across the group; every lane gets the result |
subgroup_shuffle_down(x, delta) |
type of x |
x from lane lane + delta; undefined past the last lane |
Metal, GL 4.3 with GL_KHR_shader_subgroup, and WebGPU lower these
directly. D3D11 (shader model 5.0) has no wave operations. A subgroup
call there is a compile error unless it sits under when gpu(...) with
a shared-memory branch for d3d11.
Half precision: unpack_f16x2(u32) β float2 and
pack_f16x2(float2) β u32 convert two IEEE binary16 values packed in
one word, low half first, round to nearest even. Both work on every
dialect.
The types f16, f16x2 and f16x4 exist on Metal (half) and WebGPU
(enable f16;) only. They sit under when gpu(metal) || gpu(webgpu).
On D3D11 (shader model 5.0) and GL 4.3 an ungated f16 type is a compile
error. Where the types exist, f16 supports:
- locals,
@shared f16[N], and[]f16storage buffers + - * /- swizzles of length 1, 2 and 4
- the 1-arg and 2-arg math builtins
- the subgroup reductions
Rules:
f16never mixes withf32in one expression.cast(f16, x)andcast(f32, h)convert.- A float or int literal beside an
f16adopts the type. @uniformblocks take no f16.- Half arithmetic rounds per backend. Results are not bit-identical across GPUs.
Packed 8-bit lanes: each argument is a u32 word holding four 8-bit
lanes, low byte first. All four builtins are exact on every dialect.
| Builtin | Returns | Meaning |
|---|---|---|
dot4_i8(a, b) |
i32 |
Sum of the four signed lane products |
dot4_u8(a, b) |
u32 |
Sum of the four unsigned lane products |
unpack_i8x4(w) |
float4 |
The four lanes as signed floats |
unpack_u8x4(w) |
float4 |
The four lanes as unsigned floats |
WebGPU has an instruction for each. On Metal, D3D11 and GL the dots run
a small per-shader helper; the unpacks are one reinterpret and convert
on Metal and shifts elsewhere. unpack_* is the quantized-weight decode
in one call instead of a shift pair per lane.
Loop unroll hint: @unroll before a for loop asks the shader
compiler to unroll it fully. Metal gets _Pragma("clang loop
unroll(full)") and HLSL gets [unroll]; GLSL, WGSL and native code
ignore it. Use it on constant-trip inner loops that index arrays of
simdgroup tiles, which otherwise may not stay in registers.
Simdgroup matrices (compute, Metal only, under when gpu(metal)):
simdgroup_mat8x8 (f32) and simdgroup_mat8x8_f16 are 8Γ8 tiles held
by one SIMD group. Every lane of the group must reach each call. Other
dialects reject an ungated use by name.
| Builtin | Meaning |
|---|---|
simdgroup_zero(), simdgroup_zero_f16() |
A zero tile |
simdgroup_load(m, arr, offset, stride) |
Fill m from arr, starting at element offset, stride elements per row |
simdgroup_load_t(m, arr, offset, stride) |
Same, loading the transpose |
simdgroup_mad(d, a, b, c) |
d = a * b + c |
simdgroup_store(m, arr, offset, stride) |
Write m to arr at element offset, stride elements per row |
arr is a @shared array or a storage buffer of the tile's element
type, or of its 2- or 4-vectors for staging with vector stores. Offsets
and strides count scalar elements either way. In simdgroup_mad, a
and b share one element type and c and d share one; f16 operands
into an f32 accumulator is allowed.
Atomic (compute and fragment):
| Builtin | Meaning |
|---|---|
atomic_add(buf, index, value) β T |
Add; returns the value the slot held before |
atomic_min(buf, index, value) β T |
Store the minimum; returns the previous value |
atomic_max(buf, index, value) β T |
Store the maximum; returns the previous value |
atomic_exchange(buf, index, value) β T |
Store value; returns the previous value |
atomic_cmp_exchange(buf, index, compare, value) β T |
Store value if the slot holds compare; returns the previous value, so the store happened when it equals compare |
buf is a storage buffer or a @shared array with an i32 or u32
element type T. The result may be discarded. Compare-exchange is
strong on every backend: a slot that holds compare is always
exchanged.
u32 slot = atomic_add(counter, 0, 1); // claim an output slot
out[slot] = item;
u32 old = 0; // compare-exchange retry loop
u32 seen = atomic_cmp_exchange(total, 0, old, old + 1);
while seen != old {
old = seen;
seen = atomic_cmp_exchange(total, 0, old, old + 1);
}
Using shaders with sokol
Each @shader function generates a ShaderMeta global named funcname_shader.
Create a sokol shader from a vertex + fragment pair:
sg_shader shd = sokol_make_shader(&cube_vs_shader, &cube_fs_shader);
Pass uniforms via sg_apply_uniforms(); each @uniform parameter becomes a
separate uniform block (block 0 for the first, block 1 for the second, etc.).
Cross-platform GPU targets
The shader backend is selected automatically from --target:
| Target | GPU backend | Shader language |
|---|---|---|
| Windows | D3D11 | HLSL |
| Linux | OpenGL | GLSL 410 |
| WASM/Android | OpenGL ES | GLSL ES 300 |
| macOS/iOS | Metal | MSL |
Emitted versions rise where a feature needs it: GLSL 420 for storage images and
430 for compute, GLSL ES 310 for gather or storage images. WebGL2 is GLSL ES
300 only. It cannot run the shaders that ask for 310.
Override with @gpu "target" at file scope (before shader functions):
@gpu "opengl", @gpu "d3d11", @gpu "metal", @gpu "opengles",
@gpu "webgpu". WebGPU has no --target of its own; the pragma is how a build
selects WGSL.
Build commands
minc main.mc # build release, output main.exe (or main)
minc build main.mc # same as above
minc build debug main.mc # build with debug info
minc run main.mc # build and run
minc run debug main.mc # build debug and run
minc main.mc -o custom.exe # explicit output name
minc main.mc --target linux # cross-compile to Linux
Output filename is derived from input: app.mc β app.exe (Windows) or app (Linux/macOS).
Use -o to override.
Source tags
Build options can be specified in source files using @ directives.
Tags propagate transitively through #include: if a library has @link,
any program that includes it automatically links the specified file.
@link "c_code.obj" // link external object file
@gui // set PE subsystem to GUI (no console)
@unchecked // disable bounds checking
@define "SG_D3D11" // define compile-time flag (for when defined())
@must_use i32 try_parse(...) // warn when caller discards the result
Tags are typically placed in library files:
// lib/my_library.mc
when os(windows) {
@link "c_code.obj"
@gui
}
when os(linux) {
@link "c_code.o"
}
With this, minc app.mc is all that's needed to compile a linked app.
Source location: __line__ and __file__
__line__ and __file__ fold at compile time to ordinary literals:
__line__ is the 1-based line number of the token itself, as an
integer literal; __file__ is the basename of the file being
compiled (main.mc, never a full path), as a string literal.
void check(bool ok, str file, i32 line) {
if !ok {
print("check failed at {}:{}\n", file, line);
exit(1);
}
}
check(x > 0, __file__, __line__);
The fold happens at parse time. They are usable anywhere the
equivalent written literal is: function bodies, global initializers,
shader code (__line__; __file__ fails where any string does).
Each token reports its own position: in an included file, they name
the included file's coordinates, not the includer's.
Both are reserved words: i32 __line__; is a compile error.
Compiler flags
minc [build|run] [debug] <input.mc> [options]
-o FILE Output filename (default: derived from input)
-g Emit debug info (DWARF); no codegen change (doesn't add -Og)
-Og Debug-friendly codegen: -Os + locals/params in stack slots
-Os optimize for size (skip code-expanding passes)
--target <t> Cross-compile (see target table below)
--link <file> Link external object file (also available as @link tag)
--shared Emit shared library (.dll / .so / .dylib) exporting the `export` functions
--gui Set PE subsystem to GUI (also available as @gui tag)
--def <file.def> Load additional .def file for DLL mapping (Windows)
--unchecked Disable bounds checking
--no-dce Keep all top-level functions (disable dead code elimination)
--hash Print the output's SHA-256 (a `hash` record under --agent=json)
--deps <file> Write the source files the compile read, one per line
--track-alloc Count allocations and report what is live at exit (see "Leak check")
-DFLAG Define compile-time flag
-DFLAG=value Define with value
--no-color Disable colored diagnostics
--color Keep colour and source excerpts when stderr is not a terminal
--agent Agent mode: one line per diagnostic; sets MINC_AGENT=1 for child processes
--list-builtins List built-in Windows API symbols
--version Print compiler version
Standard Windows API symbols (kernel32, user32, gdi32, ucrtbase, d3d11, ole32,
shell32) are built into the compiler; no --def flags needed for common APIs.
Run minc --list-builtins to see all 263 available symbols.
Diagnostics on a terminal, a pipe, or under an agent
On a terminal a diagnostic is coloured and shows the source line with a caret. When stderr is a pipe or a file, both are dropped and each diagnostic is one line:
app.mc:12:9: error: initializer type mismatch
1 error(s) found
NO_COLOR in the environment disables the colour only. MINC_AGENT=1
in the environment, or --agent on the command line, selects the
one-line form on any stream. --agent also sets MINC_AGENT=1 for the
processes the compiler starts (a build.mc, the test runner, the
program minc run launches). --agent=json sets MINC_AGENT=json the
same way. MINC_AGENT=0 turns agent mode off. --color keeps the
terminal form on a pipe.
In agent mode a diagnostic the compiler knows how to repair carries the repair on its line:
app.mc:3:9: error: initializer type mismatch [fix: wrap 3:13-14 in cast(i32, ...)]
app.mc:4:5: warning: result of 'f' is unused; wrap in 'ignore ...;' to silence [fix: insert "ignore " at 4:5]
app.mc:2:9: warning: unused variable 'x' [fix: delete line 2]
Under MINC_AGENT=json or --agent=json every line the compiler would
write to stderr is one JSON object. A diagnostic is
{"kind":"error","file":"app.mc","line":3,"col":9,"code":"type-mismatch",
"message":"initializer type mismatch",
"fix":{"kind":"wrap","line":3,"col":13,"end_col":14,"before":"cast(i32, ","after":")"}}
kind is error, warning or note. code is present where the
compiler classifies the diagnostic: type-mismatch, sign-mismatch,
unused-variable, must-use, unreachable, syntax. fix is present
for the mechanical repairs: wrap puts before and after around the
expression at line:col (end_col is one past the span when the
expression is a single identifier), insert puts text at line:col,
delete_line removes the line. A failed compile ends with
{"kind":"summary","errors":N}, a successful one with
{"kind":"output","file":"app.exe","bytes":N,"lines":N}. Other text
arrives as {"kind":"note","message":"..."}. minc query prints its
records in the same mode.
Running a program: minc run
minc run app.mc builds and runs the program. The binary is staged
under build/ in the working directory (-o picks another place; a
wasm run stages in the system temp directory and serves it from there).
Arguments after -- go to the program. minc run returns the program's
exit status. When the program dies of a fault, minc run names it:
crash: access violation (minc debug app.mc shows where)
--timeout N kills the program after N seconds with exit status 124.
--memory N caps it at N MB (a job object on Windows, RLIMIT_AS
elsewhere). In agent mode the run ends with a summary line:
run: exit 3 in 0.039 s, peak 3592 KB
run: exit 0xc0000005 in 0.092 s, peak 3092 KB, access violation
and under --agent=json with a record, on every run:
{"kind":"run","exit":3,"ms":39,"peak_kb":3592}
{"kind":"run","exit":124,"ms":1014,"peak_kb":3576,"timeout":true}
{"kind":"run","exit":139,"ms":8,"peak_kb":3056,"signal":"access_violation"}
{"kind":"run","exit":-1073741819,"ms":92,"peak_kb":3092,"signal":"access_violation","status":"0xc0000005"}
signal is one of access_violation, bus_error, arithmetic,
illegal_instruction, abort, breakpoint, stack_overflow,
killed, terminated, broken_pipe, or signal_N /
exception_<hex> for anything else. On Windows a fault's exit code is
the NT status: exit holds it as the number a shell sees and status
repeats it in hex.
Leak check: --track-alloc
--track-alloc builds the program with two counters, bumped by every
heap allocation and non-null free β alloc, alloc<T>, new, and
the allocations builtins make on their own (format, string(),
string builder growth) β readable in the program
as alloc_live() and alloc_total() (both 0 without the flag).
realloc moves neither. When main returns, the program reports. On a
terminal, one line on stderr when something is still live:
leak: 1 live of 2 allocations
Under --agent=json, a record on every run:
{"kind":"leak","live":1,"total":2}
The report follows the compile's mode. A program that leaves main
through exit() gives none. minc run app.mc --track-alloc and
minc test --track-alloc build tracked programs and tests; with a
build.mc the flag reaches the script's compiles too. The compiler's
own runtime allocations are not counted.
Debugging a program: minc debug
minc debug app.mc [-- args] builds with -g -Og and opens the result
in minc-dbg, installed beside the compiler: breakpoints, stepping,
bt, print. A crash stops there with the faulting context; c lets
the program die of it.
minc debug --batch app.mc runs the program to its end with no prompt.
A crash prints the fault and the backtrace and exits with the program's
status:
crash: access violation at 0x401042
#0 depth at rec.mc:2
#1 depth at rec.mc:3
#2 depth at rec.mc:3
#3 depth at rec.mc:3
#4 main at rec.mc:7
Under --agent=json the same run gives one record for the crash and
one for the exit:
{"kind":"crash","signal":"access_violation","pc":"0x401042","frames":[{"fn":"depth","file":"rec.mc","line":2},{"fn":"depth","file":"rec.mc","line":3},{"fn":"depth","file":"rec.mc","line":3},{"fn":"depth","file":"rec.mc","line":3},{"fn":"main","file":"rec.mc","line":7}]}
{"kind":"run","exit":-1073741819,"status":"0xc0000005"}
pc is the faulting instruction's address. A frame's file is the
source path as the compiler was given it. An inlined function reports
under its caller's name with its own line. The same batch mode is
minc-dbg --batch [--json] <exe> on any binary built with -g.
Profiling a program: minc profile
minc profile app.mc [--seconds N] [-- args] builds with -g (no
codegen change) and runs the program with its instruction pointer
sampled about a thousand times a second. When the program exits, or
after N seconds, the functions are listed by share, most first:
samples: 1095
79.1% 867 quick_sort
18.3% 201 insertion_sort
2.4% 27 main
Time in system code (the allocator, file reads) is listed as
<outside .text>. Under --agent=json the run gives one record:
{"kind":"profile","samples":1095,"ms":1674,"exit":0,"functions":[{"fn":"quick_sort","samples":867,"pct":79.1},{"fn":"insertion_sort","samples":201,"pct":18.3},{"fn":"main","samples":27,"pct":2.4}]}
The sampler is minc-dbg --profile [--json] [--seconds N] <exe> on
any binary built with -g; a binary without line tables reports every
sample as outside. Only the main thread is sampled.
Running the tests: minc test
In a project without a build.mc, minc test compiles every
test/*.mc, runs each, and counts an exit status of 0 as a pass. The
binaries go under build/test/. On a terminal every test gets a line;
a failing test's output follows its line:
FAIL test_math (exit 1)
expected 4, got 5
PASS test_ok
FAIL test_types (compile)
test/test_types.mc:3:9: error: initializer type mismatch
1 passed, 2 failed
In agent mode the passes are a count in the results line and only the
failures get a line. Under --agent=json every test is one record with
the compile and run time, the exit status, and the captured output of a
failure or a failing compile's diagnostic records; the run ends with a
results record:
{"kind":"test","name":"test_math","status":"fail","ms":59,"exit":1,"output":"expected 4, got 5\n"}
{"kind":"test","name":"test_ok","status":"pass","ms":353}
{"kind":"test","name":"test_types","status":"fail","reason":"compile","diagnostics":[{"kind":"error","file":"test/test_types.mc","line":3,"col":9,"code":"type-mismatch","message":"initializer type mismatch","fix":{"kind":"wrap","line":3,"col":13,"end_col":14,"before":"cast(i32, ","after":")"}},{"kind":"summary","errors":1}]}
{"kind":"results","passed":1,"failed":2,"unchanged":0}
--filter <text> runs the tests whose name contains the text.
--timeout N kills a test after N seconds; its record says
"reason":"timeout". --changed runs only the tests an edit reaches:
each compile writes the files it read to build/test/<name>.d, a pass
leaves build/test/<name>.ok, and a test runs again when it has no
stamp or any file in its list is newer than the stamp or gone. The
others count as unchanged. The same list is available to any compile
as --deps <file>.
With a build.mc the verb and its flags go to the script, which
reports in its own shape. This repository's battery prints the same
lines and records through MINC_AGENT, one record per test and one
per scope, and rejects --changed.
lib/test_framework.mc provides the above as a module:
import test_framework;
i32 main() {
str verb = str_from_cstr(get_arg(1));
if str_equal(verb, "test") { return test_run_dir("test"); }
...
}
test_run_dir compiles and runs every .mc in the directory, reads
--filter, --changed and --timeout off the command line, and
follows MINC_AGENT for the output shape. A suite that is not one
binary per file reports through test_begin, test_report,
test_report_fail and test_finish, which write the same records. The
compiler it spawns comes from MINC, then PATH.
The source index: minc query
minc query answers from the index the language server keeps: a
reference is an identifier the parser resolved, not a substring. It
scans the working directory (or --root <dir>), skipping what the
root's .gitignore lists, build, and directories starting with a
dot, and prints one tab-separated record per line, sorted by file and
position, with paths relative to the root. It needs minc-lsp beside
the compiler.
minc query def <name> definitions; Struct.field names a field
minc query refs <name> reference sites, each with the declaration holding it
minc query callers <name> functions whose bodies reference it, with site counts
minc query callees <name> functions its body references, with site counts
minc query symbols <text> definitions whose name contains the text
minc query defines <name> the files that define it
minc query closure <file> every file a compile of the file reads
| option | meaning |
|---|---|
--root <dir> |
the tree to index; default the working directory |
--target <t> |
the platform whose when arms the index holds, named as for the compiler; default this machine |
--agent=json |
one JSON record per line, before or after the verb |
$ minc query callers helper
main.mc:2:5 caller other 1
main.mc:3:5 caller main 2
$ minc query def helper --agent=json
{"kind":"def","file":"util.mc","line":2,"col":5,"symbol":"helper","node":"function"}
| record | text form | JSON fields |
|---|---|---|
definition (def, symbols) |
util.mc:2:5 function helper, then private or forward when so |
kind def, file, line, col, symbol, node, private, forward |
reference (refs) |
main.mc:2:27 ref helper other |
kind ref, file, line, col, symbol, in |
caller (callers) |
main.mc:3:5 caller main 2 |
kind caller, file, line, col, symbol, caller, sites |
callee (callees) |
util.mc:2:5 callee helper 2 |
kind callee, file, line, col, symbol, callee, sites |
file (defines, closure) |
util.mc |
kind file, file |
| nothing found, on stderr | no references to lonely (defined at util.mc:3:5), no symbol named missing |
kind empty, query, symbol, known, and the definition's file, line, col when known |
| a file the index could not parse, on stderr | note: broken.mc:2:24: not indexed (expected expression), mentions helper |
kind note, file, line, col, error, mentions |
A name defined more than once: each site binds to the definition in
its own file, else to a public one in a program the file belongs to,
and refs and callers records name it as of file:line (JSON
of); a callees record names the definition whose body holds the
sites the same way. file:name in place of the name asks about that
definition alone:
$ minc query callers src/vm.mc:def_method
src/vm.mc:3850:13 caller vm_execute 2 of src/vm.mc:2185
The exit status is 0 with results, 1 with none, and 2 for a bad
command line. The index holds top-level declarations (functions,
externs, structs, unions, enums and their members, globals, type
aliases, struct fields), not locals; the modules the tree imports from
outside it, the standard library included; and one platform's when
arms. closure, refs, callers and callees answer for that platform while
def and symbols list every file's declarations. A file with a
syntax error holds no entries, and the note names it when its text
mentions the name asked about.
The agent server: minc agent
minc agent is a Model Context Protocol server over stdio: JSON-RPC
2.0, one object per line, initialize, tools/list, tools/call,
ping. For Claude Code:
claude mcp add minc -- minc agent
The project is the client's first root when the client offers the MCP
roots capability (the server asks roots/list after the handshake
and again when the list changes), else the working directory the
server was started in. --root <dir> pins it. The server runs in
minc-lsp beside the compiler. The first call indexes the project;
later calls re-index only the files whose modification time or size
changed. A root change drops the index and scans the new one.
| tool | arguments | result |
|---|---|---|
query |
what (def, refs, callers, callees, symbols, defines, closure), name, target (as --target; a change rebuilds the index) |
the minc query records, one per line |
compile |
file, output, flags, hash |
the diagnostic records, with codes and fixes, then the summary or output record; hash adds the hash record |
run |
file, args, timeout, memory, flags |
the program's output, then the run record |
debug |
file, args |
minc debug --batch: the crash record with the backtrace, then the run record |
profile |
file, args, seconds |
minc profile: the profile record, functions by share |
test |
scope, flags |
one record per test and the results record; --filter, --changed, --timeout go in flags |
Every result is one text block. isError is set when the command's
exit status was nonzero: for compile that means diagnostics, for
run that the program failed or crashed. compile, run, debug
and profile act on the file they name even in a project with a
build.mc; test is the project's. The same server is
minc-lsp --agent [--root <dir>] [--minc-exe <path>] [--target <t>]
on its own.
Cross-compilation targets
| Target | Output format | Architecture | Example |
|---|---|---|---|
windows |
PE executable | x86-64 | minc app.mc --target windows -o app.exe |
linux |
ELF executable | x86-64 | minc app.mc --target linux -o app |
linux-arm64 |
ELF executable | ARM64 | minc app.mc --target linux-arm64 -o app |
macos |
Mach-O executable | ARM64 | minc app.mc --target macos -o app |
wasm |
WebAssembly | WASM32 | minc app.mc --target wasm -o app.wasm |
ios |
Mach-O executable | ARM64 | minc app.mc --target ios -o app |
ios-sim |
Mach-O executable | ARM64 | minc app.mc --target ios-sim -o app |
android |
ELF shared library | ARM64 | minc app.mc --target android -o libapp.so |
The default target matches the host platform. Cross-compilation produces native binaries without requiring any toolchain for the target platform.
Packaging for distribution still needs the platform's standard tools. iOS
needs Xcode (xcode-select --install) and an Apple Developer account for
on-device signing. Android needs the SDK + JDK for APK assembly, plus the
NDK if the app pulls in C code. See SETUP.md for setup details.
Numerics
How integer and floating-point operations behave in minc.
Every rule below has the same value on every target: x64 (Windows and Linux), ARM64 (macOS, iOS, Linux, Android), and wasm. The same source produces the same bit pattern.
Integer types
Signed: i8, i16, i32, i64. Two's complement.
Unsigned: u8, u16, u32, u64.
Narrow types (i8, i16, u8, u16) promote to i32 / u32
before any operation. The result carries i32 / u32 semantics.
Writing it back to a narrow slot needs an explicit cast and
truncates:
i8 a = 100;
i8 b = 100;
i8 c = cast(i8, a + b); // -56 (200 wraps to i8)
Pointers, Windows handles, array indices, and sizeof are 64-bit.
Integer arithmetic: + - *
These wrap in two's complement at the type's width. The compiler does not treat overflow as undefined and never reorders or removes an operation by assuming overflow cannot happen.
i32 a = 2147483647;
i32 b = a + 1; // -2147483648
i32 c = a * 2; // -2
i32 d = -a - 1; // -2147483648 (also -INT_MIN wraps)
u32 e = 0;
u32 f = e - 1; // 4294967295
i64 g = 1000000000000;
i64 h = g * g; // wraps in i64
Literal arithmetic takes the width of where it lands: a declaration,
an assignment, a call argument, an array element, a struct literal
field or a return of type i64 makes 1 << 33 an i64. Beside a
typed operand it takes the operand's type when that is 32 bits or
wider: with i64 x, x == 1 << 40 compares against 2^40, and a
literal that does not fit that operand's type is an error, as in a
declaration (x * 3000000000 with an i32 x). A narrower operand
adopts the literal as a bit pattern (flags & ~MASK on a u8).
A shift with a literal on the left counts as literal arithmetic even
when the count is a variable: 1 << n and (1 << n) - 1 take their
type from the declaration or the other operand, never from n. The
operands of a comparison take no type from the declaration the
comparison lands in: i64 ok = bits <= (1 << len) compares in the
type of bits.
-2147483648 is an i32 literal like -2147483647. In an
overloaded call it resolves as the computed constant would if written,
so f(1 << 33) reaches f(i64). In a
generic call the parameter's type counts once it is known, from an
explicit type argument, from the other arguments or from the expected
return type: vec_push(&v, 1 << 33) into a Vec<i64> stores 2^33,
while twice(1 << 33) on its own infers T as i32 from the
literal. A string literal argument follows the same rule: Box<str> s
= boxed("lit") infers T as str. cast(T, e) types e the way T x = e; does,
whatever surrounds the cast: cast(u64, 1 << 33) is 2^33. Where that
declaration would reject a constant that does not fit T, the cast
truncates it (cast(u8, 300) is 44).
Division and modulo
/ truncates toward zero. % carries the sign of the dividend.
Dividing by a literal zero is a compile error.
i32 a = -7 / 3; // -2 (not -3)
i32 b = -7 % 3; // -1 (sign of dividend)
i32 c = 7 % -3; // 1
INT_MIN / -1 and INT_MIN % -1 are defined:
i32 lo = (0 - 2147483647) - 1; // INT32_MIN
i32 q = lo / -1; // INT32_MIN (wraps)
i32 r = lo % -1; // 0
/ 0 and % 0 trap. There is no recovery. Check the divisor
before dividing if it is not a compile-time constant.
if d != 0 { r = n / d; }
--unchecked does not disable the divide-by-zero trap. The trap
comes from the hardware on x86 and from a runtime guard on ARM64.
Bitwise: & | ^ ~
Operate on the bit pattern. No overflow concept. The result has the type of the operands.
i32 a = 0xF0 & 0xFF; // 0xF0
u32 b = (~0) & 0xFFFF; // 0xFFFF
Same-width mixed signed/unsigned operands are accepted for &, |,
^ (and <<). The result takes the type of the left operand. These
operate on bit patterns and the sign of either operand does not
change the bits produced. Different widths still require an explicit
cast to make the widening direction clear.
i32 s = 0xF0;
u32 u = 0xFF;
i32 r = s & u; // ok: same width, result i32
Shifts: << >>
<< is the same for signed and unsigned operands.
>> is arithmetic (sign-extending) on signed types, logical
(zero-filling) on unsigned types. Mixed signed/unsigned operands
to >> are a static error; the result would depend on which
side's sign wins. Cast one side first.
i32 a = -1 >> 1; // -1 (arithmetic, sign-extends)
u32 b = cast(u32, -1) >> 1; // 2147483647 (logical, zero-fills)
Shift amount past the type's width is defined. Result on every target:
| Source | Shift amount | << |
>> logical (u32) |
>> arithmetic (i32) |
|---|---|---|---|---|
i32 / u32 |
0β31 | normal shift | normal shift | normal shift |
i32 / u32 |
β₯ 32 | 0 | 0 | 0 if positive, -1 if negative |
i64 / u64 |
0β63 | normal shift | normal shift | normal shift |
i64 / u64 |
β₯ 64 | 0 | 0 | 0 or -1 by sign |
A negative shift amount counts as past the width.
i32 v = 1;
i32 a = v << 32; // 0
i32 b = v << 33; // 0
i32 n = -5;
i32 c = n >> 32; // -1 (arithmetic; sign extended)
u32 u = cast(u32, -1);
u32 d = u >> 32; // 0 (logical; saturates)
If you want modular-count semantics (the count wraps mod 32 / 64 instead of saturating to 0), mask the count explicitly:
i32 a = 1 << 32; // 0 (default: saturates)
i32 b = 1 << (32 & 31); // 1 (mask: wraps to <<0)
Integer literals
Decimal literals are range-checked against the destination type at the assignment point:
i32 a = 2147483647; // ok
i32 b = 2147483648; // error: out of i32 range
i32 c = -2147483648; // ok (unary minus folded)
u32 d = 4294967295; // ok
u32 e = 4294967296; // error
Hex and binary literals are bit-pattern literals. They fit any type of matching width:
i32 a = 0xFFFFFFFF; // -1 (the bit pattern reinterpreted)
u32 b = 0xFFFFFFFF; // 4294967295
u8 c = 0b11111111; // 255
A hex literal wider than the destination is an error.
Casts: cast(T, expr)
cast(T, expr) is required when the conversion:
- Narrows (
i64βi32,i32βu8). - Crosses signedness at the same width (
i32βu32). - Crosses integer / float / pointer categories.
- Loses information in any other way.
Implicit widening is allowed when it loses no information:
i32 β i64, u32 β u64, u8 β i32, i32 β f64.
Narrowing keeps the low bits:
i64 v = 12030160680303; // fits in i64
i32 t = cast(i32, v); // -42715793 (low 32 bits, signed)
u8 b = cast(u8, 300); // 44
Cross-signedness keeps the bit pattern:
u32 u = cast(u32, -1); // 4294967295
i32 s = cast(i32, 4000000000); // -294967296
Float arithmetic
f32 and f64 follow IEEE 754 with default rounding (round to
nearest, ties to even). + - * / produce the same bit pattern on
every target for the same inputs.
f64 a = 1.0e308 * 10.0; // +inf
f64 b = 1.0 / 0.0; // +inf
f64 c = -1.0 / 0.0; // -inf
f64 d = 0.0 / 0.0; // NaN
f64 z = -0.0 + 0.0; // 0.0 (positive zero)
NaN compares not-equal to itself. Use that to test for NaN:
f64 nan = 0.0 / 0.0;
bool is_nan = nan != nan; // true
Float arithmetic does not trap. Overflow produces Β±inf, underflow
produces subnormals or 0.
f32 literals end with f: 0.5f, 1.5e-10f. Without the
suffix, the literal is f64.
Float literals are parsed bit-exactly. The literal you write is the value you get, with no rounding beyond what IEEE 754 requires.
Float β int cast
Saturating. Out-of-range and NaN produce defined values:
i32 a = cast(i32, 1.0e20); // 2147483647 (INT32_MAX)
i32 b = cast(i32, -1.0e20); // -2147483648 (INT32_MIN)
i32 c = cast(i32, 0.0 / 0.0); // 0 (NaN β 0)
i64 d = cast(i64, 1.0e20); // 9223372036854775807
In-range values truncate toward zero. Negative values truncate toward zero too (not toward minus-infinity):
i32 a = cast(i32, 1.7); // 1
i32 b = cast(i32, -1.7); // -1
Saturation is to the underlying conversion width: i32 for any
i32-or-narrower target, i64 for i64. Narrower destinations
(i8, i16, u8, u16, u32) take the low bits of the
saturated value. A huge float becomes INT32_MAX first and
then truncates. If you want i16/i8-range saturation, clamp the
float yourself before the cast.
i16 a = cast(i16, 1.0e20); // -1 (low 16 bits of INT32_MAX)
u8 b = cast(u8, 1.0e20); // 255 (low 8 bits)
i16 c = cast(i16, 100000.0); // -31072 (low 16 bits of 100000)
Int β float
i32 β f64 and u32 β f64 are implicit. The 53-bit mantissa
holds every 32-bit integer exactly. No cast is needed. Integer
literals within i32 range coerce directly.
f64 a = 12345678; // exact (literal coerces)
i32 n = 12345678;
f64 b = n; // exact (implicit widen)
i64 β f64 and u64 β f64 require an explicit cast. Integer
literals outside i32 range fall into this category. Conversion
rounds round-to-nearest-even when the magnitude exceeds 2^53:
f64 a = cast(f64, 9007199254740992); // exact (2^53)
f64 b = cast(f64, 9007199254740993); // rounds to 9007199254740992
// (next f64 above is 2^53 + 2)
cast(f32, integer) rounds to 24-bit mantissa precision.
f32 β f64
cast(f64, f32_val) is exact (every f32 has an exact f64
representation).
cast(f32, f64_val) rounds to nearest even. Out-of-range overflows
to Β±inf; underflow produces an f32 subnormal or zero.
f32 a = cast(f32, 1.0e100); // +inf
f32 b = cast(f32, 1.0e-50); // subnormal f32
Comparisons
==, !=, <, >, <=, >= work on integers and floats. A
struct, union or array has no ==; compare its fields. A cast()
converts scalars and pointers, never a struct or union value.
Mixed signed/unsigned operands are a static error. Cast one side explicitly:
i32 a = -1;
u32 b = 5;
// if a < b { ... } // error: mixed sign
if cast(u32, a) < b { ... } // false (-1 β 4294967295)
if a < cast(i32, b) { ... } // true
Comparison results are always bool, never the operand type.
Float comparisons follow IEEE 754: every ordered compare (<, >,
<=, >=) returns false if either operand is NaN. == returns
false for NaN, != returns true.
f64 nan = 0.0 / 0.0;
nan == nan // false
nan != nan // true
nan < 1.0 // false
1.0 < nan // false
nan <= nan // false
Constant folding
The compiler folds integer and float expressions whose operands are all literals. The folded value matches what the same expression produces at runtime. There is no folder-vs-runtime mismatch.
i32 a = (2147483647 + 1); // -2147483648 (folded as wrap)
i32 b = (1 << 32); // 0
i32 c = cast(i32, 1.0e20); // 2147483647 (saturated)
Folded floats use the same IEEE 754 rounding as runtime ops.
What traps
These are the only runtime traps from arithmetic and indexing:
| Operation | Trap |
|---|---|
n / 0, n % 0 (any integer width) |
Yes. |
| Out-of-bounds array index | Yes. Disable with --unchecked. |
| Stack overflow | Yes. |
No overflow trap. No NaN trap. No shift trap. No cast trap.
Differences from C
| Case | C | minc |
|---|---|---|
INT_MAX + 1 |
UB | Wraps to INT_MIN. |
INT_MIN / -1 |
UB (often #DE) |
Wraps to INT_MIN. |
INT_MIN % -1 |
UB | 0. |
1 << 32 (i32) |
UB | 0. |
1 >> 33 (i32, signed) |
UB | 0 or -1 by sign. |
cast(i32, NaN) |
UB | 0. |
cast(i32, 1.0e20) |
UB | INT32_MAX. |
i32 x = 4294967296 |
Truncates with warning | Error. |
i32 x = 0xFFFFFFFF |
Implementation-defined | -1 (bit pattern). |
| Mixed-sign compare | Promotes both to unsigned | Static error. |
If you bring code from C that overflows signed integers, the portable fix is to use an unsigned type, which gets the same wrap behavior on both languages without changing the comparison semantics.
Design principles
- Explicit over implicit: sizes in type names, explicit casts, no magic
- Minimal runtime: no default runtime, import what you use
- Cross-platform numerics: same code, same results everywhere
- Direct hardware access: C ABI, extern to call .dll/.so, syscalls, kernel intrinsics