cuda Zelfstudie
Aan de slag met cuda

Opmerkingen

versies

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

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.

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.