cudaAan de slag met cuda


Opmerkingen

CUDA is een eigen NVIDIA parallelle computertechnologie en programmeertaal voor hun GPU's.

GPU's zijn zeer parallelle machines die duizenden lichtgewicht threads parallel kunnen laten draaien. Elke GPU-thread is meestal langzamer in uitvoering en hun context is kleiner. Aan de andere kant kan GPU meerdere duizenden threads parallel en zelfs meer gelijktijdig uitvoeren (precieze aantallen zijn afhankelijk van het werkelijke GPU-model). CUDA is een C ++ dialect speciaal ontworpen voor NVIDIA GPU-architectuur. Vanwege de verschillen in architectuur kunnen de meeste algoritmen echter niet eenvoudigweg worden gekopieerd en geplakt vanuit C ++ - ze werken wel, maar zijn erg traag.

Terminologie

  • host - verwijst naar normale CPU-gebaseerde hardware en normale programma's die in die omgeving worden uitgevoerd
  • apparaat - verwijst naar een specifieke GPU waarin CUDA-programma's worden uitgevoerd. Een enkele host kan meerdere apparaten ondersteunen.
  • kernel - een functie op het apparaat die kan worden opgeroepen vanuit de hostcode.

Fysieke processorstructuur

De voor CUDA ingeschakelde GPU-processor heeft de volgende fysieke structuur:

  • de chip - de hele processor van de GPU. Sommige GPU's hebben er twee.
  • streaming multiprocessor (SM) - elke chip bevat maximaal ~ 100 SM's, afhankelijk van een model. Elke SM werkt bijna onafhankelijk van een andere en gebruikt alleen globaal geheugen om met elkaar te communiceren.
  • CUDA-kern - een enkele scalaire rekeneenheid van een SM. Hun precieze aantal hangt af van de architectuur. Elke kern kan een paar threads tegelijkertijd in een snelle opeenvolging verwerken (vergelijkbaar met hyperthreading in CPU).

Bovendien beschikt elke SM over een of meer warp-planners . Elke planner verzendt een enkele instructie naar verschillende CUDA-cores. Dit zorgt er effectief voor dat de SM in 32-brede SIMD- modus werkt.

CUDA-uitvoeringsmodel

De fysieke structuur van de GPU heeft directe invloed op hoe kernels op het apparaat worden uitgevoerd en hoe men ze programmeert in CUDA. Kernel wordt aangeroepen met een oproepconfiguratie die aangeeft hoeveel parallelle threads worden voortgebracht.

  • het raster - geeft alle threads weer die worden voortgebracht bij een kernelaanroep. Het wordt gespecificeerd als een één of twee dimensionale set blokken
  • het blok - is een semi-onafhankelijke reeks threads . Elk blok wordt toegewezen aan een enkele SM. Als zodanig kunnen blokken alleen communiceren via het wereldwijde geheugen. Blokken worden op geen enkele manier gesynchroniseerd. Als er te veel blokken zijn, kunnen sommige na elkaar worden uitgevoerd. Aan de andere kant, als bronnen het toelaten, kan meer dan één blok op dezelfde SM draaien, maar de programmeur kan daar niet van profiteren (behalve voor de hand liggende prestatieboost).
  • de draad - een scalaire reeks instructies uitgevoerd door een enkele CUDA-kern. Threads zijn 'lichtgewicht' met minimale context, waardoor de hardware ze snel in en uit kan wisselen. Vanwege hun aantal werken CUDA-threads met een paar registers die aan hen zijn toegewezen, en een zeer korte stapel (bij voorkeur helemaal geen!). Om die reden geeft de CUDA-compiler er de voorkeur aan om alle functieaanroepen te inline te maken om de kernel af te vlakken zodat deze alleen statische sprongen en lussen bevat. Function ponter-aanroepen en virtuele methode-aanroepen, hoewel ondersteund door de meeste nieuwere apparaten, leiden meestal tot een grote prestatieboete.

Elke thread wordt geïdentificeerd door een block index blockIdx en thread index binnen de block threadIdx . Deze getallen kunnen op elk moment worden gecontroleerd door elke lopende thread en is de enige manier om de ene thread van de andere te onderscheiden.

Bovendien zijn draden georganiseerd in scheringen , die elk exact 32 draden bevatten. Threads binnen een enkele warp worden perfect synchroon uitgevoerd, in SIMD fahsion. Threads van verschillende warps, maar binnen hetzelfde blok kunnen in elke volgorde worden uitgevoerd, maar kunnen worden gedwongen om te synchroniseren door de programmeur. Threads van verschillende blokken kunnen op geen enkele manier worden gesynchroniseerd of direct communiceren.

