Apple M3 GPU reverse engineering progress

Anna Maniscalco (Valve) @pac85@treehouse.systems

Alyssa Milburn @noopwafel@teh.entar.net

TellowKrinkle @TellowKrinkle@treehouse.systems

Viola @viola@treehouse.systems - Couldn’t be here

Three pieces

• Shader compiler
• Userspace driver
• Kernel driver

Where our story begins?

• At XDC2024 TellowKrinkle had started figuring the ISA
  • And continued to poke at it through 2025
• At the start of 2026:
  • Anna and Vio started reverse-engineering the cmdbuf
  • Alyssa bought an M3 and started poking at it

The new HW

How much could Apple possibly have changed?

Easy port?

// PBE descriptor
agx_pack(..., PBE, cfg) {
  cfg.dimension = AGX_TEXTURE_DIMENSION_2D;
  cfg.layout = AGX_LAYOUT_LINEAR;
  cfg.channels = R8G8B8A8;
  cfg.type = AGX_TEXTURE_TYPE_UNORM;
  cfg.swizzle_r = 0;
  cfg.swizzle_g = 1;
  cfg.swizzle_b = 2;
  cfg.swizzle_a = 3;
  cfg.width = width;
  cfg.height = height;
  cfg.buffer = out_bo.va->addr & 0xffffffff;
  cfg.level = 1;
  cfg.stride = (width * 4) - 16;
}

// PBE descriptor
agx_pack(..., PBE, cfg) {
  cfg.dimension = AGX_TEXTURE_DIMENSION_2D;
  cfg.layout = AGX_LAYOUT_LINEAR;
  cfg.channels = R8G8B8A8;
  cfg.type = AGX_TEXTURE_TYPE_UNORM;
  cfg.swizzle_r = 0;
  cfg.swizzle_g = 1;
  cfg.swizzle_b = 2;
  cfg.swizzle_a = 3;
  cfg.width = width;
  cfg.height = height;
  cfg.buffer = out_bo.va->addr;
  cfg.level = 1;
  cfg.stride = (width * 4) - 16;
}

Shaders

Oh no, new ISA!

M1 ISA RE

• Started with wrap.c
  • Intercepts calls from the user mode driver to the kernel mode driver
  • Gives access to almost everything
  • Apple’s kernel driver’s API is unstable

Metal Binary Archives

• Metal allows caching pipelines to disk
  • What’s inside of those?

00000000  cb fe ba be 00 00 00 02  01 00 00 13 00 00 01 42
00000010  00 00 00 30 00 00 20 a0  00 00 00 04 01 00 00 17
00000020  00 00 00 0b 00 00 20 d0  00 00 12 3c 00 00 00 04
00000030  cf fa ed fe 13 00 00 01  42 01 00 00 0d 00 00 00

Metal Binary Archives

• Metal allows caching pipelines to disk
  • What’s inside of those?

00000000  cb fe ba be 00 00 00 02  01 00 00 13 00 00 01 42
00000010  00 00 00 30 00 00 20 a0  00 00 00 04 01 00 00 17
00000020  00 00 00 0b 00 00 20 d0  00 00 12 3c 00 00 00 04
00000030  cf fa ed fe 13 00 00 01  42 01 00 00 0d 00 00 00

Metal Binary Archives

• Metal allows caching pipelines to disk
  • What’s inside of those?

00000000  cb fe ba be 00 00 00 02  01 00 00 13 00 00 01 42
00000010  00 00 00 30 00 00 20 a0  00 00 00 04 01 00 00 17
00000020  00 00 00 0b 00 00 20 d0  00 00 12 3c 00 00 00 04
00000030  cf fa ed fe 13 00 00 01  42 01 00 00 0d 00 00 00

• MacOS file utility knows about these

$ file test.bin
test.bin:  [applegpu_g13g:Mach-O 64-bit GPU executable applegpu_g13g] [air64:MetalLib executable (MacOS), version 1.2.8]
test.bin (for architecture applegpu_g13g):  Mach-O 64-bit GPU executable applegpu_g13g
test.bin (for architecture cputype (16777239) cpusubtype (11)): MetalLib executable (MacOS), version 1.2.8

Metal Binary Archives

• Contains sections like __TEXT,__vertex, __TEXT,__fragment, and __TEXT,__compute
  • What’s in those?

Metal Binary Archives

• Contains sections like __TEXT,__vertex, __TEXT,__fragment, and __TEXT,__compute
  • What’s in those?

