cudaPierwsze kroki z cuda


Uwagi

CUDA to zastrzeżona technologia obliczeń równoległych NVIDIA i język programowania dla ich układów GPU.

Procesory graficzne to wysoce równoległe maszyny, zdolne do równoległego uruchamiania tysięcy lekkich wątków. Każdy wątek GPU działa zwykle wolniej, a jego kontekst jest mniejszy. Z drugiej strony GPU jest w stanie uruchomić kilka tysięcy wątków równolegle, a nawet bardziej równolegle (dokładne liczby zależą od rzeczywistego modelu GPU). CUDA to dialekt C ++ zaprojektowany specjalnie dla architektury GPU NVIDIA. Jednak ze względu na różnice w architekturze większość algorytmów nie można po prostu wkleić z poziomu zwykłego C ++ - działałyby, ale byłyby bardzo wolne.

Terminologia

  • host - odnosi się do normalnego sprzętu opartego na procesorze i normalnych programów działających w tym środowisku
  • urządzenie - odnosi się do konkretnego procesora graficznego, w którym działają programy CUDA. Jeden host może obsługiwać wiele urządzeń.
  • jądro - funkcja znajdująca się na urządzeniu, którą można wywołać z kodu hosta.

Struktura procesora fizycznego

Procesor GPU z obsługą CUDA ma następującą strukturę fizyczną:

  • chip - cały procesor GPU. Niektóre procesory graficzne mają dwa z nich.
  • streaming multiprocessor (SM) - każdy układ zawiera do ~ 100 SM, w zależności od modelu. Każdy SM działa prawie niezależnie od siebie, wykorzystując tylko globalną pamięć do komunikowania się ze sobą.
  • Rdzeń CUDA - pojedyncza skalarna jednostka obliczeniowa SM. Ich dokładna liczba zależy od architektury. Każdy rdzeń może obsłużyć kilka wątków wykonywanych jednocześnie w krótkich odstępach czasu (podobnie jak hyperthreading w CPU).

Ponadto każde SM ma jeden lub więcej harmonogramów wypaczania . Każdy planista wysyła jedną instrukcję do kilku rdzeni CUDA. To skutecznie powoduje, że SM działa w 32- calowym trybie SIMD .

Model wykonania CUDA

Fizyczna struktura GPU ma bezpośredni wpływ na to, jak jądra są wykonywane na urządzeniu i jak programuje się je w CUDA. Jądro jest wywoływane z konfiguracją wywołania, która określa liczbę równoległych wątków.

  • siatka - reprezentuje wszystkie wątki, które pojawiają się po wywołaniu jądra. Jest określony jako jeden lub dwa wymiarowe zestawy bloków
  • blok - to pół-niezależny zestaw wątków . Każdy blok jest przypisany do jednego SM. Jako takie, bloki mogą komunikować się tylko poprzez pamięć globalną. Bloki nie są w żaden sposób synchronizowane. Jeśli bloków jest zbyt wiele, niektóre mogą być wykonywane sekwencyjnie po innych. Z drugiej strony, jeśli pozwalają na to zasoby, więcej niż jeden blok może działać na tym samym SM, ale programista nie może z tego skorzystać (z wyjątkiem oczywistego zwiększenia wydajności).
  • wątek - skalarna sekwencja instrukcji wykonywanych przez pojedynczy rdzeń CUDA. Wątki są „lekkie” z minimalnym kontekstem, co pozwala sprzętowi szybko je zamieniać i zamieniać. Ze względu na ich liczbę wątki CUDA działają z kilkoma przypisanymi do nich rejestrami i bardzo krótkim stosem (najlepiej wcale!). Z tego powodu kompilator CUDA woli wstawiać wszystkie wywołania funkcji w celu spłaszczenia jądra, tak aby zawierało tylko statyczne skoki i pętle. Wywołania funkcji Ponter i wirtualne wywołania metod, choć obsługiwane w większości nowszych urządzeń, zwykle wiążą się z dużym spadkiem wydajności.

Każdy wątek jest identyfikowany za pomocą indeksu bloku blockIdx i indeksu wątku w obrębie bloku threadIdx . Liczby te można sprawdzić w dowolnym momencie przez dowolny działający wątek i jest to jedyny sposób na odróżnienie jednego wątku od drugiego.

