xref: /aosp_15_r20/external/clpeak/src/kernels/compute_short_kernels.cl (revision 1cd03ba3888297bc945f2c84574e105e3ced3e34) 
1MSTRINGIFY(
2
3// Avoiding auto-vectorize by using vector-width locked dependent code
4
5\n#undef MAD_4
6\n#undef MAD_16
7\n#undef MAD_64
8\n
9\n#define MAD_4(x, y)     x = (y*x) + y;      y = (x*y) + x;      x = (y*x) + y;      y = (x*y) + x;
10\n#define MAD_16(x, y)    MAD_4(x, y);        MAD_4(x, y);        MAD_4(x, y);        MAD_4(x, y);
11\n#define MAD_64(x, y)    MAD_16(x, y);       MAD_16(x, y);       MAD_16(x, y);       MAD_16(x, y);
12\n
13
14__kernel void compute_short_v1(__global short *ptr, short _A)
15{
16    short x = _A;
17    short y = (short)get_local_id(0);
18
19    for(int i=0; i<64; i++)
20    {
21        MAD_16(x, y);
22    }
23
24    ptr[get_global_id(0)] = y;
25}
26
27
28__kernel void compute_short_v2(__global short *ptr, short _A)
29{
30    short2 x = (short2)(_A, (_A+1));
31    short2 y = (short2)get_local_id(0);
32
33    for(int i=0; i<32; i++)
34    {
35        MAD_16(x, y);
36    }
37
38    ptr[get_global_id(0)] = (y.S0) + (y.S1);
39}
40
41__kernel void compute_short_v4(__global short *ptr, short _A)
42{
43    short4 x = (short4)(_A, (_A+1), (_A+2), (_A+3));
44    short4 y = (short4)get_local_id(0);
45
46    for(int i=0; i<16; i++)
47    {
48        MAD_16(x, y);
49    }
50
51    ptr[get_global_id(0)] = (y.S0) + (y.S1) + (y.S2) + (y.S3);
52}
53
54
55__kernel void compute_short_v8(__global short *ptr, short _A)
56{
57    short8 x = (short8)(_A, (_A+1), (_A+2), (_A+3), (_A+4), (_A+5), (_A+6), (_A+7));
58    short8 y = (short8)get_local_id(0);
59
60    for(int i=0; i<8; i++)
61    {
62        MAD_16(x, y);
63    }
64
65    ptr[get_global_id(0)] = (y.S0) + (y.S1) + (y.S2) + (y.S3) + (y.S4) + (y.S5) + (y.S6) + (y.S7);
66}
67
68__kernel void compute_short_v16(__global short *ptr, short _A)
69{
70    short16 x = (short16)(_A, (_A+1), (_A+2), (_A+3), (_A+4), (_A+5), (_A+6), (_A+7),
71                    (_A+8), (_A+9), (_A+10), (_A+11), (_A+12), (_A+13), (_A+14), (_A+15));
72    short16 y = (short16)get_local_id(0);
73
74    for(int i=0; i<4; i++)
75    {
76        MAD_16(x, y);
77    }
78
79    short2 t = (y.S01) + (y.S23) + (y.S45) + (y.S67) + (y.S89) + (y.SAB) + (y.SCD) + (y.SEF);
80    ptr[get_global_id(0)] = t.S0 + t.S1;
81}
82
83
84)
85