cera 0.5.5

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
141
142
143
144
145
146
147
148
149
150
151

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

#line 60 "cera/src/backend/shaders/slang/overlap_add.slang"
struct KernelContext_0
{
    packed_uint4 device* par_buf_0;
    float device* hann_buf_0;
    float device* td_buf_0;
    float device* out_buf_0;
};


#line 33
[[kernel]] void overlap_add(uint3 gid_0 [[thread_position_in_grid]], packed_uint4 device* par_buf_1 [[buffer(3)]], float device* hann_buf_1 [[buffer(1)]], float device* td_buf_1 [[buffer(0)]], float device* out_buf_1 [[buffer(2)]])
{

#line 33
    thread KernelContext_0 kernelContext_0;

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

#line 33
    (&kernelContext_0)->hann_buf_0 = hann_buf_1;

#line 33
    (&kernelContext_0)->td_buf_0 = td_buf_1;

#line 33
    (&kernelContext_0)->out_buf_0 = out_buf_1;
    uint n_frames_0 = (uint4(*(par_buf_1+int(0))) ).x;
    uint n_fft_0 = (uint4(*(par_buf_1+int(0))) ).y;
    uint hop_0 = (uint4(*(par_buf_1+int(0))) ).z;

    uint g_0 = gid_0.x;
    if(g_0 >= (n_frames_0 * hop_0))
    {

#line 40
        return;
    }

    uint i_hi_0 = g_0 / hop_0;

#line 43
    uint i_hi_1;
    if(i_hi_0 >= n_frames_0)
    {

#line 44
        i_hi_1 = n_frames_0 - 1U;

#line 44
    }
    else
    {

#line 44
        i_hi_1 = i_hi_0;

#line 44
    }

#line 44
    uint i_lo_0;



    if(g_0 >= n_fft_0)
    {

#line 49
        uint _S1 = (g_0 - n_fft_0) / hop_0;

#line 49
        i_lo_0 = _S1 + 1U;

#line 48
    }
    else
    {

#line 48
        i_lo_0 = 0U;

#line 48
    }

#line 48
    uint i_0 = i_lo_0;

#line 48
    float numer_0 = 0.0f;

#line 48
    float denom_0 = 0.0f;

#line 54
    for(;;)
    {

#line 54
        if(i_0 <= i_hi_1)
        {
        }
        else
        {

#line 54
            break;
        }

#line 55
        uint local_0 = g_0 - i_0 * hop_0;
        float w_0 = (&kernelContext_0)->hann_buf_0[local_0];
        float numer_1 = numer_0 + (&kernelContext_0)->td_buf_0[i_0 * n_fft_0 + local_0] * w_0;
        float denom_1 = denom_0 + w_0 * w_0;

#line 54
        i_0 = i_0 + 1U;

#line 54
        numer_0 = numer_1;

#line 54
        denom_0 = denom_1;

#line 54
    }

#line 60
    float device* _S2 = (&kernelContext_0)->out_buf_0+g_0;

#line 60
    if(denom_0 > 9.99999993922529029e-09f)
    {

#line 60
        numer_0 = numer_0 / denom_0;

#line 60
    }

#line 60
    *_S2 = numer_0;
    return;
}