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