Come utilizzare async_work_group_copy in OpenCL?

Question

Vorrei capire come utilizzare correttamente la chiamata async_work_group_copy() in OpenCL. Diamo uno sguardo su un esempio semplificato:.

__kernel void test(__global float *x) { 
    __local xcopy[GROUP_SIZE]; 

    int globalid = get_global_id(0); 
    int localid = get_local_id(0); 
    event_t e = async_work_group_copy(xcopy, x+globalid-localid, GROUP_SIZE, 0); 
    wait_group_events(1, &e); 
} 

Il riferimento http://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/async_work_group_copy.html dice "Eseguire una copia asincrona di num_elements gentype elementi da src a dst La copia asincrona viene eseguita da tutti i work-elementi di un gruppo di lavoro e questa funzione built-in deve quindi essere incontrata da tutti gli elementi di lavoro in un gruppo di lavoro che esegue il kernel con gli stessi valori di argomento, altrimenti i risultati non sono definiti. "

Ma questo non chiarisce le mie domande ...

Vorrei sapere, se le seguenti ipotesi sono corrette:

1. La chiamata a async_work_group_copy() devono essere eseguite da tutti i lavori -Esso nel gruppo.
2. La chiamata deve essere in un modo, che l'indirizzo di origine sia identico per tutti gli elementi di lavoro e punti al primo elemento dell'area di memoria da copiare.
3. Poiché il mio indirizzo di origine è relativo in base all'ID dell'oggetto di lavoro globale del primo elemento di lavoro nel gruppo di lavoro. Quindi devo sottrarre l'id locale per avere l'indirizzo identico per tutti gli articoli di lavoro ...
4. Il terzo parametro è davvero il numero di elementi (non la dimensione in byte)?

domande bonus:

a. Posso semplicemente usare la barriera (CLK_LOCAL_MEM_FENCE) invece di wait_group_events() e ignorare il valore restituito? Se è così, sarebbe probabilmente più veloce?

b. Anche una copia locale ha senso per l'elaborazione su CPU o è un overhead perché condividono comunque una cache?

saluti, Stefan

Answer 1

Uno dei motivi principali di questa funzione esistente è quello di permettere al compilatore conducente/kernel per copiare in modo efficiente la memoria senza lo sviluppatore dover fare ipotesi circa l'hardware.

Si descrive quale memoria è necessario copiare come se si trattasse di una copia a thread singolo e async_work_group_copy viene eseguita per l'utente utilizzando l'hardware parallelo.

Per le vostre domande specifiche:

1. Non ho mai visto async_work_group_copy usato da solo alcuni degli elementi di lavoro in un gruppo. Ho sempre pensato che fosse perché era necessario. Penso che la natura di blocco di wait_group_events costringa tutti gli elementi di lavoro a far parte della copia.

2. Sì. Gli indirizzi di origine (e destinazione) devono essere gli stessi per tutti gli elementi di lavoro.

3. È possibile sottrarre il proprio ID locale per ottenere l'indirizzo corretto, ma trovo che basare l'indirizzo su groupId risolva anche questo problema. (get_group_id)

4. Sì. L'ultimo parametro è il numero di elementi, non la dimensione in byte.

a. No. In base all'evento, scoprirai che la tua barriera viene colpita quasi immediatamente dagli elementi di lavoro e che i dati non verranno necessariamente copiati. Questo ha senso perché alcuni componenti hardware opzionali potrebbero non utilizzare nemmeno le unità di calcolo per eseguire l'operazione di copia effettiva.

b. Penso che le implementazioni della cpu opencl potrebbero garantire l'utilizzo della cache L1 quando si utilizza la memoria locale. L'unico modo per sapere con certezza se questo si comporta meglio è quello di confrontare l'applicazione con varie impostazioni.