cudaErste Schritte mit cuda


Bemerkungen

CUDA ist eine proprietäre NVIDIA Parallel Computing-Technologie und Programmiersprache für ihre GPUs.

GPUs sind hochparallele Maschinen, mit denen Tausende leichter Threads parallel ausgeführt werden können. Jeder GPU-Thread wird normalerweise langsamer ausgeführt und sein Kontext ist kleiner. Auf der anderen Seite kann GPU mehrere Tausend Threads parallel und sogar parallel ausführen (genaue Zahlen hängen vom tatsächlichen GPU-Modell ab). CUDA ist ein C ++ - Dialekt, der speziell für die NVIDIA-GPU-Architektur entwickelt wurde. Aufgrund der Architekturunterschiede können die meisten Algorithmen jedoch nicht einfach aus C ++ kopiert werden - sie würden ausgeführt, wären aber sehr langsam.

Terminologie

  • Host - bezieht sich auf normale CPU-basierte Hardware und normale Programme, die in dieser Umgebung ausgeführt werden
  • Gerät - bezieht sich auf eine bestimmte GPU, in der CUDA-Programme ausgeführt werden. Ein einzelner Host kann mehrere Geräte unterstützen.
  • Kernel - Eine Funktion, die sich auf dem Gerät befindet und vom Hostcode aus aufgerufen werden kann.

Physikalische Prozessorstruktur

Der CUDA-fähige GPU-Prozessor hat die folgende physische Struktur:

  • der Chip - der gesamte Prozessor der GPU. Einige GPUs haben zwei davon.
  • Streamming-Multiprozessor (SM) - Jeder Chip enthält je nach Modell bis zu ~ 100 SMs. Jedes SM arbeitet nahezu unabhängig voneinander und verwendet nur globalen Speicher, um miteinander zu kommunizieren.
  • CUDA-Kern - eine einzelne Skalar-Recheneinheit eines SM. Ihre genaue Anzahl hängt von der Architektur ab. Jeder Kern kann einige Threads verarbeiten, die gleichzeitig ausgeführt werden (ähnlich wie Hyperthreading in CPU).

Darüber hinaus verfügt jedes SM über einen oder mehrere Warp-Scheduler . Jeder Scheduler sendet eine einzelne Anweisung an mehrere CUDA-Kerne. Dies bewirkt effektiv, dass der SM im 32-breiten SIMD- Modus arbeitet.

CUDA-Ausführungsmodell

Die physische Struktur der GPU hat direkten Einfluss darauf, wie Kernel auf dem Gerät ausgeführt werden und wie sie in CUDA programmiert werden. Der Kernel wird mit einer Aufrufkonfiguration aufgerufen, die angibt, wie viele parallele Threads erzeugt werden.

  • Das Raster - stellt alle Threads dar, die beim Kernel-Aufruf erzeugt werden. Es wird als ein oder zwei dimensionale Blöcke angegeben
  • Der Block - ist ein semi-unabhängiger Satz von Threads . Jeder Block ist einem einzelnen SM zugeordnet. Daher können Blöcke nur über den globalen Speicher kommunizieren. Blöcke werden in keiner Weise synchronisiert. Wenn zu viele Blöcke vorhanden sind, werden einige nacheinander ausgeführt. Auf der anderen Seite können, wenn die Ressourcen dies zulassen, mehr als ein Block auf demselben SM ausgeführt werden, aber der Programmierer kann davon nicht profitieren (abgesehen von der offensichtlichen Leistungssteigerung).
  • der Thread - eine skalare Folge von Anweisungen, die von einem einzelnen CUDA-Kern ausgeführt werden. Threads sind "leicht" mit minimalem Kontext, sodass die Hardware sie schnell ein- und auswechseln kann. Aufgrund ihrer Anzahl arbeiten CUDA-Threads mit einigen zugewiesenen Registern und sehr kurzen Stapeln (vorzugsweise gar nicht!). Aus diesem Grund zieht es der CUDA-Compiler vor, alle Funktionsaufrufe einzubinden, um den Kernel so zu glätten, dass er nur statische Sprünge und Schleifen enthält. Funktions-Ponter-Aufrufe und Aufrufe von virtuellen Methoden werden zwar von den meisten neueren Geräten unterstützt, sind jedoch in der Regel mit einer erheblichen Leistungseinschränkung verbunden.

