char-fork/llama.cpp

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

Masashi Yoshimuraggml-webgpu: fix CI errors from #25025 and #25262 (#26566)84f712946

char/aliceai-t5-moe
brancheschar/aliceai-t5-moeall refs →
1.9 KiB84 linesraw
1enable f16;
2
3#ifdef SRC_F32
4#define SRC_TYPE f32
5#elif defined(SRC_F16)
6#define SRC_TYPE f16
7#elif defined(SRC_I32)
8#define SRC_TYPE i32
9#endif
10
11#ifdef DST_F32
12#define DST_TYPE f32
13#elif defined(DST_F16)
14#define DST_TYPE f16
15#elif defined(DST_I32)
16#define DST_TYPE i32
17#endif
18
19@group(0) @binding(0)
20var<storage, read_write> src: array<SRC_TYPE>;
21
22@group(0) @binding(1)
23var<storage, read_write> dst: array<DST_TYPE>;
24
25struct Params{
26    ne: u32,
27    offset_src: u32,
28    offset_dst: u32,
29
30    stride_src0: u32,
31    stride_src1: u32,
32    stride_src2: u32,
33    stride_src3: u32,
34
35
36    stride_dst0: u32,
37    stride_dst1: u32,
38    stride_dst2: u32,
39    stride_dst3: u32,
40
41    src_ne0: u32,
42    src_ne1: u32,
43    src_ne2: u32,
44
45    dst_ne0: u32,
46    dst_ne1: u32,
47    dst_ne2: u32
48};
49
50@group(0) @binding(2)
51var<uniform> params: Params;
52
53@compute @workgroup_size(WG_SIZE)
54fn main(
55    @builtin(global_invocation_id) gid: vec3<u32>,
56) {
57    if (gid.x >= params.ne) {
58        return;
59    }
60
61    var i = gid.x;
62    let i3 = i / (params.src_ne2 * params.src_ne1 * params.src_ne0);
63    i = i % (params.src_ne2 * params.src_ne1 * params.src_ne0);
64    let i2 = i / (params.src_ne1 * params.src_ne0);
65    i = i % (params.src_ne1 * params.src_ne0);
66    let i1 = i / params.src_ne0;
67    let i0 = i % params.src_ne0;
68
69    var j = gid.x;
70    let j3 = j / (params.dst_ne2 * params.dst_ne1 * params.dst_ne0);
71    j = j % (params.dst_ne2 * params.dst_ne1 * params.dst_ne0);
72    let j2 = j / (params.dst_ne1 * params.dst_ne0);
73    j = j % (params.dst_ne1 * params.dst_ne0);
74    let j1 = j / params.dst_ne0;
75    let j0 = j % params.dst_ne0;
76
77    let src_idx = i0 * params.stride_src0 + i1 * params.stride_src1 +
78                  i2 * params.stride_src2 + i3 * params.stride_src3;
79
80    let dst_idx = j0 * params.stride_dst0 + j1 * params.stride_dst1 +
81                  j2 * params.stride_dst2 + j3 * params.stride_dst3;
82
83    dst[params.offset_dst + dst_idx] = DST_TYPE((src[params.offset_src + src_idx]));
84}