Lisää GPU-ohjelmoinnista, Osa 2

Säikeiden määrä ja säieryhmän koko

• Säikeiden määrän valintaan vaikuttaa useampi ristiriitainen tavoite:
  • Säikeiden määrän tulisi olla tarpeeksi suuri, jotta prosessointielementtien liukuhihnat saadaan pidettyä työllistettynä
  • Toisaalta liian suuri määrä säikeitä aiheuttaa ylimääräistä overheadia
  • Kaikille säikeille ei riitä resursseja (rekisterit jne) kerralla, joten osa säikeistä joutuu jonoon. Onko näistä säikeistä mitään hyötyä?
  • Säikeiden määrän tulisi sopia yhteen ohjelmakoodin luonteen kanssa (rinnakkaisien tehtävien määrä vs säikeiden määrä)

• Säieryhmän koon valintaan vaikuttaa vastaavasti moni asia:
  • Säieryhmän koon tulisi olla warpin/wavefrontin moninkerta
  • Säieryhmän käynnistys aiheuttaa overheadia, joten suuri säieryhmän on parempi (vähemmän käynnistettäviä säieryhmiä)
  • Samaan säieryhmään kuuluvat säikeet voivat kommunikoida keskenään, joten suuri säieryhmä vähentää globaalin kommunikoinnin tarvetta
  • Säieryhmän sisäinen synkronointi aiheuttaa enemmän overheadia mikäli säieryhmän koko on suuri (enemmän säikeitä odottamassa esteellä)

• Lisää säieryhmän kokoon vaikuttavia asioita:
  • Samaan säieryhmään kuuluvien säikeiden varaamat resurssit (rekisterit, lokaali muisti) rajoittavat säieryhmän suurinta mahdollista kokoa
  • Säieryhmä, joka varaa paljon resursseja estää muita säieryhmiä käyttämästä laskentayksikköä samanaikaisesti. Tällöin säieryhmän koon valinta on erityisen tärkeätä, jotta prosessointiyksiköt saataisiin hyödynnettyä tehokkaasti.
  • Säiryhmien määrän tulisi sopia yhteen ohjelmakoodin luonteen kanssa (toisistaan riippumattomien rinnakkaisien tehtävien määrä vs säieryhmien määrä)

• OpenCL:stä löytyy CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE-lippu, jonka avulla ytimeltä voidaan kysyä luku, jonka moninkerrat soveltuisivat säieryhmän kooksi (linkki)
• CUDA:sta löytyy cudaOccupancyMaxPotentialBlockSize() ja cudaOccupancyMaxPotentialBlockSizeVariableSMem() aliohjelmat, jotka tarjoavat heuristiikkoja säieryhmän koon valintaan (linkki)
• GPU-valmistajien dokumentaatiosta löytyy hyviä vinkkejä!
• Erilaiset profilointityökalut GPU-valmistajilta saattavat olla myöskin hyödyksi

Lisää globaalin muistin käyttämisestä

Kaksiulotteisen taulukon käyttäminen

• Väärällä tavalla varatuilla kaksiulotteisilla taulukoilla saattaa olla merkittävä vaikutus ohjelmakoodin tehokkuuteen
• Edellisen rivin loppu ja rivin alku saattavat esimerkiksi kuulua samalle välimuistin riville:

• Ongelma voidaan ratkaista varaamalla ylimääräistä tilaa rivien päähän:

• Ratkaisu hukkaa jonkin verran muistia, mutta saavutettu nopeushyöty on usein sen arvoista
• Ylimääräisen tilan määrä voidaan valita siten, että jokainen taulukon rivi alkaa välimuistirivin alusta

OpenCL

• OpenCL ei tarjoa valmiita työkaluja kaksiulotteisen taulukon varaamiseen
• Yksinkertaisin lähestymistapa on varata taulukko siten, että jokaisen rivin pituus on warpin/wavefrontin moninkerta:

#define CEIL(x,y) (((x)+(y)-1)/(y))        // x/y pyöristettynä ylöspäin
#define ROUND_TO_NEXT(x,n) ((n)*CEIL(x,n)) // Pyöristää seuraavaan moninkertaan
#define WARPFRONT 32                       // Nvidian näytönohjain

cl_uint height= 125:
cl_uint width = 432;

