OpenCL

[___]

Tested on latest 19.10.1 drivers. Windows 10 x64 1903

I attached cl file and cpp program which would launch this simple addVec kernel.

Opencl code:

#pragma OPENCL EXTENSION cl_amd_printf : enable

__attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
__kernel void addVec(__global uint *a, __global uint *b, __global uint *c)
{
  const uint gid = get_global_id(0);
  
  uint aLocal = a[gid];
  uint bLocal = b[gid];
  
  // Good
  printf("a[%u] = %u; b[%u] = %u;\n", gid, aLocal, gid, bLocal);
  
  // Bad
  /*printf("a[%u] = %u;\n", gid, aLocal);
  printf("b[%u] = %u;\n", gid, bLocal);*/

  c[gid] = aLocal + bLocal;
}‍‍‍‍‍‍‍‍‍‍‍‍‍‍‍‍‍‍‍

If you use printf from "Good" section output is fine as expected.

But if you comment first printf and uncomment 2 printfs from "Bad" section it would look like this on vega64:

a[0] = 128;
a[0] = 1024;
b[%u] = %u;

And like this on 5700XT:

a[0] = 128;
b[0] = 1024;

So it works correctly in case of navi and wrong in case of vega. It was broken on vega for a while, at least few months, maybe even before July. I don't remember when i first time noticed that more than 1 printf leads to strange output on vega.