Looking for cuda Answers? Try Ask4KnowledgeBase
Looking for cuda Keywords? Try Ask4Keywords

cudaEmpezando con cuda


Observaciones

CUDA es una tecnología de computación paralela de NVIDIA y un lenguaje de programación para sus GPU.

Las GPU son máquinas altamente paralelas capaces de ejecutar miles de subprocesos ligeros en paralelo. Cada subproceso de GPU suele ser más lento en ejecución y su contexto es más pequeño. Por otro lado, GPU puede ejecutar varios miles de subprocesos en paralelo e incluso más concurrentemente (los números precisos dependen del modelo de GPU real). CUDA es un dialecto de C ++ diseñado específicamente para la arquitectura de GPU NVIDIA. Sin embargo, debido a las diferencias de arquitectura, la mayoría de los algoritmos no se pueden copiar y pegar simplemente desde C ++ simple; se ejecutarían, pero serían muy lentos.

Terminología

  • host : se refiere al hardware normal basado en CPU y a los programas normales que se ejecutan en ese entorno
  • dispositivo : se refiere a una GPU específica en la que se ejecutan los programas CUDA. Un solo host puede admitir múltiples dispositivos.
  • kernel : una función que reside en el dispositivo que se puede invocar desde el código del host.

Estructura del procesador físico

El procesador de GPU habilitado para CUDA tiene la siguiente estructura física:

  • El chip - todo el procesador de la GPU. Algunas GPU tienen dos de ellas.
  • multiprocesador de flujo continuo (SM): cada chip contiene hasta ~ 100 SM, según el modelo. Cada SM funciona de forma casi independiente, utilizando solo memoria global para comunicarse entre sí.
  • CUDA core : una única unidad de cálculo escalar de un SM. Su número preciso depende de la arquitectura. Cada núcleo puede manejar unos pocos subprocesos ejecutados simultáneamente en una sucesión rápida (similar a un hipervínculo en la CPU).

Además, cada SM cuenta con uno o más programadores warp . Cada programador envía una única instrucción a varios núcleos CUDA. Esto hace que el SM funcione de manera efectiva en el modo SIMD de 32 anchos.

Modelo de Ejecución CUDA

La estructura física de la GPU tiene una influencia directa sobre cómo se ejecutan los núcleos en el dispositivo y cómo uno los programa en CUDA. El kernel se invoca con una configuración de llamada que especifica cuántos subprocesos paralelos se generan.

  • la cuadrícula : representa todos los subprocesos que se generan a partir de la llamada al kernel. Se especifica como un conjunto de bloques de una o dos dimensiones.
  • El bloque - es un conjunto semi-independiente de hilos . Cada bloque se asigna a un solo SM. Como tal, los bloques solo pueden comunicarse a través de la memoria global. Los bloques no están sincronizados de ninguna manera. Si hay demasiados bloques, algunos pueden ejecutarse secuencialmente después de otros. Por otro lado, si los recursos lo permiten, más de un bloque puede ejecutarse en el mismo SM, pero el programador no puede beneficiarse de que eso suceda (excepto por el aumento de rendimiento obvio).
  • el hilo : una secuencia escalar de instrucciones ejecutadas por un solo núcleo CUDA. Los hilos son 'ligeros' con un contexto mínimo, lo que permite que el hardware los intercambie rápidamente hacia adentro y hacia afuera. Debido a su número, los hilos CUDA operan con unos pocos registros asignados a ellos, y una pila muy corta (¡preferiblemente ninguno!). Por esa razón, el compilador CUDA prefiere alinear todas las llamadas de función para aplanar el kernel de modo que contenga solo saltos y bucles estáticos. Las llamadas de ponter de función y las llamadas de método virtual, aunque son compatibles con la mayoría de los dispositivos más nuevos, generalmente incurren en una mayor penalidad de rendimiento.

Cada subproceso se identifica mediante un índice de bloque blockIdx e índice de subproceso dentro del bloque threadIdx . Estos números pueden ser verificados en cualquier momento por cualquier hilo en ejecución y es la única forma de distinguir un hilo de otro.

Además, los hilos se organizan en urdimbres , cada uno de los cuales contiene exactamente 32 hilos. Los hilos dentro de una sola urdimbre se ejecutan en una sincronización perfecta, en la versión SIMD. Los hilos de diferentes deformaciones, pero dentro del mismo bloque pueden ejecutarse en cualquier orden, pero el programador puede obligarlos a sincronizar. Los hilos de diferentes bloques no se pueden sincronizar o interactuar directamente de ninguna manera.

Organización de la memoria

En la programación normal de la CPU, la organización de la memoria suele estar oculta al programador. Los programas típicos actúan como si solo hubiera RAM. Todas las operaciones de la memoria, como la gestión de registros, el uso del almacenamiento en caché L1-L2-L3, el intercambio en el disco, etc., se realizan mediante el compilador, el sistema operativo o el hardware.

