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