Ensimmäinen kosketus GPU-ohjelmointiin

OpenCL vs CUDA

Markkinoilla on tällä hetkellä kaksi rajapintaa GPU-ohjelmointiin: OpenCL ja CUDA. Valmistajien tuki eri rajapinnoille vaihtelee:

alusta \ valmistaja	Nvidia	AMD	Intel
OpenCL 1.2	Tuettu	Tuettu	Tuettu
OpenCL 2.0	Ei tuettu	Tuettu	Rajoitettu tuki
CUDA	Tuettu	Ei tuettu	Ei tuettu

GPU-ohjelmointi on myöskin mahdollista OpenACC:n kaltaisien rajapintojen avulla

Yleiskuva GPU-laitteistosta

OpenCL-terminologiassa tietokonelaitteisto jaetaan isäntälaiteeseen (host) ja OpenCL-laitteisiin (OpenCL devices). Kurssin aikana tullaan myöskin puhumaan CUDA-laitteista.

Isäntälaite

Isäntälaite on useimmissa tapauksissa tietokoneen keskusprosessori ja tämän kurssin puitteissa OpenCL-laitteiden oletetaan olevan näytönohjaimia
Laskenta on tarkoitus suorittaa pääasiassa OpenCL-laitteilla ja isäntälaitteen tehtäväksi jää laitteistoon asennettujen OpenCL-laitteiden hallinta
OpenCL-sovelluksen suorittaminen alkaa isäntälaiteella normaalin sovelluksen tavoin, mutta OpenCL-rajapinta mahdollistaa OpenCL-laitteiden resurssien hyödyntämisen
Isäntälaite voi esimerkiksi varata osan OpenCL-laiteen muistista käyttöönsä ja käskeä OpenCL-laitteita suorittamaan erinäisiä toimintoja

OpenCL-laite

Kukin OpenCL-laite sisältää yhden tai useamman laskentayksikön (computing unit) ja yksittäinen laskentayksikkö sisältää yhden tai useamman prosessointielementin (processing element)

Laskentayksiköt ja prosessointielementit

Samaan laskentayksikköön kuuluvat prosessointielementit ovat tyypillisesti läheisesti kytköksissä toisiinsa ja jakavat resursseja keskenään
Laskenta tapahtuu prosessointielementeillä, mutta yksittäistä prosessointielementtiä ei välttämättä voida suoraan rinnastaa perinteiseen prosessoriytimeen

Laskentayksiköt ja prosessointielementit (jatkuu)

OpenCL-määrityksien puolesta samaan laskentayksikköön kuuluvat prosessointielementit suorittavat ohjelmakoodia joko yhdessä SIMD-yksikkönä (eli kukin laskentayksikkö on vektoriprosessori) tai erikseen SPMD-yksikköinä (eli erillisinä ytiminä)

Varsinainen laitteistotason toteutus ei kuitenkaan välttämättä täysin sovi kumpaakaan kategoriaan. Tähän asiaan palataan myöhemmin kun puhutaan Nvidian ja AMD:n toteutustason ratkaisuista

OpenCL-sovellus

OpenCL-sovellus käyttäytyy isäntälaitteella normaalin sovelluksen tavoin, mutta OpenCL-rajanpinta mahdollistaa myöskin OpenCL-laitteiden laskentaresurssien hyödyntämisen

OpenCL-sovellus (jatkuu)

Tarkemmin sanottuna OpenCL-sovellus jakautuu
- ytimiin (kernels), jotka suoritetaan OpenCL-laitteilla ja
- isäntäohjelmaan (host program), joka suoritetaan isäntälaitteella

Ydin

Ydin voidaan mieltää aliohjelmaksi, jota voidaan tarvittaessa kutsua laskentaoperaatioiden suorittamiseksi OpenCL-laitteella
Isäntäohjelma vastaa ytimien luomisesta, käynnistämisestä ja hallinnoimisesta

Globaali indeksiavaruus

Isäntäohjelman tulee määrittää erityinen indeksiavaruus ennen jokaista ytimen käynnistystä
Ytimen käynnistyttyä tämä indeksiavaruus määrittelee globaalit indeksinumerot joukolle säikeitä (work-items)

Säikeiden suorituspolut

