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.

Measured support#

wgsl-text is a device-kernel emitter, not a general-language target. It lowers @ kernel compute kernels and related GPU views, and deliberately nothing else, so it is not scored against the general corpus — a percentage there would read as a completion score it is not. Its measured support is the device kernel support summary.

A compute kernel#

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

Faber source

reader locale

faber convert --to en — English reader surface

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

faber convert --to la — canonical Faber

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

faber convert --to th-TH — Thai

@ เคอร์เนล
ฟังก์ชัน multiplico(tf32[16, 8] a, tf32[8, 16] b, tf32[16, 16] out, u32 id) → เปล่า {
    คงที่ tf32[16, 16] c ← a.คูณเมทริกซ์(b)
}

faber convert --to zh-Hans — Simplified Chinese

@ 内核
函数 multiplico(tf32[16, 8] a, tf32[8, 16] b, tf32[16, 16] out, u32 id) → 无值 {
    常量 tf32[16, 16] c ← a.矩阵乘法(b)
}

faber convert --to zh-Hant — Traditional Chinese

@ 內核
函式 multiplico(tf32[16, 8] a, tf32[8, 16] b, tf32[16, 16] out, u32 id) → 空值 {
    定值 tf32[16, 16] c ← a.矩陣乘法(b)
}

faber convert --to vi — Vietnamese

@ hạt_nhân
hàm multiplico(tf32[16, 8] a, tf32[8, 16] b, tf32[16, 16] out, u32 id) → trống {
    hằng tf32[16, 16] c ← a.nhân_ma_trận(b)
}

faber convert --to ar — Arabic

@ نواة
دالة multiplico(tf32[16, 8] a, tf32[8, 16] b, tf32[16, 16] out, u32 id) → فراغ {
    ثابت tf32[16, 16] c ← a.ضرب_المصفوفات(b)
}

faber convert --to hi — Hindi

@ कर्नेल
फलन multiplico(tf32[16, 8] a, tf32[8, 16] b, tf32[16, 16] out, u32 id) → रिक्त {
    स्थिर tf32[16, 16] c ← a.आव्यूह_गुणन(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 targets · Measured support per term