Geheugen organisatie

Bij normaal CPU-programmeren is de geheugenorganisatie meestal verborgen voor het programmeerapparaat. Typische programma's doen alsof er alleen RAM is. Alle geheugenbewerkingen, zoals het beheren van registers, het gebruik van L1- L2- L3- caching, swapping naar disk, etc. worden afgehandeld door de compiler, het besturingssysteem of de hardware zelf.

Dit is niet het geval met CUDA. Terwijl nieuwere GPU-modellen de last gedeeltelijk verbergen, bijvoorbeeld door het Unified Memory in CUDA 6, is het toch de moeite waard om de organisatie te begrijpen om prestatieredenen. De basis CUDA-geheugenstructuur is als volgt:

  • Hostgeheugen - het normale RAM-geheugen. Meestal gebruikt door de hostcode, maar nieuwere GPU-modellen hebben er ook toegang toe. Wanneer een kernel toegang krijgt tot het hostgeheugen, moet de GPU communiceren met het moederbord, meestal via de PCIe-connector en als zodanig is het relatief langzaam.
  • Apparaatgeheugen / Wereldwijd geheugen - het belangrijkste off-chip geheugen van de GPU, beschikbaar voor alle threads.
  • Gedeeld geheugen - bevindt zich in elke SM voor veel snellere toegang dan wereldwijd. Gedeeld geheugen is privé voor elk blok. Threads binnen een enkel blok kunnen het gebruiken voor communicatie.
  • Registers - snelste, privé, niet-adresseerbaar geheugen van elke thread. Over het algemeen kunnen deze niet worden gebruikt voor communicatie, maar een paar intrinsieke functies maken het mogelijk om hun inhoud binnen een warp te schudden.
  • Lokale geheugen - private geheugen van elke draad die is geadresseerd. Dit wordt gebruikt voor morsen van registers en lokale arrays met variabele indexering. Fysiek bevinden ze zich in het wereldwijde geheugen.
  • Textuurgeheugen, constant geheugen - een deel van het globale geheugen dat is gemarkeerd als onveranderlijk voor de kernel. Hierdoor kan de GPU caches voor speciale doeleinden gebruiken.
  • L2-cache - on-chip, beschikbaar voor alle threads. Gezien de hoeveelheid threads, is de verwachte levensduur van elke cache-lijn veel lager dan op de CPU. Het wordt meestal gebruikt voor verkeerd uitgelijnde en gedeeltelijk willekeurige geheugentoegangspatronen.
  • L1-cache - bevindt zich in dezelfde ruimte als het gedeelde geheugen. Nogmaals, het aantal is vrij klein, gezien het aantal threads dat het gebruikt, dus verwacht niet dat gegevens daar lang zullen blijven. L1-caching kan worden uitgeschakeld.

versies

Berekenbaarheid architectuur GPU-codenaam Publicatiedatum
1.0 Tesla G80 2006-11-08
1.1 Tesla G84, G86, G92, G94, G96, G98, 2007-04-17
1.2 Tesla GT218, GT216, GT215 2009-04-01
1.3 Tesla GT200, GT200b 2009-04-09
2.0 Fermi GF100, GF110 2010-03-26
2.1 Fermi GF104, GF106 GF108, GF114, GF116, GF117, GF119 2010-07-12
3.0 Kepler GK104, GK106, GK107 2012-03-22
3.2 Kepler GK20A 2014/04/01
3.5 Kepler GK110, GK208 2013/02/19
3.7 Kepler GK210 2014/11/17
5.0 Maxwell GM107, GM108 2014/02/18
5.2 Maxwell GM200, GM204, GM206 2014/09/18
5.3 Maxwell GM20B 2015/04/01
6.0 Pascal GP100 2016/10/01
6.1 Pascal GP102, GP104, GP106 2016/05/27

De releasedatum markeert de release van de eerste GPU die gegeven rekenmogelijkheden ondersteunt. Sommige data zijn bij benadering, bijv. 3.2 kaart is vrijgegeven in Q2 2014.

De voorbeeldprogramma's compileren en uitvoeren

De NVIDIA-installatiehandleiding eindigt met het uitvoeren van de voorbeeldprogramma's om uw installatie van de CUDA Toolkit te verifiëren, maar geeft niet expliciet aan hoe. Controleer eerst alle vereisten. Controleer de standaard CUDA-directory voor de voorbeeldprogramma's. Als het niet aanwezig is, kan het worden gedownload van de officiële CUDA-website. Navigeer naar de map waar de voorbeelden aanwezig zijn.