Jeder Thread wird durch einen Blockindex blockIdx und einen blockIdx innerhalb des Blockes threadIdx . Diese Nummern können jederzeit von jedem laufenden Thread überprüft werden und sind die einzige Möglichkeit, einen Thread von einem anderen zu unterscheiden.

Darüber hinaus sind Threads in Warps organisiert, die jeweils genau 32 Threads enthalten. Threads innerhalb eines einzelnen Warp werden in einer SIMD-Funktion in perfekter Synchronisation ausgeführt. Threads aus verschiedenen Warps, aber innerhalb desselben Blocks, können in beliebiger Reihenfolge ausgeführt werden, können aber vom Programmierer zur Synchronisation gezwungen werden. Threads aus verschiedenen Blöcken können nicht synchronisiert werden oder in irgendeiner Weise direkt interagieren.

Speicherorganisation

Bei der normalen CPU-Programmierung ist die Speicherorganisation normalerweise vor dem Programmierer verborgen. Typische Programme verhalten sich so, als wäre nur RAM vorhanden. Alle Speicheroperationen, wie z. B. das Verwalten von Registern, das L1-L2-L3-Caching, das Wechseln auf die Festplatte usw., werden vom Compiler, dem Betriebssystem oder der Hardware selbst ausgeführt.

Dies ist bei CUDA nicht der Fall. Während neuere GPU-Modelle die Belastung teilweise überdecken, z. B. durch das Unified Memory in CUDA 6, ist es dennoch aus Gründen der Leistung sinnvoll, die Organisation zu verstehen. Die grundlegende CUDA-Speicherstruktur sieht wie folgt aus:

  • Hostspeicher - das reguläre RAM. Wird hauptsächlich vom Hostcode verwendet, aber auch neuere GPU-Modelle können darauf zugreifen. Wenn ein Kernel auf den Hostspeicher zugreift, muss die GPU normalerweise über den PCIe-Connector mit der Hauptplatine kommunizieren und ist daher relativ langsam.
  • Gerätespeicher / Globaler Speicher - der Hauptspeicher der GPU, der allen Threads zur Verfügung steht.
  • Shared Memory - in jedem SM befindet sich ein viel schnellerer Zugriff als global. Der gemeinsam genutzte Speicher ist für jeden Block privat. Threads innerhalb eines einzelnen Blocks können es für die Kommunikation verwenden.
  • Register - der schnellste, private, nicht adressierbare Speicher jedes Threads. Im Allgemeinen können diese nicht für die Kommunikation verwendet werden, aber einige intrinsische Funktionen ermöglichen das Mischen ihres Inhalts innerhalb eines Warp.
  • Die lokale Speicher - private Speicher jeden Thread, die adressierbar ist. Dies wird für Registerüberläufe und lokale Arrays mit variabler Indizierung verwendet. Physisch befinden sie sich im globalen Speicher.
  • Texture memory, Constant memory - ein Teil des globalen Speichers, der für den Kernel als unveränderlich markiert ist. Dies ermöglicht der GPU die Verwendung spezieller Caches.
  • L2-Cache - auf dem Chip, für alle Threads verfügbar. In Anbetracht der Anzahl der Threads ist die erwartete Lebensdauer jeder Cachezeile viel niedriger als bei der CPU. Es wird meistens verwendet, um falsch ausgerichtete und teilweise zufällige Speicherzugriffsmuster zu unterstützen.
  • L1-Cache - befindet sich im selben Speicherbereich wie der gemeinsam genutzte Speicher. Auch hier ist die Anzahl angesichts der Anzahl der verwendeten Threads eher gering. Erwarten Sie also nicht, dass die Daten dort lange bleiben. L1-Caching kann deaktiviert werden.

Versionen

Rechenleistung Die Architektur GPU-Codename Veröffentlichungsdatum
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 01.04.2015
6,0 Pascal GP100 2016-10-01
6.1 Pascal GP102, GP104, GP106 2016-05-27

Das Veröffentlichungsdatum kennzeichnet die Veröffentlichung der ersten GPU, die die angegebenen Berechnungsfunktionen unterstützt. Einige Daten sind ungefähr, z. B. wurde im 2. Quartal 2014 eine 3,2-Karte veröffentlicht.

Kompilieren und Ausführen der Beispielprogramme

