Metal

Apple GPU compute shader source (MSL).

Part of the GPU lane. Every panel below is compiler output.

How to read it#

The same kernel, a different ABI. Compare it against the WGSL beside it — one Faber function, two unrelated shading languages, neither written by hand.

A compute kernel#

A function marked @ nucleum. Device lanes only — this is a different kind of source, not a variant of the programs above.

Faber source

@ kernel { }
fn multiplico(tensor<f32, [16, 8]> a, tensor<f32, [8, 16]> b, tensor<f32, [16, 16]> out, int<u32> id) → void {
    const tensor<f32, [16, 16]> c ← a.matmul(b)
}

@ nucleum { }
functio multiplico(tensor<f32, [16, 8]> a, tensor<f32, [8, 16]> b, tensor<f32, [16, 16]> out, numerus<u32> id) → vacuum {
    fixum tensor<f32, [16, 16]> c ← a.matmul(b)
}

@ เคอร์เนล { }
ฟังก์ชัน multiplico(เทนเซอร์<f32, [16, 8]> a, เทนเซอร์<f32, [8, 16]> b, เทนเซอร์<f32, [16, 16]> out, จำนวน<u32> id) → เปล่า {
    คงที่ เทนเซอร์<f32, [16, 16]> c ← a.matmul(b)
}

@ 内核 { }
函数 multiplico(张量<f32, [16, 8]> a, 张量<f32, [8, 16]> b, 张量<f32, [16, 16]> out, 整数<u32> id) → 无值 {
    常量 张量<f32, [16, 16]> c ← a.matmul(b)
}

@ 內核 { }
函式 multiplico(張量<f32, [16, 8]> a, 張量<f32, [8, 16]> b, 張量<f32, [16, 16]> out, 整數<u32> id) → 空值 {
    定值 張量<f32, [16, 16]> c ← a.matmul(b)
}

@ hạt_nhân { }
hàm multiplico(ten_xo<f32, [16, 8]> a, ten_xo<f32, [8, 16]> b, ten_xo<f32, [16, 16]> out, số<u32> id) → trống {
    hằng ten_xo<f32, [16, 16]> c ← a.matmul(b)
}

@ نواة { }
دالة multiplico(موتر<f32, [16, 8]> a, موتر<f32, [8, 16]> b, موتر<f32, [16, 16]> out, عدد<u32> id) → فراغ {
    ثابت موتر<f32, [16, 16]> c ← a.matmul(b)
}

@ कर्नेल { }
फलन multiplico(टेंसर<f32, [16, 8]> a, टेंसर<f32, [8, 16]> b, टेंसर<f32, [16, 16]> out, संख्या<u32> id) → रिक्त {
    स्थिर टेंसर<f32, [16, 16]> c ← a.matmul(b)
}

Metal — 4 lines in, 36 out (9.0×)

// Generated by radix metal-text (supported-with-limitations compute source).
#include <metal_stdlib>
using namespace metal;

kernel void multiplico(
    device const float* a_in [[buffer(0)]],
    device const float* b_in [[buffer(1)]],
    device float* output [[buffer(2)]],
    uint3 id [[thread_position_in_grid]],
    uint3 local_id [[thread_position_in_threadgroup]]
) {
    uint i = id.x;
threadgroup float shared_a[64];
threadgroup float shared_b[64];
float acc = 0.0;
uint row = id.y;
uint col = id.x;
uint ty = local_id.y;
uint tx = local_id.x;
for (uint k_tile = 0u; k_tile < 1u; k_tile++) {
    uint k_start = k_tile * 8u;
    uint a_idx = row * 8u + (k_start + tx);
    if (a_idx < 16u * 8u) { shared_a[ty * 8u + tx] = a_in[a_idx]; }
    if (a_idx >= 16u * 8u) { shared_a[ty * 8u + tx] = 0.0; }
    uint b_idx = (k_start + ty) * 16u + col;
    if (b_idx < 8u * 16u) { shared_b[ty * 8u + tx] = b_in[b_idx]; }
    if (b_idx >= 8u * 16u) { shared_b[ty * 8u + tx] = 0.0; }
    threadgroup_barrier(mem_flags::mem_threadgroup);
    for (uint kk = 0u; kk < 8u; kk++) {
        acc += shared_a[ty * 8u + kk] * shared_b[kk * 8u + tx];
    }
    threadgroup_barrier(mem_flags::mem_threadgroup);
}
uint out_idx = row * 16u + col;
if (row < 16u && col < 16u) { output[out_idx] = acc; }
}

---

All lanes · Measured support per term