To replicate:
Compile the following and pass in x = 1.f. Run on GPU device.
Expected behaviour: main_float_ returns -Inf.
Observed behaviour: main_float_ returns +Inf.
I suspect this is due to a mishandling of signed zero optimsations.
By default OpenCL shouldn't ignore signed zeros.
If you do the computation with signed zeros, should get -Inf
If you ignore signed zeros, then you can optimise
0.f * (x - 2.f)
to
0.f
giving +Inf overall. But this optimisation shouldn't be done if not explicitly requested.
float main_float_(const float x)
{
return (2.f / (0.f * (x - 2.f)));
}
__kernel void main_kernel(float x, __global float * const restrict output_buffer)
{
output_buffer[0] = main_float_(x);
}
AMD Catalyst Control Center Version: 2015.0804.21.41908
Catalyst Version: 15.7.1
GPU: HD7770
Thanks for reporting this issue.
It indeed seems an optimization issue. The kernel is working as expected if it is built without any optimization [i.e. with flag "-cl-opt-disable" or "-O0" during clBuildProgram() ]
Could you please check and confirm?
Regards,
For me, it still produces the incorrect result (+Inf), even with -cl-opt-disable.
In my case, below kernel built with optimization option "-O0" prints "-infinity" when running on HD 7770 with the latest catalyst 15.9 (15.201.1151).
float main_float_(const float x)
{
return (2.f / (0.f * (x - 2.f)));
}
__kernel void main_kernel()
{
float x = main_float_(1.0f);
printf("\nx: %f\n", x);
}
Regards,
Update:
One bug report has been filed against this issue and engg. team is working on it.
Regards,