cudaKomma igång med cuda


Anmärkningar

CUDA är en egen NVIDIA parallell datorteknik och programmeringsspråk för sina GPU: er.

GPU: er är mycket parallella maskiner som kan köra tusentals lätta trådar parallellt. Varje GPU-tråd är vanligtvis långsammare i körning och deras sammanhang är mindre. Å andra sidan kan GPU köra flera tusentals trådar parallellt och ännu mer samtidigt (exakta antal beror på den faktiska GPU-modellen). CUDA är en C ++ -dialekt designad specifikt för NVIDIA GPU-arkitektur. På grund av skillnaderna i arkitektur kan de flesta algoritmer emellertid inte bara kopieras in från vanlig C ++ - de skulle köra, men skulle vara väldigt långsamma.

Terminologi

  • värd - avser normal CPU-baserad hårdvara och normala program som körs i den miljön
  • enhet - avser en specifik GPU som CUDA-program körs i. En enda värd kan stödja flera enheter.
  • kernel - en funktion som finns på enheten som kan åberopas från värdkoden.

Fysisk processorstruktur

Den CUDA-aktiverade GPU-processorn har följande fysiska struktur:

  • chipet - hela processorn på GPU. Vissa GPU: er har två av dem.
  • strömmande multiprocessor (SM) - varje chip innehåller upp till ~ 100 SM, beroende på en modell. Varje SM arbetar nästan oberoende av en annan och använder endast ett globalt minne för att kommunicera med varandra.
  • CUDA-kärna - en enda skalbar datorenhet för en SM. Deras exakta antal beror på arkitekturen. Varje kärna kan hantera några trådar som körs samtidigt i en snabb följd (liknar hypertrådning i CPU).

Dessutom har varje SM en eller flera varpschemaläggare . Varje schemaläggare skickar en enda instruktion till flera CUDA-kärnor. Detta gör att SM fungerar effektivt i 32-brett SIMD- läge.

CUDA exekveringsmodell

GPU: s fysiska struktur har direkt påverkan på hur kärnor körs på enheten och hur man programmerar dem i CUDA. Kärnan åberopas med en samtalskonfiguration som anger hur många parallella trådar som spawnas.

  • rutnätet - representerar alla trådar som spawnas vid kärnanrop. Det anges som en eller två dimensionell uppsättning block
  • blocket - är en semi-oberoende uppsättning trådar . Varje block tilldelas en enda SM. Som sådant kan block bara kommunicera genom det globala minnet. Block synkroniseras inte på något sätt. Om det finns för många block kan vissa köra sekventiellt efter andra. Å andra sidan, om resurser tillåter det, kan mer än ett block köra på samma SM, men programmeraren kan inte dra nytta av att det händer (förutom det uppenbara prestationsökningen).
  • tråden - en skalär sekvens av instruktioner som utförs av en enda CUDA-kärna. Trådar är "lätta" med minimal kontext, vilket gör att hårdvaran snabbt kan byta in och ut dem. På grund av deras antal fungerar CUDA-trådar med några få register tilldelade dem, och mycket kort stack (helst inget alls!). Av den anledningen föredrar CUDA-kompilatorn att inlineta alla funktionssamtal för att platta kärnan så att den bara innehåller statiska hopp och slingor. Funktionspontersamtal och virtuella metodsamtal, medan de stöds i de flesta nyare enheter, har vanligtvis en stor prestationsstraff.

Varje tråd identifieras av ett blockIdx och blockIdx i block threadIdx . Dessa nummer kan när som helst kontrolleras av valfri tråd och är det enda sättet att skilja en tråd från en annan.

Dessutom är trådar organiserade i varp , var och en innehåller exakt 32 trådar. Trådar inom en enda varp körs i en perfekt synk, i SIMD-fahsion. Trådar från olika varp, men inom samma block kan köras i valfri ordning, men kan tvingas synkroniseras av programmeraren. Trådar från olika block kan inte synkroniseras eller interagera direkt på något sätt.

Minneorganisation

Vid normal CPU-programmering är minnesorganisationen vanligtvis dold för programmeraren. Typiska program fungerar som om det bara fanns RAM. Alla minnesoperationer, t.ex. hantering av register, användning av L1-L2-L3-caching, byte till disk, etc. hanteras av kompilatorn, operativsystemet eller hårdvaran själv.

