diff --git a/build.zig b/build.zig
index 8fef2dc..c0ff663 100644
--- a/build.zig
+++ b/build.zig
@@ -512,7 +512,7 @@ fn customPhi(
     _: *std.Build.Module,
     _: *std.Build.Module,
     base_c_mod: *std.Build.Module,
-    _: *std.Build.Module,
+    shader_ir_mod: *std.Build.Module,
     target: std.Build.ResolvedTarget,
     optimize: std.builtin.OptimizeMode,
     use_llvm: bool,
@@ -533,6 +533,7 @@ fn customPhi(
     options.addOption([]const u8, "phi_daemon_host_prefix", daemon_host_prefix);
 
     lib_mod.addImport("phi_c", base_c_mod);
+    lib_mod.addImport("shader_ir", shader_ir_mod);
 
     const miclib = b.lazyDependency("miclib", .{
         .target = target,
diff --git a/src/phi/compiler/abi.zig b/src/phi/compiler/abi.zig
new file mode 100644
index 0000000..96ed944
--- /dev/null
+++ b/src/phi/compiler/abi.zig
@@ -0,0 +1,25 @@
+pub const version: u32 = 1;
+
+pub const BufferDescriptor = extern struct {
+    address: u64,
+    size: u64,
+};
+
+pub const KernelContext = extern struct {
+    abi_version: u32,
+    resource_count: u32,
+    resources: u64,
+    push_constants: u64,
+    push_constant_size: u32,
+    reserved: u32 = 0,
+    base_group: [3]u32,
+    group_count: [3]u32,
+    local_size: [3]u32,
+    num_workgroups: [3]u32,
+};
+
+pub const EntryPoint = *const fn (
+    context: *const KernelContext,
+    begin_workgroup: u64,
+    end_workgroup: u64,
+) callconv(.c) void;
diff --git a/src/phi/compiler/analysis.zig b/src/phi/compiler/analysis.zig
new file mode 100644
index 0000000..b6998b4
--- /dev/null
+++ b/src/phi/compiler/analysis.zig
@@ -0,0 +1,21 @@
+const std = @import("std");
+const block_layout = @import("block_layout.zig");
+const edge_copies = @import("edge_copies.zig");
+const liveness = @import("liveness.zig");
+const register_allocator = @import("register_allocator.zig");
+
+pub const Analysis = struct {
+    layout: block_layout.Layout = block_layout.Layout.empty(),
+    liveness: liveness.Analysis = liveness.Analysis.empty(),
+    allocation: register_allocator.Allocation = register_allocator.Allocation.empty(),
+    edge_copy_plans: []edge_copies.Plan = &.{},
+
+    pub fn deinit(self: *Analysis, allocator: std.mem.Allocator) void {
+        self.layout.deinit(allocator);
+        self.liveness.deinit(allocator);
+        self.allocation.deinit(allocator);
+        for (self.edge_copy_plans) |*plan| plan.deinit(allocator);
+        if (self.edge_copy_plans.len != 0) allocator.free(self.edge_copy_plans);
+        self.* = undefined;
+    }
+};
diff --git a/src/phi/compiler/artifact.zig b/src/phi/compiler/artifact.zig
new file mode 100644
index 0000000..c971795
--- /dev/null
+++ b/src/phi/compiler/artifact.zig
@@ -0,0 +1,26 @@
+const std = @import("std");
+const abi = @import("abi.zig");
+
+pub const ResourceBinding = struct {
+    set: u32,
+    binding: u32,
+};
+
+pub const KernelInfo = struct {
+    abi_version: u32 = abi.version,
+    workgroup_size: [3]u32,
+    dispatch_width: u8,
+    stack_size: u32,
+    resources: []const ResourceBinding,
+};
+
+pub const Artifact = struct {
+    code: []u8,
+    info: KernelInfo,
+
+    pub fn deinit(self: *Artifact, allocator: std.mem.Allocator) void {
+        allocator.free(self.code);
+        allocator.free(self.info.resources);
+        self.* = undefined;
+    }
+};
diff --git a/src/phi/compiler/block_layout.zig b/src/phi/compiler/block_layout.zig
new file mode 100644
index 0000000..1cbf2a2
--- /dev/null
+++ b/src/phi/compiler/block_layout.zig
@@ -0,0 +1,17 @@
+const std = @import("std");
+const shader_ir = @import("shader_ir").ir;
+
+pub const Layout = struct {
+    blocks: []shader_ir.id.BlockId,
+    positions: []?usize,
+
+    pub fn empty() Layout {
+        return .{ .blocks = &.{}, .positions = &.{} };
+    }
+
+    pub fn deinit(self: *Layout, allocator: std.mem.Allocator) void {
+        if (self.blocks.len != 0) allocator.free(self.blocks);
+        if (self.positions.len != 0) allocator.free(self.positions);
+        self.* = undefined;
+    }
+};
diff --git a/src/phi/compiler/code_buffer.zig b/src/phi/compiler/code_buffer.zig
new file mode 100644
index 0000000..d275e24
--- /dev/null
+++ b/src/phi/compiler/code_buffer.zig
@@ -0,0 +1,67 @@
+const std = @import("std");
+const shader_ir = @import("shader_ir").ir;
+const Error = @import("errors.zig").Error;
+
+pub const Label = enum(u32) { _ };
+
+pub const Fixup = struct {
+    displacement_offset: usize,
+    instruction_end: usize,
+    target: Label,
+};
+
+pub const BlockLabel = struct {
+    block: shader_ir.id.BlockId,
+    label: Label,
+};
+
+pub const CodeBuffer = struct {
+    allocator: std.mem.Allocator,
+    bytes: std.ArrayList(u8) = .empty,
+    label_offsets: std.ArrayList(?usize) = .empty,
+    fixups: std.ArrayList(Fixup) = .empty,
+
+    pub fn init(allocator: std.mem.Allocator) CodeBuffer {
+        return .{ .allocator = allocator };
+    }
+
+    pub fn deinit(self: *CodeBuffer) void {
+        self.bytes.deinit(self.allocator);
+        self.label_offsets.deinit(self.allocator);
+        self.fixups.deinit(self.allocator);
+        self.* = undefined;
+    }
+
+    pub fn offset(self: *const CodeBuffer) usize {
+        return self.bytes.items.len;
+    }
+
+    pub fn emitByte(self: *CodeBuffer, byte: u8) std.mem.Allocator.Error!void {
+        try self.bytes.append(self.allocator, byte);
+    }
+
+    pub fn emitBytes(self: *CodeBuffer, bytes: []const u8) std.mem.Allocator.Error!void {
+        try self.bytes.appendSlice(self.allocator, bytes);
+    }
+
+    pub fn createLabel(self: *CodeBuffer) std.mem.Allocator.Error!Label {
+        const label: Label = @enumFromInt(self.label_offsets.items.len);
+        try self.label_offsets.append(self.allocator, null);
+        return label;
+    }
+
+    pub fn bindLabel(self: *CodeBuffer, label: Label) Error!void {
+        const index = @intFromEnum(label);
+        if (index >= self.label_offsets.items.len or self.label_offsets.items[index] != null)
+            return error.EncodingFailed;
+        self.label_offsets.items[index] = self.offset();
+    }
+
+    pub fn resolveFixups(_: *CodeBuffer) Error!void {
+        return error.CodeGenerationNotImplemented;
+    }
+
+    pub fn toOwnedSlice(self: *CodeBuffer) std.mem.Allocator.Error![]u8 {
+        return self.bytes.toOwnedSlice(self.allocator);
+    }
+};
diff --git a/src/phi/compiler/codegen.zig b/src/phi/compiler/codegen.zig
new file mode 100644
index 0000000..7ef5991
--- /dev/null
+++ b/src/phi/compiler/codegen.zig
@@ -0,0 +1,17 @@
+const shader_ir = @import("shader_ir").ir;
+const Analysis = @import("analysis.zig").Analysis;
+const Error = @import("errors.zig").Error;
+const Encoder = @import("imci/encoder.zig").Encoder;
+
+pub const Options = struct {
+    dispatch_width: u8 = 16,
+};
+
+pub fn emitComputeKernel(
+    _: *Encoder,
+    _: *const shader_ir.module.Module,
+    _: *const Analysis,
+    _: Options,
+) Error!void {
+    return error.CodeGenerationNotImplemented;
+}
diff --git a/src/phi/compiler/compiler.zig b/src/phi/compiler/compiler.zig
new file mode 100644
index 0000000..bd1fcf1
--- /dev/null
+++ b/src/phi/compiler/compiler.zig
@@ -0,0 +1,60 @@
+const std = @import("std");
+const shader_ir = @import("shader_ir").ir;
+
+pub const abi = @import("abi.zig");
+pub const analysis = @import("analysis.zig");
+pub const artifact = @import("artifact.zig");
+pub const block_layout = @import("block_layout.zig");
+pub const code_buffer = @import("code_buffer.zig");
+pub const codegen = @import("codegen.zig");
+pub const control_flow = @import("control_flow.zig");
+pub const edge_copies = @import("edge_copies.zig");
+pub const errors = @import("errors.zig");
+pub const imci = @import("imci/imci.zig");
+pub const liveness = @import("liveness.zig");
+pub const register_allocator = @import("register_allocator.zig");
+
+pub const Artifact = artifact.Artifact;
+pub const Error = errors.Error;
+
+pub const Options = struct {
+    dispatch_width: u8 = 16,
+};
+
+pub fn compileCompute(allocator: std.mem.Allocator, module: *shader_ir.module.Module, options: Options) Error!Artifact {
+    _ = allocator;
+
+    if (module.stage != .compute)
+        return error.UnsupportedStage;
+    if (module.entry_point == null)
+        return error.MissingEntryPoint;
+    if (module.execution_modes.workgroup_size == null)
+        return error.MissingWorkgroupSize;
+    if (!module.properties.structured_control_flow)
+        return error.UnstructuredControlFlow;
+    if (options.dispatch_width != 16)
+        return error.UnsupportedType;
+
+    return error.CodeGenerationNotImplemented;
+}
+
+test "[compiler] foundation declarations compile" {
+    std.testing.refAllDecls(abi);
+    std.testing.refAllDecls(analysis);
+    std.testing.refAllDecls(artifact);
+    std.testing.refAllDecls(block_layout);
+    std.testing.refAllDecls(code_buffer);
+    std.testing.refAllDecls(codegen);
+    std.testing.refAllDecls(control_flow);
+    std.testing.refAllDecls(edge_copies);
+    std.testing.refAllDecls(imci);
+    std.testing.refAllDecls(liveness);
+    std.testing.refAllDecls(register_allocator);
+}
+
+test "[compiler] rejects non-compute modules before code generation" {
+    var module = shader_ir.module.Module.init(std.testing.allocator, .vertex);
+    defer module.deinit();
+
+    try std.testing.expectError(error.UnsupportedStage, compileCompute(std.testing.allocator, &module, .{}));
+}
diff --git a/src/phi/compiler/control_flow.zig b/src/phi/compiler/control_flow.zig
new file mode 100644
index 0000000..18d4d8c
--- /dev/null
+++ b/src/phi/compiler/control_flow.zig
@@ -0,0 +1,27 @@
+const shader_ir = @import("shader_ir").ir;
+const registers = @import("imci/registers.zig");
+
+pub const ActiveMask = union(enum) {
+    full,
+    register: registers.Mask,
+    spilled: u32,
+};
+
+pub const Region = union(enum) {
+    block: shader_ir.id.BlockId,
+    selection: struct {
+        header: shader_ir.id.BlockId,
+        merge: shader_ir.id.BlockId,
+    },
+    loop: struct {
+        header: shader_ir.id.BlockId,
+        merge: shader_ir.id.BlockId,
+        continue_block: shader_ir.id.BlockId,
+    },
+};
+
+pub const State = struct {
+    active_mask: ActiveMask = .full,
+    loop_depth: u16 = 0,
+    selection_depth: u16 = 0,
+};
diff --git a/src/phi/compiler/edge_copies.zig b/src/phi/compiler/edge_copies.zig
new file mode 100644
index 0000000..a521de8
--- /dev/null
+++ b/src/phi/compiler/edge_copies.zig
@@ -0,0 +1,18 @@
+const std = @import("std");
+const shader_ir = @import("shader_ir").ir;
+
+pub const Copy = struct {
+    source: shader_ir.id.ValueId,
+    destination: shader_ir.id.ValueId,
+};
+
+pub const Plan = struct {
+    predecessor: shader_ir.id.BlockId,
+    successor: shader_ir.id.BlockId,
+    copies: []Copy,
+
+    pub fn deinit(self: *Plan, allocator: std.mem.Allocator) void {
+        allocator.free(self.copies);
+        self.* = undefined;
+    }
+};
diff --git a/src/phi/compiler/errors.zig b/src/phi/compiler/errors.zig
new file mode 100644
index 0000000..9d4c83a
--- /dev/null
+++ b/src/phi/compiler/errors.zig
@@ -0,0 +1,15 @@
+const std = @import("std");
+
+pub const Error = std.mem.Allocator.Error || error{
+    MissingEntryPoint,
+    MissingWorkgroupSize,
+    InvalidModule,
+    UnsupportedStage,
+    UnsupportedType,
+    UnsupportedOperation,
+    UnstructuredControlFlow,
+    RegisterAllocationFailed,
+    EncodingFailed,
+    BranchOutOfRange,
+    CodeGenerationNotImplemented,
+};
diff --git a/src/phi/compiler/imci/encoder.zig b/src/phi/compiler/imci/encoder.zig
new file mode 100644
index 0000000..e77983b
--- /dev/null
+++ b/src/phi/compiler/imci/encoder.zig
@@ -0,0 +1,46 @@
+const std = @import("std");
+const CodeBuffer = @import("../code_buffer.zig").CodeBuffer;
+const Error = @import("../errors.zig").Error;
+const encoding = @import("encoding.zig");
+const registers = @import("registers.zig");
+
+pub const Encoder = struct {
+    code: CodeBuffer,
+
+    pub fn init(allocator: std.mem.Allocator) Encoder {
+        return .{ .code = CodeBuffer.init(allocator) };
+    }
+
+    pub fn deinit(self: *Encoder) void {
+        self.code.deinit();
+        self.* = undefined;
+    }
+
+    pub fn prologue(_: *Encoder, _: u32) Error!void {
+        return Error.CodeGenerationNotImplemented;
+    }
+
+    pub fn epilogue(_: *Encoder) Error!void {
+        return Error.CodeGenerationNotImplemented;
+    }
+
+    pub fn moveVector(_: *Encoder, _: registers.Zmm, _: encoding.VectorSource, _: ?registers.Mask) Error!void {
+        return Error.CodeGenerationNotImplemented;
+    }
+
+    pub fn addVector(_: *Encoder, _: encoding.VectorElement, _: registers.Zmm, _: registers.Zmm, _: encoding.VectorSource, _: ?registers.Mask) Error!void {
+        return Error.CodeGenerationNotImplemented;
+    }
+
+    pub fn compareVector(_: *Encoder, _: encoding.VectorElement, _: encoding.Condition, _: registers.Mask, _: registers.Zmm, _: encoding.VectorSource, _: registers.Mask) Error!void {
+        return Error.CodeGenerationNotImplemented;
+    }
+
+    pub fn gather(_: *Encoder, _: encoding.VectorElement, _: registers.Zmm, _: registers.Gpr, _: registers.Zmm, _: registers.Mask) Error!void {
+        return Error.CodeGenerationNotImplemented;
+    }
+
+    pub fn scatter(_: *Encoder, _: encoding.VectorElement, _: registers.Gpr, _: registers.Zmm, _: registers.Zmm, _: registers.Mask) Error!void {
+        return Error.CodeGenerationNotImplemented;
+    }
+};
diff --git a/src/phi/compiler/imci/encoding.zig b/src/phi/compiler/imci/encoding.zig
new file mode 100644
index 0000000..c9eb0a2
--- /dev/null
+++ b/src/phi/compiler/imci/encoding.zig
@@ -0,0 +1,38 @@
+const registers = @import("registers.zig");
+
+pub const VectorElement = enum {
+    i32,
+    u32,
+    f32,
+    i64,
+    u64,
+    f64,
+};
+
+pub const Immediate = union(enum) {
+    u8: u8,
+    u32: u32,
+    i32: i32,
+    u64: u64,
+};
+
+pub const Memory = struct {
+    base: ?registers.Gpr = null,
+    index: ?registers.Gpr = null,
+    scale: enum(u2) { one, two, four, eight } = .one,
+    displacement: i32 = 0,
+};
+
+pub const VectorSource = union(enum) {
+    register: registers.Zmm,
+    memory: Memory,
+};
+
+pub const Condition = enum {
+    equal,
+    not_equal,
+    less,
+    less_equal,
+    greater,
+    greater_equal,
+};
diff --git a/src/phi/compiler/imci/imci.zig b/src/phi/compiler/imci/imci.zig
new file mode 100644
index 0000000..e5e48a6
--- /dev/null
+++ b/src/phi/compiler/imci/imci.zig
@@ -0,0 +1,5 @@
+pub const encoder = @import("encoder.zig");
+pub const encoding = @import("encoding.zig");
+pub const registers = @import("registers.zig");
+
+pub const Encoder = encoder.Encoder;
diff --git a/src/phi/compiler/imci/registers.zig b/src/phi/compiler/imci/registers.zig
new file mode 100644
index 0000000..0fe3f2c
--- /dev/null
+++ b/src/phi/compiler/imci/registers.zig
@@ -0,0 +1,64 @@
+pub const Gpr = enum(u5) {
+    rax,
+    rcx,
+    rdx,
+    rbx,
+    rsp,
+    rbp,
+    rsi,
+    rdi,
+    r8,
+    r9,
+    r10,
+    r11,
+    r12,
+    r13,
+    r14,
+    r15,
+};
+
+pub const Zmm = enum(u5) {
+    zmm0,
+    zmm1,
+    zmm2,
+    zmm3,
+    zmm4,
+    zmm5,
+    zmm6,
+    zmm7,
+    zmm8,
+    zmm9,
+    zmm10,
+    zmm11,
+    zmm12,
+    zmm13,
+    zmm14,
+    zmm15,
+    zmm16,
+    zmm17,
+    zmm18,
+    zmm19,
+    zmm20,
+    zmm21,
+    zmm22,
+    zmm23,
+    zmm24,
+    zmm25,
+    zmm26,
+    zmm27,
+    zmm28,
+    zmm29,
+    zmm30,
+    zmm31,
+};
+
+pub const Mask = enum(u3) {
+    k0,
+    k1,
+    k2,
+    k3,
+    k4,
+    k5,
+    k6,
+    k7,
+};
diff --git a/src/phi/compiler/liveness.zig b/src/phi/compiler/liveness.zig
new file mode 100644
index 0000000..7f7233f
--- /dev/null
+++ b/src/phi/compiler/liveness.zig
@@ -0,0 +1,25 @@
+const std = @import("std");
+const shader_ir = @import("shader_ir").ir;
+
+pub const Position = u32;
+
+pub const LiveRange = struct {
+    value: shader_ir.id.ValueId,
+    first: Position,
+    last: Position,
+};
+
+pub const Analysis = struct {
+    ranges: []LiveRange,
+    value_ranges: []?usize,
+
+    pub fn empty() Analysis {
+        return .{ .ranges = &.{}, .value_ranges = &.{} };
+    }
+
+    pub fn deinit(self: *Analysis, allocator: std.mem.Allocator) void {
+        if (self.ranges.len != 0) allocator.free(self.ranges);
+        if (self.value_ranges.len != 0) allocator.free(self.value_ranges);
+        self.* = undefined;
+    }
+};
diff --git a/src/phi/compiler/register_allocator.zig b/src/phi/compiler/register_allocator.zig
new file mode 100644
index 0000000..b2c2d62
--- /dev/null
+++ b/src/phi/compiler/register_allocator.zig
@@ -0,0 +1,61 @@
+const std = @import("std");
+const shader_ir = @import("shader_ir").ir;
+const registers = @import("imci/registers.zig");
+
+pub const StackSlot = struct {
+    offset: u32,
+    size: u32,
+    alignment: u32,
+};
+
+pub const VectorComponent = struct {
+    chunks: []const registers.Zmm,
+};
+
+pub const VectorLocation = struct {
+    components: []const VectorComponent,
+};
+
+pub const Immediate = union(enum) {
+    integer: u64,
+    float: u64,
+};
+
+pub const Location = union(enum) {
+    immediate: Immediate,
+    vector: VectorLocation,
+    mask: registers.Mask,
+    stack: StackSlot,
+};
+
+pub const Constraints = struct {
+    temporary_vectors: u8 = 0,
+    temporary_masks: u8 = 0,
+    temporary_gprs: u8 = 0,
+};
+
+pub const Allocation = struct {
+    value_locations: []?Location,
+    stack_size: u32,
+    scratch_vector: ?registers.Zmm,
+    scratch_mask: ?registers.Mask,
+
+    pub fn empty() Allocation {
+        return .{
+            .value_locations = &.{},
+            .stack_size = 0,
+            .scratch_vector = null,
+            .scratch_mask = null,
+        };
+    }
+
+    pub fn location(self: *const Allocation, value: shader_ir.id.ValueId) ?Location {
+        if (value.index() >= self.value_locations.len) return null;
+        return self.value_locations[value.index()];
+    }
+
+    pub fn deinit(self: *Allocation, allocator: std.mem.Allocator) void {
+        if (self.value_locations.len != 0) allocator.free(self.value_locations);
+        self.* = undefined;
+    }
+};
diff --git a/src/phi/lib.zig b/src/phi/lib.zig
index 42d492d..b237d17 100644
--- a/src/phi/lib.zig
+++ b/src/phi/lib.zig
@@ -7,6 +7,7 @@ pub const proto = @import("phi_protocol_c");
 pub const config = base.config;
 pub const mic = @import("miclib");
 pub const scif = @import("scif.zig");
+pub const compiler = @import("compiler/compiler.zig");
 
 pub const PhiInstance = @import("PhiInstance.zig");
 pub const PhiDevice = @import("PhiDevice.zig");
@@ -75,6 +76,7 @@ test {
     std.testing.refAllDecls(PhiPhysicalDevice);
     std.testing.refAllDecls(PhiTransport);
     std.testing.refAllDecls(scif);
+    std.testing.refAllDecls(compiler);
     std.testing.refAllDecls(PhiPipeline);
     std.testing.refAllDecls(PhiPipelineCache);
     std.testing.refAllDecls(PhiPipelineLayout);