// Lasketaan seuraava 32 moninkerta eli tässä tapauksessa 448
cl_uint ldf = ROUND_TO_NEXT(width,WARPFRONT);

// Varataan height x ldf int-taulukko
cl::Buffer deviceBuffer(context, CL_MEM_READ_WRITE, height*ldf*sizeof(int));

// Asetetaan ydin komentojonoon
kernel(cl::EnqueueArgs(...), deviceBuffer, height, width, ldf);

• Muuttuja ldf (rivin pituus alkioina) välitetään ytimelle, joka hyödyntää sitä taulukon käsittelyssä:

__kernel void kernel(__global int *buffer, uint height, uint width, uint ldf) {
    ...

    int value = buffer[i*ldf+j]; // Sama kuin buffer[i][j]

    ...
}

CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE

• Välimuistin rivin pituus voidaan myöskin kysyä cl::Device::getInfo-jäsenfunktiolla, jolloin taulukon rivin pituudeksi voidaan asettaa välimuistin rivin pituuden moninkerta:

#define CEIL(x,y) (((x)+(y)-1)/(y))        // x/y pyöristettynä ylöspäin
#define ROUND_TO_NEXT(x,n) ((n)*CEIL(x,n)) // Pyöristää seuraavaan moninkertaan

cl_uint cacheLine;
device.getInfo(CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE, &cacheLine);

if(cacheLine < 1)
    cacheLine = 1;

cl_uint height = 125:
cl_uint width = 432;

// Lasketaan seuraava välimuistin rivin pituuden moninkerta
cl_uint ldf = ROUND_TO_NEXT(width*sizeof(int), cacheLine);

// Varataan taulukko, jossa jokaisen rivin pituus on ldf tavua
cl::Buffer deviceBuffer(context, CL_MEM_READ_WRITE, height*ldf);

kernel(cl::EnqueueArgs(...), deviceBuffer, height, width, ldf);

• Ytimen lähdekoodissa tulee huomioida, että muuttuja ldf välittää nyt tiedon rivin pituudesta tavuina. Huomaa myös, että buffer-taulukko välitetään char-tyyppisenä.

__kernel void kernel(__global char *buffer, uint height, uint width, uint ldf) {
    ...
    
    // Lasketaan rivi; buffer[i][*]
    __global int *line = (__global int*)(buffer + i*ldf);
    
    int value1 = line[j];                       // Sama kuin buffer[i][j]
    
    int value2 = *((int*)(buffer + i*ldf) + j); // Sama kuin buffer[i][j]

    ...
}

CUDA

• CUDA tarjoaa cudaMallocPitch-aliohjelman kaksiulotteisien taulukoiden varaamiseen:

cudaError_t cudaMallocPitch(
    void   **devPtr, // Osoitin taulukon alkuun
    size_t  *ldf,    // Optimaalinen arvo taulukon rivin pituudeksi tavuina
    size_t   height, // Taulukon korkeus alkioina
    size_t   width   // Taulukon leveys _tavuina_
    );

size_t height = 125:
size_t width = 432;

size_t ldf;
char* deviceBuffer;

// Varataan taulukko, jossa jokaisen rivin pituus on ldf tavua
cudaMallocPitch(&deviceBuffer, &ldf, width * sizeof(int), height);

kernel<<<...>>>(deviceBuffer, height, width, ldf);

• Siirto isäntälaitteen muistista globaaliin muistiin onnistuu kätevästä cudaMemcpy2DAsync-aliohjelmalla:

int *hostBuffer;
cudaMallocHost(&hostBuffer, height * width * sizeof(int));

// hostBuffer täytetään datalla tässä välillä...

err = cudaMemcpy2DAsync(
    deviceBuffer,           // Kohde
    ldf,                    // Riviä varten varattu tila tavuina kohteessa
    hostBuffer,             // Lähde
    width * sizeof(int),    // Riviä varten varattu tila tavuina lähteessä
    width * sizeof(int),    // Taulukon leveys tavuina
    height,                 // Taulukon korkeus
    cudaMemcpyHostToDevice, // Siirron suunta
    0);                     // Komentojono

• Ytimen lähdekoodissa tulee huomioida, että muuttuja ldf välittää tiedon rivin pituudesta tavuina:

