Hi all,
I have a severe problem when using one of my OpenCL kernels in HSA mode.
When HSA is disabled, everything computes just fine, but by enabling HSA, a certain kernel produces a wrong output!
For demonstration purposes I wrote a little test class (see attachment: HSATest.cpp). Note that the kernel is much shorter than the one I use but results in the same misbehavior!
kernel code part:
(...)
__kernel void Test(__global int* left,
__global int* under,
const int start_pos)
{
if (start_pos > 0) {
int left_child = ((start_pos + 1) / 2) - 1;
left[start_pos] = left_child;
if ((left_child % 2) != 0) { // this is false! WHY?!
under[left[start_pos]] = start_pos;
}
}
}
(...)
'left' and 'under' are both arrays of the length 4 filled with -1 values. 'start_pos' is a simple integer of value 3.
The problem lies in the line "if ((left_child % 2) != 0)" which results in being 'false' even if 'left_child' is 1 (which is wrong, because 1%2 = 1 and therefore != 0)!
If I replace this line with "if (left_child == 1)" or even "if ((left_child % 2) == 1)" it results in 'true' (which is correct)!
Please read the comments in the (kernel) code for more information on this problem.
The output should be:
under result: [ -1 3 -1 -1 ]
left result: [ -1 -1 -1 1 ]
but instead I get:
under result: [ -1 -1 -1 -1 ]
left result: [ -1 -1 -1 1 ]
because of the "if ((left_child % 2) != 0)" line being interpreted as 'false' and therefore it skips the part which sets the second entry of the 'under'-Array to 3.
My system specs are:
Windows 7 Professional with Service Pack 1 (64-bit)
CPU: Kaveri AMD A10-7850K
Motherboard: A88X-PRO from ASUSTeK COMPUTER INC.
BIOS version: 0904
I use Microsoft Visual C++ 2010 Express SP1 with Windows SDK 7.1.
I also attached the output from clinfo for more information.
The thing which also bothers me is that clinfo shows my device as being a HSA Device but "cl_amd_hsa" is missing from the Platform Extensions!
Update:
After i couldn't determine the cause of this problem, I did a fresh OS install with Windows 8.1 Pro (64-bit). I redownloaded and installed the drivers from http://developer.amd.com/tools-and-sdks/opencl-zone/opencl-1-2-beta-driver/ and did everything as noted in the readme file.
The installation went smooth and when running my device in HSA mode, the clinfo finally shows the "cl_amd_hsa" Platform Extension. I again attached the clinfo output.
Sadly the problem remains!
Even after compiling my code with Visual Studio 2013 Express with the Visual Studio 2013 (v120) Platformtoolset and compiling for 64bit or 32bit.
Am I the only one having issues with this kernel?