Kukin säie aloittaa ohjelmakoodin suorittamisen ytimen alusta, mutta indeksinumerointi mahdollistaa sen, että eri säikeiden suorituspolut pääsevät haarautumaan:

Säieryhmät

Säikeet jaetaan isäntäohjelman määrittämällä tavalla säieryhmiin (work-groups) ja kukin säieryhmä saa oman säieryhmäindeksinumeronsa
Tämän lisäksi kukin säie saa oman lokaalin indeksinumeron säieryhmän sisällä.
Indeksointi voi olla yksi-, kaksi- tai kolmiulotteinen

Esimerkki säikeiden indeksoimisesta

Alla esimerkki tilanteessa, jossa 16 säiettä on jaettu neljään säieryhmään, joissa kussakin on neljä säiettä:

Säieryhmät ja osatehtävät

Säikeiden jakaminen säieryhmiin ohjaa ohjelmoijan pilkkomaan tehtävän osatehtäviin, joka voidaan ratkaista toisistaan riippumattomasti rinnakkain:

Säiryhmät ja laskentayksiköt

Vain ajonaikaisen järjestelmä tulee tietää OpenCL-laitteen todellisen rakenteen:

Muistimalli

Yksittäisellä säikeellä on käytettävissään neljä erilaista muistialuetta:

Ohjelmoijan tulee määritellä missä muistialueessa viitattu data sijaitsee. OpenCL 2.0 helpottaa tätä jonkin verran.

Globaali muisti (`__global`)

Globaali muisti (global memory) on muistialue, johon jokaisella säikeellä on luku- ja kirjoitusoikeus riippumatta siitä mihin säieryhmään ne kuuluvat
Näytönohjaimien tapauksessa suurin osa videomuistista on käytettävissä globaalin muistin muodossa
Globaali muisti on useimmassa tapauksissa toteutettu näytönohjainpiirin ulkopuolella, joten muistikaista on rajoitettu ja muistin käyttämiseen liittyvät latenssiajat ovat satojen kellojaksojen luokkaa
Modernien näytönohjaimien tapauksessa globaalia muistia käytetään välimuistin kautta, mutta OpenCL-määritykset eivät vaadi tätä

Lokaali muisti (`__local`)

Lokaali muisti (local memory) on muistialue, joka näkyy pelkästään saman säieryhmän säikeille
Lokaalia muistia käytetään tyypillisesti tilanteissa, joissa samaan säieryhmään kuuluvat säikeet haluavat jakaa dataa keskenään
Moderneissa näytönohjaimissa lokaali muisti on totetettu osana näytönohjainpiiriä (tyypillisesti jokaisella laskentayksiköllä on oma lokaali muisti), joten lokaali muisti on noin kertaluokkaa nopeampi kuin globaali muisti
CUDA-terminologiassa lokaalia muistia nimitetään jaetuksi muistiksi (shared memory)

Vakiomuisti (`__constant`)

Vakiomuisti (constant memory) on kaikille säikeille näkyvä muistialue, jonka sisältö pysyy samana ytimen suorituksen ajan
Vakiomuistiin osoittavat muistipyynnöt kulkevat usein välimuistin kautta. OpenCL-määritykset eivät kuitenkaan vaadi tätä laitteistolta.

Yksityinen muisti `(__private`)

Yksityinen muisti (private memory) on muistialue, joka näkyy pelkästään yksittäiselle säikeelle
Kääntäjä saattaa tallentaa dataa yksityiseen muistiin tilanteissa, joissa laskentayksikön resurssit ovat vähissä
Toteutetty tyypillisesti osana globaalia muistia, joten yksityisen muistin eksplisiittinen käyttäminen ei ole suositeltavaa

Relaksoitu konsistenssi (relaxed consistency)

OpenCL:n muistimalli on konsistentti yksittäisen säikeen näkökulmasta
Tässä yhteydessä konsistentti tarkoittaa mm. sitä, että muistioperaatiot tapahtuvat ennalta määrätyssä järjestyksessä
OpenCL:n muistimalli ei ole konsistentti eri säikeiden välillä
Tämä tarkoittaa mm. sitä, että eri säikeiden tekemien muistioperaatioiden suoritusjärjestystä ei ole määritelty
Muistin saattaminen konsistenttiin tilaan onnistuu joissakin tilanteissa, mutta vaatii erityisen synkronointikomennon