__global__ void kernel(char *buffer, size_t height, size_t width, size_t ldf) {
    ...
    
    int *line = (int*)(buffer + i*ldf);          // Lasketaan rivin
    int value1 = line[j];                        // Sama kuin buffer[i][j]
    
    int value2 = *((int*)(buffer + i*ldf) + j);  // Sama kuin buffer[i][j]

    ...
}

Kolmiulotteiset taulukot

• CUDA tarjoaa cudaMalloc3D-aliohjelman kolmiulotteisien taulukoiden varaamiseen:

cudaError_t cudaMalloc3D(
    struct cudaPitchedPtr *pitchedDevPtr, // Varattu 3-ulotteinen taulukko
    struct cudaExtent      extent         // Tieto taulukon koosta
    );

struct cudaExtent make_cudaExtent(
    size_t w,    // Taulukon leveys tavuina
    size_t h,    // Taulukon korkeus alkioina
    size_t d);	 // Taulukon syvyys alkioina

struct cudaPitchedPtr {
    size_t 	pitch; // Yhdelle taulukon riville varattu tila tavuina
    void * 	ptr;   // Osoitin taulukon alkuun
    size_t 	xsize; 
    size_t 	ysize;
};

struct cudaExtent {
    size_t 	depth;
    size_t 	height;
    size_t 	width;
};

• Seuraava esimerkki varaa 54 \(\times\) 23 \(\times\) 5 -taulukon ja kutsuu ydintä:

size_t width = 54, height = 23, depth = 5;

// Luodaan cudaExtent structi 3D-taulukon varaamista varten
cudaExtent extent = make_cudaExtent(width * sizeof(int), height, depth);

// Varataan 3D-taulukko edellä luodun cudaExtentin perusteella
cudaPitchedPtr pitchPointer;
cudaMalloc3D(&pitchPointer, extent);

// Kutsutaan ydintä
kernel<<<...>>>(pitchPointer, width, height, depth);

__global__ void kernel(
    cudaPitchedPtr pitchPointer, size_t width, size_t height, size_t depth) {
    
    ...
    
    // Lasketaan rivin (i,j,*) alku
    int *line = (int*)(
        (char*)pitchPointer.ptr + (i*height + j)*pitchPointer.pitch);
    
    int value = line[k]; // Sama kuin buffer[i][j][k]
}

Huomautus

• Huomaa, että 2- ja 3-ulotteisia taulukoita voidaan edelleen käsitellä yksiulotteisena taulukkona
• Säieryhmien indeksointia voidaan hyödyntää taulukoiden käsittelyssä:

__kernel void (__global float *A, uint height, uint width) {
    const int i = get_group_id(1);  // height säieryhmää suuntaan 1
    const int j = get_global_id(0); // Vähintään width säiettä suuntaan 0
    
    float value = 0.0;
    if(j < width)
        value = A[i][j]; // Jokainen säie lukee yhden luvun taulukosta 
    
    ...
}

Muistipuskurin jakaminen alipuskureihin

• CUDA:n GPU:n muistiin osoittavat osoittimen käyttäytyvät normaalien osoittimien tavoin
• Niiden kanssa voidaan siis harrastaa osoitinaritmetiikkaa alipuskureiden luomiseksi:

int *buffer;
cudaMalloc((void **)&buffer, N*sizeof(int));

// Alipuskuri, joka alkaa indeksista N/3
int *subBuffer = buffer + N/3;

kernel<<<...>>>(buffer, ...);    // Suorittaa operaation koko buffer-taulukolle
kernel<<<...>>>(subBuffer, ...); // Suorittaa operaation pelkästään osalle
                                 // buff-taulukkoa

OpenCL

• OpenCL:ssä alipuskurit täytyy luoda cl::Buffer-luokan createSubBuffer-jäsenfunktiolla:

cl::Buffer cl::Buffer::createSubBuffer(
    cl_mem_flags           flags,
    cl_buffer_create_type  buffer_create_type,
    const void            *buffer_create_info,
    cl_int                *err = NULL)

• Esimerkki:

cl::Buffer buffer(context, CL_MEM_READ_WRITE, N*sizeof(int));

cl_buffer_region region;
region.origin = N/3;   // Alipuskuri alkaa indeksistä N/3
region.size = N/3;     // Alipuskurin koko on N/3

