SwiftUI52 views · 75 mods Stdlib586 types · 200+ fn Foundation87+ functions Metal106 kw · 90+ builtins
Standalone library: /Users/2n/dev/miniswift/metal — links as libMetal.a

Metal — MSL → WGSL Transpiler

A four-pass C library that takes Apple Metal Shading Language source and produces WebGPU Shading Language plus a JSON pipeline-metadata sidecar. No LLVM, no codegen runtime — just lexer → parser → lower → emit. 224 tests, 100% pass.

106
MSL keywords lexed
90+
Type mappings (scalar · vector · matrix · texture)
90+
Builtin function rewrites
224
Tests passing (lexer · parser · MIR · WGSL · editor)
Full — translates to WGSL · Partial — translates with helper · Reject — diagnostic + migration hint
Pipeline Lexer (106 kw) Types (90+) Builtins (90+) Attributes Emit (WGSL + JSON) Loud rejects

The Pipeline

Four passes — each surfaces structured diagnostics on failure, never NULL or a corrupt result.

1. Lexer

Tokenises MSL — 106 keywords, full C99 operator set, double-bracket attributes [[position]] / [[buffer(0)]], preprocessor directives. Single-pass, malloc-free hot loop.

2. Parser

Recursive descent for declarations + statements, Pratt precedence climbing for expressions. Handles the full MSL surface: structs, function templates (single-param + multi-param + non-type int), vertex/fragment/kernel qualifiers, attributes, default arguments, function overloading.

3. Lower (AST → MIR)

Type mapping (90+ scalar/vector/matrix/texture entries), attribute → builtin/location resolution, ~90 builtin function rewrites, template monomorphization, address-space tracking, function-constant detection.

4. Emit

MIR → WGSL text + JSON pipeline metadata. Operator-precedence-aware parenthesization, struct entry-point I/O detection, storage texture format/access threading.

Preprocessor

#include resolution + simple #define substitution. 21 predefined macros (__METAL__, __METAL_VERSION__, …).

Editor API

msl_lex + msl_semantic_tokens for IDE highlighting; LSP-compatible 1-based line/column diagnostics.

Lexer Keywords

All 106 reserved words tokenised correctly. Grouped here by purpose — each item lexes; what happens after lex is governed by parser/lower coverage.

Function qualifiers5
vertex / fragment / kernel
Entry-point qualifiers → @vertex / @fragment / @compute in WGSL.
constant / device / threadgroup
Address-space qualifiers — tracked through MIR; emitted as var<uniform> / var<storage> / var<workgroup>.
thread / threadgroup_imageblock
Stack + tile-memory qualifiers.
constexpr
Compile-time evaluable; folded where possible.
Scalar types11
int / uint / short / ushort / char / uchar / long / ulong
Integer scalars — i32 / u32 / i32 / u32 / i32 / u32 / i32 / u32 in WGSL.
bool
Boolean scalar.
float / half
f32 / f16 (with half emulation when WGSL impl lacks half).
bfloat
Brain-float — recognised; widens to f32 if WGSL host doesn't expose bf16.
Vector types32
float2 · float3 · float4
vec2f / vec3f / vec4f.
half2 · half3 · half4
vec2<f16> / vec3<f16> / vec4<f16>.
int2 · int3 · int4 (+ uint / short / ushort / char / uchar)
Full N×{2,3,4} integer vector grid.
bool2 · bool3 · bool4
Boolean vectors.
Matrix types18
float2x2 · float2x3 · float2x4 · float3x2 · float3x3 · float3x4 · float4x2 · float4x3 · float4x4
All 9 column-major float matrices → mat<C>x<R>f.
half2x2 … half4x4
Same 9-cell grid for half precision.
Texture types14
texture1d / texture2d / texture3d
Standard sampleable textures.
texture1d_array / texture2d_array / texturecube / texturecube_array
Array + cubemap variants.
texture2d_ms / texture2d_ms_array
Multi-sampled textures.
depth2d / depth2d_array / depth2d_ms / depthcube / depthcube_array
Depth texture flavours.
Atomic types2
atomic_int / atomic_uint
→ atomic<i32> / atomic<u32>. Full atomic_* op family rewritten.
Control flow + storage14
if / else / for / while / do / switch / case / default / break / continue
Standard C-style control flow → WGSL equivalents.
return / discard_fragment
Return / fragment kill.
struct / enum / class / typedef / using
Type declarations — class restricted to value-style use.
const / static
Mutability + linkage.