$ cd /path/to/samples/
$ ls
 

U zou een uitvoer moeten zien die lijkt op:

0_Simple     2_Graphics  4_Finance      6_Advanced       bin     EULA.txt
1_Utilities  3_Imaging   5_Simulations  7_CUDALibraries  common  Makefile
 

Zorg ervoor dat de Makefile aanwezig is in deze map. De opdracht make in op UNIX gebaseerde systemen bouwt alle voorbeeldprogramma's. Of navigeer naar een submap waar een andere Makefile aanwezig is en voer het make commando van daar uit om alleen dat monster te bouwen.

Voer de twee voorgestelde voorbeeldprogramma's uit - deviceQuery en bandwidthTest :

$ cd 1_Utilities/deviceQuery/
$ ./deviceQuery 
 

De uitvoer is vergelijkbaar met die hieronder:

./deviceQuery Starting...

 CUDA Device Query (Runtime API) version (CUDART static linking)

Detected 1 CUDA Capable device(s)

Device 0: "GeForce GTX 950M"
  CUDA Driver Version / Runtime Version          7.5 / 7.5
  CUDA Capability Major/Minor version number:    5.0
  Total amount of global memory:                 4096 MBytes (4294836224 bytes)
  ( 5) Multiprocessors, (128) CUDA Cores/MP:     640 CUDA Cores
  GPU Max Clock rate:                            1124 MHz (1.12 GHz)
  Memory Clock rate:                             900 Mhz
  Memory Bus Width:                              128-bit
  L2 Cache Size:                                 2097152 bytes
  Maximum Texture Dimension Size (x,y,z)         1D=(65536), 2D=(65536, 65536), 3D=(4096, 4096, 4096)
  Maximum Layered 1D Texture Size, (num) layers  1D=(16384), 2048 layers
  Maximum Layered 2D Texture Size, (num) layers  2D=(16384, 16384), 2048 layers
  Total amount of constant memory:               65536 bytes
  Total amount of shared memory per block:       49152 bytes
  Total number of registers available per block: 65536
  Warp size:                                     32
  Maximum number of threads per multiprocessor:  2048
  Maximum number of threads per block:           1024
  Max dimension size of a thread block (x,y,z): (1024, 1024, 64)
  Max dimension size of a grid size    (x,y,z): (2147483647, 65535, 65535)
  Maximum memory pitch:                          2147483647 bytes
  Texture alignment:                             512 bytes
  Concurrent copy and kernel execution:          Yes with 1 copy engine(s)
  Run time limit on kernels:                     Yes
  Integrated GPU sharing Host Memory:            No
  Support host page-locked memory mapping:       Yes
  Alignment requirement for Surfaces:            Yes
  Device has ECC support:                        Disabled
  Device supports Unified Addressing (UVA):      Yes
  Device PCI Domain ID / Bus ID / location ID:   0 / 1 / 0
  Compute Mode:
     < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >

deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 7.5, CUDA Runtime Version = 7.5, NumDevs = 1, Device0 = GeForce GTX 950M
Result = PASS
 

De instructie Result = PASS aan het einde geeft aan dat alles correct werkt. Voer nu het andere voorgestelde voorbeeldprogramma van bandwidthTest op dezelfde manier uit. De uitvoer is vergelijkbaar met:

[CUDA Bandwidth Test] - Starting...
Running on...

 Device 0: GeForce GTX 950M
 Quick Mode

 Host to Device Bandwidth, 1 Device(s)
 PINNED Memory Transfers
   Transfer Size (Bytes)    Bandwidth(MB/s)
   33554432            10604.5

 Device to Host Bandwidth, 1 Device(s)
 PINNED Memory Transfers
   Transfer Size (Bytes)    Bandwidth(MB/s)
   33554432            10202.0

 Device to Device Bandwidth, 1 Device(s)
 PINNED Memory Transfers
   Transfer Size (Bytes)    Bandwidth(MB/s)
   33554432            23389.7

Result = PASS

NOTE: The CUDA Samples are not meant for performance measurements. Results may vary when GPU Boost is enabled.
 

Nogmaals, de instructie Result = PASS geeft aan dat alles correct is uitgevoerd. Alle andere voorbeeldprogramma's kunnen op dezelfde manier worden uitgevoerd.

Laten we een enkele CUDA-thread starten om hallo te zeggen

Dit eenvoudige CUDA-programma demonstreert hoe een functie te schrijven die op de GPU wordt uitgevoerd (ook bekend als "apparaat"). De CPU, of "host", creëert CUDA-threads door speciale functies aan te roepen die "kernels" worden genoemd. CUDA-programma's zijn C ++ -programma's met extra syntaxis.

