llama-cpp-sys-4 0.9.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

#pragma OPENCL EXTENSION cl_khr_fp16 : enable

// Quantize a contiguous K-major [N, K] f32 activation buffer (K contiguous) into 32-value Q8_1 SoA blocks.
// Store int8 quants, per-block scale d, and scaled sum s = d * sum(qs); one work-item handles each block.
__kernel void kernel_quant_a_q8_1(
        __global const float * src,   // [N * K]
        __global       char  * qa,    // [N * K]
        __global       half  * da,    // [N * (K/32)]
        __global       half  * sa,    // [N * (K/32)]
        int total_blocks              // N * (K/32)
) {
    const int blk = get_global_id(0);
    if (blk >= total_blocks) {
        return;
    }

    const int base = blk * 32;

    float v[32];
    float amax = 0.0f;
    #pragma unroll
    for (int i = 0; i < 32; ++i) {
        v[i] = src[base + i];
        amax = fmax(amax, fabs(v[i]));
    }

    const float d  = amax / 127.0f;
    const float id = (amax > 0.0f) ? (127.0f / amax) : 0.0f;

    int sum = 0;
    #pragma unroll
    for (int i = 0; i < 32; ++i) {
        const int q = (int)rint(v[i] * id);
        qa[base + i] = (char)q;
        sum += q;
    }

    da[blk] = (half)d;
    sa[blk] = (half)(d * (float)sum);
}