Detta är inte fallet med CUDA. Medan nyare GPU-modeller delvis döljer bördan, t.ex. genom Unified Memory i CUDA 6, är det fortfarande värt att förstå organisationen av prestandaskäl. Den grundläggande CUDA-minnesstrukturen är följande:

  • Värdminne - det vanliga RAM- minnet . Vanligtvis används av värdkoden, men nyare GPU-modeller kan också komma åt den. När en kärna får åtkomst till värdminnet måste GPU kommunicera med moderkortet, vanligtvis via PCIe-anslutningen och som sådan är det relativt långsamt.
  • Enhetsminne / Globalt minne - GPU: s huvudminne utanför chip, tillgängligt för alla trådar.
  • Delat minne - som finns i varje SM möjliggör mycket snabbare åtkomst än global. Delat minne är privat för varje block. Trådar inom ett enda block kan använda det för kommunikation.
  • Registreringar - snabbaste, privata, oadresserbara minne för varje tråd. I allmänhet kan dessa inte användas för kommunikation, men några få funktioner innebär att blanda innehållet i ett varp.
  • Lokalt minne - privat minne för varje tråd som är adresserbar. Detta används för registerutsläpp och lokala matriser med variabel indexering. De är fysiskt bosatta i det globala minnet.
  • Texturminne, konstant minne - en del av det globala minnet som är markerat som oföränderligt för kärnan. Detta tillåter GPU att använda cacheminnesinlägg för speciella ändamål.
  • L2-cache - on-chip, tillgängligt för alla trådar. Med tanke på mängden trådar är den förväntade livslängden för varje cachelinje mycket lägre än på CPU. Det används mestadels missanpassade och delvis slumpmässiga minnesåtkomstmönster.
  • L1-cache - ligger i samma utrymme som delat minne. Återigen är mängden ganska litet med tanke på antalet trådar som använder det, så förvänta dig inte att data kommer att stanna där länge. L1-cachning kan inaktiveras.

versioner

Beräkningsförmåga Arkitektur GPU-kodnamn Utgivningsdatum
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

Släppdatumet markerar frisläppandet av den första GPU som stöder en given beräkningsförmåga. Vissa datum är ungefärliga, till exempel släpptes 3,2 kort under andra kvartalet 2014.

Kompilera och köra provprogrammen

NVIDIA installationsguide slutar med att köra provprogrammen för att verifiera din installation av CUDA Toolkit, men anger inte uttryckligen hur. Kontrollera först alla förutsättningar. Kontrollera standardkatalogen för CUDA för exempelprogrammen. Om den inte finns kan den laddas ner från den officiella CUDA-webbplatsen. Navigera till katalogen där exemplen finns.

$ cd /path/to/samples/
$ ls
 

Du bör se en utgång som liknar:

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

Se till att Makefile finns i den här katalogen. Kommandot make i UNIX-baserade system kommer att bygga alla exempelprogram. Alternativt navigerar du till en underkatalog där en annan Makefile är närvarande och kör make kommandot därifrån för att bara bygga det exemplet.

Kör de två föreslagna provprogrammen - deviceQuery och bandwidthTest :

$ cd 1_Utilities/deviceQuery/
$ ./deviceQuery 
 

Utgången kommer att likna den som visas nedan:

./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
 

Uttalandet Result = PASS i slutet indikerar att allt fungerar korrekt. Kör nu det andra föreslagna provprogrammet bandwidthTest på liknande sätt. Utgången kommer att likna:

[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.
 

Återigen indikerar Result = PASS att allt kördes korrekt. Alla andra provprogram kan köras på liknande sätt.

Låt oss lansera en enda CUDA-tråd för att säga hej

Detta enkla CUDA-program visar hur man skriver en funktion som kommer att köras på GPU (alias ”enhet”). CPU, eller "värd", skapar CUDA-trådar genom att kalla speciella funktioner som kallas "kärnor". CUDA-program är C ++ -program med ytterligare syntax.

För att se hur det fungerar, lägg följande kod i en fil med namnet 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;
}
 

(Observera att för att kunna använda printf funktionen på enheten behöver du en enhet som har en datorkapacitet på minst 2,0. Se versionens översikt för mer information.)

Låt oss nu sammanställa programmet med NVIDIA-kompilatorn och köra det:

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

Några ytterligare information om exemplet ovan:

  • nvcc står för "NVIDIA CUDA Compiler". Den separerar källkoden i värd- och enhetskomponenter.
  • __global__ är ett CUDA-nyckelord som används i funktionsdeklarationer som indikerar att funktionen körs på GPU-enheten och kallas från värden.
  • Trippelvinklar ( <<< , >>> ) markerar ett samtal från värdkod till enhetskod (även kallad "kärnlansering"). Siffrorna inom dessa treparenteser anger antalet gånger som ska köras parallellt och antalet trådar.

