WGSL

WebGPU compute shader source.

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

How to read it#

Bindings, workgroup declarations, and bounds guards that the kernel never spells out. This is the case for writing kernels in Faber: the source stays about the computation.

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)
}

WGSL — 4 lines in, 34 out (8.5×)

// Generated by radix wgsl-text (supported-with-limitations compute source).

var<workgroup> shared_a: array<f32, 64u>;
var<workgroup> shared_b: array<f32, 64u>;

@group(0) @binding(0) var<storage, read> a_in: array<f32>;
@group(0) @binding(1) var<storage, read> b_in: array<f32>;
@group(0) @binding(2) var<storage, read_write> output: array<f32>;

@compute @workgroup_size(8, 8, 1)
fn multiplico(@builtin(global_invocation_id) id: vec3<u32>, @builtin(local_invocation_id) local_id: vec3<u32>) {
  let i: u32 = id.x;
var acc: f32 = 0.0;
let row = id.y;
let col = id.x;
let ty = local_id.y;
let tx = local_id.x;
for (var k_tile: u32 = 0u; k_tile < 1u; k_tile++) {
    let k_start = k_tile * 8u;
    let 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; }
    let 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; }
    workgroupBarrier();
    for (var kk: u32 = 0u; kk < 8u; kk++) {
        acc += shared_a[ty * 8u + kk] * shared_b[kk * 8u + tx];
    }
    workgroupBarrier();
}
let out_idx = row * 16u + col;
if (row < 16u && col < 16u) { output[out_idx] = acc; }
}

---

All lanes · Measured support per term