cl::Buffer subBuffer = buffer.createSubBuffer(
    CL_MEM_READ_WRITE, CL_BUFFER_CREATE_TYPE_REGION, &region);

Siirrot RAM <-> VRAM

• Erityyppisien muistisiirtojen nopeudet:
  • Lokaali muisti: \(\mathcal{O}(10^{12})\) B/s
  • Globaali muisti: \(\mathcal{O}(10^{11})\) B/s
  • RAM-muisti: \(\mathcal{O}(10^{10})\) B/s
  • RAM-VRAM-siirto: \(\mathcal{O}(10^9)\) B/s
• Datan siirtäminen RAM-muistista GPU:n globaaliin muistiin on siis erittäin kallis operaatio ja sitä tulisi välttää!

• Yksinkertaisien laskentaoperaatioiden tapauksessa (esim. axpy) suurin osa ajasta saattaa kulua siirtoon RAM-muistin ja globaalin muistin välillä:

• Myöskin datan siirtäminen laskennan välissä hidastaa laskentaa merkittävästi:

• Pyri siis
  • siirtämään kaikki tarvittava data kerralla globaaliin muistiin,
  • laskemaan mahdollisimman paljon ja
  • siirtämään tulos vasta lopuksi takaisin RAM-muistiin:

Page-locked / pinned / kiinnitetty muisti

• Page-locked / pinned -tyyppisen eli kiinnitetyn muistin varaaminen nopeuttaa datasiirtoa huomattavasti, koska PCIe-väylän DMA-ohjain voi toimia tällöin tehokkaammin!
• Muistutus: CUDA:ssa kiinnitettyä muistia voidaan varata ja vapauttaa cudaMallocHost ja cudaFreeHost aliohjelmilla:

cudaError_t cudaMallocHost(void **ptr, size_t size);	

cudaError_t cudaFreeHost(void *ptr);	

• OpenCL-määritykset eivät aseta vaatimuksia kiinnitetyn muistin käytölle
• AMD:n toteutus toteuttaa suuret datasiirrot kiinnittämällä isäntälaitteen puoleisen muistialueen paloittain siirron aikana
• NVidia:n OpenCL-toteutuksen dokumentaatio on vajanaisempi, mutta se toimii todennäköisesti samalla tavalla

Muistin mappaaminen

• Muistiavaruuksien ns. "mappaaminen" yksinkertaistaa ongelmakoodia ja parantaa suorituskykyä joissakin tapauksissa
• Käytännössä CPU ja GPU jakavat (ainakin hetkellisesti) saman muistialueen
• Mappaaminen mahdollistaa periaatteessa sen, että data siirretään vasta sitten kun sitä tarvitaan

Periaate

• Ennen muistin mappausta CPU:n ja GPU:n muistiavaruuden ovat fyysisesti ja loogisesti erilliset:

• Muistin mappaaminen yhdistää CPU:n ja GPU:n muistiavaruudet loogisella tasolla:

• CPU voi kirjoittaa laskentaan käytettävän datan mapattuun muistialueeseen:

• GPU näkee mapatun muistialueen sisällön ytimen käynnistyttyä:

• GPU voi tallentaa laskennan tuloksen mapattuun muistialueeseen:

• CPU näkee näkee mapatun muistialueen sisällön kun ytimen suoritus valmistuu:

Käytäntö

• Käytännössä laskentaan käytettävä data saatetaan siirtää fyysisesti globaaliin muistiin ennen kuin ydin käynnistyy
• Periaatteessa on myöskin mahdollista, että data siirretään fyysisesti vasta kun GPU tarvitsee sitä laskennassa:

• Vastaavasti laskennan tulos voidaan siirtää fyysisesti RAM-muistiin kerralla tai paloissa silloin kun CPU tarvitsee osaa datasta:

OpenCL

• OpenCL-tarjoaa kolme erilaista tapaa käyttää mapattua muistia:
  • CL_MEM_ALLOC_HOST_PTR annettu cl::Buffer-olion luonnin yhteydessä
  • CL_MEM_USE_HOST_PTR annettu cl::Buffer-olion luonnin yhteydessä
  • cl::Buffer-olio luotu ilman ylimääräisiä lippuja
• AMD laajentaa tarjontaa omalla CL_MEM_USE_PERSISTENT_MEM_AMD-lipullaan

CL_MEM_ALLOC_HOST_PTR