Ensimmäinen OpenCL 1.2 -ohjelma

Tällä kurssilla OpenCL-rajapintaa käytetään pääasiassa C++ -kääreen kautta. Tämä valinta selkeyttää ohjelmakoodia jonkin verran ja yksinkertaistaa erityisesti muistinhallintaa.
Mikäli tietokoneeseen on asennettu toimiva OpenCL-toteutus, tarvitsee C++ -kääre toimiakseen pelkästään CL/cl.hpp -otsikkotiedoston
Tarvittaessa kääre on ladattavissa Khronos groupin verkkosivulta. Tarkempi dokumentaatio löytyy täältä.

Tutustutaan seuraavaksi seuraavaan esimerkkiohjelmaan: git / ensimmainen-opencl-ohjelma
Esimerkkiohjelma sisältää laajennetun version (main.cpp), jossa tarkistetaan käytettyjen parametrien sopivuus käytetylle OpenCL-laitteella ja käsitellään kaikki virhetilanteet
Tämän lisäksi paketti sisältää yksinkertaistetun version, jossa demotaan C++ -kääreen C++11 -ominaisuuksia ja poikkeuksien käsittelyä
Lopuksi mukana on myöskin ns. barebones-versio, jossa on mukana vain absoluuttinen minimi määrä koodia (ei argumenttien validointia ja virhekoodin käsittelyä).

main.cpp-tiedosto alkaa CL/cl.hpp -otsikkotiedoston sisällyttämisellä:

#include <CL/cl.hpp>

Kaikki C++-kääreeseen liittyvät luokat ja funktiot kuuluvat cl-nimiavaruuteen
Nimiavaruuden etuliitettä ei tarvitsisi mainita erikseen mikäli käyttäisime using namespace cl -komentoa. Otamme sen kuitankin mukaan tässä esimerkissä selvyyden vuoksi.
C++kääreessä esiintyvät VECTOR_CLASS ja STRING_CLASS viittaavat oletuksena std::vector ja std::string tietotyyppeihin

Virheiden käsittely

Suurin osa OpenCL-kirjaston kutsuista joko
- palauttaa cl_int-tyyppisen virhekoodin tai
- ottaa argumenttinaan osoittimen cl_int-tyyppiseen virhekoodimuuttujaan, johon OpenCL-kutsun status kirjoitetaan
CL_SUCCESS (0) on geneerinen kaikki ok -paluuarvo
Muista tyypillisiä virhekoodeja:
- CL_DEVICE_NOT_FOUND / -1,
- CL_OUT_OF_RESOURCES / -5,
- CL_INVALID_COMMAND_QUEUE / -37,
- CL_INVALID_KERNEL_ARGS / -52
Virhekoodeja vastaavat numeroarvot löytyvät CL/cl.h-tiedostosta

C++-kääre osaa myöskin ilmoittaa virhetilanteista poikkeuksien avulla
Poikkeukset ovat oletuksena pois käytöstä, mutta ne voidaan aktivoida seuraavasti:

#define __CL_ENABLE_EXCEPTIONS
#include <CL/cl.hpp>

Poikkeukset otetaan kiinni try-catch-rakenteella:

try {
    ...
} catch(const cl::Error &err) {
    std::cerr << "Virheviesti: " << err.what() << std::endl;
}

OpenCL-alusta

Tietokoneelle asennettuun OpenCL-toteutuksen hyödyntämiseksi meidän täytyy ensin hankkia käyttöömme OpenCL-alusta (OpenCL platform layer)
OpenCL-alustan kautta pääsemme edelleen käsiksi laitteistoon asennettuihin OpenCL-laitteisiin ja niiden resursseihin
C++ -kääreessä OpenCL-alusta on paketoitu cl::Platform -luokan sisälle, joka tarjoaa staattisen jäsenfunktion OpenCL-alustojen hakemiseen:

static cl_int cl::Platform::get(VECTOR_CLASS<Platform> * platforms)

cl_int err;
	
std::vector<cl::Platform> platforms;
err = cl::Platform::get(&platforms);