Type Mapping

90+ scalar / vector / matrix / texture / sampler entries cross-walked from MSL to WGSL. Highlights:

Scalar widening

char / short → i32; uchar / ushort → u32. WGSL only has 32-bit integer scalars, so narrower MSL ints widen with explicit conversion at boundaries.

Half precision

half / halfN → f16 / vecN<f16>. Pack/unpack helpers (pack_half_to_unorm2x16 etc.) emitted as helper functions.

Matrix layout

MSL matrices are column-major; WGSL matrices are also column-major. floatCxR → matCxRf with the same indexing.

Texture formats

Read-only texture2d<float> → texture_2d<f32>; access mode (access::read / ::write / ::read_write) determines storage texture vs sampleable texture.

Address spaces

Tracked end-to-end: constant & → var<uniform>, device * → var<storage>, threadgroup → var<workgroup>, thread → var<function>.

Sampler

sampler + access::sample mapped to WGSL sampler; sampler descriptor metadata routed into pipeline JSON.

bfloat

Recognised; widens to f32 when target environment lacks bf16.

imageblock

Tile-memory threadgroup_imageblock tokenises but lowers conservatively.

Builtin Functions

~90 builtin rewrites — direct passthrough where WGSL has the same name, helper-fn emission where it doesn't. Helpers are declared once at the top of the WGSL output, never inlined per call.

Math & trig25+
abs / sign / clamp / min / max / mix / step / smoothstep
Direct WGSL passthrough.
sin / cos / tan / asin / acos / atan / atan2 / sinh / cosh / tanh / asinh / acosh / atanh
Trig family — direct.
exp / exp2 / log / log2 / pow / sqrt / rsqrt / inversesqrt
Exponential / logarithm / sqrt family.
ceil / floor / round / trunc / fract / fmod / modf
Rounding + remainder.
sincos
Emitted as __sincos(x, &s, &c) helper; calls sin + cos under the hood.
min3 / max3 / median3 / fmedian3
Three-arg variants — emitted as helpers.
Vector & geometric12
dot / cross / length / normalize / distance / reflect / refract / faceforward
Standard vector ops — direct.
distance_squared / length_squared
Squared variants — emitted as helpers.
all / any / select
Boolean vector reductions + branchless select.
Matrix5
transpose / determinant / inverse
Standard matrix ops.
m * v / m * m
Operator overloading honours MSL semantics in WGSL.
Atomic operations18
atomic_load / atomic_store / atomic_exchange
Plain + _explicit variants → atomicLoad / atomicStore / atomicExchange.
atomic_fetch_add / sub / and / or / xor / min / max
Full RMW family — direct WGSL atomic rewrites.
atomic_compare_exchange_weak / strong (+ _explicit)
CAS variants → atomicCompareExchangeWeak.
atomic_thread_fence
→ storageBarrier() / workgroupBarrier() by scope.
Pack / unpack10+
pack2x16snorm / pack2x16unorm / pack4x8snorm / pack4x8unorm
→ WGSL pack2x16snorm et al.
pack_half_to_snorm2x16 / unorm2x16 / snorm4x8 / unorm4x8
Half-format pack/unpack — emitted as helpers.
Saturating conversion6
convert_char_saturate / short / int / uchar / ushort / uint
Clamping conversions — emitted as helpers (WGSL has no native saturating cast).
add_sat / sub_sat / abs_diff / hadd
Saturating arithmetic — emitted as helpers.
SIMD-group / quad10+
quad_min / quad_max / quad_and
Quad-scope reductions — emitted as helper subgroup ops; depend on WGSL subgroup extension support.
simd_sum / simd_min / simd_max / simd_and / simd_or / simd_xor / simd_product
SIMD-group reductions — same caveat.
Texture sampling10+
.sample(sampler, coord) / .sample(... level: ...) / .sample(... grad2d: ...)
→ textureSample / textureSampleLevel / textureSampleGrad.
.read(coord) / .write(value, coord)
Storage-texture access → textureLoad / textureStore.
.get_width() / .get_height() / .get_depth() / .get_num_levels()
→ textureDimensions / textureNumLevels.
.gather(...) / .gather_compare(...)
→ textureGather / textureGatherCompare.