• CL_MEM_ALLOC_HOST_PTR-lipun antaminen cl::Buffer-luokan muodostinfunktiolle kertoo OpenCL-toteutukselle, että OpenCL-toteutuksen tuli varata muisti siten, että isäntälaitteella on pääsy siihen:

cl::Buffer deviceBuffer = cl::Buffer(context, 
    CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, N*sizeof(int));

• AMD:n OpenCL-toteutuksessa mappaukseen liittyvä data sijaitsee fyysisesti kiinnitetyssä RAM-muistissa

• Muistin mappaaminen tapahtuu komentojonon enqueueMapBuffer-jäsenfunktiolla:

void * cl::CommandQueue::enqueueMapBuffer(
    const Buffer&              buffer,
    cl_bool                    blocking_map,
    cl_map_flags               flags,
    ::size_t                   offset,
    ::size_t                   size,
    const VECTOR_CLASS<Event> *events = NULL,
    Event                     *event = NULL,
    cl_int                    *err = NULL);

• Paluuarvona palautetaan mapatun muistialueen alun osoite isäntäohjelman muistiavaruudessa
• blocking_map-lipun asettaminen arvoon CL_TRUE tekee mappaamiseoperaatiosta blockkaavan
• flags-lippu määrittää onko muistipuskuri tarkoitus mapata luettavaksi (CL_MAP_READ) vai kirjoitettavaksi (CL_MAP_WRITE)

• Mapataan seuraavaksi deviceBuffer osaksi isäntälaitteen muistia blockkaavalla operaatiolla siten, että CPU voi kirjoittaa mapattuun muistialueeseen:

int *hostPointer = 0;
hostPointer = (int*) queue.enqueueMapBuffer(
    deviceBuffer, CL_TRUE, CL_MAP_WRITE, 0, N*sizeof(int));

• Isäntäohjelma voi nyt kirjoittaa mapattuun muistialueeseen (hostPointer)

for(int i = 0; i < N; i++)
    hostPointer[i] = checkBuffer[i];

• Muistialueen mappaus täytyy purkaa ennen kuin ydin voi käyttää mappauksen avulla jaettua dataa
• Purkaminen tapahtuu komentojonon enqueueUnmapMemObject-jäsenfunktiolla:

cl_int cl::CommandQueue::enqueueUnmapMemObject(
    const Memory              &memory,
    void                      *mapped_ptr,
    const VECTOR_CLASS<Event> *events = NULL,
    Event                     *event = NULL)

• Puramme siis aikaisemmin luovamme mappauksen, jossa deviceBuffer mapattiin isäntäohjelman muistiavaruudessa sijaitsevaan muistialueseen (hostPointer):

queue.enqueueUnmapMemObject(deviceBuffer, hostPointer);

• Mappauksen purkamisen jälkeen voimme kutsua ydintä normaalisti:

kernel(cl::EnqueueArgs(...), deviceBuffer, ...);

• Seuraavaksi asetamme komentojonoon blockkaavat käskyn mapata deviceBuffer osaksi isäntäohjelman muistiavaruutta siten, että isäntäohjelma voi lukea mapatusta muistialueesta:

hostPointer = (int*) queue.enqueueMapBuffer(
    deviceBuffer, CL_TRUE, CL_MAP_READ, 0, N*sizeof(int));

• Mappauksen jälkeen isäntäohjelma voi esimerkiksi tarkistaa GPU:n suorittaman tuloksen oikeellisuuden:

int rightValues = 0;
for(int i = 0; i < N; i++)
    if(hostPointer[i] == checkBuffer[i]+1)
        rightValues++;

• Lopuksi meidän täytyy vielä muistaa purkaa mappaus:

queue.enqueueUnmapMemObject(deviceBuffer, hostBuffer);

CL_MEM_USE_HOST_PTR

• Mappaukseen käytettävä isäntälaitteen puoleinen muistialue voidaan myöskin varata etukäteen:

int *hostBuffer = new int[N];

• Tällöin hostBuffer täytyy olla varattu siten, että muistialueen alun osoite on OpenCL-laitteen puolella käytettävän tietotyypin moninkerta!
• new-operaattorin palauttama osoitin on usein valmiiksi sopiva perustietotyyppien tapauksessa. Muissa tapauksissa joudutaan ehkä soveltamaan hieman osoitinaritmetiikkaa.

