OpenCL, Некорректная работа на AMD A6-6420K (HD8470D) при корректной на Polaris12. Косяк у меня или в llvm?

Question

Yupa20171123 @Yupa20171123

OpenCL

OpenCL, Некорректная работа на AMD A6-6420K (HD8470D) при корректной на Polaris12. Косяк у меня или в llvm?

sources TestGpu.tar.gz: https://yadi.sk/d/Bzkxr-nz3Q3zCp

Вычисленные значения температуры заливались в массив пикселей и выводились на экран (через socket, на просмотровшик).
Иногда (около 20% случаев) на экране были полосы или области с шумом. На процессоре прогонка работала. На GPU - нет.

Стал тестировать отдельные операции (работа с uchar - не работает, ...)

Выделяю кусок памяти в GPU.
Заливаю туда такой же по размеру залитый нулем.
Десять раз: { Читаю оттуда 256*256*4 и сохраняю в *.bmp. Запускаю kernel. }

Каждый поток (item) gpu: (их 256)

n= get_global_id(0);
 //if (n & 1 ) - блокирует все.  
 //if n in { 1,2,4,8,...} return; (used switch)  - должен дать полоски черного цвета. 
for (val = 0; val<256;val++){
   a=n;
   while (a<END){ mem[a]=val; a+=256; }
   a=n;
   while  (a<END){ if (mem[a]!=val){ goto ERROR; } a+=256; }
}

На llvm 3.* на встроенном в CPU программа собиралась. На RX550 нет (unsupported call of get_local_size,...)
На ~amd64, llvm5 - polaris12 заработал, (пока кажется) корректно. А встроенный все также.

Драйвера: Встроенный - radeon. Polaris (rx550) - amdgpu.

Если вставить сразу за get_global_id (перед проверкой памяти) чтото вроде:

for (i=0;i<n*1000;i++){   for (a=0;a<256;a++){ pix1[a]=pix2[a]; }   }

Тогда (они пишут память по очереди) количество ошибок сильно меньше. Операция "pix1[...." - чтоб не зарезал оптимизатор (компилятора).

Продолжение: https://yadi.sk/d/vjGHTbQR3Q9bgW
pfTDMF залит 0.0, [117] и [138] = 0.5.

На Polaris12 это работает:

while (a<e){
    if (pfTDMF[a]==0){ a++; continue; }
    b=a+1;
    while (b<e){
      if (pfTDMF[b]!=0.0f){ b++; continue; }
      break;
    }  
    if (a<e) pfUF[a]= 0.60f;
    if (b<e) pfUF[b]=-0.60f;    
    a=b;        
  }

Получается две точки.
А вот так:

FindStart:;  
  if (pfTDMF[a]==0.0f){//ucModeL[i] & 4 
    a+=uiStep;  
    if (a<e) goto FindStart;
    goto End; 
  }
//  c++; //<--- need for work :(  
  b = a+uiStep;   
  if (b>=e){
//    pfUF[a]=pfTDMF[a]/pfTDMB[a];
    goto End;
  }
FindEnd:;  
  if (pfTDMF[b]!=0.0f){
    b+=uiStep;  
    if (b<e) goto FindEnd;
  }  
if (a<e) pfUF[a]= 0.6f;
if (b<e) pfUF[b]=-0.6f;
if (b>=e) goto End; 
a = b; 
//  c++; //<--- need for work :(    Compilator error (?) 
goto FindStart;
End:  
  return;

не работает.
Если только один поток (127 к примеру) - работает - две точки.
Два и больше - не работает - второй точки нет.
Если же добавлю "с++;" (или что то другое) - работает.
Надеюсь косяк у меня. Но может быть кэш. Может нет команды ожидания.

Разбросал на отдельные ядра, не помогло. (https://yadi.sk/d/XwJXXd5e3QFsQC)
Прогонка, шаг 3 . Так как ожидается несколько слоев, хотел их паралельно запустить. Поэтому цикл.

Так работает:

if (n & 1 )    Work(n*256+80,64,1,pfTDMAX,pfTDMBX,pfTDMCX,pfTDMFX,pfAlphaX,pfBetaX,pfUFX,puiFlagsX);
  barrier(CLK_GLOBAL_MEM_FENCE);  
  if (!(n & 1) )   Work(n*256+80,64,1,pfTDMAX,pfTDMBX,pfTDMCX,pfTDMFX,pfAlphaX,pfBetaX,pfUFX,puiFlagsX);

n*256+80 - чтоб исключить выход за диапозон если вдруг. (n = get_global_id(0);).
Можно 0x1 или 0x4. Но 0х8 не работает тоже:
Так не работает:

Work(n*256,256,1,pfTDMAX,pfTDMBX,pfTDMCX,pfTDMFX,pfAlphaX,pfBetaX,pfUFX,puiFlagsX);

Похоже на то что часть потоков прекрашают работу. Возможно всетаки не могут закачать в модуль нужную память.
Ядрами читается [n*256+i] для i от 0 до 255.