00000000  cf fa ed fe 13 00 00 01  42 01 00 00 01 00 00 00
00000010  03 00 00 00 90 02 00 00  00 00 00 00 00 00 00 00
00000020  19 00 00 00 28 02 00 00  00 00 00 00 00 00 00 00
00000030  00 00 00 00 00 00 00 00  00 00 00 00 00 00 00 00

Metal Binary Archives

• Contains sections like __TEXT,__vertex, __TEXT,__fragment, and __TEXT,__compute
  • What’s in those?

00000000  cf fa ed fe 13 00 00 01  42 01 00 00 01 00 00 00
00000010  03 00 00 00 90 02 00 00  00 00 00 00 00 00 00 00
00000020  19 00 00 00 28 02 00 00  00 00 00 00 00 00 00 00
00000030  00 00 00 00 00 00 00 00  00 00 00 00 00 00 00 00

• Second level mach-o contains more normal sections like __TEXT,__text, and has symbols like _agc.main and _agc.main.constant_program, which contain GPU ISA

Compiler Explorer

kernel void test(uint pos [[thread_position_in_grid]], device uint4* out, const device uint4* in0, device uint4* in1) {
    out[pos] = in0[pos] + in1[pos];
}
$ pbpaste | python3 compiler_explorer.py -
compute shader:
   0: 72111004             get_sr           r4, sr80 (thread_position_in_grid.x)
   4: 0501840e00c8f200     device_load      0, i32, xyzw, r0_r1_r2_r3, u2_u3, r4, unsigned, lsl 2
   c: 0529880e00c8f200     device_load      0, i32, xyzw, r5_r6_r7_r8, u4_u5, r4, unsigned, lsl 2
  14: 3800                 wait             0
  16: 0e01ca022c000000     iadd             r0, r5.discard, r0.discard
  1e: 0e05cc222c000000     iadd             r1, r6.discard, r1.discard
  26: 0e09ce422c000000     iadd             r2, r7.discard, r2.discard
  2e: 0e0dd0622c000000     iadd             r3, r8.discard, r3.discard
  36: 4501800e00c8f200     device_store     0, i32, xyzw, r0_r1_r2_r3, u0_u1, r4, unsigned, lsl 2, 0
  3e: 8800                 stop

Running Custom Assembly

• Binary archives have no integrity checks, so you can just modify the code and run it
  • The driver deduplicates them on load

Running Custom Assembly

$ python3 hwtestbed.py "isub.sat r0, r0, r1" -r{1,-8},{2,0x7fffffff}
Thread 0: r0: 00000000 (0         ), input: 1,2
Thread 1: r0: fffffff6 (-10       ), input: -8,2
Thread 2: r0: 00000000 (0         ), input: 1,0x7fffffff
Thread 3: r0: 7ffffff9 (2147483641), input: -8,0x7fffffff
$ python3 hwtestbed.py "isub.sat r0, r0, r1.sx" -r{1,-8},{2,0x7fffffff}
Thread 0: r0: ffffffff (-1        ), input: 1,2
Thread 1: r0: 7fffffff (2147483647), input: -8,2
Thread 2: r0: 80000002 (-2147483646), input: 1,0x7fffffff
Thread 3: r0: 7ffffff9 (2147483641), input: -8,0x7fffffff

Running Custom Assembly

$ python3 hwtestbed.py -b 4e01402a24000000 -r{1,-8},{2,0x7fffffff}
Thread 0: r0: 00000000 (0         ), input: 1,2
Thread 1: r0: fffffff6 (-10       ), input: -8,2
Thread 2: r0: 00000000 (0         ), input: 1,0x7fffffff
Thread 3: r0: 7ffffff9 (2147483641), input: -8,0x7fffffff
$ python3 hwtestbed.py -b 4e01402a64000000 -r{1,-8},{2,0x7fffffff}
Thread 0: r0: ffffffff (-1        ), input: 1,2
Thread 1: r0: 7fffffff (2147483647), input: -8,2
Thread 2: r0: 80000002 (-2147483646), input: 1,0x7fffffff
Thread 3: r0: 7ffffff9 (2147483641), input: -8,0x7fffffff

M3 ISA RE

• compiler_explorer.py works without modification
  • (If you ignore the broken disassembly)

compute shader prolog:
   0: 0300                 <disassembly failed>
   2: 0700                 <disassembly failed>
TODO: 1000

M1 vs M3

M3 ISA RE