if(err != CL_SUCCESS) {
	std::cerr << "OpenCL-alustojen hakeminen epäonnistui." << 
	    std::endl;
	std::cerr << "OpenCL-virhekoodi: " << err << std::endl;
	return 1;
}

if(platforms.size() < 1) {
	std::cerr << "OpenCL-alustoja ei löytynyt." << std::endl;
	return 1;
}

Seuraavaksi voisimme kysellä OpenCL-alustalta eri asioita getInfo -jäsenfunktion avulla:

cl_int cl::Platform::getInfo(cl_platform_info name,
                             STRING_CLASS * param)

OpenCL-alustan tukeman OpenCL-version voi tarkistaa CL_PLATFORM_VERSION -lipulla ja valmistajan nimen saa selville CL_PLATFORM_VENDOR - lipulla
OpenCL-alusta saattaa tukea myöskin erinäisiä standardeja ja toteutuskohtaisia laajennuksia. CL_PLATFORM_EXTENSIONS -lippu palauttaa listan tuetuista laajennuksista.

OpenCL-laitteiden hakeminen

OpenCL-alustaan saattaa liittyä useampi OpenCL-laite
C++ -kääreessä OpenCL-laite on paketoitu cl::Device -luokan sisään ja cl::Platform -luokka sisältää getDevices-jäsenfunktion OpenCL-laitteiden kyselemistä varten:

cl_int cl::Platform::getDevices(cl_device_type type,
                                VECTOR_CLASS<Device> * devices)

Yllä type -argumentti voi saada esimerkiksi arvot CL_DEVICE_TYPE_CPU, CL_DEVICE_TYPE_GPU tai CL_DEVICE_TYPE_ACCELERATOR. Tämän kurssin puitteissa CL_DEVICE_TYPE_GPU on tietenkin kaikkein mielenkiintoisin vaihtoehto.

std::vector<cl::Device> devices;
err = platforms[0].getDevices(CL_DEVICE_TYPE_GPU, &devices);

if(err != CL_SUCCESS) {
	std::cout << "OpenCL-laitteiden hakeminen epäonnistui." << 
	    std::endl;
	std::cerr << "OpenCL-virhekoodi: " << err << std::endl;
	return 1;
}

if(devices.size() < 1) {
	std::cout << "OpenCL-laitteita ei löytynyt." << std::endl;
	return 1;
}

Voisimme jo tuttuun tapaan kysellä cl::Device -oliolta eri asioita:

template <typename T>
cl_int cl::Device::getInfo(cl_device_info name,
                           T * param)

CL_DEVICE_MAX_WORK_GROUP_SIZE -lippu palauttaa säieryhmän suurimman mahdollisen koon
CL_DEVICE_MAX_MEM_ALLOC_SIZE -lippu palauttaa suurimman varattavissa olevan globaalin muistialueen koon
CL_DEVICE_LOCAL_MEM_SIZE-lippu palauttaa suurimman varattavissa olevan lokaalin muistialueen koon
CL_DEVICE_EXTENSIONS -lippu palauttaa OpenCL-laitteen tukeman laajennukset

OpenCL-konteksti

OpenCL-laitteiden resurssien hallinta tapahtuu OpenCL-kontekstin kautta
C++ -kääreessä OpenCL-konteksti on paketoitu cl::Context -luokan sisään, jonka muodostinfunktio ottaa argumenttinaan mukaan liitettävät OpenCL-laiteet:

cl::Context::Context(VECTOR_CLASS<Device>& devices,
                     cl_context_properties * properties = NULL,
                     void (CL_CALLBACK * pfn_notify)(
                         const char * errorinfo,
                         const void * private_info,
                         ::size_t cb,
                         void * user_data) = NULL,
                     void * user_data = NULL,
                     cl_int * err = NULL)

std::vector<cl::Device> oneDevice;
oneDevice.push_back(devices[0]);
cl::Context context(oneDevice, 0, 0, 0, &err);

if(err != CL_SUCCESS) {
	std::cout << "OpenCL-kontekstin luominen epäonnistui." << 
	    std::endl;
	std::cerr << "OpenCL-virhekoodi: " << err << std::endl;
	return 1;
}

