Akari is the engine’s pure shader language: one .akari source emits Metal (MSL), D3D12 (HLSL) and Vulkan (HLSL in its Vulkan dialect, compiled to SPIR-V), then platform compilers produce bytecode. Product packaging and pass roles live in Shader authoring. GPU packing rules: docs/reference/akari-abi.md. Tool layout: src/akari/. Editor highlighting: tools/vscode-extensions/akari (VS Code / Cursor); tools/nvim-plugins/akari (Neovim / Vim).
This page is the language surface for authors. It describes implemented syntax and checks only. Concrete syntax: src/akari/syntax.ebnf (W3C EBNF aligned with akari_syntax).
Mental model
.akari → parse → HIR → IR → pure MSL | HLSL | Vulkan HLSL → metalc | DXC | DXC -spirv → .metallib | .cso | .spvVulkan is the HLSL emitter in a second dialect, not a separate text backend: math, texture, ray-query and wave lowerings are shared with D3D12, and only the binding model differs ([[vk::binding(binding, set)]] from the manifest, runtime descriptor arrays for the bindless heaps, explicit [[vk::location]] stage I/O, [[vk::image_format]] on storage textures).
- Packages (
#shader) are one pass / PSO family (often multi-stage). - Modules (
#module) are shared helpers imported by name. - Bindings use
@bind(group.name); slots come frombindings.kaji.json, not hardcoded registers. - No dual native
metal//d3d12//vulkan/authoring trees. - Array lengths are integer constant expressions: literals or
constnames (uint[HIKARI_BIN_BUCKETS],float[N * 2u]). A length that does not evaluate, or evaluates to zero, is a compile error.
File structure
One root per file. Headers are #… directives. Body items are struct / enum / bitfield / fn / const, plus package bindings (@bind(…) let, #shader only) and #require.
#shader composite
#version 2
#layout fullscreen_texture_table
#import preamble, fullscreen_triangle
#require ray_query
struct V {
@position position: float4;
uv: float2;
}
const PI: float = 3.14159265;
@vertex
fn composite_vertex(@vertex_id vid: uint) -> V {
return V {
position: hikari_full_screen_position(vid),
uv: hikari_full_screen_uv(vid),
};
}
@fragment
fn composite_fragment(
@stage_in inn: V,
@bind(post.hdr) hdr: Tex2D,
@sampler(linear, clamp) samp,
) -> float4 {
return hdr.sample(samp, inn.uv);
}Headers
| Directive | Where | Role |
|---|---|---|
#shader name | package root | Pass package |
#module name | module root | Shared import unit |
#version N | either | Syntax version (default 2) |
#layout name | #shader only | Binding layout contract (standard, ui, …) |
#import a, b | either | Load pure modules from akariImportPaths / engine SDK |
#require a, b | either | Explicit capabilities (optional; also inferred from use) |
Paths:
| Tree | Role |
|---|---|
src/hikari/shaders/akari/packages/*.akari | Engine packages |
src/hikari/shaders/akari/modules/*.akari | Engine modules |
<game>/assets/shaders/akari/packages/ | Game packages |
Every statement, declaration and struct field ends with ;. # directives (#import, #require, …) take none; enum cases, bitfield fields, arguments and literals are separated by , (a trailing , is allowed).
Types
Scalars and vectors
| Family | Forms |
|---|---|
| Bool | bool, bool2…bool4 |
| Signed int | int, int2…int4 |
| Unsigned | uint, uint2…uint4 |
| Float | float, float2…float4 |
| Half | half, half2…half4 |
| Wide int | uint64 (bindless addresses) |
Matrices
Column-major. Named float{cols}x{rows}:
float2x2, float2x3, float2x4, float3x2, float3x3, float3x4, float4x2, float4x3, float4x4.
* between matrix and vector is matrix–vector mul. Prefer mul(m, v) when you want an explicit call. transpose(m) is first-class.
Arrays
let planes: float4[6] = …Length must be a constant integer. Uniform array stride follows ABI (std140-like); storage buffers pack tightly.
Resources
| Type | Meaning |
|---|---|
Uniform<T> | Constant / cbuffer |
Buf<T> | Read-only structured / device buffer |
RwBuf<T> | Read–write structured buffer |
AtomicBuf<T> | Atomic buffer (Metal atomic_T*, HLSL UAV interlocked) |
Tex2D / Tex2DArray / Tex3D / TexCube / TexCubeArray | Sampled textures (texel float4) |
TexDepth2D | Depth texture for sample_cmp (MSL depth2d<float>, HLSL Texture2D<float>; load → float) |
Tex2D<uint> / Tex2D<int> (and the same on the other four) | Integer textures — load only, texel uint4 / int4 |
RwTex2D / RwTex2D<rgba32f|r16f|r32f|rg16f|r11g11b10f|rgba16f|r32u|r32i> | Storage image (default rgba32f, texel float4) |
RwTex3D / RwTex3D<r16f|r32f|rg16f|r11g11b10f|rgba16f|r32u|r32i> | Storage volume — same formats; store takes a uint3 coord |
Sampler | Filtering sampler |
SamplerCmp | Comparison sampler (shadows) |
AccelStruct | TLAS handle (Metal IAS / DXR) |
Structs
struct SurfaceOut {
@location(0) color: float4;
@location(1) normal: float4;
}
@packed
struct Tight {
a: float;
b: float3; // storage packing even under uniform-style rules
}
struct Tap {
weight: float = 1.0; // used when a literal omits the field
offset: float2 = float2(0.0);
count: uint;
}A field default is part of the struct: it may name module constants and functions, never a local of the literal's scope. Defaults do not change layout.
Field attributes:
| Attr | Role |
|---|---|
@location(N) | Colour target N ([[color(N)]] / SV_TargetN) |
@position | Clip position |
@clip_distance | Homogeneous user clip distances (float[1..4] vertex output); negative distances clip primitives |
@flat | No interpolation |
@noperspective | Linear (no perspective) |
@target(group.name) | MRT color via HK_TARGET |
Literals and operators
let i: int = -3;
let u: uint = 0x80000000u;
let f: float = 1.5;
let b: bool = true;
let t: float = cond ? a : b; // right-associativeArithmetic: + - * / %
Compare: == != < <= > >=
Logic: && \|\| !
Bitwise: & \| ^ ~ << >>
Precedence, weakest to strongest: ?:, ||, &&, == !=, < <= > >=, |, ^,
&, << >>, + -, * / %, unary, as, postfix. Bitwise operators bind tighter
than comparisons, so mask & MaterialFlags.Emissive != 0u tests the masked bits
and a << 2u == b compares the shifted value.
Inc/dec: ++x --x x++ x--
Assign: = and += -= *= /= %= &= \|= ^= <<= >>=
Cast / construct:
let h: half = float(i) as half;
let v: float3 = float3(1.0, 0.0, 0.0);
let z: float3 = float3(0.0); // splat: one scalar fills every component
let u: uint = as_uint(f); // bitcast
let f2: float = as_float(u);Swizzles: .xyzw / .rgba (length 1–4).
Struct literals
Named fields. A missing field takes its declared default; otherwise it is an error unless the literal ends with .. (zero-fill) or ..base (copy remaining fields from base). Zero-fill is scalars/vectors/arrays/nested structs; not resources.
let s = VsOut { position: p, uv: uv }
let rest = VsOut { position: p, .. }
let skinned = GBufferVertex { lodFade: OPAQUE, ..common }Positional VsOut(p, uv) still works (declaration order, all fields required).
Array literals
let words = [0u, 1u, 2u, 3u];
let zeros = [0u; 4]; // repeat; count is a constant integerModule-level const arrays also accept list and repeat literals. They emit native
constant aggregate initializers on Metal and HLSL, so lookup tables do not require
per-thread initialization helpers.
const ORDER: uint[4] = [3u, 0u, 2u, 1u];Compile-time layout
Folded to uint literals from the GPU ABI (storage packing by default; sizeof(Uniform<T>) uses uniform packing):
const N: uint = sizeof(MeshletWork);
const O: uint = offsetof(MeshletWork, flags);Declarations
const K: float = 16.0; // module-scope constant
fn helper(x: float3) -> float3 {
const SCALE: float = 2.0; // immutable local (init required)
let a: float = 1.0; // immutable local
var b: float = 0.0; // mutable
return normalize(x) * a * SCALE;
}| Form | Scope | Mutable | Init |
|---|---|---|---|
const (item) | module | no | required |
const (stmt) | local | no | required |
let | local | no | optional |
var | local | yes | optional |
let _ = value; evaluates value and discards it (no local is declared, so it may repeat).
Assign / ++ / -- to const or let bindings (including fields/swizzles of them) is a type error. Module const and local const share the keyword; locals are not merged into the module constant table.
Local bindings belong to their lexical block. A nested block may shadow an outer
binding; leaving the block restores the outer binding. A for initializer is
visible in its condition, step, and body, and ends with the loop. Duplicate
declarations within one block are errors. Compound assignments such as
values[index++] += amount evaluate the destination once.
Logical && / || short-circuit, and condition ? yes : no evaluates only the
selected expression. These rules also apply to atomics and texture queries.
Loop conditions are evaluated on each iteration; continue still evaluates a
for step or a do / while condition before the next iteration.
Enums
Typed closed sets of constants (bitflags and ordinals).
enum MapMask: uint {
Albedo = 1u,
Normal = 2u,
MetallicRoughness = 4u,
// omitted discriminant → previous + 1
}
fn sample_maps(mask: uint) {
if ((mask & MapMask.Albedo) != 0u) { /* … */ }
}| Rule | Detail |
|---|---|
| Underlying | Required: uint or int |
| Access | Name.Case (dot path) |
| Emit | Module consts Name_Case as underlying type (MSL constant / HLSL static const) |
| Typing | A case reads as its underlying integer wherever one is expected (bitwise, compare, indices); an integer becomes an enum only through x as Name |
| Auto values | First omitted case is 0; each following omitted case is previous + 1 |
| Explicit values | Checked integer constant expressions; named constants and earlier enum cases are allowed. Values must fit the underlying type; uint bitwise complement uses 32 bits. |
Prefer shared module enums over pasting local let/const flag bits in every package.
For a counted set, end the enum with an explicit Count case and size arrays with it.
Bitfields
Packed fields over a uint, least-significant field first (the same order as a
Zig packed struct). The value is its uint: same layout in buffers and
uniforms, and as converts both ways without touching bits.
bitfield VisibleMeshletWord: uint {
work: 22, // bits 0..21
fade: 8, // bits 22..29
target: 1,
late: 1, // bit 31
}
let w = records[i] as VisibleMeshletWord;
let fade = w.fade; // (w >> 22u) & 255u
var next = VisibleMeshletWord { work: index, fade: fade, .. }; // `..` zeros the rest
next.late = 1u; // read-modify-write of one field
records[i] = next as uint;| Rule | Detail |
|---|---|
| Storage | uint only; widths are constant integers 1..=32 summing to at most 32 |
| Read | w.field → uint, shift and mask (a field ending at bit 31 needs no mask) |
| Literal | Every field, or .. for zeros; values are masked to their width; literal values fold at compile time |
| Write | w.field = v (plain = only) |
| Conversion | A bitfield reads as its uint wherever one is expected; a uint becomes one only through raw as Word |
Unknown, misplaced, malformed, duplicate, and conflicting attributes are errors. Compute thread tuples contain exactly three positive integer dimensions.
fn without a stage attribute is a free function (shared helper). Stage entries use attributes below.
Visibility (private)
Unmarked items stay import-visible (same as today). Prefix private to keep a
symbol local to its defining #module / #shader — importers cannot call or
name it; the defining file still can. Stored on struct / fn / enum /
const / bitfield / instances. Private types remain usable in the defining
module's helpers, including signatures and local annotations. A parameter or
local binding with the same name as a private value shadows it normally.
#module brdf
private fn schlick_weight(n_dot_v: float) -> float { … }
fn hikari_fresnel_schlick_nov(f0: float3, n_dot_v: float) -> float3 {
return … schlick_weight(n_dot_v) …
}Default arguments, named arguments, overloads
fn eval_ibl(..., reflections_own_specular: bool = true) -> float3 { … }
// Trailing defaults may be omitted; names reorder before arity checks.
let a = eval_ibl(u, sky, …, N, V, …);
let b = eval_ibl(…, reflections_own_specular: false);
// Same source name, different parameter types → overload (emit uses distinct symbols).
fn fresnel(f0: float3, n_dot_v: float) -> float3 { … }
fn fresnel(f0: float, n_dot_v: float) -> float { … }A scalar user function also lifts to vectors when no exact overload exists: srgb(float3) becomes float3(srgb(x), srgb(y), srgb(z)). An explicit vector overload still wins.
ref / inout
Passing convention, written after the colon. Handles (Buf, textures, samplers) ignore it — they are already cheap.
fn bump(quota: inout uint) { quota = quota - 1u }
fn read_row(row: ref VsOut) -> float4 { return row.position }inout is a mutable borrow (writes go back). ref is read-only. Both require an lvalue. Buffer elements (buf[i]) cannot be borrowed — pass Buf<T> and the index (a ref T of a float4x4[6] row would copy on HLSL).
A context is a struct of resource handles. It is a compile-time parameter pack — backends flatten it into per-handle arguments, because ordinary GPU structs cannot hold textures, samplers, or device pointers:
struct LightCtx {
lights: Buf<GpuLightEntry>;
atlas: TexDepth2D;
samp: SamplerCmp;
}
fn evaluate(ctx: LightCtx, ...) { … hikari_load_light(ctx.lights, i) … }
// emits as evaluate(ctx_lights, ctx_atlas, ctx_samp, ...)Function values
Compile-time only. Parameters of type fn(T) -> R specialize the callee at each call site (no GPU function pointers):
fn apply(x: float, f: fn(float) -> float) -> float { return f(x) }
let y = apply(3.0, scale);The argument must be a free-function name or a closure. Closures capture outer locals (including resource packs) and are specialized the same way — captures become extra parameters, not GPU function pointers:
fn apply(x: float, f: fn(float) -> float) -> float { return f(x) }
let y = apply(3.0, |v| v * s);
let vis = evaluate_lights(ctx, |light, kind, L, pos, depth, dist, NoL, shadowed| {
var v = 1.0;
if (shadowed) { v = raster_vis(light, kind, L, pos, depth) }
return v;
})Captures of @groupshared locals are read-only borrows that retain their threadgroup address space; they do not copy the shared array into each invocation. Metal retains a threadgroup reference; HLSL reads the owning function's qualified global storage directly. Synchronize the producing lanes with barrier() before reading through a closure.
Parameter types may be omitted when the expected fn(...) type supplies them. |x: inout uint| { x = x - 1u } and || expr (no parameters) are valid.
A closure may be bound to a let and used several times. It is lowered where
it is used and captures by reference, seeing the bindings of its definition
(later shadowing does not leak in). Calling one directly needs typed parameters
and a return type:
let tap = |offset: float2| -> float4 { return history.sample(samp, uv + offset * step); };
let sum = tap(float2(0.0)) + filter(uv, tap);A fn-typed parameter can be passed straight on to another fn-typed
parameter, and may have a default (a function name or closure):
fn atrous(uv: float2, receiver: fn(float2) -> float = receiver_at_texel) -> float { … }Compile-time parameters (const, optional resources) and instances
A const parameter is a compile-time value; a resource parameter written T?
may be none. Like fn parameters, they make the function a template: it
is never emitted as written, only as a specialization per distinct argument set
(pick__mirrored_true). Inside the specialization a const is its literal, so
it can size arrays, and any if / ?: whose condition folds — const values,
literals, == none / != none — keeps only the taken arm. The other arm is
still type-checked.
fn reserve(request: uint, const mirrored: bool = false) -> uint {
if (mirrored) { … } // folded per specialization
…
}
fn shade(scene: AccelStruct?, p: float3) -> float {
if (scene != none) { return rt_visibility(scene, p, up, 0.01, far, 255u); }
return 1.0;
}
let a = shade(none, p); // `shade__scene_none` has no TLAS parameterfn name = target(args) names a specialization. Its arguments fix only const,
fn and optional parameters (positional ones fill those in order); the runtime
parameters and the stage carry over, so it is how one source produces several
entry points. An instance is an ordinary function otherwise: code anywhere in
the package can call it by name with the runtime arguments. Instances take the
template's place in emission order.
@fragment
fn ambient_occlusion_atrous(const step: float, @stage_in inn: V) -> float4 { … step … }
fn ambient_occlusion_atrous_1 = ambient_occlusion_atrous(step: 1.0);
fn ambient_occlusion_atrous_2 = ambient_occlusion_atrous(step: 2.0);Const values are literals, consts, enum cases and + - * / / && || ! over
them. none is valid only as an optional argument and in == none / != none.
When the host loads each choice as its own package, the template lives in a
module and every package file is one instance line — tonemap_pass.akari holds
the pass with const operator: TonemapOperator, and tonemap.akari,
tonemap_reinhard.akari, … each contain
fn tonemap_fragment = tonemap_fragment_for(operator: TonemapOperator.Reinhard);.
Stages and system values
@vertex
fn vs(…) -> VsOut { … }
@fragment
fn ps(…) -> float4 { … }
@compute(threads = (8, 8, 1))
fn cs(…) { … }| Attribute | Typical type | Role |
|---|---|---|
@vertex_id | uint | Vertex index |
@instance_id | uint | Instance index |
@dispatch_id | uint / uint3 | Global thread id |
@group_id | uint / uint3 | Workgroup id |
@local_id | uint3 | Thread in group |
@local_index | uint | Flat index in group |
@front_facing | bool | Front face |
@stage_in | struct | Interpolated inputs to fragment |
There is no vertex-input attribute. A vertex stage fetches its own data by @vertex_id / @instance_id — scene meshes through bindless_geometry (hikari_load_mesh_vertex), other draws from a bound Buf<T> — and no pipeline declares a vertex layout.
Resources and bindings
Prefer semantic binds:
@bind(frame.uniforms) uniforms: Uniform<HikariUniforms>;
@bind(draw.object) objects: Buf<HikariObjectData>;
@bind(debug.albedo) albedo: Tex2D;
@bind(debug.shadow_depth) atlas: TexDepth2D;
@sampler(linear, repeat) samp
@sampler(comparison, greater_equal, linear, clamp) shadow_s: SamplerCmp;| Form | Notes |
|---|---|
@bind(group.name) | Resolved via HK_BUFFER / HK_TEXTURE in the manifest |
@bind(group.name) tex: TexCube[N] / Tex2D[N] / RwTex2D[N] | Resource array: base key + N consecutive texture slots (Metal array<…, N> [[texture(base)]], HLSL Texture… name[N] : register(tBase)). No new host bind keys — declare only the base (e.g. depth_pyramid_compute.mips). |
@slot(N) | Raw slot (rare; compute experiments) |
@sampler(filter, address) | linear/point + clamp/repeat/mirror. Must be a state of the manifest's static_samplers set, whose row gives its HLSL register / Vulkan binding; Metal bakes a constexpr sampler. |
@sampler(comparison, filter, address) | Comparison sampler; compare func defaults to less_equal |
@sampler(comparison, compare_func, filter, address) | Explicit compare (greater_equal, less_equal, less, greater, equal, not_equal, always, never) — pick to match the host depth convention (Hikari reverse-Z → greater_equal) |
@bind(depth_pyramid_compute.mips) mips: RwTex2D<rg16f>[8],
// mips[i].store(px, value)Unknown binds error when the manifest is loaded. Soft/unit tests may keep empty manifests.
Package bindings
A #shader may declare its resources once at package scope instead of on every
entry:
@bind(vsm.frame) let frame: Uniform<VsmFrame>;
@bind(vsm.pages) let pages: RwBuf<uint>;
@sampler(point, clamp) let point_s;
fn mark(page: uint) { pages[page] = frame.epoch; }
@compute(threads = (64, 1, 1))
fn vsm_mark(@dispatch_id id: uint) { mark(id); }Every function of the package — and the closures and specializations made from
it — can name them. After specialization, each entry receives exactly the
bindings it can reach (through calls too), appended after its own parameters in
declaration order; helpers that reach one get it as a trailing parameter. A
binding behind a folded const / none branch is not part of that entry's
interface. A local or parameter with the same name shadows the binding.
Modules cannot declare bindings — their layout belongs to the package that
imports them.
@sampler registers come from the sampler's state (static_samplers in the
manifest), so neither declaration order nor which samplers an entry reaches
changes them.
MRT outputs:
struct GBufferOutput {
@target(gbuffer.albedo) albedo: float4;
@target(gbuffer.normal) normal: float4;
// …
}Control flow
if (cond) { … } else if (other) { … } else { … }
for (var i: uint = 0u; i < n; i = i + 1u) { … }
while (cond) { … }
do { … } while (cond)
switch (x) {
case 0:
…
break;
case 1:
case 2:
…
break;
default:
…
break;
}
break;
continue;
discard; // fragment killFlow hints
On if / for / while / do:
| Akari | HLSL | Metal |
|---|---|---|
@unroll / @unroll(N) | [unroll] / [unroll(N)] | [[unroll]] (N dropped) |
@loop | [loop] | [[dont_unroll]] |
@branch | [branch] | (no-op) |
@flatten | [flatten] | (no-op) |
@unroll
for (var i: uint = 0u; i < 6u; i = i + 1u) { … }
@branch
if ((mask & 1u) != 0u) { … }Textures
Texture operations are methods on the texture value:
| Op | Example | Notes |
|---|---|---|
| sample | tex.sample(samp, uv) | VS/CS force lod 0 on HLSL |
| sample_level | tex.sample_level(samp, uv, mip) | Explicit lod |
| sample_cmp | depth.sample_cmp(cmp_samp, uv, ref) | Returns float |
| sample_grad | tex.sample_grad(samp, uv, ddx, ddy) | Explicit gradients |
| gather | tex.gather(samp, uv) / tex.gather(samp, uv, c) | One component (c = constant 0..3, default 0) of the four bilinear-footprint texels; Tex2D, Tex2DArray, TexCube, TexDepth2D (one component). Texel order (−,+) (+,+) (+,−) (−,−) on every backend |
| gather_cmp | depth.gather_cmp(cmp_samp, uv, ref) | Four comparison results; TexDepth2D only |
| load | tex.load(coord) / tex.load(coord, mip) | Integer coords; result follows the texel type |
| store | rw.store(coord, value) | RwTex2D (uint2) / RwTex3D (uint3) |
| dimensions | tex.dimensions() / tex.dimensions(mip) | uint2 or uint3 |
| mip_levels | tex.mip_levels() | Mip count |
Array / cube packing
Tex2DArray: UV + layer asfloat3(uv, layer)(Metal splits.xy+uint(.z)).TexCubeArray: direction + layer asfloat4where applicable.
Store arity
store must supply at least the target format's channel count — float2 for
rg16f, float3 for r11g11b10f, a scalar for r16f/r32f/r32u/r32i,
float4 otherwise. Extra components are dropped (HLSL swizzles the assignment
down; Metal's hardware ignores them), so a float4 store into any format stays
valid. A short store is a compile error at lowering rather than a backend
error: HLSL rejects it while Metal's write — which only has a vec<T, 4>
overload the emitter zero-fills to reach — would have accepted it, so the check
keeps the two backends agreeing on what compiles.
Integer textures
No hardware filters integer texels, so the whole sample family is rejected on
Tex2D<uint> / Tex2D<int> at compile time rather than surfacing as a DXC or
Metal error against generated source. Read them with load and integer
coordinates. Use them for bitfields — the G-buffer material target is one —
where a filtered read would silently blend unrelated flags across an edge.
Depth textures
sample_cmp requires TexDepth2D (Metal only exposes sample_compare on
depth2d). A colour Tex2D is a compile error. load, sample,
sample_level, and sample_grad on a depth texture return a scalar float
(not float4) — matching Metal depth2d and HLSL Texture2D<float>.
let c = arr.sample(samp, float3(uv, float(layer)));
let shadow = atlas.sample_cmp(shadow_s, uv, ref_z);
let d = atlas.load(px); // float
rw.store(uint2(id.xy), float4(h, 0.0, 0.0, 0.0));Compute
@compute(threads = (64, 1, 1))
fn meshlet_cull(
@bind(meshlet_cull.frame) frame: Uniform<CullFrame>,
@bind(meshlet_cull.arguments) arguments: AtomicBuf<uint>,
@dispatch_id index: uint,
@local_id lid: uint3,
@local_index lix: uint,
) {
@groupshared var tile: float[64];
// …
barrier();
let slot = atomic_add(arguments, batch * 8u + 1u, 1u);
}| Feature | Notes |
|---|---|
@groupshared var name: T | Threadgroup / groupshared storage |
barrier() | Group memory + execution barrier |
| Atomics | See below (compute entries) |
Atomics
On AtomicBuf / compatible UAV storage:
| Call | Args |
|---|---|
atomic_add / atomic_sub | (buf, index, value) |
atomic_min / max / and / or / xor | same |
atomic_exchange | (buf, index, value) |
atomic_compare_exchange | (buf, index, expected, desired) → previous |
Return type is typically the previous value (uint path). Stage: compute entries only (free helpers allowed; stage checked on entries).
Indexing an AtomicBuf (buf[i], buf[i] = v) is a relaxed atomic load / store
on Metal and a plain UAV access on HLSL. Where another thread group's writes in
the same dispatch must be observed on D3D12, read with an RMW such as
atomic_add(buf, i, 0u).
Lane groups (wave)
Ops across a SIMD-group / wave. The point is to replace N same-address atomics with one: have every lane vote, prefix-sum the votes for a local offset, and let a single lane take the group's base.
| Call | Metal | HLSL |
|---|---|---|
wave_prefix_sum(uint) -> uint | simd_prefix_exclusive_sum | WavePrefixSum |
wave_sum(uint) -> uint | simd_sum | WaveActiveSum |
wave_is_first() -> bool | simd_is_first | WaveIsFirstLane |
wave_read_first(T) -> T | simd_broadcast_first | WaveReadLaneFirst |
wave_any(bool) -> bool / wave_all(bool) -> bool | simd_any / simd_all | WaveActiveAnyTrue / WaveActiveAllTrue |
wave_count(bool) -> uint | simd_sum(b ? 1u : 0u) | WaveActiveCountBits |
wave_prefix_count(bool) -> uint | simd_prefix_exclusive_sum(b ? 1u : 0u) | WavePrefixCountBits |
Prefix sums are exclusive on both backends. Stage: compute and fragment — Metal has no SIMD-group functions in a vertex stage, so the language rejects it rather than emitting something only D3D12 can run. D3D12 needs SM 6.0, which is already the floor.
Both backends define these over the active lanes, so calling them inside
divergent control flow gives the divergent answer. The portable idiom is to vote
0/1 from every lane instead of branching around the op:
var vote = 0u;
if (survived) { vote = 1u }
let offset = wave_prefix_sum(vote);
let total = wave_sum(vote);
var base = 0u;
if (wave_is_first() && total > 0u) { base = atomic_add(counters, 0u, total) }
base = wave_read_first(base);
if (survived) { out_list[base + offset] = value }Math and builtins
Portable math (renamed per backend as needed; Akari spells them mix, fract):
min max clamp abs floor ceil round trunc sign fract saturate
sin cos tan asin acos atan atan2
sqrt rsqrt exp log exp2 log2 pow fma
normalize length distance dot cross reflect refract
mix step smoothstep any all select isnan isinf
min3 max3 median3
transpose mul
countbits firstbitlow firstbithigh reversebits extract_bits(x, offset, count) insert_bits(base, bits, offset, count) (integer bit ops on uint/uintN; firstbithigh(0u) is 0xFFFFFFFFu on every backend)
ddx ddy (fragment entries)
remap_clip_z (engine clip-Z convention)
Packing, one formula on every backend (x in the lowest bits):
| Builtin | Formula |
|---|---|
pack_unorm4x8(float4) -> uint / unpack_unorm4x8(uint) -> float4 | uint(clamp(v, 0, 1) * 255.0 + 0.5) per byte / float(byte) / 255.0 |
pack_unorm2x16(float2) -> uint / unpack_unorm2x16(uint) -> float2 | the same with 65535.0 |
pack_snorm2x16(float2) -> uint / unpack_snorm2x16(uint) -> float2 | int(floor(clamp(v, -1, 1) * 32767.0 + 0.5)) / max(float(int16) / 32767.0, -1.0) |
pack_half2x16(float2) -> uint / unpack_half2x16(uint) -> float2 | IEEE half, round to nearest even |
pack_uint4x8(uint4) -> uint / unpack_uint4x8(uint) -> uint4 | each component masked to 8 bits |
Engine helpers live in modules (hikari_*, lighting, etc.) and are normal functions after #import.
Ray tracing
Language intrinsics (rt_*, bindless_*). Engine helpers (hikari_rt_*,
hikari_rt_candidate_accepts, …) live in modules and are normal functions after
#import.
let vis: float = rt_visibility(scene, origin, dir, t_min, t_max, mask);
// The same any-hit query reporting where it stopped: the accepted hit's ray
// parameter, or -1.0 when the segment is clear. Any hit, not the nearest.
let occluder_t: float = rt_occluder_distance(scene, origin, dir, t_min, t_max, mask);
let hit: RtTraceHit = rt_trace(scene, origin, dir, t_min, t_max, mask);
// hit.hit, hit.distance, hit.instance_id, hit.primitive_id, hit.barycentric
// Alpha-aware: non-opaque candidates run hikari_rt_candidate_accepts
// (engine helper; import raytraced_scene) inside the traversal.
let vis_m: float = rt_visibility_masked(
scene, instances, materials, buffer_table, texture_table,
instance_count, origin, dir, t_min, t_max, mask);
let occluder_t_m: float = rt_occluder_distance_masked(
scene, instances, materials, buffer_table, texture_table,
instance_count, origin, dir, t_min, t_max, mask);
let hit_m: RtTraceHit = rt_trace_masked(
scene, instances, materials, buffer_table, texture_table,
instance_count, origin, dir, t_min, t_max, mask);
let f = bindless_load_f32(table, buf_index, elem);
let f2 = bindless_load_f32x2(table, buf_index, elem);
let f3 = bindless_load_f32x3(table, buf_index, elem);
let f4 = bindless_load_f32x4(table, buf_index, elem);
let u = bindless_load_u32(table, buf_index, elem);
let t2d = bindless_tex2d(texture_table, index);
let cube = bindless_texcube(texture_table, index);
let vol = bindless_tex3d(texture_table, index);
let depth = bindless_texdepth(texture_table, index);- Metal:
#include <metal_raytracing>+ helper bodies when RT is used. - D3D12: DXC shader model ≥ 6_5 when
ray_queryis required (auto if the spec omitsshaderModel).
Emit hygiene (authors)
- Only what entries reach is emitted. A package file carries the functions,
constants and structs its entry points use, transitively; everything else an
import brings in is dropped. On D3D12 and Vulkan each entry is its own file
and carries only its own closure; Metal's one file per package carries the
union. A
#moduletranspiled on its own (no entries) emits everything. - Metal: non-entry functions get forward prototypes, so helper order in a file does not matter.
- HLSL: reserved identifiers (e.g. field
sample) are renamed on both declaration and access (_sample). Prefer non-reserved names in new code.
Capabilities (#require)
Optional file header. Also inferred from use.
| Cap | Triggered by | Stage rule (entries) |
|---|---|---|
derivatives | ddx / ddy | @fragment only |
atomics | atomic_* | @compute only |
threadgroup | barrier, @groupshared | @compute only |
wave | wave_* | @compute / @fragment |
ray_query | rt_*, AccelStruct | any |
rw_texture | .store / RwTex2D / RwTex3D | any |
#require ray_query, atomics- D3D12:
ray_query⇒ SM 6_5 minimum (auto or reject too-low explicit SM). - Metal: no flag matrix; RT packages emit raytracing includes. Stage gates still apply.
- Unknown names → error at the misspelled name:
unknown capability…``.
You rarely need #require if you call the ops correctly; it documents intent and fails early on typos.
Type checking
At IR lower (before metalc/DXC):
- Unknown symbols / functions / fields / methods
- Wrong arity (user fns, math, texture ops)
- Assign / init / return type mismatch
- Conditions (
if, loops,?:) and&&/||operands must bebool;?:also takes aboolvector, selecting per lane - Arithmetic, bitwise and comparison operands must be scalars, vectors or
matrices, and two vectors must have the same width (
float3 * float2is an error; a scalar still broadcasts) - Capability stage violations
Every parse and type error names its file, line and column, quotes the source line, and — inside a specialized template — says where the specialization was requested:
error: modules/m.akari:3:57: unknown field `nope` on `uint`
| fn scale(const k: uint, x: float) -> float { return x * k.nope; }
= note: in instance `triple` at packages/p.akari:5:1The CLI also prints each as @akari diag JSON on stderr for tooling. Library
callers get AkariError::Diagnostics(sources, diagnostics);
diagnostic_messages() renders them as above without the severity.
When metalc/DXC fails, diagnostics are remapped through the structured emit line map to .akari path and line when possible. Toolchain “candidate function” notes and known Metal attribute noise are stripped so the first error is actionable.
Complete mini examples
Vertex + fragment (fullscreen texture copy)
This is a complete language example for a fullscreen pass. Scene materials must also satisfy the engine's G-buffer contract; use the current gbuffer.akari package as their starting point. Rigid scene meshes use meshlet attribute pulling rather than hardware vertex input.
#shader texture_copy
#layout fullscreen_texture_table
#import preamble, fullscreen_triangle
struct VsOut {
@position position: float4;
uv: float2;
}
@vertex
fn texture_copy_vertex(@vertex_id vid: uint) -> VsOut {
return VsOut {
position: hikari_full_screen_position(vid),
uv: hikari_full_screen_uv(vid),
};
}
@fragment
fn texture_copy_fragment(
@stage_in inn: VsOut,
@bind(post.hdr) source: Tex2D,
@sampler(linear, clamp) samp,
) -> float4 {
return source.sample(samp, inn.uv);
}Compute + atomics + groupshared
Dispatch a multiple of 64 threads and provide at least that many source elements plus one initialized counter. Every lane must reach the barrier; a partial final group needs an explicit bounds-aware load rather than an early return.
#shader reduce
#layout standard
#require atomics, threadgroup
@compute(threads = (64, 1, 1))
fn reduce(
@bind(meshlet_cull.source) source: Buf<float>,
@bind(meshlet_cull.arguments) counters: AtomicBuf<uint>,
@dispatch_id gid: uint,
@local_index li: uint,
) {
@groupshared var partial: float[64];
partial[li] = source[gid];
barrier();
if (li == 0u) {
var s: float = 0.0;
@unroll
for (var i: uint = 0u; i < 64u; i = i + 1u) {
s = s + partial[i];
}
if (s > 0.0) {
let _ = atomic_add(counters, 0u, 1u);
}
}
}Content compute write (RwTex2D)
#shader water_sim
#layout standard
#import content_compute
@compute(threads = (8, 8, 1))
fn water_sim(
@bind(content_compute.output) output: RwTex2D<r32f>,
@bind(content_compute.frame) frame: Uniform<HikariContentComputeFrame>,
@dispatch_id id: uint3,
) {
if (id.x >= uint(frame.resolution) || id.y >= uint(frame.resolution)) { return; }
let uv = float2(id.xy);
let h = sin(uv.x * 0.1 + frame.simTime) * 0.5;
output.store(id.xy, float4(h, 0.0, 0.0, 0.0));
}Tooling (authors)
# Transpile one package (--target=metal | d3d12 | spirv)
bin/akari/akari transpile --input=path.akari --target=metal --out=/tmp/out \
--import-dir=src/hikari/shaders/akari/modules --bindings=src/hikari/shaders/bindings.kaji.json
# Full spec compile (Metal, D3D12 or Vulkan — build.akari.{metal,d3d12,vulkan}.json)
bin/akari/akari compile \
--spec=src/hikari/shaders/build.akari.metal.json \
--out=/tmp/shaders --cache=/tmp/akari-cache \
--engine-sdk=. --config=debug
# Layout dump (CPU mirror checks)
bin/akari/akari transpile … --dump-layoutProduct builds go through Kaji / Shinra (--shaders:config=debug|release → Shinra
--shaders-config= → akari compile --config=). Binary:
flat bin/akari/akari.
Authoring rules (product)
- Pure
.akarionly — no hand-edited MSL/HLSL SoT. @bind/@target/@samplerover raw registers (except deliberate@slotexperiments).- Prefer
#importmodules over copy-paste. - Keep CPU struct layouts aligned with ABI (
--dump-layout). - Material features → runtime map masks, not Cartesian permutations. Entry points
that differ by a constant are one
const-parameter template plusfninstances. - Resources several entries share are package bindings, not repeated parameters.
- Add new binds to
bindings.kaji.jsonand shaders.md together.
Related docs
| Doc | Contents |
|---|---|
src/akari/syntax.ebnf | Formal grammar (W3C EBNF; akari_syntax) |
| Shader authoring | Passes, layouts, manifests, specs, validation |
| Rendering | Graph order, composability contracts |
| ABI | Offsets, packing, @packed |