kernel void test(uint pos [[thread_position_in_grid]], device uint4* out, const device uint4* in0, device uint4* in1) {
    out[pos] = in0[pos] + in1[pos];
}
compute shader:
00000000: 8C A0 10 06 06 04 06 00  06 00 06 00 67 10 54 00  ............g.T.
00000010: 01 08 20 00 57 01 00 40  40 00 67 00 44 08 02 08  .. .W..@@.g.D...
00000020: 20 00 57 01 00 40 40 00  06 80 06 00 06 00 06 00   .W..@@.........
00000030: 9F 01 56 00 02 10 00 A8  17 05 9F 01 54 02 02 14  ..V.........T...
00000040: 08 A8 17 05 9F 01 54 04  02 18 10 A8 17 05 9F 01  ......T.........
00000050: 54 06 02 1C 18 A8 17 05  E7 00 54 00 00 08 21 00  T.........T...!.
00000060: 17 00 00 10 10 00 0E 00  00 00                    ..........

M3 ISA RE

kernel void test(uint pos [[thread_position_in_grid]], device uint4* out, const device uint4* in0, device uint4* in1) {
    out[pos] = in0[pos] + in1[pos];
}
compute shader:
00000000: 8C A0 10 06 06 04 06 00  06 00 06 00 67 10 54 00  ............g.T.
00000010: 01 08 20 00 57 01 00 40  40 00 67 00 44 08 02 08  .. .W..@@.g.D...
00000020: 20 00 57 01 00 40 40 00  06 80 06 00 06 00 06 00   .W..@@.........
00000030: 9F 01 56 00 02 10 00 A8  17 05 9F 01 54 02 02 14  ..V.........T...
00000040: 08 A8 17 05 9F 01 54 04  02 18 10 A8 17 05 9F 01  ......T.........
00000050: 54 06 02 1C 18 A8 17 05  E7 00 54 00 00 08 21 00  T.........T...!.
00000060: 17 00 00 10 10 00 0E 00  00 00                    ..........

M3 ISA RE

kernel void test(uint pos [[thread_position_in_grid]], device uint4* out, const device uint4* in0, device uint4* in1) {
    out[pos] = in0[pos] + in1[pos];
}
compute shader:
   0: 8C A0 10 06
   4: 06 04 06 00 06 00 06 00
   c: 67 10 54 00 01 08 20 00 57 01 00 40 40 00
  1a: 67 00 44 08 02 08 20 00 57 01 00 40 40 00
  28: 06 80 06 00 06 00 06 00
  30: 9F 01 56 00 02 10 00 A8 17 05
  3a: 9F 01 54 02 02 14 08 A8 17 05
  44: 9F 01 54 04 02 18 10 A8 17 05
  4e: 9F 01 54 06 02 1C 18 A8 17 05
  58: E7 00 54 00 00 08 21 00 17 00 00 10 10 00
  66: 0E 00 00 00

M3 ISA RE

kernel void test(uint pos [[thread_position_in_grid]], device uint4* out, const device uint4* in0, device uint4* in1) {
    out[pos] = in0[pos] + in1[pos];
}
compute shader:
   0: 8C A0 10 06                                get_sr           0, r8, sr32 (thread_position_in_grid.x)
   4: 06 04 06 00 06 00 06 00                    wait_long        0
   c: 67 10 54 00 01 08 20 00 57 01 00 40 40 00  device_load      5, i32, xyzw, r0_r1_r2_r3, u2_u3, r8, lsl 4, 0, wait 0, 1, 1, 0
  1a: 67 00 44 08 02 08 20 00 57 01 00 40 40 00  device_load      5, i32, xyzw, r4_r5_r6_r7, u4_u5, r8, lsl 4, 0, 1, 0, 0
  28: 06 80 06 00 06 00 06 00                    wait_long        5
  30: 9F 01 56 00 02 10 00 A8 17 05              iadd             r0, r4.discard, r0.discard, wait 5
  3a: 9F 01 54 02 02 14 08 A8 17 05              iadd             r1, r5.discard, r1.discard
  44: 9F 01 54 04 02 18 10 A8 17 05              iadd             r2, r6.discard, r2.discard
  4e: 9F 01 54 06 02 1C 18 A8 17 05              iadd             r3, r7.discard, r3.discard
  58: E7 00 54 00 00 08 21 00 17 00 00 10 10 00  device_store     0, i32, xyzw, r0_r1_r2_r3, u0_u1, r8.discard, lsl 4, 0, 1, 1
  66: 0E 00 00 00                                stop

M3 ISA RE

