cera 0.5.2

Rust-native LLM inference engine
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

@binding(2) @group(0) var<storage, read> par_buf_0 : array<vec2<u32>>;

@binding(0) @group(0) var<storage, read_write> a_buf_0 : array<f32>;

@binding(1) @group(0) var<storage, read> b_buf_0 : array<f32>;

@compute
@workgroup_size(256, 1, 1)
fn add_inplace(@builtin(global_invocation_id) gid_0 : vec3<u32>)
{
    var i_0 : u32 = gid_0.x;
    if(i_0 >= (par_buf_0[i32(0)].x))
    {
        return;
    }
    a_buf_0[i_0] = a_buf_0[i_0] + b_buf_0[i_0];
    return;
}

@compute
@workgroup_size(256, 1, 1)
fn scaled_add_inplace(@builtin(global_invocation_id) gid_1 : vec3<u32>)
{
    var i_1 : u32 = gid_1.x;
    if(i_1 >= (par_buf_0[i32(0)].x))
    {
        return;
    }
    a_buf_0[i_1] = a_buf_0[i_1] + (bitcast<f32>((par_buf_0[i32(0)].y))) * b_buf_0[i_1];
    return;
}

@compute
@workgroup_size(256, 1, 1)
fn mul_inplace(@builtin(global_invocation_id) gid_2 : vec3<u32>)
{
    var i_2 : u32 = gid_2.x;
    if(i_2 >= (par_buf_0[i32(0)].x))
    {
        return;
    }
    a_buf_0[i_2] = a_buf_0[i_2] * b_buf_0[i_2];
    return;
}

@compute
@workgroup_size(256, 1, 1)
fn silu_mul_inplace(@builtin(global_invocation_id) gid_3 : vec3<u32>)
{
    var i_3 : u32 = gid_3.x;
    if(i_3 >= (par_buf_0[i32(0)].x))
    {
        return;
    }
    var g_0 : f32 = clamp(a_buf_0[i_3], -80.0f, 80.0f);
    a_buf_0[i_3] = g_0 / (1.0f + exp(- g_0)) * b_buf_0[i_3];
    return;
}