Lue ensin tällä sivulla olevat yleisohjeet. Tarkemmat tiedot löytyvät tehtäväkohtaisista ohjeista. Jokaiseen tehtävään on saatavilla oma zip-tiedosto, josta löytyvää koodia voit käyttää ratkaisujesi pohjana.
| Tehtävä | Palautuksia | Pisteet | Max. | Taso | Suos. | Määräaika |
|---|---|---|---|---|---|---|
| I8MM2: perustason CPU-toteutus | ||||||
Toteuta yksinkertainen peräkkäinen perusratkaisu. Älä yritä vielä käyttää minkäänlaista rinnakkaisuutta; yritä ensin saada toteutus toimimaan oikein. |
||||||
| – | – | 3 + 0 | ★ | Suos. | 2026-10-07 klo 23:59:59 | |
| I8MM3: nopea CPU-toteutus | ||||||
Toteuta nopea CPU-ratkaisu hyödyntäen monisäikeistystä. Käytä käskytason rinnakkaisuutta, vektorointia sekä rekisterien ja välimuistin uudelleenkäyttöä. |
||||||
| – | – | 5 + 0 | ★+ | Suos. | 2026-10-21 klo 23:59:59 | |
| I8MM4: perustason GPU-toteutus | ||||||
Toteuta perustason GPU-ratkaisu. |
||||||
| – | – | 3 + 0 | ★ | Suos. | 2026-11-04 klo 23:59:59 | |
| I8MM5: nopea GPU-toteutus | ||||||
Toteuta nopea GPU-ratkaisu käyttäen tavallista (siirrettävää) CUDAa. |
||||||
| – | – | 3 + 0 | ★★ | Suos. | 2026-11-18 klo 23:59:59 | |
| I8MM6: perustason Tensor Core -toteutus | ||||||
Toteuta Tensor Core -perusratkaisu. Älä yritä vielä käyttää datan uudelleenkäyttöä. Kyseessä on tekniikkaharjoitus, joten hyväksyttävän ratkaisun täytyy hyödyntää tensoriytimiä (tensor cores). |
||||||
| – | – | 4 + 0 | ★★ | Suos. | 2026-12-02 klo 23:59:59 | |
| I8MM9a: CPU AVX512-VNNI | ||||||
Toteuta nopea CPU-ratkaisu käyttäen AVX512-VNNI-käskyjä. |
||||||
| – | – | 4 + 2 | ★★ | Suos. | 2026-12-02 klo 23:59:59 | |
| I8MM9b: SIMD GPU | ||||||
Toteuta nopea GPU-ratkaisu hyödyntäen Turing-arkkitehtuurin SIMD-erityiskäskyjä.
Kyseessä on tekniikkaharjoitus, joten hyväksyttävän ratkaisun täytyy käyttää |
||||||
| – | – | 4 + 2 | ★★ | Suos. | 2026-12-02 klo 23:59:59 | |
| I8MM9c: nopea Tensor Core -toteutus | ||||||
Toteuta nopea Tensor Core -ratkaisu. |
||||||
| – | – | 7 + 2 | ★★★ | Suos. | 2026-12-02 klo 23:59:59 | |
Saat m × k -matriisin ja k × n -matriisin, jotka koostuvat 8-bittisistä kokonaisluvuista. Tehtävänäsi on laskea näiden kahden matriisin m × n -kokoinen tulo.
Toteuta seuraava funktio:
void gemm(int m, int n, int k, const int8_t* A, const int8_t* B, int32_t* C);
Tässä A ja B ovat osoittimia syötematriiseihin. Matriisissa A on m riviä ja k saraketta, ja matriisissa B on
k riviä ja n saraketta.
Kaikilla 0 <= y < m ja 0 <= x < k matriisin A rivin y ja sarakkeen x alkio on tallennettu kohtaan A[x + y*k].
Kaikilla 0 <= y < k ja 0 <= x < n matriisin B rivin y ja sarakkeen x alkio on tallennettu kohtaan B[x + y*n].
Funktion tulee ratkaista seuraava tehtävä:
kaikilla i ja j, joilla 0 <= i < m ja 0 <= j < n,
laske matriisin A rivin i ja matriisin B sarakkeen j välinen sisätulo ja tallenna tulos kohtaan C[j + i*n].
Taulukot A, B ja C on jo varattu valmiiksi funktiota kutsuvassa koodissa; sinun ei tarvitse tehdä näihin taulukoihin liittyvää muistinhallintaa.
CPU:lla ratkaistavissa tehtävissä A, B ja C osoittavat CPU-muistiin,
ja GPU:lla ratkaistavissa tehtävissä ne osoittavat laitemuistiin.
Älä oleta, että C sisältää kutsuhetkellä kelvollisia arvoja. Erityisesti ei ole taattua, että se olisi alustettu nollilla.
Summausdimensio k on taatusti pienempi kuin 65536, joten kaikki tulokset voidaan esittää 32-bittisinä etumerkillisinä kokonaislukuina.
Vaikka liukuluku- ja kokonaislukumatriisien kertolasku näyttävät hyvin samankaltaisilta, mikroarkkitehtuurin tasolla niissä on yksi ratkaiseva ero: kun kerrotaan kaksi 32-bittistä liukulukua, tulos on jälleen 32-bittinen liukuluku, joka voidaan lisätä 32-bittiseen liukulukuun. Sen sijaan kahden 8-bittisen kokonaisluvun tulo on 16-bittinen kokonaisluku, ja jos tällaisia tuloja halutaan laskea yhteen useita, akkumulointiin tarvitaan 32-bittinen kokonaisluku.
Ei voi olla sellaista SIMD-käskyä, joka ottaisi kaksi pakattuja 8-bittisiä kokonaislukuja sisältävää vektorirekisteriä ja
akkumuloisi tuloksen kolmanteen rekisteriin (kuten esimerkiksi VFMADDPS tekee); kohderekisteri on yksinkertaisesti liian pieni
kaikkien 64 tulon akkumuloimiseen, jos oletetaan 512-bittiset rekisterit. Sen sijaan laitteisto toteuttaa sisätulon kaltaisia operaatioita: otetaan yhdestä operandista 8-bittisiä kokonaislukuja pareittain (tai neljän ryhmissä), kerrotaan kukin vastaavalla
toisen operandin 8-bittisellä kokonaisluvulla, summataan yksittäiset tulot ja akkumuloidaan ne kohdeoperandiin. Näin
kohde voi sisältää vähemmän kokonaislukuja, mutta suuremmilla bittileveyksillä.
Yleisessä AVX-512:ssa on yksi käsky 8-bittisen sisätulon laskemiseen lukupareista 16-bittisellä akkumuloinnilla.
Tästä ei ole erityisesti hyötyä, koska akkumulointi täytyy tehdä 32 bitillä ylivuotojen estämiseksi.
Samankaltainen käsky on kuitenkin olemassa kahden 16-bittisen luvun sisätulolle, jossa akkumulointi tehdään 32 bitillä.
8-bittisten lukujen laajentaminen 16-bittisiksi ja _mm512_madd_epi16-käskyn käyttäminen voi olla toimiva strategia.
Huomaa VNNI-tehtävässä, että käytettävissä oleva käsky _mm512_dpbusds_epi32 sallii vain
yhden etumerkillisen ja yhden etumerkittömän operandin kertomisen. Jotta saisit tämän käskyn nopeushyödyt käyttöön,
sinun täytyy siis toteuttaa esi- ja jälkikäsittely, joka muuntaa etumerkillisten kokonaislukumatriisien kertolaskun
matriisikertolaskuksi, jossa kerrotaan etumerkillisiä ja etumerkittömiä kokonaislukuja. CUDA:n __dp4a-intrinsic-funktio tukee suoraan etumerkillisten kokonaislukujen kertolaskua.
Yksinkertainen ajatusmalli tensoriytimen perusoperaatiolle on, että se laajentaa tavallisen SIMD-käsittelyn vektorioperaatiot käskyiksi, jotka operoivat kiinteänkokoisilla matriisifragmenteilla. Kun käytetään 8-bittisiä kokonaislukuoperandeja, GPU:n kukin warp voi käsitellä kahden 16 × 16 -fragmentin tulon yhdellä käskyllä. Näiden käskyjen C++-rajapinta on dokumentoitu CUDA Programming Guide-oppaassa.
Voit siis ajatella syöte- ja tulosmatriisien koostuvan 16 × 16 -laatoista, ja algoritmi voidaan toteuttaa samalla tavalla kuin skalaarinen matriisikertolasku, paitsi että jokainen alkio on nyt matriisifragmentti. Erityisesti optimoinnit, kuten rekisterien uudelleenkäyttö (nyt kokonaisten fragmenttien tasolla), jaettu muisti ja oikean data-asettelun valinta, ovat edelleen ratkaisevan tärkeitä hyvän suorituskyvyn kannalta.