Om te zien hoe het werkt, plaats je de volgende code in een bestand met de naam hello.cu :

#include <stdio.h>

// __global__ functions, or "kernels", execute on the device
__global__ void hello_kernel(void)
{
  printf("Hello, world from the device!\n");
}

int main(void)
{
  // greet from the host
  printf("Hello, world from the host!\n");

  // launch a kernel with a single thread to greet from the device
  hello_kernel<<<1,1>>>();

  // wait for the device to finish so that we see the message
  cudaDeviceSynchronize();

  return 0;
}
 

(Let op: om de printf functie op het apparaat te gebruiken, hebt u een apparaat nodig met een rekencapaciteit van minimaal 2.0. Zie het versieoverzicht voor meer informatie.)

Laten we nu het programma compileren met de NVIDIA-compiler en het uitvoeren:

$ nvcc hello.cu -o hello
$ ./hello
Hello, world from the host!
Hello, world from the device!
 

Enkele aanvullende informatie over het bovenstaande voorbeeld:

  • nvcc staat voor "NVIDIA CUDA Compiler". Het scheidt de broncode in host- en apparaatcomponenten.
  • __global__ is een CUDA-sleutelwoord dat wordt gebruikt in functieverklaringen die aangeven dat de functie wordt uitgevoerd op het GPU-apparaat en wordt aangeroepen door de host.
  • Drievoudige punthaken ( <<< , >>> ) markeren een aanroep van hostcode naar apparaatcode (ook wel "kernel launch" genoemd). De nummers tussen deze drievoudige haakjes geven het aantal keren dat parallel moet worden uitgevoerd en het aantal threads aan.

voorwaarden

Download en installeer de CUDA Toolkit en het ontwikkelaarstuurprogramma om te beginnen met programmeren met CUDA. De toolkit bevat nvcc , de NVIDIA CUDA Compiler en andere software die nodig is om CUDA-toepassingen te ontwikkelen. Het stuurprogramma zorgt ervoor dat GPU-programma's correct worden uitgevoerd op CUDA-compatibele hardware , die u ook nodig hebt.

U kunt bevestigen dat de CUDA Toolkit correct op uw machine is geïnstalleerd door nvcc --version uit te voeren vanaf een opdrachtregel. Op een Linux-machine bijvoorbeeld

$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2016 NVIDIA Corporation
Built on Tue_Jul_12_18:28:38_CDT_2016
Cuda compilation tools, release 8.0, V8.0.32
 

voert de compilerinformatie uit. Als de vorige opdracht niet is geslaagd, is de CUDA Toolkit waarschijnlijk niet geïnstalleerd of maakt het pad naar nvcc ( C:\CUDA\bin op Windows-machines, /usr/local/cuda/bin op POSIX-besturingssystemen) geen deel uit van uw PATH omgevingsvariabele.

Bovendien hebt u ook een host-compiler nodig die met nvcc werkt om CUDA-programma's te compileren en te bouwen. In Windows is dit cl.exe , de Microsoft-compiler, die wordt geleverd met Microsoft Visual Studio. Op POSIX-besturingssystemen zijn andere compilers beschikbaar, waaronder gcc of g++ . De officiële CUDA Quick Start Guide kan u vertellen welke compilerversies worden ondersteund op uw specifieke platform.

Laten we, om ervoor te zorgen dat alles correct is ingesteld, een triviaal CUDA-programma compileren en uitvoeren om ervoor te zorgen dat alle tools correct samenwerken.

__global__ void foo() {}

int main()
{
  foo<<<1,1>>>();

  cudaDeviceSynchronize();
  printf("CUDA error: %s\n", cudaGetErrorString(cudaGetLastError()));

  return 0;
}
 

Om dit programma te compileren, kopieert u het naar een bestand met de naam test.cu en compileert u het vanaf de opdrachtregel. Op een Linux-systeem zou bijvoorbeeld het volgende moeten werken:

$ nvcc test.cu -o test
$ ./test
CUDA error: no error
 

Als het programma zonder fouten slaagt, laten we beginnen met coderen!

Som twee arrays op met CUDA

Dit voorbeeld illustreert hoe u een eenvoudig programma maakt dat twee int arrays met CUDA optelt.

Een CUDA-programma is heterogeen en bestaat uit onderdelen die zowel op CPU als GPU worden uitgevoerd.

