memoria fisica su dispositivi AMD: locale vs privato

Sto scrivendo un algoritmo in OpenCL in cui avrei bisogno di ogni unità di lavoro per ricordare una buona parte dei dati, diciamo qualcosa tra uno long[70] e uno long[200] o giù di lì per kernel.memoria fisica su dispositivi AMD: locale vs privato

I dispositivi AMD recenti hanno 32 memoria KiB __local, che è (per la quantità data di dati per kernel) sufficiente per memorizzare le informazioni per 20-58 unità di lavoro. Tuttavia, da quello che ho capito dall'architettura (e in particolare da this drawing), ogni core dello shader ha anche una quantità dedicata di memoria privata. Tuttavia non riesco a trovare la sua dimensione.

Qualcuno può dirmi come scoprire quanta memoria privata ha ogni kernel?

Sono particolarmente curioso dell'HD7970, dal momento che ho in programma di acquistarne presto alcuni.

Edit: Problema risolto, la risposta è here nell'Appendice D.

fonte

2012-02-17 user1111929

Non credo che la memoria privata sia dedicata per core: si associa al file di registro, che è per risorsa dell'unità di calcolo. Ogni elemento di lavoro ottiene i registri allocati dal file di registro dell'unità di calcolo, quanti sono necessari per determinare il numero di fronti d'onda in volo in un dato istante. – talonmies

Dal famoso disegno in tutto il mondo http://www.codeproject.com/KB/showcase/Memory-Spaces/image001.jpg Ho concluso che la memoria privata è fisicamente diversa dalla memoria __locale, no? – user1111929

Sì, sono fisicamente diversi. La memoria privata si associa al file di registro dell'unità di calcolo, la memoria locale per calcolare la memoria condivisa a livello di unità nella maggior parte dei dispositivi AMD moderni. Alcune prime GPU OpenCL compatibili non avevano memoria condivisa, e la memoria locale era solo SDRAM. Nessuno dei due è per core e quanto si utilizza per articolo di lavoro per gruppo privato e per gruppo di lavoro per effetti locali il numero di fronti d'onda simultanei in esecuzione per unità di calcolo. – talonmies

La risposta è stata data dai commenti degli utenti nei commenti, quindi scriverò una nuova risposta qui per chiudere la domanda.

Questi valori sono disponibili nell'Appendice D dell'API AMD Guida alla programmazione OpenCL http://developer.amd.com/sdks/amdappsdk/assets/amd_accelerated_parallel_processing_opencl_programming_guide.pdf (esiste un documento simile per nVidia). Apparentemente un registro è a 128 bit (4x32) per dispositivi AMD e ci sono 16384 registri per tutti i moderni dispositivi di fascia alta, quindi si tratta di un notevole 256KB per unità di calcolo.

fonte

2012-03-01 12:16:47 user1111929

Penso che si sta cercando per la memoria __local. Questo è il riferimento a 32 KB di memoria locale. Non penso che tu possa eseguire il polling del dispositivo per ottenere la quantità di memoria privata.

È possibile passare un riferimento NULL lungo * cl_mem per allocare la memoria. Penso che sia meglio usare una quantità di memoria statica per WI. Supponendo che sia richiesto un tempo [200] per ogni oggetto di lavoro, si utilizzerà il codice seguente. Sarebbe anche una buona idea dividere il lavoro in gruppi che hanno le stesse (o simili) esigenze di memoria, al fine di ottenere il massimo dalla memoria LDS.

void __kernel(__local long* localMem, const int localMemPerItem 
     //more args... 
     ) 
{ 
    //host has 'passed' localMemPerItem*get_local_size() long values in as locamMem 
    //this work item has access to all of it, but can choose to restrict 
    //itself to only the portion it needs. 
    //work group size will be limited to CL_DEVICE_LOCAL_MEM_SIZE/(8*localMemPerItem) 
    int startIndex=localMemPerItem*get_local_id(0); 
    //use localMem[startIndex+ ... ] 
}

fonte

2012-02-17 18:41:18 mfa

Non è possibile eseguire il polling, ma esiste? Dal famoso disegno in tutto il mondo http://www.codeproject.com/KB/showcase/Memory-Spaces/image001.jpg ho ipotizzato che ci sia un insieme fisicamente separato di registri privati su ciascuna unità di lavoro. No? Speravo in qualche modo di fare meglio di una limitazione CL_DEVICE_LOCAL_MEM_SIZE/(8 * localMemPerItem), poiché lascia approssimativamente metà dei core inutilizzati. L'accesso alla memoria globale sarebbe probabilmente troppo lento, anche se è solo l'incremento di un contatore. – user1111929

Ho trovato alcune ulteriori informazioni sulle dimensioni dei registri di cipresso, cayman e fermi qui: http://www.realworldtech.com/page.cfm?ArticleID=RWT121410213827&p=11 Dovresti essere in grado di modificare alcune vars private di dimensioni adeguate in quelle dimensioni . Penso che la LDS sarà comunque la scelta migliore. – mfa

Per rispondere a quanto grande è registrare file in una scheda della serie 79xx, sin dal suo basata su architettura GCN è 64KB come da immagine in AnandTech: http://www.anandtech.com/print/5261

Per rispondere alla tua domanda come per scoprire come molta memoria ogni kernel usa .. puoi guardare eseguire AMD APP Profiler sul tuo kernel, ti dirà nella sezione di occupazione del kernel quanto spazio è utilizzato dal kernel.

fonte

2012-02-20 14:50:24 kiranputtur

Oh davvero? Quello è strano. Ho pensato di aver trovato la risposta, ma è diversa. Nella guida alla programmazione AMD OpenCL http://developer.amd.com/sdks/amdappsdk/assets/amd_accelerated_parallel_processing_opencl_programming_guide.pdf nell'Appendice D, c'è la dimensione totale del file di registro, ed è elencata come 256 KB per tutti i dispositivi moderni. Quale è corretto ora? : S – user1111929

Credo che entrambi siano corretti. A quanto ho capito, nell'architettura GCN, un'unità SIMD ha un file di registro da 64 kb, e ci sono 4 unità SIMD per unità di calcolo, vale a dire. 4 * 64kb = 256kb del file di registro totale per unità di calcolo. – talonmies

memoria fisica su dispositivi AMD: locale vs privato

risposta

Problemi correlati