kernel void test(uint pos [[thread_position_in_grid]], device float4* out, const device float4* in0, device float4* in1) {
    out[pos] = in0[pos] + in1[pos];
}
compute shader:
   0: 8C A0 10 06                                get_sr           0, r8, sr32 (thread_position_in_grid.x)
   4: 06 04 06 00 06 00 06 00                    wait_long        0
   c: 67 10 54 00 01 08 20 00 57 01 00 40 40 00  device_load      5, i32, xyzw, r0_r1_r2_r3, u2_u3, r8, lsl 4, 0, wait 0, 1, 1, 0
  1a: 67 00 44 08 02 08 20 00 57 01 00 40 40 00  device_load      5, i32, xyzw, r4_r5_r6_r7, u4_u5, r8, lsl 4, 0, 1, 0, 0
  28: 06 80 06 00 06 00 06 00                    wait_long        5
  30: 09 09 1C 01 00 C0                          fadd             r0, r4.discard, r0.discard, wait 5
  36: 19 0B 18 03                                fadd             r1, r5.discard, r1.discard
  3a: 29 0D 18 05                                fadd             r2, r6.discard, r2.discard
  3e: 39 0F 18 07                                fadd             r3, r7.discard, r3.discard
  42: E7 00 54 00 00 08 21 00 17 00 00 10 10 00  device_store     0, i32, xyzw, r0_r1_r2_r3, u0_u1, r8.discard, lsl 4, 0, 1, 1
  50: 0E 00 00 00                                stop

M1 vs M3

Load Immediate

M1

Load Immediate

M1

Load Immediate

M1

M3

Enums

• Load/Store instructions have a mask, to say which elements of a vector to load
  • e.g. x, z, xz, xyzw
• On M1, this is a four-bit bitfield, with one byte each for x, y, z, w

Enums

Enums

Enums

Texture one is at bits 27, 28, 37, and 86

Compare Select

• Both M1 and M3 have a compare select instruction, allowing you to calculate things like D = A > B ? X : Y
  • 5 registers is a lot
  • M1’s version is 8-10 bytes, and can’t use input modifiers (negate, abs)
  • M3 has a compressed form that aliases (A, B) with (X, Y), allowing 6-14 byte variants
    • Compare and select inputs are interpreted differently

Control Flow

• M1 control flow used instructions with a built-in compare, like if_icmp
• M3 adds new condition code registers to pass flags to if
  • Those registers are written by a variant of cmpsel
    • It comes in a simpler cmp and full cmpselcc variants

Compare Select CC

• cmp is encoded as a cmpsel with the destination size set to 32 but with a non-32-bit-aligned register
• There’s a bit that Apple’s compiler always clears for cmp and sets for cmpselcc
  • If you flip it yourself, you’ll find that it inverts the comparison result
  • When left unset, cmpselcc calculates A > B ? Y : X

Puzzling out the rest of userspace

• First step was to revive wrap.c
• Apple’s kernel driver’s API is unstable

Staring at bytes

• Pointers are easy to spot
• We can poison them from the shim
• Diffing entire runs to find a particular thing
• Then chasing our way back through pointers

Road to demo

• Once we understood enough we started replacing thins
• Eventually there was little enough left that it could be factored out
• -> stand-alone compute demo

Apple’s way of doing drivers

• Normally kernel API is fairly slim
• Usermode driver produces commands that setup all HW state

Think different

• Apple has a thicker interface
• Mostly sane for compute, pointer to texture heap and cdm

CDM?

• Maps to a command buffer in vulkan…mostly
• Has commands to set state and dispatch

What about graphics

• Way less sane, framebuffer state is passed to the kernel and handheld there
• Commands contained in VDM

VDM?

• State setting and draw commands
• Maps to commands inside of a render pass in vulkan

How does it map to vulkan?

• A command buffer might need more than one submissions
• Each render pass needs it’s own
• Switching between compute and gfx requires splitting submissions

The big news

• USC gone
• Pipeline points to a shader directly

How it used to work

• USC contained two pointers
• One, a pre shader, would run with 1x1 wg size
• The other is the actual shader
• Also points the GPU at user data

Preshader?

• Also known as preamble
• The idea is to split out scalar parts of a shader

The new HW

• Pipeline only points to a pre shader
• Pre shader uses a special instruction to set the actual shader
• Pre shader dynamically allocates registers with an instruction
• Similarly shared memory is allocaed with another instruction
• Any costant is hardcoded or loaded from this shader (usually pointers)
• If there’s a pre shader it jumps to it at the end

This shader is treated like a descriptor

