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.
#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];
printf("a[%u] = %u; b[%u] = %u;\n", gid, aLocal, gid, bLocal);
/*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 = 128;
a = 1024;
b[%u] = %u;
And like this on 5700XT:
a = 128;
b = 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.