Perché questo codice opencl non è deterministico?

Question

Il seguente codice python utilizza PyOpenCL per riempire l'array a_plus_b con la somma degli elementi nell'array b (questo non è il mio obiettivo reale, ma è il codice più semplice che riesca a trovare che mostri ancora il problema).

import pyopencl as cl 
import numpy as np 
import numpy.linalg as la 

height = 50 
width = 32 

b = np.arange(width,dtype=np.int32) 

ctx = cl.create_some_context() 
queue = cl.CommandQueue(ctx) 

mf = cl.mem_flags 
b_buf = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=b) 
dest_buf = cl.Buffer(ctx, mf.WRITE_ONLY, height*4) 

prg = cl.Program(ctx, """ 
    __kernel void sum(__global const int *b, __global int *c) 
    { 
     int x = get_global_id(1); 
     int y; 
     c[x] = 0; 
     for(y=0;y<get_global_size(0);y++) { 
      c[x] += b[y]; 
     } 
    } 
    """).build() 

prg.sum(queue, (width,height), None, b_buf, dest_buf) 

a_plus_b = np.empty(height,dtype=np.int32) 
cl.enqueue_copy(queue, a_plus_b, dest_buf) 

print(np.sum(b)) 
print(a_plus_b) 
print(np.sum(a_plus_b-np.sum(b))) 

concede l'output:

496 
[496 496 496 496 496 496 496 496 496 496 496 496 496 496 496 496 496 496 
496 496 496 496 496 496 496 496 496 496 496 496 496 496 496 496 496 496 
496 496 496 496 496 496 496 496 496 496 496 496 496 496] 
0 

Tuttavia, se cambio larghezza da 32 a 33, la matrice non è più lo stesso elemento più e più volte.

528 
[555 557 555 559 560 528 560 560 528 528 528 528 528 528 528 528 528 528 
528 528 528 531 540 569 581 528 705 591 560 560 545 560 560 528 560 528 
528 528 528 528 528 528 528 528 528 528 528 532 533 535] 
752 

Infatti, ogni volta che il codice viene eseguito, produce un risultato diverso.

528 
[560 560 559 560 560 560 560 528 528 528 528 528 528 528 528 528 528 528 
528 528 528 560 528 514 565 553 621 650 560 560 560 560 560 528 528 528 
528 528 528 528 528 528 528 528 549 528 528 544 528 537] 
724 

Che cosa causa la differenza? Cosa non sono

Answer 1

Si eseguono articoli di lavoro LARGHEZZA x ALTEZZA. Per ogni valore di X nel tuo kernel, ci saranno gli elementi di lavoro di WIDTH che fanno esattamente la stessa cosa in parallelo: impostando C [X] a 0 e quindi aggiornandolo nel ciclo Y. Tutti questi articoli di lavoro di LARGHEZZA leggeranno C [X] e quindi lo aggiorneranno più o meno allo stesso tempo. Questo "più o meno" è la causa delle variazioni che osservate.

L'algoritmo è 1D, ed è necessario eseguire solo HEIGHT-item e passare WIDTH come argomento del kernel. Sostituisci C [X] con un registro "SUM" e fai un singolo C [X] = SUM alla fine.