cera 0.5.3

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
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
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140

#include <metal_stdlib>
#include <metal_math>
#include <metal_texture>
using namespace metal;

#line 41 "cera/src/backend/shaders/slang/conv1d_fused.slang"
struct KernelContext_0
{
    uint device* par_buf_0;
    float device* proj_buf_0;
    float device* rbuf_0;
    float device* w_buf_0;
    float device* out_buf_0;
};


#line 24
[[kernel]] void conv1d_fused(uint3 gid_0 [[thread_position_in_grid]], uint device* par_buf_1 [[buffer(4)]], float device* proj_buf_1 [[buffer(0)]], float device* rbuf_1 [[buffer(1)]], float device* w_buf_1 [[buffer(2)]], float device* out_buf_1 [[buffer(3)]])
{

#line 24
    thread KernelContext_0 kernelContext_0;

#line 24
    (&kernelContext_0)->par_buf_0 = par_buf_1;

#line 24
    (&kernelContext_0)->proj_buf_0 = proj_buf_1;

#line 24
    (&kernelContext_0)->rbuf_0 = rbuf_1;

#line 24
    (&kernelContext_0)->w_buf_0 = w_buf_1;

#line 24
    (&kernelContext_0)->out_buf_0 = out_buf_1;
    uint ch_0 = gid_0.x;
    uint hs_0 = par_buf_1[int(0)];
    uint ks_0 = par_buf_1[int(1)];
    uint d_conv_0 = par_buf_1[int(2)];

    if(ch_0 >= hs_0)
    {

#line 31
        return;
    }


    float c_val_0 = (&kernelContext_0)->proj_buf_0[hs_0 + ch_0];

    float bx_0 = (&kernelContext_0)->proj_buf_0[ch_0] * (&kernelContext_0)->proj_buf_0[2U * hs_0 + ch_0];

#line 37
    uint k_0 = 0U;

#line 37
    float sum_0 = 0.0f;


    for(;;)
    {

#line 40
        if(k_0 < d_conv_0)
        {
        }
        else
        {

#line 40
            break;
        }

#line 41
        float sum_1 = sum_0 + *((&kernelContext_0)->rbuf_0+(k_0 * hs_0 + ch_0)) * (&kernelContext_0)->w_buf_0[ch_0 * ks_0 + k_0];

#line 40
        k_0 = k_0 + 1U;

#line 40
        sum_0 = sum_1;

#line 40
    }


    float sum_2 = sum_0 + bx_0 * (&kernelContext_0)->w_buf_0[ch_0 * ks_0 + d_conv_0];

    if(d_conv_0 > 1U)
    {

#line 45
        k_0 = 0U;
        for(;;)
        {

#line 46
            if(k_0 < (d_conv_0 - 1U))
            {
            }
            else
            {

#line 46
                break;
            }

#line 47
            uint _S1 = k_0 + 1U;

#line 47
            *((&kernelContext_0)->rbuf_0+(k_0 * hs_0 + ch_0)) = *((&kernelContext_0)->rbuf_0+(_S1 * hs_0 + ch_0));

#line 46
            k_0 = _S1;

#line 46
        }

#line 45
    }

#line 50
    if(d_conv_0 > 0U)
    {

#line 51
        *((&kernelContext_0)->rbuf_0+((d_conv_0 - 1U) * hs_0 + ch_0)) = bx_0;

#line 50
    }



    *((&kernelContext_0)->out_buf_0+ch_0) = c_val_0 * sum_2;
    return;
}