Just change REAL4 to double4. The following code should be suffice, if the problem is not in the unused kernel parameters. I try to check it now, both with -O0 switch.
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
//2 component vector to hold the real and imaginary parts of a complex number:
typedef double2 cREAL;
/* Multiply two complex numbers */
inline cREAL cmult(cREAL a, cREAL b) {
return (cREAL)(fma(a.x, b.x, -a.y * b.y), fma(a.x, b.y, a.y * b.x));
}
__kernel void p3loop(__global double * const results) {
const uint
ind_gl = ((get_global_id(0) * get_global_size(1) + get_global_id(1)) * get_global_size(2) + get_global_id(2)); // global index
double4 ptemp;
cREAL alpha[6];
alpha[0] = (cREAL)(0, -0.2);
alpha[1] = (cREAL)(0, 0.05);
alpha[2] = (cREAL)(0.6, 0.0005);
alpha[3] = (cREAL)(-0.1, 0.75);
alpha[4] = (cREAL)(0.9, -.03);
alpha[5] = (cREAL)(0.2, 0.1);
ptemp.lo = (cREAL)(1, 0);
for (volatile int j = 0; j < 6; ++j) {
ptemp.hi = (cREAL)(1, 0) - alpha;
ptemp.hi = cmult(ptemp.hi, ptemp.hi);
ptemp.lo = cmult(ptemp.hi, ptemp.lo);
/*
if (j == 5) {
results[ind_gl] = ptemp.lo.y;
return;
}
*/
}
results[ind_gl] = ptemp.lo.y;
return;
} // end of kernel