cl::Context -olio voidaan luoda myöskin ilman cl::Device -oliota:

cl::Context::Context(cl_device_type type,
                     cl_context_properties * properties = NULL,
                     void (CL_CALLBACK * pfn_notify)(
                         const char * errorinfo,
                         const void * private_info,
                         ::size_t cb,
                         void * user_data) = NULL,
                     void * user_data = NULL,
                     cl_int * err = NULL)

Tällöin mukaan liitetään kaikki type argumenttia vastaavat laitetyypit

OpenCL-kontekstin tietoihin päästään käsiksi getInfo -jäsenfunktiolla:

template <typename T>
cl_int cl::Context::getInfo(cl_context_info name,
                            T * param)

Erityisesti CL_CONTEXT_DEVICES -lippu, joka palauttaa kontekstiin liitetyt OpenCL-laiteet, on hyödyllinen mikäli cl::Context -olio luotiin ilman cl::Device -oliota.

Ytimen lähdekoodi

Ytimet kirjoitetaan (rajoitetulla) C-kielellä
Tarkastellaan seuraavaksi yksinkertaista ydintä joka ottaa argumenttinaan taulukon (buffer) kokonaislukuja (n kpl) ja lisää jokaiseen alkioon luvun yksi:

__kernel void add_one(__global int *buffer, int n) {
    const int global_id = get_global_id(0);

    if(global_id < n)
        buffer[global_id]++; 
}

__kernel -avainsana kertoo, että kyseessä on ydin, jota voidaan kutsua isäntäohjelman puolelta
__global -avainsana kertoo, että buffer-taulukko on tallennettu globaaliin muistiin

Globaali indeksinumero saadaan käyttöön get_global_id -aliohjelman avulla. Aliohjelmalle annettu argumentti määrää indeksi dimension:

const int global_id = get_global_id(0);

Kukin säie kasvattaa omaa globaalia indeksinumeroaan vastaavaa taulukon alkiota yhdellä:

if(global_id < n)
    buffer[global_id]++;

Yllä oletetaan, että säikeitä on enemmän kuin taulukossa alkioita

Yksinkertaisissa tapauksessa ytimet voidaan kirjoittaa isäntäohjelman sisään merkkijonoiksi:

const char *kernel =
"__kernel void add_one(__global int *buffer, int n) {  \n" \
"    const int global_id = get_global_id(0);           \n" \
"    if(global_id < n)                                 \n" \
"        buffer[global_id]++;                          \n" \
"}                                                     \n";

Tämän jälkeen ytimien lähdekoodit voidaan pakata cl::Program::Sources -tyyppisen olion sisälle:

cl::Program::Sources sources;
sources.push_back(
    cl::Program::Sources::value_type(kernel, strlen(kernel)));

Ytimien lähdekoodi voidaan myöskin tallentaa erilliseen tiedostoon. OpenCL ei tue tätä suoraan, mutta esimerkiksi linux-ympäristöstä löytyvä xxd -työkalu kelpaa tähän tarkoitukseen loistavasti
Oletetaan seuraavaksi, että ytimen lähdekoodi on kirjoitettu erilliseen kernel.cl -tiedostoon. Tällöin komento xxd -i kernel.cl > kernel.cl.dat tuottaisi seuraavan tiedoston:

unsigned char kernel_cl[] = {
  0x0a, 0x2f, 0x2f, 0x20, 0x59, 0x64, 0x69, 0x6e, 0x2c, ...
  0x6b, 0x61, 0x20, 0x6f, 0x74, 0x74, 0x61, 0x61, 0x20, ...
  ...
  0x2b, 0x3b, 0x20, 0x0a, 0x7d
};
unsigned int kernel_cl_len = 857;

Nyt voisimme sisällyttää lähdekoodin osaksi isäntäohjelmaa:

#include "kernel.cl.dat"

...

cl::Program::Sources sources;
sources.push_back(cl::Program::Sources::value_type(
    (const char*) kernel_cl, kernel_cl_len));

Isäntäohjelman lähdekoodiin upotetun merkkijonon ja xxd -työkalun hyödyntämisen lisäksi ytimien lähdekoodi voitaisiin ladata erillisestä tekstitiedostosta ajonaikaisesti merkkijonotaulukkon