Ponadto wątki są zorganizowane w osnowy , z których każdy zawiera dokładnie 32 wątki. Wątki w ramach jednego warpu są wykonywane w doskonałej synchronizacji, w trybie SIMD. Wątki z różnych wypaczeń, ale w tym samym bloku, mogą być wykonywane w dowolnej kolejności, ale mogą być zmuszone do synchronizacji przez programistę. Wątków z różnych bloków nie można synchronizować ani bezpośrednio oddziaływać w żaden sposób.

Organizacja pamięci

W normalnym programowaniu procesora organizacja pamięci jest zwykle ukryta przed programistą. Typowe programy działają tak, jakby była tylko pamięć RAM. Wszystkie operacje pamięci, takie jak zarządzanie rejestrami, używanie buforowania L1-L2-L3-, zamiana na dysk itp. Są obsługiwane przez kompilator, system operacyjny lub sam sprzęt.

Nie jest tak w przypadku CUDA. Podczas gdy nowsze modele GPU częściowo ukrywają to obciążenie, np. Poprzez Ujednoliconą Pamięć w CUDA 6, nadal warto zrozumieć organizację ze względu na wydajność. Podstawowa struktura pamięci CUDA jest następująca:

  • Pamięć hosta - zwykła pamięć RAM. Używany głównie przez kod hosta, ale nowsze modele GPU również mogą uzyskać do niego dostęp. Gdy jądro uzyskuje dostęp do pamięci hosta, GPU musi komunikować się z płytą główną, zwykle przez złącze PCIe i jako takie jest stosunkowo wolne.
  • Pamięć urządzenia / pamięć globalna - główna pamięć poza układem graficznym, dostępna dla wszystkich wątków.
  • Pamięć współdzielona - umiejscowiona w każdym SM pozwala na znacznie szybszy dostęp niż globalny. Pamięć współdzielona jest prywatna dla każdego bloku. Wątki w jednym bloku mogą wykorzystywać go do komunikacji.
  • Rejestry - najszybsza, prywatna, nieadresowalna pamięć każdego wątku. Zasadniczo nie można ich używać do komunikacji, ale kilka wewnętrznych funkcji pozwala tasować ich zawartość w warp.
  • Pamięć lokalna - prywatna pamięć każdego wątku, który można adresować. Służy to do rejestrowania wycieków i lokalnych tablic ze zmiennym indeksowaniem. Fizycznie znajdują się w pamięci globalnej.
  • Pamięć tekstur, Pamięć stała - część pamięci globalnej oznaczona jako niezmienna dla jądra. Dzięki temu GPU może używać pamięci podręcznych specjalnego przeznaczenia.
  • Pamięć podręczna L2 - na chipie, dostępna dla wszystkich wątków. Biorąc pod uwagę liczbę wątków, oczekiwany czas życia każdej linii pamięci podręcznej jest znacznie krótszy niż na procesorze. Jest to najczęściej stosowane w celu niedopasowania i częściowo losowych wzorców dostępu do pamięci.
  • Pamięć podręczna L1 - znajduje się w tym samym miejscu, co pamięć współdzielona. Ponownie, ilość jest raczej niewielka, biorąc pod uwagę liczbę używających ją wątków, więc nie oczekuj, że dane pozostaną tam długo. Buforowanie L1 można wyłączyć.

Wersje

Możliwości obliczeniowe Architektura Nazwa kodowa GPU Data wydania
1.0 Tesla G80 2006-11-08
1.1 Tesla G84, G86, G92, G94, G96, G98, 17.04.2007
1.2 Tesla GT218, GT216, GT215 01.04.2009
1.3 Tesla GT200, GT200b 2009-04-09
2.0 Fermi GF100, GF110 26.03.2010
2.1 Fermi GF104, GF106 GF108, GF114, GF116, GF117, GF119 2010-07-12
3.0 Kepler GK104, GK106, GK107 22.03.2012
3.2 Kepler GK20A 01.04.2014
3.5 Kepler GK110, GK208 19.02.2013
3.7 Kepler GK210 17.11.2014
5.0 Maxwell GM107, GM108 18.02.2014
5.2 Maxwell GM200, GM204, GM206 2014-09-18
5.3 Maxwell GM20B 2015-04-01
6.0 Pascal GP100 01.10.2016
6.1 Pascal GP102, GP104, GP106 27.05.2016