De hoofdonderdelen van een programma dat CUDA gebruikt, zijn vergelijkbaar met CPU-programma's en bestaan uit

  • Geheugentoewijzing voor gegevens die op de GPU worden gebruikt
  • Gegevens kopiëren van hostgeheugen naar GPU's-geheugen
  • Het oproepen van de kernelfunctie om gegevens te verwerken
  • Kopieer het resultaat naar het CPU-geheugen

Om apparatengeheugen toe te wijzen, gebruiken we de cudaMalloc functie. Om gegevens tussen het apparaat en de host te cudaMemcpy , kan de cudaMemcpy functie worden gebruikt. Het laatste argument van cudaMemcpy geeft de richting van de kopieerbewerking aan. Er zijn 5 mogelijke types:

  • cudaMemcpyHostToHost - Host -> Host
  • cudaMemcpyHostToDevice - Host -> Apparaat
  • cudaMemcpyDeviceToHost - Apparaat -> Host
  • cudaMemcpyDeviceToDevice - Apparaat -> Apparaat
  • cudaMemcpyDefault - Standaardgebaseerde verenigde virtuele adresruimte

Vervolgens wordt de kernelfunctie aangeroepen. De informatie tussen de drievoudige chevrons is de uitvoeringsconfiguratie, die bepaalt hoeveel apparaatthreads de kernel parallel uitvoeren. Het eerste nummer ( 2 in het voorbeeld) geeft het aantal blokken aan en het tweede ( (size + 1) / 2 in het voorbeeld) - aantal threads in een blok. Merk op dat we in dit voorbeeld 1 aan de grootte toevoegen, zodat we één extra thread aanvragen in plaats van dat een thread verantwoordelijk is voor twee elementen.

Omdat kernel-aanroep een asynchrone functie is, wordt cudaDeviceSynchronize opgeroepen om te wachten tot de uitvoering is voltooid. Resultaatmatrices worden naar het cudaFree gekopieerd en alle geheugen dat op het apparaat is toegewezen, wordt vrijgemaakt met cudaFree .

Om de functie als kernel te definiëren, wordt de __global__ aangifte gebruikt. Deze functie wordt door elke thread opgeroepen. Als we willen dat elke thread een element van de resulterende array verwerkt, dan hebben we een manier nodig om elke thread te onderscheiden en identificeren. CUDA definieert de variabelen blockDim , blockIdx en threadIdx . De vooraf gedefinieerde variabele blockDim bevat de dimensies van elk blockDim zoals opgegeven in de tweede configuratieparameter voor de uitvoering van de kernel. De vooraf gedefinieerde variabelen threadIdx en blockIdx bevatten respectievelijk de index van de thread in het threadIdx en het blockIdx in het raster. Merk op dat, omdat we mogelijk om één thread meer vragen dan elementen in de arrays, we in size moeten doorgeven om ervoor te zorgen dat we geen toegang hebben voorbij het einde van de array.

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>

__global__ void addKernel(int* c, const int* a, const int* b, int size) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < size) {
        c[i] = a[i] + b[i];
    }
}

// Helper function for using CUDA to add vectors in parallel.
void addWithCuda(int* c, const int* a, const int* b, int size) {
    int* dev_a = nullptr;
    int* dev_b = nullptr;
    int* dev_c = nullptr;

    // Allocate GPU buffers for three vectors (two input, one output)
    cudaMalloc((void**)&dev_c, size * sizeof(int));
    cudaMalloc((void**)&dev_a, size * sizeof(int));
    cudaMalloc((void**)&dev_b, size * sizeof(int));

    // Copy input vectors from host memory to GPU buffers.
    cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);

    // Launch a kernel on the GPU with one thread for each element.
    // 2 is number of computational blocks and (size + 1) / 2 is a number of threads in a block
    addKernel<<<2, (size + 1) / 2>>>(dev_c, dev_a, dev_b, size);
    
    // cudaDeviceSynchronize waits for the kernel to finish, and returns
    // any errors encountered during the launch.
    cudaDeviceSynchronize();

    // Copy output vector from GPU buffer to host memory.
    cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);

    cudaFree(dev_c);
    cudaFree(dev_a);
    cudaFree(dev_b);
}

int main(int argc, char** argv) {
    const int arraySize = 5;
    const int a[arraySize] = {  1,  2,  3,  4,  5 };
    const int b[arraySize] = { 10, 20, 30, 40, 50 };
    int c[arraySize] = { 0 };

    addWithCuda(c, a, b, arraySize);

    printf("{1, 2, 3, 4, 5} + {10, 20, 30, 40, 50} = {%d, %d, %d, %d, %d}\n", c[0], c[1], c[2], c[3], c[4]);

    cudaDeviceReset();

    return 0;
}