char-fork/llama.cpp

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

lhezopencl: add neg, exp and diag (#20127)6c97bffd6

char/aliceai-t5-moe
brancheschar/aliceai-t5-moeall refs →
834 B27 linesraw
1kernel void kernel_diag_f32(
2    global const char * src0,
3    ulong               offset0,
4    global       char * dst,
5    ulong               offsetd,
6    ulong               nb01,
7    ulong               nb02,
8    ulong               nb03,
9    int                 ne0,
10    ulong               nb0,
11    ulong               nb2,
12    ulong               nb3
13) {
14    src0 = src0 + offset0;
15    dst  = dst + offsetd;
16
17    int i3 = get_group_id(2);
18    int i2 = get_group_id(1);
19    int i1 = get_group_id(0);
20
21    global const float * src0_ptr = (global const float *)(src0 +           i2*nb02 + i3*nb03);
22    global       float * dst_ptr  = (global       float *)(dst  + i1*nb01 + i2*nb2  + i3*nb3);
23
24    for (int i0 = get_local_id(0); i0 < ne0; i0 += get_local_size(0)) {
25        dst_ptr[i0] = i0 == i1 ? src0_ptr[i0] : 0.0f;
26    }
27}