Entry shader

   0: 2C A0 02 00 12 08 78 01                    mov_imm          r2, 3085472
   8: 3C 80 02 00 04 00 00 00                    mov_imm          r3, 256
  10: 6C E0 42 00 06 00 20 00                    mov_imm          r22, 262624
  18: 7C 80 42 00 04 00 00 00                    mov_imm          r23, 256
  20: 67 00 54 24 02 00 00 00 57 00 00 40 26 00  device_load      1, i32, xyzw, r18_r19_r20_r21, r2_r3, 0, lsl 2, 0, 1, 1, 0
  2e: 67 00 54 34 16 00 00 00 57 00 02 40 26 00  device_load      1, i32, xyzw, r26_r27_r28_r29, r22_r23, 0, lsl 2, 16, 1, 1, 0
  3c: 67 00 54 2C 16 00 00 00 57 04 00 40 26 00  device_load      1, i32, xyzw, r22_r23_r24_r25, r22_r23, 0, lsl 2, 0, 1, 1, 1
  4a: 77 00 2A 41 00 00 00 00                    load_shader.todo 130
  52: 77 01 AA A6 00 00 00 02                    load_shader.todo 17179869517
  5a: 06 00                                      nop              
  5c: F7 00 2A 00 00 10 20 10 20 00              ralloc           1026, 1026
  66: 1C 80 02 00 00 00 00 00                    mov_imm          r1, 0
  6e: 14 81 11 06 00 00 00 00                    get_sr           0, r1l, sr1
  76: 0C 80 02 00 04 00 00 00                    mov_imm          r0, 256
  7e: 9F 11 54 00 02 00 08 A8 10 05              iadd             r0, r0, r1, lsl 4, wait 0
  88: 1C 80 02 00 04 00 00 00                    mov_imm          r1, 256
  90: 0F 12 54 00 4C 00 06 08 06 00              <disassembly failed>
  9a: 06 00                                      nop              
  9c: 06 00                                      nop              
  9e: 0B 24 09 04                                mov              u0, r18.discard
  a2: 1B 26 09 04                                mov              u1, r19.discard
  a6: 2B 28 09 04                                mov              u2, r20.discard
  aa: 3B 2A 09 04                                mov              u3, r21.discard
  ae: 8B 34 09 04                                mov              u8, r26.discard
  b2: 9B 36 09 04                                mov              u9, r27.discard
  b6: AB 38 09 04                                mov              u10, r28.discard
  ba: BB 3A 09 04                                mov              u11, r29.discard
  be: 4B 2C 09 04                                mov              u4, r22.discard
  c2: 5B 2E 09 04                                mov              u5, r23.discard
  c6: 6B 30 09 04                                mov              u6, r24.discard
  ca: 7B 32 09 04                                mov              u7, r25.discard
  ce: 0E 00 00 00                                stop 

Dynamic registers

• An instruction must appear in the preamble to allocate registers
• If using less than 128 regs, 0 can be specified
• If using more they must be specified

More than 128?

• In fact, up to around 32k
• Everything goes in regs, indirect indexing is present
• No more stack
• No more spilling (Metal compiler crashes ungracefully if out of registers)

Wtf?

Wtf?

• We don’t know how this works in HW
• Are that many registers actually available on chip?
• Is it spilling them to memory automatically?

.. if only we had some way to observe this?

What does this mean for the compiler?

• Most instructions have the same semantics
• Work is mostly on the backend side of the compiler
• Register allocations mostly works the same, just no more spilling

A catch

• preshader only has 128 regs available
• stack seems to be gone -> spilling not possible
• What do we do if the pre shader needs more regs?

A catch

• Apple just seems to never use more than 128 in the preshader
• Code is moved in the main shader as needed

Bad for Mesa

• We split out the preshader in nir_opt_preamble
• Happens way before regalloc

So what do we do?

• Just disable pre shader if it uses more than 128 regs
• Can be achieved by concatenating it at the beginning of the shader
• Can be done last moment
• Negates the benefits of a preshader, hopefully won’t happen too often

Current status

We’ve implemented

• control flow
• bitop and mov

Current status

We can compile this shader

   nir_builder b = nir_builder_init_simple_shader(MESA_SHADER_COMPUTE,
                                                  &agx_nir_options, "agx_troll");
   
   nir_def *l =  nir_load_global(&b, 1, 32, nir_imm_int64(&b, iova));
   nir_def *cond = nir_ilt_imm(&b, nir_ior_imm(&b, l, 1), 10);
   nir_if * if_ = nir_push_if(&b, cond);
   nir_push_loop(&b);
   nir_store_global(&b, nir_ior(&b, l, nir_imm_int(&b, 67)), nir_imm_int64(&b, iova));
   nir_break_if(&b, cond);
   nir_pop_loop(&b, NULL);
   nir_pop_if(&b, if_); 