Attributes & Builtins

Double-bracket attributes [[...]] resolve to WGSL @builtin(...) / @location(...) / @group / @binding.

Vertex / fragment builtins12
[[position]]
→ @builtin(position).
[[vertex_id]] / [[instance_id]]
→ @builtin(vertex_index) / @builtin(instance_index).
[[point_size]]
Point primitive size; emitted where supported.
[[point_coord]]
Point fragment coord.
[[front_facing]]
→ @builtin(front_facing).
[[sample_id]] / [[sample_mask]]
MSAA-related; tracked through MIR.
[[stage_in]]
Marks struct as entry-point input — fields lift to @location(N).
Compute builtins7
[[thread_position_in_grid]]
→ @builtin(global_invocation_id).
[[thread_position_in_threadgroup]]
→ @builtin(local_invocation_id).
[[threadgroup_position_in_grid]]
→ @builtin(workgroup_id).
[[thread_index_in_threadgroup]]
→ @builtin(local_invocation_index).
[[threads_per_threadgroup]] / [[threads_per_grid]]
Workgroup + grid size queries.
Resource bindings5
[[buffer(N)]]
Buffer index → @group(0) @binding(N). Group/binding numbers tracked in pipeline JSON.
[[texture(N)]]
Texture binding slot.
[[sampler(N)]]
Sampler binding slot.
[[threadgroup(N)]]
Threadgroup memory binding.
[[function_constant(N)]]
Specialisation constant → @id(N).
User-data attributes3
[[user(name)]]
Custom semantic on struct fields.
[[attribute(N)]]
Vertex attribute slot — folds into @location(N).
[[color(N)]]
Fragment colour output → @location(N).

Emit Pass — WGSL + JSON

The fourth pass turns MIR into two artifacts: a WGSL shader text and a JSON pipeline-metadata sidecar.

WGSL text emission

Operator-precedence-aware parenthesization (no over-parenthesizing nor missing parens). Struct entry-point I/O is detected and rewritten.

Helper functions

Operations WGSL doesn't expose (sincos, quad reductions, saturating conv, half pack/unpack) are emitted as __name(...) helpers declared once at the top.

Storage texture metadata

Format + access mode threaded into the WGSL declaration: texture_storage_2d<rgba8unorm, write>.

Pipeline JSON

Lists entry points (vertex / fragment / compute), each with its bindings (group / binding / kind), vertex-attribute layout (location → format), and any specialisation constants.

Diagnostics

1-based line/column ranges (LSP-compatible). Each diagnostic carries severity (ERROR / WARNING / INFO / HINT) + structured message.

Arena allocation

AST + MIR allocate from chunk-based arenas; msl_free(result) releases everything at once. Zero ref-counting overhead in the hot loop.

Loud Rejects

MSL features without a WGSL equivalent emit structured diagnostics with migration hints — never silently produce gibberish WGSL.

Ray tracing

raytracing::* namespace, ray_data, intersection_query. WebGPU doesn't expose ray tracing.

Mesh shaders

mesh, object, amplification qualifiers + mesh_grid_properties. No WGSL equivalent.

Inline samplers

constexpr sampler s(...) at file scope. WGSL samplers must be uniform-bound.

Lambdas / function pointers

Closures, visible_function_table, indirect dispatch. WGSL has no first-class functions.

Recursive functions

Detected during lower; rejected with cycle trace. WGSL is non-recursive by spec.

Simdgroup matrices

simdgroup_matrix + matrix multiply ops. WGSL subgroup matrix extension not yet shipping.

Imageblock data

imageblock<T> tile-memory storage. Tokenises but lower returns a TODO diagnostic.

Argument buffers

Tier-2 argument buffers with descriptor indexing. WebGPU supports limited indexing only.

Visible function table

Indirect callables for ray-tracing intersection / visibility programs.

Tessellation

Hull / domain shaders + patch primitives. WebGPU has no tessellation pipeline.

Pulled vertex fetch

[[stage_in]] with no buffer attribute → vertex-fetch attribute pulling. Falls back to per-attribute bindings.

Performance counters

counter, visibility_buffer. No corresponding WGSL feature.