Este no es el caso con CUDA. Si bien los modelos de GPU más recientes ocultan parcialmente la carga, por ejemplo, a través de la Memoria unificada en CUDA 6, todavía vale la pena entender a la organización por razones de rendimiento. La estructura básica de la memoria CUDA es la siguiente:

  • Memoria de host - la memoria RAM regular. Principalmente utilizado por el código del host, pero los modelos más nuevos de GPU también pueden acceder a él. Cuando un kernel accede a la memoria del host, la GPU debe comunicarse con la placa base, generalmente a través del conector PCIe y, como tal, es relativamente lento.
  • Memoria del dispositivo / Memoria global: la memoria principal fuera de chip de la GPU, disponible para todos los subprocesos.
  • Memoria compartida : ubicada en cada SM permite un acceso mucho más rápido que el global. La memoria compartida es privada para cada bloque. Los hilos dentro de un solo bloque pueden usarlo para la comunicación.
  • Registros : la memoria más rápida, privada y no direccionable de cada hilo. En general, no se pueden usar para la comunicación, pero algunas funciones intrínsecas permiten barajar su contenido dentro de una deformación.
  • Memoria local - memoria privada de cada hilo que es direccionable. Esto se usa para registros de derrames y arreglos locales con indexación variable. Físicamente, residen en la memoria global.
  • Memoria de textura, Memoria constante : una parte de la memoria global que está marcada como inmutable para el kernel. Esto permite que la GPU use cachés de propósito especial.
  • L2 cache - on-chip, disponible para todos los hilos. Dada la cantidad de subprocesos, la vida útil esperada de cada línea de caché es mucho menor que en la CPU. Se utiliza principalmente para ayudar a los patrones de acceso de memoria desalineados y parcialmente aleatorios.
  • Caché L1 : ubicado en el mismo espacio que la memoria compartida. Nuevamente, la cantidad es bastante pequeña, dado el número de subprocesos que la utilizan, por lo que no espere que los datos permanezcan allí por mucho tiempo. El almacenamiento en caché de L1 se puede deshabilitar.

Versiones

Capacidad de cálculo Arquitectura Nombre en clave de GPU Fecha de lanzamiento
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

La fecha de lanzamiento marca el lanzamiento de la primera GPU que admite la capacidad de cálculo dada. Algunas fechas son aproximadas, por ejemplo, la tarjeta 3.2 se lanzó en el segundo trimestre de 2014.

Compilando y ejecutando los programas de muestra

La guía de instalación de NVIDIA termina con la ejecución de los programas de ejemplo para verificar su instalación del kit de herramientas de CUDA, pero no indica explícitamente cómo. En primer lugar, compruebe todos los requisitos previos. Compruebe el directorio predeterminado de CUDA 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ía 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 el Makefile esté presente en este directorio. El comando make en los sistemas basados ​​en UNIX construirá todos los programas de ejemplo. Alternativamente, navegue a un subdirectorio donde esté presente otro Makefile y ejecute el comando make desde allí para compilar solo esa muestra.

Ejecute los dos programas de ejemplo 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 instrucción Result = PASS al final indica que todo funciona correctamente. Ahora, ejecute la otra prueba de bandwidthTest de bandwidthTest programa de muestra sugerida de una 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.
 

Nuevamente, la instrucción Result = PASS indica que todo se ejecutó correctamente. Todos los demás programas de ejemplo se pueden ejecutar de una manera similar.

Vamos a lanzar un solo hilo CUDA para saludar

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

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

(Tenga en cuenta que para utilizar 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 las versiones para obtener más información).

Ahora compilemos el programa usando el compilador NVIDIA y ejecutémoslo:

$ 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 "Compilador NVIDIA CUDA". Separa el código fuente en componentes de host y dispositivo.
  • __global__ es una palabra clave CUDA utilizada en las declaraciones de funciones que indica 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 "lanzamiento del kernel"). Los números entre estos tres paréntesis indican el número de veces que se ejecutan en paralelo y el número de subprocesos.

Prerrequisitos

Para comenzar a programar con CUDA, descargue e instale el kit de herramientas CUDA y el controlador del desarrollador . El kit de herramientas incluye nvcc , el compilador NVIDIA CUDA y otro software necesario para desarrollar aplicaciones CUDA. El controlador garantiza que los programas de GPU se ejecuten correctamente en un hardware compatible con CUDA , que también necesitará.

Puede confirmar que el kit de herramientas CUDA está instalado correctamente en su máquina ejecutando nvcc --version desde una línea de comandos. 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 el kit de herramientas CUDA 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 sea 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 se incluye 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 de compilador son compatibles con su plataforma en particular.

Para asegurarnos de que todo esté configurado correctamente, compilemos y ejecutemos un programa trivial de CUDA para asegurarnos de 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 error, ¡entonces comencemos a codificar!

Suma dos matrices con CUDA.

Este ejemplo ilustra cómo crear un programa simple que sumará dos arrays 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
  • Invocando la función del kernel para procesar datos
  • Copie el resultado a la memoria de la CPU

Para asignar memoria a los dispositivos cudaMalloc 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 unificadas basadas en cudaMemcpyDefault predeterminados

A continuación se invoca la función del núcleo. La información entre los chevrones triples es la configuración de ejecución, que determina cuántos subprocesos de dispositivo ejecutan el kernel en paralelo. El primer número ( 2 en el ejemplo) especifica el número de bloques y el segundo ( (size + 1) / 2 en el ejemplo) - número de hilos en un bloque. Tenga en cuenta que en este ejemplo agregamos 1 al tamaño, de modo que solicitamos un hilo adicional en lugar de tener un hilo 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 __global__ se usa el especificador de declaración. Esta función será invocada por cada hilo. 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 hilo dentro de su bloque de hilo y el bloque del hilo dentro de la cuadrícula, respectivamente. Tenga en cuenta que dado que potencialmente solicitamos un subproceso más que los elementos en las matrices, debemos pasar de size para asegurarnos de que no accedamos 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;
}