Current status

output

   0: 2C 80 02 00 00 00 0C 00                    mov_imm          r2, 98304
   8: 3C 80 02 00 04 00 00 00                    mov_imm          r3, 256
  10: 67 00 44 02 02 00 00 00 31 00 00 00 20 00  device_load      0, i32, x, r1, r2_r3, 0, lsl 4, 0, 0, 0, 0
  1e: 06 04 06 00 06 00 06 00                    wait_long        0
  26: 4B 03 07 80 12 0A 00 00 00 00              or               r4, r1, 1
  30: 0A 09 0D 8A 22 81 07 02 20 80              cmpsel           slt, r0h, r4.discard, 10, 1, 0
  3a: 1A 80 23 00 06 00 97 00 00 00              cmp              ine, cc0, 0, r0h
  44: 0F 05 54 01                                if               cc0, 1
  48: 0A 80 23 80 06 00 87 02 00 00              cmp              ieq, cc0, 0, 0
  52: 0F 05 54 01                                if               cc0, 1
  56: 4B 03 07 C2 12 0A 00 00 00 00              or               r4, r1, 67
  60: E7 00 44 08 02 00 00 00 31 00 00 00 08 00  device_store     0, i32, x, r4, r2_r3, 0, lsl 4, 0, 0, 0
  6e: 06 04 06 00 06 00 06 00                    wait_long        0
  76: 1A 80 23 00 06 00 97 00 00 00              cmp              ine, cc0, 0, r0h
  80: 0F 05 54 01                                if               cc0, 1
  84: 8F 05 54 02 00 06                          <disassembly failed>
  8a: 0F 06 04 01 00 00                          pop_exec         1
  90: 0F 00 54 E2 FF FF FF FF FF 00              jmp_exec_any     0x72
  9a: 0F 06 04 01 00 00                          pop_exec         1
  a0: 0F 06 04 01 00 00                          pop_exec         1
  a6: 0E 00 00 00                                stop

...
Output buffer after dequeue, 67

Where do we go from there

• More work on the compiler needed
• Adapting the rest of the driver

None of this is hard, mostly just significant work

Kernel driver