• Varattu muistipuskuri voidaan nyt antaa OpenCL-toteutukselle cl::Buffer-olion luomisen yhteydessä:

cl::Buffer deviceBuffer = cl::Buffer(context, 
    CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR, N*sizeof(int), hostBuffer);

• CL_MEM_USE_HOST_PTR-lippu kertoo, että OpenCL-toteutus voi käyttää hostBuffer-osoittimen viittaamaa muistialuetta vapaasti mappausoperatioissa. Puskuria ei voi käyttää enää muihin tarkoituksiin!

• Komentojonon enqueueMapBuffer-aliohjelma palauttaa käytännössä cl::Buffer-olion luonnin yhteydessä annetun osoitteen. OpenCL-määritykset eivät kuitenkaan takaa tätä, joten mapatun muistialueen osoite kannattaa tallentaa aina erilliseen osoitinmuuttujaan:

int *hostPointer = 0;
hostPointer = (int*) queue.enqueueMapBuffer(
    deviceBuffer, CL_TRUE, CL_MAP_WRITE, 0, N*sizeof(int));

cl::Buffer-olio luotu ilman ylimääräisiä lippuja

• Mappaukseen käytetty cl::Buffer-olio voidaan luoda myöskin ilman CL_MEM_ALLOC_HOST_PTR ja CL_MEM_USE_HOST_PTR lippuja
• AMD:n OpenCL-toteutuksessa mappauksen purkaminen käynnistää siirtooperaation RAM-muistin ja globaalin muistin välillä

CUDA

• Muistin mappaaminen vaatii CPU:n ja GPU:n suorituksen synkronointia

• Seuraavassa esimerkissä synkronointi on toteutettu komentojonon avulla:

cudaStream_t queue;
cudaStreamCreate(&queue);

• Synkronoinnin voisi hoitaa myöskin ilman komentojonoa cudaDeviceSynchronize()-aliohjelmalla

• Varataan aluksi kiinnitettyä muistia isäntäohjelman muistiavaruudesta:

int *hostPointer;
cudaMallocHost(&hostPointer, N*sizeof(int), cudaHostAllocMapped);

• cudaHostAllocMapped-lippu kertoo CUDA:lle, että varattu muistialue näkyy CUDA-laitteelle

• Isäntäohjelma voi kirjoittaa laskentaan käytettävän datan varattuun muistialueeseen normaalisti:

for(int i = 0; i < N; i++)
    hostPointer[i] = (i & 0x1) * i;

• Synkronointi ei ole tässä vaiheessa pakollista kunhan ydintä ei laiteta komentojonoon ennen kuin isäntäohjelma on kirjoittanut kaiken vaadittavan datan mapattuun muistialueeseen

• Vanhempien GPU:iden tapauksessa (CC < 2.0, ei UVA-tukea) meidän täytyy ensin hankkia tietoomme GPU:n puolelta varatun mapattavat muistialueen osoite:

int *devicePointer;
cudaHostGetDevicePointer(&devicePointer, hostPointer, 0);

• Ydin voidaan nyt asettaa komentojonoon joka käyttäen edellä hankittua devicePointer-osoitinta:

kernel<<<..., queue>>>(devicePointer, ...);

• Uudenpien GPU:den tapauksessa voimme käyttää suoraan hostPointer-osoitinta:

kernel<<<..., queue>>>(hostPointer, ...);

• Seuraavaksi isäntäohjelman täytyy synkronoida suorituksensa GPU:n kanssa:

cudaStreamSynchronize(queue);

• Muussa tapauksessa hostPointer ei olisi vielä välttämättä valmis käytettäväksi

• Synkronoinnin jälkeen isäntäohjelma voi käyttää hostPointer-osoitinta normaalisti:

int rightValues = 0;
for(int i = 0; i < N; i++)
    if(hostPointer[i] == (i & 0x1) * i + 1)
        rightValues++;

Useamman komentojonon käyttäminen

• Osa RAM-muistin ja GPU:n muistin välisiin siirtoihin liittyvästä ajasta voidaan peittää käyttäen useampaa komentojonoa:

• Tällöin GPU:n prosessointielementit voivat tehdä laskentaa datalla yhdessä komentojonossa ja DMA-ohjain voi siirtää dataa toisessa komentojonossa