Comenzando con cuda

Iniciemos un solo hilo de CUDA para saludar

Este sencillo programa CUDA demuestra cómo escribir una función que se ejecutará en la GPU (también conocida como “dispositivo”). La CPU, o “host”, crea subprocesos CUDA llamando a funciones especiales llamadas “núcleos”. Los programas CUDA son programas C++ con sintaxis adicional.

Para ver cómo funciona, coloca el siguiente código en un archivo llamado hola.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;
}

(Tenga en cuenta que para usar la función printf en el dispositivo, necesita un dispositivo que tenga una capacidad de cómputo de al menos 2.0. Consulte la [descripción general de versiones] (https://www.wikiod.com/es/cuda/comenzando-con-cuda /introduction-to-cuda#t=20160918151348344166&a=versions) para obtener más detalles).

Ahora compilemos el programa usando el compilador de NVIDIA y ejecútelo:

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

Alguna información adicional sobre el ejemplo anterior:

  • nvcc significa “NVIDIA CUDA Compiler”. Separa el código fuente en componentes de host y dispositivo.
  • __global__ es una palabra clave de CUDA utilizada en las declaraciones de funciones que indican que la función se ejecuta en el dispositivo GPU y se llama desde el host.
  • Los corchetes angulares triples (<<<,>>>) marcan una llamada desde el código del host al código del dispositivo (también llamado “inicio del núcleo”). Los números dentro de estos corchetes triples indican la cantidad de veces que se ejecutará en paralelo y la cantidad de subprocesos.

Requisitos previos

Para comenzar a programar con CUDA, descargue e instale CUDA Toolkit and Developer Driver. El kit de herramientas incluye nvcc, NVIDIA CUDA Compiler y otro software necesario para desarrollar aplicaciones CUDA. El controlador garantiza que los programas de la GPU se ejecuten correctamente en hardware compatible con CUDA, que también necesitará.

Puede confirmar que CUDA Toolkit está correctamente instalado en su máquina ejecutando nvcc --version desde una línea de comando. Por ejemplo, en una máquina 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

genera la información del compilador. Si el comando anterior no tuvo éxito, es probable que CUDA Toolkit no esté instalado o que la ruta a nvcc (C:\CUDA\bin en máquinas Windows, /usr/local/cuda/bin en sistemas operativos POSIX ) no es parte de su variable de entorno PATH.

Además, también necesitará un compilador host que funcione con nvcc para compilar y crear programas CUDA. En Windows, este es cl.exe, el compilador de Microsoft, que viene con Microsoft Visual Studio. En los sistemas operativos POSIX, hay otros compiladores disponibles, incluidos gcc o g++. La Guía de inicio rápido oficial de CUDA puede decirle qué versiones del compilador son compatibles con su plataforma en particular.

Para asegurarnos de que todo esté configurado correctamente, compilemos y ejecutemos un programa CUDA trivial para garantizar que todas las herramientas funcionen juntas correctamente.

__global__ void foo() {}

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

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

  return 0;
}

Para compilar este programa, cópielo en un archivo llamado test.cu y compílelo desde la línea de comandos. Por ejemplo, en un sistema Linux, lo siguiente debería funcionar:

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

Si el programa tiene éxito sin errores, ¡entonces comencemos a codificar!

Compilar y ejecutar los programas de muestra

La guía de instalación de NVIDIA finaliza con la ejecución de los programas de muestra para verificar la instalación de CUDA Toolkit, pero no indica explícitamente cómo. Primero verifique todos los requisitos previos. Verifique el directorio CUDA predeterminado para los programas de muestra. Si no está presente, se puede descargar desde el sitio web oficial de CUDA. Navegue hasta el directorio donde están presentes los ejemplos.

$ cd /path/to/samples/
$ ls

Deberías ver una salida similar a:

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

Asegúrese de que Makefile esté presente en este directorio. El comando make en los sistemas basados ​​en UNIX construirá todos los programas de muestra. Alternativamente, navegue a un subdirectorio donde esté presente otro Makefile y ejecute el comando make desde allí para construir solo esa muestra.

Ejecute los dos programas de muestra sugeridos: deviceQuery y bandwidthTest:

$ cd 1_Utilities/deviceQuery/
$ ./deviceQuery

La salida será similar a la que se muestra a continuación:

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

La declaración Resultado = APROBADO al final indica que todo funciona correctamente. Ahora, ejecute el otro programa de muestra sugerido bandwidthTest de manera similar. La salida será similar a:

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

De nuevo, la sentencia Resultado = APROBADO indica que todo se ejecutó correctamente. Todos los demás programas de muestra se pueden ejecutar de manera similar.

Suma dos arreglos con CUDA

Este ejemplo ilustra cómo crear un programa simple que sumará dos matrices int con CUDA.

Un programa CUDA es heterogéneo y consta de partes que se ejecutan tanto en la CPU como en la GPU.

Las partes principales de un programa que utiliza CUDA son similares a los programas de CPU y consisten en

  • Asignación de memoria para los datos que se utilizarán en la GPU
  • Copia de datos desde la memoria del host a la memoria de las GPU
  • Invocar la función del kernel para procesar datos
  • Copiar resultado a la memoria de la CPU

Para asignar memoria a los dispositivos usamos la función cudaMalloc. Para copiar datos entre el dispositivo y el host, se puede utilizar la función cudaMemcpy. El último argumento de cudaMemcpy especifica la dirección de la operación de copia. Hay 5 tipos posibles:

  • cudaMemcpyHostToHost - Host -> Host
  • cudaMemcpyHostToDevice - Host -> Dispositivo
  • cudaMemcpyDeviceToHost - Dispositivo -> Host
  • cudaMemcpyDeviceToDevice - Dispositivo -> Dispositivo
  • cudaMemcpyDefault - Espacio de direcciones virtuales unificado basado en predeterminado

A continuación, se invoca la función del kernel. La información entre los cheurones triples es la configuración de ejecución, que dicta cuántos subprocesos del dispositivo ejecutan el kernel en paralelo. El primer número (2 en el ejemplo) especifica el número de bloques y el segundo ((tamaño + 1) / 2 en el ejemplo) - el número de subprocesos en un bloque. Tenga en cuenta que en este ejemplo agregamos 1 al tamaño, por lo que solicitamos un subproceso adicional en lugar de tener un subproceso responsable de dos elementos.

Dado que la invocación del kernel es una función asíncrona, se llama a cudaDeviceSynchronize para esperar hasta que se complete la ejecución. Las matrices de resultados se copian en la memoria del host y toda la memoria asignada en el dispositivo se libera con cudaFree.

Para definir la función como kernel, se utiliza el especificador de declaración __global__. Esta función será invocada por cada subproceso. Si queremos que cada subproceso procese un elemento de la matriz resultante, entonces necesitamos un medio para distinguir e identificar cada subproceso. CUDA define las variables blockDim, blockIdx y threadIdx. La variable predefinida blockDim contiene las dimensiones de cada bloque de subprocesos como se especifica en el segundo parámetro de configuración de ejecución para el lanzamiento del kernel. Las variables predefinidas threadIdx y blockIdx contienen el índice del subproceso dentro de su bloque de subprocesos y el bloque de subprocesos dentro de la cuadrícula, respectivamente. Tenga en cuenta que dado que potencialmente solicitamos un subproceso más que elementos en las matrices, debemos pasar tamaño para asegurarnos de no acceder más allá del final de la matriz.

#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;
}
Reducción paralela (por ejemplo, cómo sumar una matriz) →