AnsweredAssumed Answered

OpenCL driver bug

Question asked by rosenrodt on Jan 30, 2019
Latest reply on Apr 1, 2019 by dipak

EDIT: reformat 

EDIT 2: correct driver version


Found a weird behavior in AMD's OpenCL compiler. Code taken straight from Boost library:


__kernel void serial_adjacent_find(const uint size, __global uint* output, __global char* buf0) 
  uint result = size;
  for (uint i = 0; i<size - 1; i++) {
    if (buf0[i+1] < buf0[i]) {
      result = i;
  *output = result; }

The kernel checks if input array `buf0` is monotonically increasing. If any element in the array `buf0` is greater than the next element, it breaks out of the loop and alters the output value, which the host side checks against. When the input array is of type `char`, the kernel will break out of the loop prematurely. However with other types such as int or float this kernel runs fine.

I can repro the bug on AMD RX500/Vega series. Running on Windows 10 + Adrenalin 18.2.2 18.12.2/19.1.1.

I don't see this exact kernel causing an issue with Intel/Nvidia/Qualcomm OpenCL driver compilers.


I have attached a reproducible CMake project here for your quick reference