Das NVIDIA-Installationshandbuch endet mit der Ausführung der Beispielprogramme, um die Installation des CUDA Toolkit zu überprüfen, gibt jedoch nicht explizit die Vorgehensweise an. Überprüfen Sie zunächst alle Voraussetzungen. Überprüfen Sie das Standard-CUDA-Verzeichnis für die Beispielprogramme. Wenn es nicht vorhanden ist, kann es von der offiziellen CUDA-Website heruntergeladen werden. Navigieren Sie zu dem Verzeichnis, in dem die Beispiele vorhanden sind.

$ cd /path/to/samples/
$ ls
 

Sie sollten eine Ausgabe ähnlich der folgenden sehen:

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

Stellen Sie sicher, dass das Makefile in diesem Verzeichnis vorhanden ist. Mit dem Befehl make in UNIX-basierten Systemen werden alle Beispielprogramme erstellt. Navigieren Sie alternativ zu einem Unterverzeichnis, in dem sich ein anderes Makefile befindet, und führen Sie den Befehl make von dort aus aus, um nur dieses Beispiel zu erstellen.

Führen Sie die zwei empfohlenen Beispielprogramme aus - deviceQuery und bandwidthTest :

$ cd 1_Utilities/deviceQuery/
$ ./deviceQuery 
 

Die Ausgabe ähnelt der unten gezeigten:

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

Die Anweisung Result = PASS am Ende zeigt an, dass alles korrekt funktioniert. Führen Sie nun das andere vorgeschlagene Beispielprogramm bandwidthTest auf ähnliche Weise aus. Die Ausgabe wird ähnlich sein:

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

Die Anweisung Result = PASS zeigt erneut an, dass alles ordnungsgemäß ausgeführt wurde. Alle anderen Beispielprogramme können auf ähnliche Weise ausgeführt werden.

Lassen Sie uns einen einzelnen CUDA-Thread starten, um Hallo zu sagen

Dieses einfache CUDA-Programm zeigt, wie eine Funktion geschrieben wird, die auf der GPU (auch als "Gerät" bezeichnet) ausgeführt wird. Die CPU oder "Host" erstellt CUDA-Threads durch Aufrufen spezieller Funktionen, die als "Kernels" bezeichnet werden. CUDA-Programme sind C ++ - Programme mit zusätzlicher Syntax.

Um zu sehen, wie es funktioniert, hello.cu den folgenden Code in eine Datei namens 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;
}
 

(Beachten Sie, dass Sie zur Verwendung der printf Funktion auf dem Gerät ein Gerät mit einer Rechenkapazität von mindestens 2,0 benötigen. Weitere Informationen finden Sie in der Versionsübersicht .)

Lassen Sie uns nun das Programm mit dem NVIDIA-Compiler kompilieren und ausführen:

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

Einige zusätzliche Informationen zum obigen Beispiel:

  • nvcc steht für "NVIDIA CUDA Compiler". Es trennt den Quellcode in Host- und Gerätekomponenten.
  • __global__ ist ein CUDA-Schlüsselwort, das in Funktionsdeklarationen verwendet wird, um __global__ , dass die Funktion auf dem GPU-Gerät ausgeführt wird und vom Host aufgerufen wird.
  • Dreifache spitze Klammern ( <<< , >>> ) kennzeichnen einen Aufruf vom Hostcode zum Gerätecode (auch als "Kernel-Start" bezeichnet). Die Zahlen in diesen dreifachen Klammern geben die Anzahl der parallel auszuführenden Zeiten und die Anzahl der Threads an.

Voraussetzungen

Laden Sie den CUDA Toolkit und den Entwicklertreiber herunter, um mit der Programmierung mit CUDA zu beginnen. Das Toolkit enthält nvcc , den NVIDIA CUDA Compiler und andere Software, die zur Entwicklung von CUDA-Anwendungen erforderlich ist. Der Treiber stellt sicher, dass GPU-Programme auf CUDA-fähiger Hardware ordnungsgemäß ausgeführt werden, was auch erforderlich ist.

Sie können bestätigen, dass das CUDA Toolkit auf Ihrem Computer ordnungsgemäß installiert ist, indem Sie nvcc --version über eine Befehlszeile nvcc --version . Zum Beispiel auf einem Linux-Rechner

$ 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
 

gibt die Compiler-Informationen aus. Wenn der vorherige Befehl nicht erfolgreich war, ist das CUDA Toolkit wahrscheinlich nicht installiert, oder der Pfad zu nvcc ( C:\CUDA\bin auf Windows-Computern, /usr/local/cuda/bin unter POSIX-Betriebssystemen) gehört nicht zu Ihrem Umgebungsvariable PATH

