cancel
Showing results for 
Search instead for 
Did you mean: 

Archives Discussions

Ono-Sendai
Journeyman III

Issue with signed zero handling

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

0 Likes
4 Replies
dipak
Big Boss

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,

0 Likes

For me, it still produces the incorrect result (+Inf), even with -cl-opt-disable.

0 Likes

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,

0 Likes

Update:

One bug report has been filed against this issue and engg. team is working on it.

Regards,

0 Likes