Hoofdstuk 11: GPU Onderzoeksrichtingen op Scalarisatie en Affiene Uitvoering

Zoals beschreven in Hoofdstuk 2, hebben GPU-rekenprogramma's, zoals CUDA en OpenCL, een MIMD-achtig programmamodel dat de programmeur in staat stelt om een groot aantal scalaire threads op de GPU te lanceren. Hoewel elk van deze scalaire threads zijn eigen unieke uitvoeringspad kan volgen en willekeurige geheugenlocaties kan openen, volgen ze in de meeste gevallen een klein aantal uitvoeringspaadjes en voeren ze soortgelijke bewerkingen uit.

De convergente controle-stroom tussen GPU-threads wordt geëxploiteerd op de meeste, zo niet alle, moderne GPU's via het SIMT-uitvoeringsmodel, waarbij scalaire threads worden gegroepeerd in warps die worden uitgevoerd op SIMD-hardware (zie Sectie 3.1.1). Dit hoofdstuk vat een reeks onderzoeken samen die de gelijkenis van deze scalaire threads verder exploiteren via scalarisatie en affiene uitvoering.

Het belangrijkste inzicht van dit onderzoek ligt in de observatie van de waardestructuur [Kim et al., 2013] over threads die dezelfde rekenkernel uitvoeren. De twee typen waardestructuur, uniform en affien, worden geïllustreerd in de rekenkernel in Voorbeeld 11.1.

Uniforme variabele Een variabele die dezelfde constante waarde heeft voor elke thread in de rekenkernel. In Algoritme 11.1 hebben de variabele a, evenals de literals THRESHOLD en Y_MAX_VALUE, allemaal uniforme waarden over alle threads in de rekenkernel. Een uniforme variabele kan worden opgeslagen in één scalair register en door alle threads in de rekenkernel worden hergebruikt.

Affiene variabele Een variabele met waarden die een lineaire functie zijn van de thread-ID voor elke thread in de rekenkernel. In Algoritme 11.1 kan het geheugenadres van de variabele y[idx] worden voorgesteld als een affiene transformatie van de thread-ID threadIdx.x:

&(y[idx]) = &(y[0]) + sizeof(int) * threadIdx.x;

Deze affiene representatie kan worden opgeslagen als een paar scalaire waarden, een basis en een stride, wat veel compacter is dan de volledig uitgevouwen vector.

__global__ void vsadd( int y[], int a ) {
    // Vertaal alleen de commentaren, niet de code
    // Voer de berekening uit op de GPU
}
```Hier is de Nederlandse vertaling van het Markdown-bestand, waarbij de code-opmerkingen zijn vertaald:
 
```c
int idx = threadIdx.x;
    y[idx] = y[idx] + a;
    if ( y[idx] > THRESHOLD )
        y[idx] = Y_MAX_VALUE;
}

Algoritme 11.1: Voorbeeld van schaal- en affiene bewerkingen in een compute-kernel (uit [Kim et al., 2013]).

Er zijn verschillende onderzoeksvoorstellen over hoe uniforme of affiene variabelen in GPU's kunnen worden gedetecteerd en benut. De rest van dit hoofdstuk vat deze voorstellen samen op deze twee aspecten.

Detectie van uniforme of affiene variabelen

Er zijn twee hoofdbenaderingen om het bestaan van uniforme of affiene variabelen in een GPU-compute-kernel te detecteren: Compiler-Driven Detection en Detection via Hardware.

Compiler-Driven Detection

Een manier om het bestaan van uniforme of affiene variabelen in een GPU-compute-kernel te detecteren, is door middel van een speciale compileranalyse. Dit is mogelijk omdat de bestaande GPU-programmeringsmodellen, CUDA en OpenCL, al middelen bieden voor de programmeur om een variabele als constant door de hele compute-kernel te verklaren, evenals het bieden van een speciale variabele voor de thread-ID. De compiler kan een controle-afhankelijkheidsanalyse uitvoeren om variabelen te detecteren die puur afhankelijk zijn van constanten en thread-ID's, en deze als uniform/affien markeren. Bewerkingen die uitsluitend op uniforme/affiene variabelen werken, komen vervolgens in aanmerking voor schalarisatie.

AMD GCN [AMD, 2012] vertrouwt op de compiler om uniforme variabelen en scalaire bewerkingen te detecteren die kunnen worden opgeslagen en verwerkt door een speciale scalaire processor.

Asanovic et al. [2013] introduceren een gecombineerde convergente en variante analyse waarmee de compiler bewerkingen in een willekeurige compute-kernel kan bepalen die in aanmerking komen voor schalarisatie en/of affiene transformatie. Instructies binnen de convergente regio's van een compute-kernel kunnen worden omgezet in scalaire/affiene instructies. Bij elke overgang van divergente naar convergente regio's van een compute-kernel voegt de compiler een syncwarp-instructie in om controle-flow-geïnduceerde registervafhankelijkheden tussen de twee regio's te verwerken. Asanovic et al. [2013] hebben deze aanpak overgenomenHier is de Nederlandse vertaling van het bestand, waarbij de code-opmerkingen zijn vertaald:

Deze analyse om scalaire bewerkingen te genereren voor de Temporal-SIMT-architectuur [Keckler et al., 2011, Krashinsky, 2011].

Ontkoppelde Affiene Berekening (DAC) [Wang en Lin, 2017] maakt gebruik van een soortgelijke compilatoranalyse om scalaire en affiene kandidaten te extraheren die in een afzonderlijke warp moeten worden ontkoppeld. Wang en Lin [2017] breiden het proces uit met een divergente affiene analyse, met als doel instructiestromen te extraheren