Sto sviluppando un'applicazione OpenCL 1.2 che si occupa di immagini di grandi dimensioni. Al momento, l'immagine che sto testando è 16507x21244 pixel. Il mio kernel viene eseguito in un ciclo che opera su blocchi dell'immagine. Il kernel prende i pezzi 32bpp (rgba) dell'immagine e passa i pezzi di float4 pixel.Come determinare la dimensione massima delle immagini che posso tranquillamente passare al/dal kernel OpenCL?
Definiamo un lato di un blocco (quadrato) in pixel come dimensione del blocco. Vale a dire, un quadrato di pixel di 8192x8192 ha una dimensione del blocco 8192. Ovviamente, sul lato destro e sul lato inferiore abbiamo blocchi rettangolari più piccoli se l'immagine non è chiaramente divisibile per la dimensione del blocco. Il mio codice gestisce questo, ma per il resto di questo post, ignoriamolo per semplicità.
Sono cercando di determinare la dimensione massima del blocco su cui posso operare in ogni iterazione del mio ciclo, nonché la dimensione del blocco ottimale (che potrebbe non essere la dimensione massima del blocco).
Per riferimento, here is the information reported by the clinfo utility on my machine. Sto usando il mio kernel su Geforce GTX 560 Ti con la piattaforma Nvidia usando i loro driver linux proprietari.
La mia ipotesi iniziale ingenua era che potevo operare sulla dimensione massima dell'immagine 2d. Tuttavia, ciò restituisce clEnqueueNDRangeKernel restituendo un codice di errore di -4 (CL_MEM_OBJECT_ALLOCATION_FAILURE).
Pensandoci, questo ha senso per me. Con 1 GB di memoria video, ci si aspetterebbe di poter contenere una singola trama di 16384x16384 pixel (32 bpp) o una trama di 8192x8192 pixel (float4). Se entrambi devono essere memorizzate nella cache sulla carta mentre il kernel funziona, potremmo aspettarci di utilizzare la seguente quantità di memoria:
4 bytes-per-pixel * chunk size^2 (input image)
+ 16 bytes-per-pixel * chunk size^2 (output image)
= 1 GiB total video memory
Risolvendo per dimensione del blocco otteniamo
chunk size = sqrt(1GiB/20)
Collegando la quantità di memoria riportato da OpenCL (che è un po 'meno di 1 GiB - 1023 MiB) e pavimentazione il risultato, otteniamo:
floor(sqrt(1072889856/20)) = 7324
Tuttavia, una dimensione pezzo di 7324 risultati ancora in CL_MEM_OBJECT_ALLOCATION_FAILURE.
La mia prossima ipotesi è stata che non possiamo passare un'immagine più grande della dimensione di allocazione massima , che OpenCL riporta come 268222464 byte per la mia scheda. Poiché la mia immagine di output ha una larghezza di pixel maggiore, determinerebbe la mia dimensione del blocco.
floor(sqrt(268222464/16)) = 4094
Ehi, funziona davvero! Ora, se provassimo ad aumentare? Con mia sorpresa, non fallisce. Attraverso tentativi ed errori, ho ristretto il 6784 come dimensione massima effettiva del blocco. A 6785, inizia a lamentarsi con CL_MEM_OBJECT_ALLOCATION_FAILURE. Non so perché il max sembra essere 6784, e non so se questo è ripetibile o se il valore oscilla (come altri stati esistenti nella memoria video che influenzano quanto può contenere). Trovo anche che correndo con una dimensione del blocco di 6784 è inferiore di alcuni secondi rispetto alla dimensione basata sull'assegnazione massima. Mi chiedo se questo è dovuto al fatto che OpenCL ha bisogno di eseguire più (costose) allocazioni sotto il cofano?
Ho anche notato la "dimensione massima dell'argomento del kernel" che OpenCL è in grado di segnalare (
CL_DEVICE_MAX_PARAMETER_SIZE). Tuttavia, quel valore sembra fasullo.Se potessi passare solo 4096 byte, ciò mi limiterà a 16x16 pixel!
Quindi sono lasciati con due domande fondamentali:
- Come faccio a determinare la dimensione massima del blocco assoluto?
- Come determinare la dimensione del blocco più veloce? (C'è un metodo diverso da prova ed errore?)
Come una domanda bonus, ci sono buone risorse ho potuto girare a per le future domande di questa natura per quanto riguarda basso livello interazioni OpenCL-hardware ?
Infine, fornirò alcuni frammenti di codice per la peer-review; Sarei estremamente grato per qualsiasi critica costruttiva!
Come sempre, grazie in anticipo per qualsiasi aiuto!
Ho un piccolo suggerimento: aggiungi 'assert (err == CL_SUCCESS);' alla riga 112 di 'rgbtolab.d'. So che il buffer dovrebbe essere uguale a NULL se l'allocazione di memoria fallisce, ma forse questo è un bug nel driver di Nvidia. Per quanto riguarda 'CL_DEVICE_MAX_PARAMETER_SIZE', l'oggetto immagine passato al kernel è come un riferimento alla matrice che contiene l'immagine. Non conosco l'esatto modulo interno, cioè struct o pointer, ma è per questo che puoi superare i 16x16 pixel. – chippies
@chippies Se l'errore può verificarsi in un programma ben funzionante (un errore di memoria esaurita è uno di questi errori), non utilizzare assert, in quanto non saranno inclusi nelle versioni di rilascio. Usa il meccanismo di eccezione ('std.exception.enforce' è un modo conveniente per farlo). –
@chippies, grazie, volevo affermarlo ma apparentemente lo mancava! Per quanto riguarda CL_DEVICE_MAX_PARAMETER_SIZE, non so perché non ci ho pensato, heh; cancellando la sezione pertinente nel post. –