Hugging Face's logo

Spaces:
natasa365
/
whisper.cpp
Running

App Files Community
whisper.cpp
File size: 4,406 Bytes 
5ff8785
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
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
 
kernel void kernel_concat_f32_contiguous(
    global const char * p_src0, ulong off_src0,
    global const char * p_src1, ulong off_src1,
    global char * p_dst, ulong off_dst,
    int d_ne00, int d_ne01, int d_ne02, // src0->ne[0..2] for the slice
    int d_ne10, int d_ne11, int d_ne12, // src1->ne[0..2] for the slice (d_ne1X must match d_ne0X on non-concat axes)
    int d_ne0,  int d_ne1,  int d_ne2,  // dst->ne[0..2] for the slice
    int dim
) {
    global const float * src0 = (global const float*)((global char*)p_src0 + off_src0);
    global const float * src1 = (global const float*)((global char*)p_src1 + off_src1);
    global float * dst        = (global float*)((global char*)p_dst + off_dst);

    int i0 = get_global_id(0); // Index along dst's 0th dimension
    int i1 = get_global_id(1); // Index along dst's 1st dimension
    int i2 = get_global_id(2); // Index along dst's 2nd dimension

    if (i0 >= d_ne0 || i1 >= d_ne1 || i2 >= d_ne2) {
        return;
    }

    ulong dst_idx = (ulong)i2 * d_ne0 * d_ne1 + (ulong)i1 * d_ne0 + i0;
    ulong src_idx;

    if (dim == 0) {
        if (i0 < d_ne00) { // Data from src0
            src_idx = (ulong)i2 * d_ne00 * d_ne01 + (ulong)i1 * d_ne00 + i0;
            dst[dst_idx] = src0[src_idx];
        } else { // Data from src1
            src_idx = (ulong)i2 * d_ne10 * d_ne11 + (ulong)i1 * d_ne10 + (i0 - d_ne00);
            dst[dst_idx] = src1[src_idx];
        }
    } else if (dim == 1) {
        if (i1 < d_ne01) { // Data from src0
            src_idx = (ulong)i2 * d_ne00 * d_ne01 + (ulong)i1 * d_ne00 + i0;
            dst[dst_idx] = src0[src_idx];
        } else { // Data from src1
            src_idx = (ulong)i2 * d_ne10 * d_ne11 + (ulong)(i1 - d_ne01) * d_ne10 + i0;
            dst[dst_idx] = src1[src_idx];
        }
    } else if (dim == 2) {
        if (i2 < d_ne02) { // Data from src0
            src_idx = (ulong)i2 * d_ne00 * d_ne01 + (ulong)i1 * d_ne00 + i0;
            dst[dst_idx] = src0[src_idx];
        } else { // Data from src1

            src_idx = (ulong)(i2 - d_ne02) * d_ne10 * d_ne11 + (ulong)i1 * d_ne10 + i0;
            dst[dst_idx] = src1[src_idx];
        }
    }
}

kernel void kernel_concat_f32_non_contiguous(
    global const char * p_src0, ulong off_src0,
    global const char * p_src1, ulong off_src1,
    global char * p_dst, ulong off_dst,

    long ne00, long ne01, long ne02, long ne03,
    ulong nb00, ulong nb01, ulong nb02, ulong nb03,

    ulong nb10, ulong nb11, ulong nb12, ulong nb13, // Strides for src1

    long d_ne0, long d_ne1, long d_ne2, long d_ne3,
    ulong d_nb0, ulong d_nb1, ulong d_nb2, ulong d_nb3,
    int dim
) {
    global const char * src0_base = p_src0 + off_src0;
    global const char * src1_base = p_src1 + off_src1;
    global char * dst_base        = p_dst + off_dst;

    long current_i1 = get_global_id(0); // Index for dst_dim_1
    long current_i2 = get_global_id(1); // Index for dst_dim_2
    long current_i3 = get_global_id(2); // Index for dst_dim_3

    if (current_i1 >= d_ne1 || current_i2 >= d_ne2 || current_i3 >= d_ne3) {
        return;
    }

    global const float * x_val_ptr;
    global float * y_val_ptr;

    for (long current_i0 = 0; current_i0 < d_ne0; ++current_i0) {
        bool use_src0;
        long s_i0 = current_i0, s_i1 = current_i1, s_i2 = current_i2, s_i3 = current_i3;

        if (dim == 0) {
            use_src0 = (current_i0 < ne00);
            if (!use_src0) { s_i0 = current_i0 - ne00; }
        } else if (dim == 1) {
            use_src0 = (current_i1 < ne01);
            if (!use_src0) { s_i1 = current_i1 - ne01; }
        } else if (dim == 2) {
            use_src0 = (current_i2 < ne02);
            if (!use_src0) { s_i2 = current_i2 - ne02; }
        } else { // dim == 3
            use_src0 = (current_i3 < ne03);
            if (!use_src0) { s_i3 = current_i3 - ne03; }
        }

        if (use_src0) {
            x_val_ptr = (global const float *)(src0_base + (ulong)s_i3*nb03 + (ulong)s_i2*nb02 + (ulong)s_i1*nb01 + (ulong)s_i0*nb00);
        } else {
            x_val_ptr = (global const float *)(src1_base + (ulong)s_i3*nb13 + (ulong)s_i2*nb12 + (ulong)s_i1*nb11 + (ulong)s_i0*nb10);
        }

        y_val_ptr = (global float *)(dst_base + (ulong)current_i3*d_nb3 + (ulong)current_i2*d_nb2 + (ulong)current_i1*d_nb1 + (ulong)current_i0*d_nb0);
        *y_val_ptr = *x_val_ptr;
    }
}