Upload shaders/matmul_gemv.wgsl with huggingface_hub

shaders/matmul_gemv.wgsl

@@ -0,0 +1,63 @@
+// matmul_gemv.wgsl — M=1 specialized matrix-vector product.
+// Ported from gfx1151_inference/shaders/matmul_gemv.comp.
+//
+// Computes c[n] = sum_k(a[k] * B[n, k]) for n in [0, N).
+//   A is a vector [K]
+//   B is row-major [N, K]
+//   C is a vector [N]
+//
+// One WORKGROUP per output element n. Dispatch (N, 1, 1).
+// Each workgroup has 64 threads cooperating on one dot product.
+// Threads stride over K: thread t handles k = t, t+64, t+128, ...
+// Tree reduction in shared memory yields c[n].
+
+const WG_SIZE: u32 = 64u;
+
+struct Params {
+    N: u32,
+    K: u32,
+}
+
+@group(0) @binding(0) var<storage, read>       a_buf: array<f32>;  // [K]
+@group(0) @binding(1) var<storage, read>       b_buf: array<f32>;  // [N, K]
+@group(0) @binding(2) var<storage, read_write> c_buf: array<f32>;  // [N]
+
+@group(1) @binding(0) var<uniform> params: Params;
+
+var<workgroup> partial: array<f32, 64>;  // WG_SIZE
+
+@compute @workgroup_size(WG_SIZE)
+fn main(
+    @builtin(workgroup_id) wg_id: vec3<u32>,
+    @builtin(local_invocation_id) lid: vec3<u32>
+) {
+    let n   = wg_id.x;
+    let tid = lid.x;
+    if (n >= params.N) { return; }
+
+    let b_base = n * params.K;
+    var acc: f32 = 0.0;
+    var k: u32 = tid;
+    loop {
+        if (k >= params.K) { break; }
+        acc = acc + a_buf[k] * b_buf[b_base + k];
+        k = k + WG_SIZE;
+    }
+    partial[tid] = acc;
+    workgroupBarrier();
+
+    // Tree reduction
+    var off: u32 = WG_SIZE / 2u;
+    loop {
+        if (off == 0u) { break; }
+        if (tid < off) {
+            partial[tid] = partial[tid] + partial[tid + off];
+        }
+        workgroupBarrier();
+        off = off >> 1u;
+    }
+
+    if (tid == 0u) {
+        c_buf[n] = partial[0u];
+    }
+}