Archives Discussions

[doqtor]

The problem is when using logarithms. See below example:

__kernel __attribute__((reqd_work_group_size(256, 1, 1)))

void test(__global double *a, __global double *c)

{

    int local_id = get_local_id(0);

    c[local_id] = log(a[local_id]);

}

Array "a" is filled with numbers from 0 to 255:

double a[256] = {0...255};

Then the first 3 results are OK but then the rest is wrong:

wrong result [3]: -1.74441e+136 <> 1.09861

wrong result [5]: 1.37262 <> 1.60944

wrong result [6]: -1.74441e+136 <> 1.79176

wrong result [7]: -2.68156e+154 <> 1.94591

wrong result [9]: 2.07944 <> 2.19722

wrong result [10]: 2.06577 <> 2.30259

...

wrong result [255]: 4.84812 <> 5.54126

The same kernel works fine when built online or when built as OpenCL1.x online or loaded from binary.

[dipak]

It would be helpful if you could provide the complete project with the host-code such that we can check it here. BTW, did you specify the OpenCL 2.0 flag correctly during the kernel program building?

[doqtor]

There you go. Attached just quickly made example reproducing the problem.

First run: "./oclLoadBinaryKernel" to output binary kernel.

Then: "./oclLoadBinaryKernel 1" to run kernel loaded from binary. It should produce totally wrong results which will be outputted to the console.

My environment:

Ubuntu 14.04

AMD R9 295

fglrx 15.20.3

AMD APP SDK 3.0

[dipak]

Thanks for sharing the code. We'll check and get back to you shortly.

Regards,

[dipak]

Hi doqtor,

As I checked with a Hawaii card, I indeed got a similar observation. However, I don't have any explanation right now. I'll check with some relevant folks and come back to you.

BTW, in my case, the kernel was running fine if :

1) replace the log(x) function by native_log(x)

OR

2) change the data-type from double to float

You may use one of the above cases as a temporary workaround, if it works for you.

Regards,

[doqtor]

Hi dipak​,

Thanks for the suggestions but none of them will work for me.

I found that exp suffers with the same issue so there may be more built-in math functions affected.

Regards,

[dipak · Accepted solution]

Hi,

Here is an update. It seems that the above issue is no longer reproducible with the latest internal build of the Catalyst driver. So, please keep your patience till that driver version is publicly available.

Regards,

[doqtor]

Hi dipak​,

Do you know if the driver with the bug fix has already been publicly released?

Regards,

[dipak]

Sorry, I haven't tested it with latest pubic Catalyst driver 15.9 (15.201.1151) [Desktop ]. So, I'm not sure.

However, the internal driver (from 15.30 series) where I didn't observe the bug is yet to be released.

Regards,

[doqtor]

I tested 201 and bug is not fixed in that release. I thought the bug fix may be released because pinform​ marked your answer as a correct answer few days ago.

[dipak]

Sorry for the inconvenience.