char-fork/llama.cpp

git clone https://git.t4t.associates/char-fork/llama.cpp

Reese Levineggml webgpu: ops support for qwen3.5 (SET, TRI_SOLVE, SSM_CONV, GATED_DELTA_NET) + GET_ROWS optimization (#20687)c1258830b

char/aliceai-t5-moe
brancheschar/aliceai-t5-moeall refs →
1.7 KiB65 linesraw
1@group(0) @binding(0)
2var<storage, read_write> src0: array<f32>;
3
4@group(0) @binding(1)
5var<storage, read_write> src1: array<f32>;
6
7@group(0) @binding(2)
8var<storage, read_write> dst: array<f32>;
9
10struct Params {
11    offset_src0: u32,
12    offset_src1: u32,
13    offset_dst: u32,
14
15    stride_src01: u32,
16    stride_src02: u32,
17    stride_src11: u32,
18
19    stride_dst0: u32,
20    stride_dst1: u32,
21    stride_dst2: u32,
22
23    nc: u32,
24    nr: u32,
25    n_t: u32,
26    n_s: u32,
27    token_tiles: u32,
28};
29
30@group(0) @binding(3)
31var<uniform> params: Params;
32
33@compute @workgroup_size(BLOCK_SIZE, TOKENS_PER_WG)
34fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
35    let i1 = gid.x;
36    let tile_y = gid.y / TOKENS_PER_WG;
37    let local_token = gid.y % TOKENS_PER_WG;
38    let i3 = tile_y / params.token_tiles;
39    let token_tile = tile_y % params.token_tiles;
40    let i2 = token_tile * TOKENS_PER_WG + local_token;
41
42    if (i1 >= params.nr || i2 >= params.n_t || i3 >= params.n_s) {
43        return;
44    }
45
46    let src0_base = params.offset_src0 + i3 * params.stride_src02 + i2 + i1 * params.stride_src01;
47    let src1_base = params.offset_src1 + i1 * params.stride_src11;
48
49    var sum = 0.0;
50
51#ifdef VECTORIZED
52    sum =
53        src0[src0_base + 0u] * src1[src1_base + 0u] +
54        src0[src0_base + 1u] * src1[src1_base + 1u] +
55        src0[src0_base + 2u] * src1[src1_base + 2u] +
56        src0[src0_base + 3u] * src1[src1_base + 3u];
57#else
58    for (var i0 = 0u; i0 < params.nc; i0++) {
59        sum += src0[src0_base + i0] * src1[src1_base + i0];
60    }
61#endif
62
63    let dst_idx = params.offset_dst + i3 * params.stride_dst2 + i2 * params.stride_dst1 + i1 * params.stride_dst0;
64    dst[dst_idx] = sum;
65}