basic_dsp_vector 0.10.2

Digital signal processing based on real or complex vectors in time or frequency domain.
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

pub const CONV_KERNEL: &str = r#"
    #define mulc32(a,b) ((float4)((float)mad(-(a).y, (b).y, (a).x * (b).x), (float)mad((a).y, (b).x, (a).x * (b).y), (float)mad(-(a).z, (b).z, (a).w * (b).w), (float)mad((a).z, (b).w, (a).w * (b).z)))

    __kernel
    void conv_vecs_r(
                __global float4 const* const src,
                __global float4 const* const conv,
                __global float4* const res)
    {
        ulong const idx = get_global_id(0);

        float4 sum1 = (float4)0.0;
        float4 sum2 = (float4)0.0;
        float4 sum3 = (float4)0.0;
        float4 sum4 = (float4)0.0;
        long start = idx - FILTER_LENGTH / 2;
        long end = start + FILTER_LENGTH;
        long j = 0;
        for (long i = start; i < end; i++) {
            float4 current = src[i];
            sum1 += current * conv[j];
            j++;
            sum2 += current * conv[j];
            j++;
            sum3 += current * conv[j];
            j++;
            sum4 += current * conv[j];
            j++;
        }

        res[idx] = (float4)
                   (sum1.x+sum1.y+sum1.z+sum1.w,
                    sum2.x+sum2.y+sum2.z+sum2.w,
                    sum3.x+sum3.y+sum3.z+sum3.w,
                    sum4.x+sum4.y+sum4.z+sum4.w);
    }

    __kernel
    void conv_vecs_c(
                __global float4 const* const src,
                __global float4 const* const conv,
                __global float4* const res)
    {
        ulong const idx = get_global_id(0);

        float4 sum1 = (float4)0.0;
        float4 sum2 = (float4)0.0;
        long start = idx - FILTER_LENGTH / 2;
        long end = start + FILTER_LENGTH;
        long j = 0;
        for (long i = start; i < end; i++) {
            float4 current = src[i];
            sum1 += mulc32(current, conv[j]);
            j++;
            sum2 += mulc32(current, conv[j]);
            j++;
        }

        // float4 order is: xyzw
        res[idx] = (float4)
                   (sum1.x+sum1.z,
                    sum1.y+sum1.w,
                    sum2.x+sum2.z,
                    sum2.y+sum2.w);
    }
"#;

pub static MUL_KERNEL: &'static str = r#"
    #define mulc32(a,b) ((float2)((float)mad(-(a).y, (b).y, (a).x * (b).x), (float)mad((a).y, (b).x, (a).x * (b).y)))

    __kernel void multiply_vector(
                __global float2 const* const coeff,
                __global float2* const srcres)
    {
        uint const idx = get_global_id(0);
        srcres[idx] = mulc32(srcres[idx], coeff[idx]);
    }
"#;