Data wydania oznacza wydanie pierwszego procesora graficznego obsługującego dane możliwości obliczeniowe. Niektóre daty są przybliżone, np. Wydano kartę 3.2 w drugim kwartale 2014 roku.

Kompilowanie i uruchamianie przykładowych programów

Przewodnik instalacji NVIDIA kończy się uruchomieniem przykładowych programów w celu zweryfikowania instalacji CUDA Toolkit, ale nie określa wyraźnie, w jaki sposób. Najpierw sprawdź wszystkie wymagania wstępne. Sprawdź domyślny katalog CUDA dla przykładowych programów. Jeśli go nie ma, można go pobrać z oficjalnej strony internetowej CUDA. Przejdź do katalogu, w którym znajdują się przykłady.

$ cd /path/to/samples/
$ ls
 

Powinieneś zobaczyć wynik podobny do:

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

Upewnij się, że Makefile znajduje się w tym katalogu. Komenda make w systemach UNIX zbuduje wszystkie przykładowe programy. Ewentualnie przejdź do podkatalogu, w którym obecny jest inny plik Makefile i stamtąd uruchom komendę make aby zbudować tylko tę próbkę.

Uruchom dwa sugerowane programy przykładowe - deviceQuery i bandwidthTest deviceQuery :

$ cd 1_Utilities/deviceQuery/
$ ./deviceQuery 
 

Dane wyjściowe będą podobne do pokazanych poniżej:

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

Instrukcja Result = PASS na końcu wskazuje, że wszystko działa poprawnie. Teraz uruchom inny sugerowany test bandwidthTest przykładowego programu w podobny sposób. Dane wyjściowe będą podobne do:

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

Ponownie instrukcja Result = PASS wskazuje, że wszystko zostało wykonane poprawnie. Wszystkie inne przykładowe programy można uruchamiać w podobny sposób.

Uruchommy pojedynczy wątek CUDA, aby się przywitać

Ten prosty program CUDA pokazuje, jak napisać funkcję, która zostanie wykonana na GPU (inaczej „urządzenie”). CPU lub „host” tworzy wątki CUDA, wywołując specjalne funkcje zwane „jądrem”. Programy CUDA to programy C ++ z dodatkową składnią.

Aby zobaczyć, jak to działa, umieść następujący kod w pliku o nazwie 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;
}
 

(Należy pamiętać, że do korzystania z funkcji printf na urządzeniu potrzebne jest urządzenie, które ma zdolność obliczeniową co najmniej 2,0. Szczegółowe informacje zawiera przegląd wersji ).

Teraz skompilujmy program za pomocą kompilatora NVIDIA i uruchommy go:

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

Kilka dodatkowych informacji o powyższym przykładzie:

  • nvcc oznacza „NVIDIA CUDA Compiler”. Oddziela kod źródłowy na elementy hosta i urządzenia.
  • __global__ to słowo kluczowe CUDA używane w deklaracjach funkcji, wskazujące, że funkcja działa na urządzeniu GPU i jest wywoływana z hosta.
  • Trójkątne nawiasy ( <<< , >>> ) oznaczają wywołanie z kodu hosta na kod urządzenia (zwane również „uruchomieniem jądra”). Liczby w tych potrójnych nawiasach wskazują liczbę uruchomień równoległych i liczbę wątków.

Wymagania wstępne

Aby rozpocząć programowanie w CUDA, pobierz i zainstaluj CUDA Toolkit i sterownik programisty . Zestaw zawiera nvcc , kompilator NVIDIA CUDA i inne oprogramowanie niezbędne do tworzenia aplikacji CUDA. Sterownik zapewnia, że programy GPU działają poprawnie na sprzęcie obsługującym CUDA , który również będzie ci potrzebny.

Możesz potwierdzić, że CUDA Toolkit jest poprawnie zainstalowany na twoim komputerze, uruchamiając nvcc --version z wiersza poleceń. Na przykład na komputerze z systemem Linux

$ 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
 

wyświetla informacje o kompilatorze. Jeśli poprzednie polecenie nie powiodło się, prawdopodobnie CUDA Toolkit nie jest zainstalowany lub ścieżka do nvcc ( C:\CUDA\bin na komputerach z systemem Windows, /usr/local/cuda/bin na systemach POSIX) nie jest częścią twojego Zmienna środowiskowa PATH .

