Come si usa memoria locale in OpenCL?

Question

Ho giocato con OpenCL di recente, e sono in grado di scrivere i kernel semplici che utilizzano la memoria solo globale. Ora mi piacerebbe iniziare a utilizzare memoria locale, ma io non riesco a capire come utilizzare get_local_size() e get_local_id() per calcolare un "pezzo" di uscita alla volta.

Per esempio, diciamo che ho voluto convertire esempio kernel OpenCL Ciao Mondo di Apple per qualcosa che gli usi della memoria locale. Come lo faresti? Ecco i sorgenti del kernel originale:

__kernel square(
    __global float *input,
    __global float *output,
    const unsigned int count)
{
    int i = get_global_id(0);
    if (i < count)
        output[i] = input[i] * input[i];
}

Se questo esempio non può essere facilmente convertita in qualcosa che mostra come fare uso della memoria locale, qualsiasi altro semplice esempio farà.

Answer 1

Controlla i campioni nei SDK NVIDIA o AMD, dovrebbero puntare nella giusta direzione. Matrix trasposizione userebbe memoria locale, per esempio.

Utilizzando il kernel squadratura, si potrebbe organizzare i dati in un buffer intermedio. Ricordarsi di passare nel parametro aggiuntivo.

__kernel square(
    __global float *input,
    __global float *output,
    __local float *temp,
    const unsigned int count)
{
    int gtid = get_global_id(0);
    int ltid = get_local_id(0);
    if (gtid < count)
    {
        temp[ltid] = input[gtid];
        // if the threads were reading data from other threads, then we would
        // want a barrier here to ensure the write completes before the read
        output[gtid] =  temp[ltid] * temp[ltid];
    }
}

Answer 2

C'è un'altra possibilità per fare questo, se la dimensione della memoria locale è costante. Senza l'utilizzo di un puntatore nella lista dei parametri kernel, il buffer locale può essere dichiarata all'interno del kernel semplicemente dichiarandolo __local:

__local float localBuffer[1024];

Questo elimina il codice a causa di chiamate clSetKernelArg meno.

Answer 3

In memoria locale OpenCL è destinato a condividere i dati su tutti gli elementi di lavoro in un gruppo di lavoro. E di solito richiede di fare una chiamata barriera prima che i dati della memoria locale possono essere utilizzati (ad esempio, un elemento di lavoro vuole leggere una memoria locale di dati che è stato scritto da altri elementi di lavoro). Barriera è costosa in hardware. Tenete a mente, memoria locale deve essere utilizzato per i dati ripetuti di lettura / scrittura. conflitto banca dovrebbe essere evitato il più possibile.

Se non si è attenti con memoria locale, si può finire con le prestazioni peggiori po 'di tempo rispetto all'utilizzo di memoria globale.