Außerdem benötigen Sie einen Host-Compiler, der mit nvcc , um CUDA-Programme zu kompilieren und zu erstellen. Unter Windows ist dies cl.exe , der Microsoft-Compiler, der mit Microsoft Visual Studio cl.exe wird. Unter POSIX-Betriebssystemen sind andere Compiler verfügbar, einschließlich gcc oder g++ . Im offiziellen CUDA Quick Start Guide erfahren Sie, welche Compilerversionen auf Ihrer jeweiligen Plattform unterstützt werden.

Um sicherzustellen, dass alles korrekt eingerichtet ist, lassen Sie uns ein triviales CUDA-Programm kompilieren und ausführen, um sicherzustellen, dass alle Tools ordnungsgemäß zusammenarbeiten.

__global__ void foo() {}

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

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

  return 0;
}
 

Um dieses Programm zu kompilieren, kopieren Sie es in eine Datei namens test.cu und kompilieren Sie es über die Befehlszeile. Auf einem Linux-System sollte beispielsweise Folgendes funktionieren:

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

Wenn das Programm ohne Fehler erfolgreich ist, beginnen wir mit der Codierung!

Addieren Sie zwei Arrays mit CUDA

Dieses Beispiel zeigt, wie Sie ein einfaches Programm erstellen, das zwei int Arrays mit CUDA summiert.

Ein CUDA-Programm ist heterogen und besteht aus Teilen, die auf CPU und GPU laufen.

Die Hauptteile eines Programms, die CUDA verwenden, sind ähnlich wie CPU-Programme und bestehen aus

  • Speicherzuordnung für Daten, die auf der GPU verwendet werden
  • Daten werden vom Hostspeicher in den GPU-Speicher kopiert
  • Aufrufen der Kernel-Funktion zum Verarbeiten von Daten
  • Ergebnis in den CPU-Speicher kopieren

Um Speicherplatz für Geräte zuzuweisen, verwenden wir die cudaMalloc Funktion. Zum Kopieren von Daten zwischen Gerät und Host cudaMemcpy Funktion cudaMemcpy verwendet werden. Das letzte Argument von cudaMemcpy gibt die Richtung des Kopiervorgangs an. Es gibt 5 mögliche Typen:

  • cudaMemcpyHostToHost - Host -> Host
  • cudaMemcpyHostToDevice - Host -> Gerät
  • cudaMemcpyDeviceToHost - Gerät -> Host
  • cudaMemcpyDeviceToDevice - Gerät -> Gerät
  • cudaMemcpyDefault - Standardbasierter, einheitlicher virtueller Adressraum

Als nächstes wird die Kernel-Funktion aufgerufen. Die Information zwischen den dreifachen Chevrons ist die Ausführungskonfiguration, die bestimmt, wie viele Gerätethreads den Kernel parallel ausführen. Die erste Anzahl ( 2 im Beispiel) gibt die Anzahl der Blöcke und die zweite ( (size + 1) / 2 im Beispiel) die Anzahl der Threads in einem Block an. Beachten Sie, dass wir in diesem Beispiel die Größe um 1 erhöhen, sodass wir einen zusätzlichen Thread anfordern, anstatt einen Thread für zwei Elemente verantwortlich zu machen.

Da der cudaDeviceSynchronize eine asynchrone Funktion ist, wird cudaDeviceSynchronize aufgerufen, um zu warten, bis die Ausführung abgeschlossen ist. Ergebnis-Arrays werden in den cudaFree kopiert und der gesamte auf dem Gerät zugewiesene Speicher wird mit cudaFree .

Um die Funktion als Kernel zu definieren, wird der __global__ Deklarationsbezeichner verwendet. Diese Funktion wird von jedem Thread aufgerufen. Wenn jeder Thread ein Element des resultierenden Arrays verarbeiten soll, brauchen wir ein Mittel, um jeden Thread zu unterscheiden und zu identifizieren. CUDA definiert die Variablen blockDim , blockIdx und threadIdx . Die vordefinierte Variable blockDim enthält die Dimensionen jedes Thread-Blocks, wie im zweiten Konfigurationsparameter für den Kernel-Start angegeben. Die vordefinierten Variablen threadIdx und blockIdx enthalten den Index des Threads in seinem threadIdx bzw. den blockIdx innerhalb des Gitters. Da wir möglicherweise einen Thread mehr anfordern als Elemente in den Arrays, müssen wir die size um sicherzustellen, dass wir nicht über das Ende des Arrays hinaus zugreifen.

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