drivers/gpu/drm/asahi/driver.rs
@@ -145,6 +145,11 @@ impl drm::driver::Driver for AsahiDriver {
+        (
+            of::DeviceId::new(c_str!("apple,agx-t8122")),
+            &hw::t8122::HWCONFIG
+        ),
+

@@ -205,6 +210,9 @@ fn probe(
+            (hw::GpuGen::G15, _, &[14, 7, 0]) => {
+                gpu::GpuManagerG15V14_7::new(&drm.clone(), &res, cfg)? as Arc<dyn gpu::GpuManager>
+            }

Kernel driver

arch/arm64/boot/dts/apple/t8122.dtsi
@@ -374,6 +408,65 @@ soc {
+   gpu: gpu@292400000 {
+     compatible = "apple,agx-t8122", "apple,agx-g15g";
+     reg = <0x2 0x92400000 0 0x70000>,

[1.245075] asahi 292400000.gpu: Probing...
[1.271715] asahi 292400000.gpu: GPU Information:
[1.272328] asahi 292400000.gpu:   Type: G15G
[1.276848] asahi 292400000.gpu:   Cores: 10 (10)
[5.139875] [drm] Initialized asahi 0.0.0 for 292400000.gpu on minor 1

Easy mode

Three small tweaks:

• New firmware version:
  • 14.x
• New GPU generation:
  • G15 (M3)
• New hardware:
  • T8122 (M3 Macbook Air)

How to port?

Tracing

CBW. WC3D job_params3 =
CBW. WC3D   Start3DJobParameters3:
CBW. WC3D.S3DP3 scissor_array = Start3DArrayAddr:
CBW. WC3D.S3DP3   ptr = 0x10001c50000
CBW. WC3D.S3DP3 visibility_result_buffer = 0x10001af8000
CBW. WC3D.S3DP3 unk_partial_bg = Start3DBackgroundProgram:
CBW. WC3D.S3DP3   pipeline_bind = 0x7800000000040
CBW. WC3D.S3DP3   address = 0x10001ae8100
CBW. WC3D.S3DP3 partial_bg = Start3DBackgroundProgram:
CBW. WC3D.S3DP3   pipeline_bind = 0xffff800000000000
CBW. WC3D.S3DP3   address = 0x10001ae8000
CBW. WC3D.S3DP3 depth_flags = 0x0
CBW. WC3D.S3DP3 unk_290 = 0x4040404

Tracing

[gfx-asc] frag work item
[gfx-asc] vert work item
[gfx-asc] kick~!
[gfx-asc] stats: Power ON: spent 0.5544s powered off (13305045 ticks)
[gfx-asc] stats: PowerState:         ps=1 100 0 0 3 0 0 0 0
[gfx-asc] stats: Power: 22.291 mW
[gfx-asc] stats: TempSensor: #3 40.19°C s=4 tmin=32 tmax=40
[gfx-asc] stats: Temp: 40.61°C s=4 tmin=32 tmax=40
[gfx-asc] stats: Power OFF: spent 0.003242s powered on (77810 ticks)
[gfx-asc] pong~!
regionB.page_mgr_info (0xfffffc20000e0000..0xfffffc20000e01ff)
  0003ae00 10000002 00004a00 0000bca0 00005c00 00000000 00000000 00000000
[gfx-asc] event fired: 1, 2

14.x Firmware

• m1n1 and kernel driver already support:
  • 12.x and 13.x firmware
  • G13 (M1) and G14 (M2)
• Versioning macros/helpers designed to make this easy
• Not a stable ABI - but how much could they have changed?

(Someone maintained the m1n1 debug tooling, right?)

Rearranging structures

     pub(crate) struct RunWorkQueueMsg {
+    #[ver(V < V14_7)]
     pub(crate) pipe_type: PipeType,
     pub(crate) work_queue: Option<GpuWeakPointer<super::workqueue::...>>,
+    #[ver(V >= V14_7)]
+    pub(crate) pipe_type: PipeType,
+    #[ver(V < V14_7)]
     pub(crate) wptr: u32,
+    #[ver(V >= V14_7)]
+    pub(crate) wptr: u16,
+    #[ver(V < V14_7)]
     pub(crate) event_slot: u32,
+    #[ver(V >= V14_7)]
+    pub(crate) event_slot: u8,
     ...

driver: G15 specifics

Surprisingly little on top of G14X:

• New address space bits; now 43 bits (from 39)
  • Some (render pipeline) addresses are still 32-bit..
• Fault information
• Many scattered struct fields
  • Yet more InitData churn
• Page manager support

Tracing

Ver("G >= G15", "g15_unk_8" / Int64ul), # 0
Ver("G >= G15", "g15_unk_10" / Int64ul), # 0
Ver("G >= G15", "g15_ptr_18" / Int64ul), # blockmgr buf
Ver("G >= G15", "g15_ptr_20" / Int64ul), # 0
Ver("G >= G15", "g15_gptr_28" / Int64ul), # pmgr buf1
Ver("G >= G15", "g15_gptr_30" / Int64ul), # pmgr buf2
Ver("G >= G15", "g15_gptr_38" / Int64ul), # pmgr buf3
Ver("G >= G15", "g15_gptr_40" / Int64ul), # pmgr buf4

Tracing

CBW. WCCP. SCC.[0x164.8] g15_unk_8 = 0x0
CBW. WCCP. SCC.[0x16c.8] g15_unk_10 = 0x0
CBW. WCCP. SCC.[0x174.8] g15_ptr_18 = 0xfffffc2000190000
CBW. WCCP. SCC.[0x17c.8] g15_ptr_20 = 0x0
CBW. WCCP. SCC.[0x184.8] g15_gptr_28 = 0x1000c5d0000
CBW. WCCP. SCC.[0x18c.8] g15_gptr_30 = 0x1000c5df400
CBW. WCCP. SCC.[0x194.8] g15_gptr_38 = 0x1000c5ee800
CBW. WCCP. SCC.[0x19c.8] g15_gptr_40 = 0x1000c5ef000

driver: Page manager

• 4MB chunks of memory
• Allocated by the hardware on-demand
  • Faults come from hardware, not firmware!

.. huh, why would the HW be allocating memory?

driver: Page manager

  ralloc           2052

  mov_imm          r6, 458855      # 0x67 | (7 << 16)
  mov              r7, r6.cache, group 4
  ...
  mov_imm          r6, 134086759   # 0x67 | (2046 << 16)
  mov              r2046, r6.cache, group 4
  mov_imm          r6, 134152295   # 0x67 | (2047 << 16)
  mov              r2047, r6.cache, group 4

what is GPU

      agx_push(out, CDM_GLOBAL_SIZE, cfg) {
         cfg.x = 8;
         cfg.y = 8;
         cfg.z = 4;
      }

driver: Page manager

>>> chexdump(agx.uat.ioread(1, 0x10080000000 + 0x17000, 0x1000))
00000000  67 00 39 00 67 00 39 00  67 00 39 00 67 00 39 00
00000010  *
00000080  67 00 3a 00 67 00 3a 00  67 00 3a 00 67 00 3a 00
00000090  *
00000100  67 00 3b 00 67 00 3b 00  67 00 3b 00 67 00 3b 00
00000110  *
000001c0  67 00 3c 00 67 00 3c 00  67 00 3c 00 67 00 3c 00
000001d0  *

UAPI

For now .. kind of works as-is?

(Assuming you don’t want more than 4GB.)

So.. Linux driver?

[8.086635] asahi 292400000.gpu: GPU: run_job
[8.087234] asahi 292400000.gpu: GPU: run_job: pipe_type=Compute
[8.087993] asahi 292400000.gpu: GPU: run_job: run()
[8.088669] asahi 292400000.gpu: Pipe: RunWorkQueueMsgG15V14_7 { timestamp: 0, work_queue:
  Some(0xfffffc0800267e80 (asahi::fw::workqueue::QueueInfoG15V14_7)), pipe_type: Compute,
  wptr: 1, event_slot: 0, is_new: true, __pad: <pad> }
[8.097272] asahi 292400000.gpu: GPU: run_job: ring doorbell
[8.098834] asahi 292400000.gpu: Stats: PowerOn { off_time: 0x1d604a9f }
[8.099680] asahi 292400000.gpu: Stats: PowerState { timestamp: 0x1d604a9f, last_busy_ts: 0x24edb549,
  active: 0, poweroff: 0, unk1: 0, pstate: 1, unk2: 100, unk3: 0 }
[8.101344] asahi: Unknown stats message: (15, [159, 74, 96, 29, 0, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0])
[8.103587] asahi 292400000.gpu: Stats: PowerOff { on_time: 0x5a96 }

Status

March 2026

commit b5076d68c85a4ca1df054c1f9a7234b31a64fbf4
Date: Sat, 7 Mar 2026

    demo: it's metalover

commit a1006e523146717b4e75078a25c67389b41bf2ce
Date: Sat, 28 Mar 2026

    M3 agx test: compute shader runs

May-August 2026

Status: September

• userspace driver is TODO
  • mesa compiler is WIP
• m1n1 side is .. kind of done?
• Kernel side is still WIP
  • Decided not to speedrun-hack the driver for XDC

Status: coming soon™

Lies

[    6.101783] asahi 292400000.gpu: Stats: TempSensor { __pad: <pad>, sensor_id: 11, raw_value: 21346, scale: 8, tmin: 34, tmax: 41 }
[    6.103294] asahi 292400000.gpu: Event: Timeout { counter: 0, unk_8: 0, event_slot: 0 }
[    6.104420] asahi 292400000.gpu:  (\________/) 
[    6.105180] asahi 292400000.gpu:   |        |  
[    6.105964] asahi 292400000.gpu: '.| \  , / |.'
[    6.106741] asahi 292400000.gpu: --| / (( \ |--
[    6.107513] asahi 292400000.gpu: .'|  _-_-  |'.
[    6.108287] asahi 292400000.gpu:   |________|  
[    6.109060] asahi 292400000.gpu: ** GPU timeout nya~!!!!! **
[    6.109940] asahi 292400000.gpu:   Event slot: 0
[    6.110712] asahi 292400000.gpu:   Timeout count: 0
[    6.111514] asahi 292400000.gpu:   Unk: 0
[    6.112231] asahi 292400000.gpu:   Pending events:
[    6.113030] asahi 292400000.gpu:   Halt count: 0
[    6.113808] asahi 292400000.gpu:   Halted: 0

What about the M4?

• ISA seems (largely?) unchanged
• Just waiting for Sven to add m1n1 support..

A word on AI

• We are passionate volunteers
• We do this for fun
• LLMs take out the fun, hence we didn’t use them
• They also go against many principles we share in this team

A word on AI

• We plan on upstreaming
• Upstream has no AI policy

A word on AI

• Some people out there have taken another route
• They might have gotten further than us..
  • but we’ll keep going at our own pace

Thank you

We’d like to thank Vio for her contributions

(As well as all the Asahi folks)

Tank you

And the audience for listening

Let’s see something now

demo time!