char-fork/llama.cpp

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

Masashi Yoshimuraggml-webgpu: add support for f16 repeat (#26307)7a2db1a0c

char/aliceai-t5-moe
brancheschar/aliceai-t5-moeall refs →
1.6 KiB70 linesraw
1enable f16;
2
3struct Params {
4    ne: u32,
5
6    offset_src0: u32,
7    offset_dst: u32,
8
9    stride_src0_0: u32,
10    stride_src0_1: u32,
11    stride_src0_2: u32,
12    stride_src0_3: u32,
13
14    a_ne0: u32,
15    a_ne1: u32,
16    a_ne2: u32,
17    a_ne3: u32,
18
19    ne0: u32,
20    ne1: u32,
21    ne2: u32,
22};
23
24#ifdef TYPE_F32
25#define DataType f32
26#endif
27#ifdef TYPE_I32
28#define DataType i32
29#endif
30#ifdef TYPE_F16
31#define DataType f16
32#endif
33#ifdef TYPE_I16
34// same size (16-bit) is sufficient for repeat
35#define DataType f16
36#endif
37
38@group(0) @binding(0)
39var<storage, read_write> src0: array<DataType>;
40
41@group(0) @binding(1)
42var<storage, read_write> dst: array<DataType>;
43
44@group(0) @binding(2)
45var<uniform> params: Params;
46
47@compute @workgroup_size(WG_SIZE)
48fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
49    if (gid.x < params.ne) {
50        var i = gid.x;
51        let i3 = i / (params.ne2 * params.ne1 * params.ne0);
52        i = i % (params.ne2 * params.ne1 * params.ne0);
53        let i2 = i / (params.ne1 * params.ne0);
54        i = i % (params.ne1 * params.ne0);
55        let i1 = i / params.ne0;
56        let i0 = i % params.ne0;
57
58        let a_i0 = i0 % params.a_ne0;
59        let a_i1 = i1 % params.a_ne1;
60        let a_i2 = i2 % params.a_ne2;
61        let a_i3 = i3 % params.a_ne3;
62
63        let a_index = a_i0 * params.stride_src0_0 +
64                           a_i1 * params.stride_src0_1 +
65                           a_i2 * params.stride_src0_2 +
66                           a_i3 * params.stride_src0_3;
67
68        dst[params.offset_dst + gid.x] = src0[params.offset_src0 + a_index];
69    }
70}