Ytimen lähdekoodin kääntäminen

Ytimien lähdekoodin kääntämistä varten tarvitsemme Ohjelma-objektin (Program object), joka kapseloi sisäänsä liittyvän OpenCL-kontekstin, ytimien lähdekoodit/binäärit, käännetyn version ytimien lähdekoodista, käännöslogit ja ytimien "kahvat"
C++ -kääreessä ohjelma-objekti on pakattu cl::Program -luokan sisälle, jonka muodostinfunktio ottaa argumenttinaan liittyvän OpenCL-kontekstin ja ytimien lähdekoodit:

cl::Program::Program(const Context& context,
                     const Sources& sources,
                     cl_int * err = NULL)

cl::Program program(context, sources, &err);

if(err != CL_SUCCESS) {
	std::cout << "Ohjelma-objektin luominen epäonnistui." << 
	    std::endl;
	std::cerr << "OpenCL-virhekoodi: " << err << std::endl;
	return 1;
}

Ohjelma-objektiin liitetyt ytimien lähdekoodit voidaan kääntää cl::Program -luokan build -jäsenfunktion avulla:

cl_int cl::Program::build(const VECTOR_CLASS<Device>& devices,
                          const char * options = NULL,
                          void (CL_CALLBACK * pfn_notify)(
                              cl_program,
                              void * user_data) = NULL,
                          void * data = NULL)

Käännösprosessin tuloste ja mahdolliset virheet saadaan selville antamalla CL_PROGRAM_BUILD_LOG -lippu ´getBuildInfo´-jäsenfunktiolle:

template <typename T>
cl_int cl::Program::getBuildInfo(cl_program_build_info name,
                                 T * param)

err = program.build(oneDevice, 0);

if(err != CL_SUCCESS) {

	std::string log;
	program.getBuildInfo(oneDevice[0], 
	    CL_PROGRAM_BUILD_LOG, &log);
	std::cout << 
		"OpenCL-kääntäjän tuloste:" << std::endl << 
		log << std::endl;
	
	std::cout << 
		"Ytimien lähdekoodin kääntäminen epäonnistui." << 
		std::endl;
	std::cerr << "OpenCL-virhekoodi: " << err << std::endl;
	return 1;
}

Mikäli ytimien lähdekoodit on tallennettu yhteen merkkijonoon, voisimme jättää cl::Program::Sources -olion luomisen kokonaan pois antamalla lähdekoodit suoraan merkkijonona cl::Program -luokan muodostinfunktiolle:

cl::Program::Program(const Context& context,
                     const STRING_CLASS& source,
                     bool build,
                     cl_int * err = NULL)

Mikäli tämän lisäksi kutsuisimme muodostinfunktiota build argumentin arvolla CL_TRUE, käännettäisiin ytimien lähdekoodit jo cl::Program -olion luonnin yhteydessä, jolloin meidän ei tarvitsisi kutsua build -jäsenfunktiota erikseen

Ydin-objektin luominen

Ydin-objekti kapseloi sisäänsä yksittäiseen ytimeen liittyvät ominaisuudet
C++ -kääreessä ydin-objekti on kapseloitu cl::Kernel -luokan sisälle, jonka muodostinfunktio ottaa argumenttinaan Ohjelma-objektin ja ytimen nimen:

cl::Kernel::Kernel(const Program& program,
                   const char * name,
                   cl_int * err = NULL)

cl::Kernel kernel(program, "add_one", &err);

if(err != CL_SUCCESS) {
	std::cout << "Ytimen luominen epäonnistui." << std::endl;
	std::cerr << "OpenCL-virhekoodi: " << err << std::endl;
	return 1;
}

Voisimme myöskin hyödyntää C++11 -standardista löytyviä funktoreita ydin-objektin luomisessa:

template<typename T0, 
         typename T1 = detail::NullType, 
         ..., 
         typename T31 = detail::NullType>
struct make_kernel::detail::functionImplementation<T0, T1, ..., T31>
cl::make_kernel::make_kernel(const Program &program,
                             const STRING_CLASS name,
                             cl_int *err = NULL)

Esimerkkiohjelmassa ydin-objekti voitaisiin luoda seuraavasti:

auto kernel = cl::make_kernel<cl::Buffer&, int>(program, "add_one");

Komentojonon luominen

Isäntäohjelma asettaa ytimien käynnistyskomennot OpenCL-laitekohteiseen komentojonoon (Command Queue)
Ajonaikainen järjestelmä suorittaa ytimet automaattisesti annetussa järjestyksessä
Komentojonoon voidaan asettaa myöskin esimerkiksi datasiirtokomentoja ja ns. eventtejä (Event), joiden avulla komentojonon tilannetta voidaan seurata reaaliajassa

Komentojono on kapseloitu cl::CommandQueue-luokan sisälle, jonka muodostinfunktio ottaa argumenttinaan OpenCL-kontekstin ja OpenCL-laitteen, joihin komentojono on tarkoitus liittää:

cl::CommandQueue::CommandQueue(
    const Context& context,
    const Device& device,
    cl_command_queue_properties properties = 0,
    cl_int * err = NULL)

cl::CommandQueue queue(context, oneDevice[0], 0, &err);

if(err != CL_SUCCESS) {
	std::cout << "Komentojonon luominen epäonnistui." << std::endl;
	std::cerr << "OpenCL-virhekoodi: " << err << std::endl;
	return 1;
}

Puskuri-objektin luominen

OpenCL-laitteen globaalia ja vakiomuistia hallitaan Puskuri-objektien (Buffer object) avulla
Puskuri-objekti on kapseloitu cl::Buffer-luokan sisälle, joka ottaa argumenttinaan liittyvän OpenCL-kontekstin, flags-lippumuuttujan ja puskurin koon:

cl::Buffer::Buffer(
    const Context& context,
    cl_mem_flags flags,
    ::size_t size,
    void * host_ptr = NULL,
    cl_int * err = NULL)

cl::Buffer deviceBuffer(context, CL_MEM_READ_WRITE, N*sizeof(int), 0, &err);

if(err != CL_SUCCESS) {
	std::cout << "Muistin varaaminen epäonnistui epäonnistui." << std::endl;
	std::cerr << "OpenCL-virhekoodi: " << err << std::endl;
	return 1;
}

Datan kirjoittaminen puskuriin

Komentojonon enqueueWriteBuffer-jäsenfunktio asetaa komentojonoon käskyn kirjoittaa dataa isäntälaiteen muistista Puskuri-objektiin
blocking_write-lipun asettaminen arvoon CL_TRUE tekee kutsusta blockaavan eli aliohjelmasta palataan vasta kun siirto on suoritettu loppuun

cl_int cl::CommandQueue::enqueueWriteBuffer(
    const Buffer& buffer,
    cl_bool blocking_write,
    ::size_t offset,
    ::size_t size,
    const void * ptr,
    const VECTOR_CLASS<Event> * events = NULL,
    Event * event = NULL)

err = queue.enqueueWriteBuffer(deviceBuffer, CL_FALSE, 0, N*sizeof(int), 
	hostBuffer, 0, 0);

if(err != CL_SUCCESS) {
	std::cout << "Isäntälaite -> OpenCL-laite -siirtokäskyn asettaminen " \
		"komentojonoon epäonnistui." << std::endl;
	std::cerr << "OpenCL-virhekoodi: " << err << std::endl;
}

Ytimen argumenttia asettaminen

Ytimen argumentit asetataan yksitellen setArg-jäsenfunktiolla:

template <typename T>
cl_int cl::Kernel::setArg(cl_uint index, T value)

kernel.setArg(0, deviceBuffer);

if(err != CL_SUCCESS) {
	std::cout << "Ytimen 1. argumentin asettaminen epäonnistui." << 
		std::endl;
	std::cerr << "OpenCL-virhekoodi: " << err << std::endl;

	return 1;
}

kernel.setArg(1, N);

if(err != CL_SUCCESS) {
	std::cout << "Ytimen 2. argumentin asettaminen epäonnistui." << 
		std::endl;
	std::cerr << "OpenCL-virhekoodi: " << err << std::endl;

	return 1;
}

Ytimen käynnistäminen

Ytimen käynnistyskäsky asetetaan komentojonoon enqueueNDRangeKernel-jäsenfunktiolla
Globaalin indeksiavaruuden koko asetetaan global-argumentilla ja säieryhmän koko vastaavasti local-argumentilla

