llama-cpp-sys-4 0.10.0

Low Level Bindings to llama.cpp
Documentation 
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93

#pragma OPENCL EXTENSION cl_khr_fp16 : enable

#ifdef cl_qcom_reqd_sub_group_size
#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
__attribute__((qcom_reqd_sub_group_size("full")))
#endif
kernel void kernel_gemm_noshuffle_iq1_m_f32(
    global const uchar * src0_qs,
    global const uchar * src0_qh,
    global const ushort * src0_sc,
    read_only image1d_buffer_t src1,
    global float * dst,
    ulong offsetd,
    int m,
    int n,
    int k,
    int n_no_padding
) {
    dst = (global float *)((global char *)dst + offsetd);
    const int gy = get_global_id(0);
    const int row = get_global_id(1) * 4;
    const int n_4 = n / 4;
    if (row >= m) {
        return;
    }

    half8 c0 = 0, c1 = 0, c2 = 0, c3 = 0;
    for (int i = 0; i < k; i += 16) {
        const uint4 w0 = convert_uint4(vload4(0, src0_sc + row + (4*(i / 256) + 0) * m));
        const uint4 w1 = convert_uint4(vload4(0, src0_sc + row + (4*(i / 256) + 1) * m));
        const uint4 w2 = convert_uint4(vload4(0, src0_sc + row + (4*(i / 256) + 2) * m));
        const uint4 w3 = convert_uint4(vload4(0, src0_sc + row + (4*(i / 256) + 3) * m));
        const ushort4 bits = convert_ushort4((w0 >> (uint4)(12)) | ((w1 >> (uint4)(8)) & (uint4)(0x00f0))
                                           | ((w2 >> (uint4)(4)) & (uint4)(0x0f00)) | (w3 & (uint4)(0xf000)));
        const half4 d = as_half4(bits);
        const uint4 sc = convert_uint4(vload4(0, src0_sc + row + (i / 64) * m));
        const half4 scale = d * convert_half4(2*((sc >> (uint4)(3*((i / 16) & 3))) & (uint4)(7)) + 1);
        const uint4 qh = convert_uint4(vload4(0, src0_qh + row + (i / 16) * m));

        for (int l = 0; l < 16; l += 4) {
            const int ki = i + l;
            const uint4 gi = convert_uint4(vload4(0, src0_qs + row + (ki / 8) * m))
                           | (((qh >> (uint4)(4*((ki / 8) & 1))) & (uint4)(7)) << (uint4)(8));
            const uint4 grid = (uint4)(iq1_s_grid[gi.s0], iq1_s_grid[gi.s1], iq1_s_grid[gi.s2], iq1_s_grid[gi.s3]);
            const float4 delta = select((float4)(-0.875f), (float4)(-1.125f), (qh & (uint4)(0x08 << (4*((ki / 8) & 1)))) != 0);
            #pragma unroll
            for (int j = 0; j < 4; ++j) {
                half8 b;
                b.s0123 = read_imageh(src1, gy * 2     + (ki + j) * n_4);
                b.s4567 = read_imageh(src1, gy * 2 + 1 + (ki + j) * n_4);
                const float4 magnitude = convert_float4((grid >> (uint4)(8*j + (ki & 4))) & (uint4)(15));
                const half4 w = convert_half4(magnitude + delta) * scale;
                c0 += b * w.s0;
                c1 += b * w.s1;
                c2 += b * w.s2;
                c3 += b * w.s3;
            }
        }
    }

    int idx = gy * 8 * m + row;
    if (idx + 3 < m * n_no_padding) {
        vstore4((float4)(c0.s0, c1.s0, c2.s0, c3.s0), 0, dst + idx);
        idx += m;
    }
    if (idx + 3 < m * n_no_padding) {
        vstore4((float4)(c0.s1, c1.s1, c2.s1, c3.s1), 0, dst + idx);
        idx += m;
    }
    if (idx + 3 < m * n_no_padding) {
        vstore4((float4)(c0.s2, c1.s2, c2.s2, c3.s2), 0, dst + idx);
        idx += m;
    }
    if (idx + 3 < m * n_no_padding) {
        vstore4((float4)(c0.s3, c1.s3, c2.s3, c3.s3), 0, dst + idx);
        idx += m;
    }
    if (idx + 3 < m * n_no_padding) {
        vstore4((float4)(c0.s4, c1.s4, c2.s4, c3.s4), 0, dst + idx);
        idx += m;
    }
    if (idx + 3 < m * n_no_padding) {
        vstore4((float4)(c0.s5, c1.s5, c2.s5, c3.s5), 0, dst + idx);
        idx += m;
    }
    if (idx + 3 < m * n_no_padding) {
        vstore4((float4)(c0.s6, c1.s6, c2.s6, c3.s6), 0, dst + idx);
        idx += m;
    }
    if (idx + 3 < m * n_no_padding) {
        vstore4((float4)(c0.s7, c1.s7, c2.s7, c3.s7), 0, dst + idx);
    }
}