Dodatkowo będziesz potrzebował kompilatora hosta, który współpracuje z nvcc do kompilacji i budowy programów CUDA. W systemie Windows jest to cl.exe , kompilator Microsoft, który jest dostarczany z Microsoft Visual Studio. W systemach POSIX dostępne są inne kompilatory, w tym gcc lub g++ . Oficjalny Skrócony przewodnik CUDA informuje, które wersje kompilatora są obsługiwane na konkretnej platformie.

Aby upewnić się, że wszystko jest poprawnie skonfigurowane, skompilujmy i uruchommy trywialny program CUDA, aby wszystkie narzędzia działały poprawnie.

__global__ void foo() {}

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

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

  return 0;
}
 

Aby skompilować ten program, skopiuj go do pliku o nazwie test.cu i skompiluj z wiersza poleceń. Na przykład w systemie Linux powinny działać następujące elementy:

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

Jeśli program powiedzie się bezbłędnie, zacznijmy kodować!

Zsumuj dwie tablice za pomocą CUDA

Ten przykład ilustruje, jak utworzyć prosty program, który zsumuje dwie tablice int z CUDA.

Program CUDA jest heterogeniczny i składa się z części działających zarówno na CPU, jak i na GPU.

Główne części programu wykorzystujące CUDA są podobne do programów CPU i składają się z

  • Przydział pamięci dla danych, które będą używane na GPU
  • Kopiowanie danych z pamięci hosta do pamięci GPU
  • Wywoływanie funkcji jądra do przetwarzania danych
  • Skopiuj wynik do pamięci procesorów

Do alokacji pamięci urządzeń używamy funkcji cudaMalloc . Aby skopiować dane między urządzeniem a hostem można cudaMemcpy funkcji cudaMemcpy . Ostatni argument cudaMemcpy określa kierunek operacji kopiowania. Istnieje 5 możliwych typów:

  • cudaMemcpyHostToHost - Host -> Host
  • cudaMemcpyHostToDevice - Host -> Urządzenie
  • cudaMemcpyDeviceToHost - Urządzenie -> Host
  • cudaMemcpyDeviceToDevice - Urządzenie -> Urządzenie
  • cudaMemcpyDefault - Domyślna ujednolicona wirtualna przestrzeń adresowa

Następnie wywoływana jest funkcja jądra. Informacje między potrójnymi szewronami to konfiguracja wykonania, która określa, ile wątków urządzenia wykonuje jądro równolegle. Pierwsza liczba (w przykładzie 2 ) określa liczbę bloków, a druga ( (size + 1) / 2 w przykładzie (size + 1) / 2 ) - liczbę wątków w bloku. Zauważ, że w tym przykładzie dodajemy 1 do rozmiaru, dlatego żądamy jednego dodatkowego wątku zamiast jednego wątku odpowiedzialnego za dwa elementy.

Ponieważ wywołanie jądra jest funkcją asynchroniczną, cudaDeviceSynchronize jest cudaDeviceSynchronize aby poczekać na zakończenie wykonywania. Tablice wyników są kopiowane do pamięci hosta, a cała pamięć przydzielona na urządzeniu jest zwalniana dzięki cudaFree .

Aby zdefiniować funkcję jako jądro __global__ używany jest specyfikator deklaracji. Ta funkcja będzie wywoływana przez każdy wątek. Jeśli chcemy, aby każdy wątek przetwarzał element wynikowej tablicy, potrzebujemy sposobu na odróżnienie i zidentyfikowanie każdego wątku. CUDA definiuje zmienne blockDim , blockIdx i threadIdx . Predefiniowana zmienna blockDim zawiera wymiary każdego bloku wątku, jak określono w drugim parametrze konfiguracyjnym wykonania dla uruchomienia jądra. Predefiniowane zmienne threadIdx i blockIdx zawierają odpowiednio indeks wątku w jego bloku wątku i blok wątku w siatce. Zauważ, że ponieważ potencjalnie żądamy więcej wątków niż elementów w tablicach, musimy przekazać size aby upewnić się, że nie uzyskamy dostępu poza koniec tablicy.

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