cl_int cl::CommandQueue::enqueueNDRangeKernel(
    const Kernel& kernel,
    const NDRange& offset,
    const NDRange& global,
    const NDRange& local,
    const VECTOR_CLASS<Event> * events = NULL,
    Event * event = NULL)

On hyvä idea tarkistaa säieryhmän maksimikoko:

::size_t maxWorkGroupSize;
oneDevice[0].getInfo(CL_DEVICE_MAX_WORK_GROUP_SIZE, &maxWorkGroupSize);

if(err != CL_SUCCESS) {
	std::cout << "OpenCL-laitteen suurimman mahdollinen säieryhmän koon "\
	"kysyminen epäonnistui." << std::endl;
	std::cerr << "OpenCL-virhekoodi: " << err << std::endl;

	return 1;
}

Globaalin indeksiavaruuden ja säieryhmn koon laskeminen:

::size_t workGroupCount = N/maxWorkGroupSize+1;

cl::NDRange globalDim(workGroupCount*maxWorkGroupSize, 1, 1);
cl::NDRange localDim(maxWorkGroupSize, 1, 1);

err = queue.enqueueNDRangeKernel(kernel, cl::NullRange, 
	globalDim, localDim, 0, 0);

if(err != CL_SUCCESS) {
	std::cout << "Ytimen käynnistyskäskyn asettaminen komentojonoon " \
		"epäonnistui." << std::endl;
	std::cerr << "OpenCL-virhekoodi: " << err << std::endl;
	
	return 1;
}

Datan lukeminen puskurista

Komentojonon enqueueReadBuffer-jäsenfunktio asetaa komentojonoon käskyn lukea dataa Puskuri-objektista isäntälaiteen muistiin
blocking_read-lipun asettaminen arvoon CL_TRUE tekee kutsusta blockaavan eli aliohjelmasta palataan vasta kun siirto on suoritettu loppuun

cl_int cl::CommandQueue::enqueueReadBuffer(
    const Buffer& buffer,
    cl_bool blocking_read,
    ::size_t offset,
    ::size_t size,
    const void * ptr,
    const VECTOR_CLASS<Event> * events = NULL,
    Event * event = NULL)

err = queue.enqueueReadBuffer(deviceBuffer, CL_TRUE, 0, N*sizeof(int), 
	hostBuffer, 0, 0);

if(err != CL_SUCCESS) {
	std::cout << "OpenCL-laite -> Isäntälaite -siirtokäskyn asettaminen " \
		"komentojonoon epäonnistui." << std::endl;
	std::cerr << "OpenCL-virhekoodi: " << err << std::endl;
	
	return 1;
}

Ensimmäinen kosketus GPU-ohjelmointiin

OpenCL vs CUDA

Yleiskuva GPU-laitteistosta

Isäntälaite

OpenCL-laite

Laskentayksiköt ja prosessointielementit

Laskentayksiköt ja prosessointielementit (jatkuu)

OpenCL-sovellus

OpenCL-sovellus (jatkuu)

Ydin

Globaali indeksiavaruus

Säikeiden suorituspolut

Säieryhmät

Esimerkki säikeiden indeksoimisesta

Säieryhmät ja osatehtävät

Säiryhmät ja laskentayksiköt

Muistimalli

Globaali muisti (__global)

Lokaali muisti (__local)

Vakiomuisti (__constant)

Yksityinen muisti (__private)

Relaksoitu konsistenssi (relaxed consistency)

Ensimmäinen OpenCL 1.2 -ohjelma

Virheiden käsittely

OpenCL-alusta

OpenCL-laitteiden hakeminen

OpenCL-konteksti

Ytimen lähdekoodi

Ytimen lähdekoodin kääntäminen

Ydin-objektin luominen

Komentojonon luominen

Puskuri-objektin luominen

Datan kirjoittaminen puskuriin

Ytimen argumenttia asettaminen

Ytimen käynnistäminen

Datan lukeminen puskurista

Globaali muisti (`__global`)

Lokaali muisti (`__local`)

Vakiomuisti (`__constant`)

Yksityinen muisti `(__private`)