förutsättningar

För att komma igång med programmeringen med CUDA, ladda ner och installera CUDA Toolkit och utvecklardrivrutin . Verktygssatsen inkluderar nvcc , NVIDIA CUDA Compiler och annan programvara som är nödvändig för att utveckla CUDA-applikationer. Drivrutinen ser till att GPU-program körs korrekt på CUDA-kapabel hårdvara , vilket du också behöver.

Du kan bekräfta att CUDA Toolkit är korrekt installerat på din maskin genom att köra nvcc --version från en kommandorad. Till exempel på en Linux-maskin,

$ 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
 

matar ut kompilatorinformationen. Om det föregående kommandot inte lyckades är CUDA Toolkit sannolikt inte installerat, eller sökvägen till nvcc ( C:\CUDA\bin på Windows-maskiner, /usr/local/cuda/bin på POSIX OSes) är inte en del av din PATH miljövariabel.

Dessutom behöver du också en värdkompilerare som fungerar med nvcc att kompilera och bygga CUDA-program. På Windows är detta cl.exe , Microsoft-kompilatorn, som levereras med Microsoft Visual Studio. På POSIX-operativsystem finns andra kompilatorer tillgängliga, inklusive gcc eller g++ . Den officiella CUDA Quick Start Guide kan berätta vilka kompilatorversioner som stöds på din specifika plattform.

För att se till att allt är korrekt konfigurerat, låt oss sammanställa och köra ett triviellt CUDA-program för att säkerställa att alla verktyg fungerar korrekt.

__global__ void foo() {}

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

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

  return 0;
}
 

För att kompilera detta program, kopiera det till en fil som heter test.cu och kompilera det från kommandoraden. Till exempel på ett Linux-system bör följande fungera:

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

Om programmet lyckas utan fel, låt oss börja kodningen!

Sammanfatta två matriser med CUDA

Detta exempel illustrerar hur man skapar ett enkelt program som kommer att sammanfatta två int arrays med CUDA.

Ett CUDA-program är heterogent och består av delar körs både på CPU och GPU.

Huvuddelen av ett program som använder CUDA liknar CPU-program och består av

  • Minneallokering för data som kommer att användas på GPU
  • Datakopiering från värdminne till GPU: s minne
  • Åkallar kärnfunktionen för att bearbeta data
  • Kopiera resultatet till CPU: s minne

För att tilldela cudaMalloc använder vi cudaMalloc funktionen. För att kopiera data mellan enhet och värd kan cudaMemcpy funktion användas. Det sista argumentet från cudaMemcpy anger riktningen för kopiering. Det finns 5 möjliga typer:

  • cudaMemcpyHostToHost - Host -> Host
  • cudaMemcpyHostToDevice - Host -> Enhet
  • cudaMemcpyDeviceToHost - Enhet -> Värd
  • cudaMemcpyDeviceToDevice - Enhet -> Enhet
  • cudaMemcpyDefault - Standardbaserat enhetligt virtuellt adressutrymme

Därefter aktiveras kärnfunktionen. Informationen mellan trippelkedjorna är exekveringskonfigurationen, som dikterar hur många enhetstrådar som kör kärnan parallellt. Det första numret ( 2 i exempel) anger antalet block och det andra ( (size + 1) / 2 i exempel) - antal trådar i ett block. Observera att i detta exempel lägger vi till 1 till storleken, så att vi begär en extra tråd snarare än att ha en tråd ansvarig för två element.

Eftersom kärnanropning är en asynkron funktion cudaDeviceSynchronize för att vänta tills körningen är klar. Resultatmatriser kopieras till värdminnet och allt minne som tilldelats på enheten frigörs med cudaFree .

För att definiera funktion som kärnan används __global__ deklarationsspecifikation. Denna funktion kommer att påkallas av varje tråd. Om vi vill att varje tråd ska bearbeta ett element i den resulterande matrisen, behöver vi ett sätt att skilja och identifiera varje tråd. CUDA definierar variablerna blockDim , blockIdx och threadIdx . Den fördefinierade variablen blockDim innehåller dimensionerna för varje trådblock som anges i den andra exekveringskonfigurationsparametern för kärnlanseringen. De fördefinierade variablerna threadIdx och blockIdx innehåller blockIdx index inom dess threadIdx respektive blockIdx i rutnätet. Observera att eftersom vi potentiellt begär ytterligare en tråd än element i matriserna, måste vi passera i size att säkerställa att vi inte kommer förbi slutet av matrisen.

#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;
}