Bloque iv
This presentation is the property of its rightful owner.
Sponsored Links
1 / 48

Bloque IV PowerPoint PPT Presentation


  • 69 Views
  • Uploaded on
  • Presentation posted in: General

Bloque IV. Prácticas de programación en CUDA. David Miraut Marcos García Ricardo Suárez. Código. http://dl.dropbox.com/u/7996589/C%C3%B3digo.zip. Instalación en Windows. Documentación: http://developer.nvidia.com/nvidia-gpu-computing-documentation

Download Presentation

Bloque IV

An Image/Link below is provided (as is) to download presentation

Download Policy: Content on the Website is provided to you AS IS for your information and personal use and may not be sold / licensed / shared on other websites without getting consent from its author.While downloading, if for some reason you are not able to download a presentation, the publisher may have deleted the file from their server.


- - - - - - - - - - - - - - - - - - - - - - - - - - E N D - - - - - - - - - - - - - - - - - - - - - - - - - -

Presentation Transcript


Bloque iv

Bloque IV

Prácticas de programación en CUDA

David Miraut

Marcos García

Ricardo Suárez


C digo

Código

http://dl.dropbox.com/u/7996589/C%C3%B3digo.zip


Instalaci n en windows

Instalación en Windows

  • Documentación:

    • http://developer.nvidia.com/nvidia-gpu-computing-documentation

    • http://developer.download.nvidia.com/compute/DevZone/docs/html/C/doc/CUDA_C_Getting_Started_Windows.pdf

  • Tarjetas compatibles:

    • Tarjetas desde la serie 8

    • Algunas Quadro

    • Teslas

    • http://www.nvidia.com/object/cuda_gpus.html

  • Página de Nvidia (Windows|Linux |Mac OS):

    • http://developer.nvidia.com/cuda-toolkit-40

    • Compatibilidad hacia atrás:

      • http://developer.nvidia.com/cuda-toolkit-archive


  • Instalaci n en windows1

    Instalación en Windows

    • Descargar el driver

      • Driver de desarrollo

        • XP (32|64)

        • Vista y Windows 7 (32|64)

        • Vista y Windows 7 (Notebooks) (32|64)

      • Driver de desarrollo vs driver gráfico

        • Valen los 2 (270.81 | 280.26)

        • Desarrollo:

          • El más antiguo en el que funciona el Toolkit

          • Soporta más dispositivos

          • Basado en la versión release

      • Instalar Driver


    Instalaci n en windows2

    Instalación en Windows

    • Toolkit 4.0

      • Contiene:

        • Cabeceras

        • Librerías

          • GPU-accelerated BLAS library

          • GPU-accelerated FFT library

          • GPU-acceleratedSparseMatrixlibrary

          • GPU-accelerated RNG library

        • Herramientas

          • Visual Profiler

        • “Integración con Visual Studio”

          • Variables de entorno

          • .rules

          • nvcc

        • Otros recursos


    Instalaci n en windows3

    Instalación en Windows

    • Toolkit 4.0

      • Versiones de 64 y 32 bits

      • Carpetas (C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.0)

        • Ejecutables para el compilador y herramientas (bin)

        • Cabeceras (include)

        • Ficheros de enlazado (bin)

        • Documentación (doc)

      • Instalar Toolkit

      • Pasos

        • Registro

        • Instalación: típica, completa, personalizada


    Instalaci n en windows4

    Instalación en Windows

    • SDK

      • Proyectos listos para funcionar en Visual Studio

      • Directorio

        • C:\Documents and Settings\AllUsers\Application Data\NVIDIA Corporation\NVIDIA GPU Computing SDK

        • %ProgramData%\NVIDIA Corporation\NVIDIA GPU Computing SDK

      • Acceso Online: http://developer.nvidia.com/gpu-computing-sdk

      • CUTIL: librería de utilidades (Fuentes)

      • Instalar SDK

      • Pasos

        • Registro

        • Instalación: típica, completa, personalizada

        • Crear acceso directo


    Instalaci n en windows5

    Instalación en Windows

    • Test

      • Ejecutar bandwidthTest

        • C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\C\bin\win32|64\Release

        • C:\Documents and Settings\All Users\Application Data\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\C\bin\win32|64\Release

      • Proyectos [X]

        • C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\C\src\bandwidthTest

        • C:\Documents and Settings\All Users\Application Data\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\C\src\bandwidthTest


    Instalaci n en windows6

    Instalación en Windows

    • Instalación en VS

      • Highlighting

        • Copiar: usertype.dat

          • De [SDK_DIR]\NVIDIA GPU Computing SDK 4.0\C\doc\syntax_highlighting\visual_studio_8

          • A [VISUAL_DIR]\Microsoft Visual Studio 8\Common7\IDE

        • En Visual Studio: Herramientas -> Opciones -> Editor de Texto -> Extensión de archivo -> agregar .cu como MSVS C++

        • Reiniciar MSVS


    Instalaci n en windows7

    Instalación en Windows

    • Instalación VS

      • CUDA VS Wizard para VS2008

        • Crea el proyecto automáticamente

        • http://sourceforge.net/projects/cudavswizard/develop

        • 32 y 64 bits

          • http://sourceforge.net/projects/cudavswizard/files/CUDA_VS_Wizard_2.2%20Beta/

        • No actualizado a la versión 4.0 [X]

          • Hay que compilar CUTIL (Portabilidad y control de errores – No oficial)

            • Disponible en el SDK ($(NVSDKCOMPUTE_ROOT)\common\)

            • Mover las DLLs a la ruta ($(NVSDKCOMPUTE_ROOT)\common\bin)

          • Cambiar las propiedades del proyecto si se trabaja en 32 bits tanto en Release como en Debug (en todos los proyectos!!!!)

            • Vinculador -> Directorios de bibliotecas adicionales

              • $(NVSDKCOMPUTE_ROOT)\common\lib\ por $(NVSDKCOMPUTE_ROOT)\common\lib\Win32


    Instalaci n en windows8

    Instalación en Windows

    • Instalación VS

      • Configuración de CUDA en VS fromscratch(V2010)

        • Crear un proyecto vacío (Win32 de consola vacío p.e.)

        • Añadir las reglas de compilación a los archivos .CU

          • Botón derecho sobre el proyecto -> añadir reglas de generación

          • La primer vez:

            • Buscar existentes:

              • $(CUDA_PATH)\extras\visual_studio_integration\rules

            • Añadir una asociada a *.cu (Runtime)

          • E.O.C.

            • Marcar la regla

            • Puede marcarse o utilizarse la regla del CUDA VS Wizard


    Instalaci n en windows9

    Instalación en Windows

    • Configuración de CUDA en VS fromscratch(V2010)

      • Añadir ficheros de inclusión:

        • Herramientas -> Opciones -> Proyectos y soluciones -> Directorios de VC++ -> En Archivos de inclusión

          • $(CUDA_INC_PATH)

          • $(NVSDKCOMPUTE_ROOT)\common\Inc\

        • También se puede hacer en propiedades del proyecto -> CUDA

      • Añadir librerías

        • Propiedades del proyecto -> Vinculador -> General -> Directorios de bibliotecas adicionales

          • $(CUDA_LIB_PATH)

          • $(NVSDKCOMPUTE_ROOT)\common\lib\Win32

      • En VS2010 Añadir el parche


    01 hellocuda

    01-HelloCuda

    • Ejemplo de CUDA VS Wizard

      • Crear un proyecto

      • No marcar la opción de cabeceras precompiladas

      • Abrir el fichero principal

    • Funciones

      • InitCUDA:

        • Cuenta el número de dispositivos

        • Busca el primer dispositivo compatible con CUDA

      • HelloCUDA

        • Kernel Global

        • Copia una frase en un lugar de la tarjeta gráfica

      • Main

        • Se reserva espacio para el resultado

        • Se lanza el kernel

        • Se espera a que termine el kernel

        • Se copian los resultados a memoria principal

        • Se libera el contexto de CUDA

    • Tareas

      • Añadir la función “getchar()” al final de la función principal

      • Añadir las modificaciones necesarias para poder compilar el código

      • Enlazar el proyecto en Release


    00 proyecto base

    00-Proyecto base

    • Cuatro ficheros

      • Main.c

        • Se ejecuta en el host

        • Llama a la función encargada de ejecutar el kernel

        • Se encarga de la medición de tiempo

      • mi.h:

        • Cabecera de la función que llama al kernel

      • mi.cu

        • Se ejecuta en el host

        • Fichero encargado de llamar al kernel

        • Selecciona el número de hilos y bloques

        • Se encarga de la trasferencia de los datos

        • Espera a que finalice el kernel

      • mi_kernel.cu:

        • Se ejecuta en el device

        • Implementación del kernel


    00 proyecto base1

    00-Proyecto base

    • Tareas

      • Deshabilitar la compilación de mi_kernel.cu

        • Utilizar un .cuh

      • Indicar al compilador que debe mostrar la información necesaria para calcular la ocupación

      • Activar las optimizaciones


    000 rc4 cpu

    000-RC4 CPU

    • Contenido

      • Implementación del RC4 en CPU

      • El fichero “main.c” contiene múltiples llamadas al RC4 que deberán paralelizarse


    02 rc4 sin sm

    02-RC4 SIN SM

    • Ejercicio

      • Implemetarunafunciónquellame al kernel y copiarsudescripción en el archivo de cabecera:

        int rc4_call_kernel(unsigned char *key, unsigned intlKey, unsigned char *text, unsigned intlText, unsigned intnKey, unsigned char *cypherT)

        • Reservar espacio de los parámetros tanto de entrada como de salida

          cutilSafeCall(cudaMalloc((void**)& puntero, elementos * sizeof(tipo)));

        • Copiar los vectores de entrada

          cutilSafeCall(cudaMemcpy(destino, origen, tamaño * sizeof(tipo), cudaMemcpyHostToDevice));


    02 rc4 sin sm1

    02-RC4 SIN SM

    • Ejercicio

      • Implementar una función que llame al kernel:

        • Dividir en bloques e hilos de forma que cada hilo procese una clave

          • Determinar el número de hilos por bloque

            blockDim.x = Número de hilos por bloque

          • Determinar el número de bloques

            blocks= nkeys/blockDim.x

            gridDim.x = (blocks*blockDim.x < nkeys)?blocks+1:blocks;


    02 rc4 sin sm2

    02-RC4 SIN SM

    • Ejercicio

      • Implementar una función que llame al kernel:

        • Llamar al kernel

          rc4_kernel_noSM<<<gridDim, blockDim>>>(d_key, lKey, d_text, lText, nKey, d_cypherT);

        • Sincronización

          cudaThreadSynchronize();


    02 rc4 sin sm3

    02-RC4 SIN SM

    • Ejercicio

      • Implementar una función que llame al kernel

        • Copia de los resultados

          cutilSafeCall(cudaMemcpy(destino, origen, tamaño * sizeof(tipo), cudaMemcpyDeviceToHost));

        • Liberar Recursos

          cutilSafeCall(cudaFree(d_cypherT));


    02 rc4 sin sm4

    02-RC4 SIN SM

    • Ejercicio

      • Implementar un conjunto de kernels que implemente el RC4

        • Implementar las funciones auxiliares como __device__

        • Implementar el kernel principal __global__

          • Determinar a qué clave se va a acceder

            unsignedintdimXxIdxX = blockDim.x * blockIdx.x;

            unsignedintpKey = dimXxIdxX * lKey + threadIdx.x * lKey;

            unsignedintpText = dimXxIdxX * lText + threadIdx.x * lText;

          • Controlar datos que no son múltiplos de 32

            • Solución 1:

              if ((dimXxIdxX + threadIdx.x) >= nKey)

              return;

            • Solución 2: rellenar con datos basura


    02 rc4 sin sm5

    02-RC4 SIN SM

    • Ejercicio

      • Implementar la toma de tiempos en el fichero main.c

        • Asignar valor a las variables de entrada

          lKey = 6;

          lText = 100;

          nKey = 1000000;

          key = (unsigned char *)malloc(lKey * nKey * sizeof(unsigned char));

          cypherT = (unsignedchar *)malloc(lText * nKey * sizeof(unsignedchar));

          text = (unsigned char *)malloc(lText * sizeof(unsigned char));

        • Calcular el tiempo medio

          for (i = 0; i < 10; ++i)

          rc4_call_kernel(key, lKey, text, lText, nKey, cypherT);


    02 rc4 sin sm6

    02-RC4 SIN SM

    • Ejercicio

      • Implementar un conjunto de kernels que implemente el RC4

        • Implementar las funciones auxiliares como __device__

        • Implementar el kernel principal __global__

          • Determinar a qué clave se va a acceder

            unsignedintdimXxIdxX = blockDim.x * blockIdx.x;

            unsignedintpKey = dimXxIdxX * lKey + threadIdx.x * lKey;

            unsignedintpText = dimXxIdxX * lText + threadIdx.x * lText;

          • Controlar datos que no son múltiplos de 32

            • Solución 1:

              if ((dimXxIdxX + threadIdx.x) >= nKey)

              return;

            • Solución 2: rellenar con datos basura


    02 rc4 sin sm7

    02-RC4 SIN SM

    • Análisis

      • Fermi (256 – 100%-441ms)


    02 rc4 sin sm8

    02-RC4 SIN SM

    • Análisis

      • Fermi (128 – 66% - 416ms)


    02 rc4 sin sm9

    02-RC4 SIN SM

    • Análisis

      • Fermi (50% - 96 – 397MS)


    03 rc4 sm

    03-RC4 SM

    • Ejercicio

      • El vector de S pasa a memoria compartida

        • Declaración

          __shared__ unsigned char S[tamaño* número de hilosporbloque];

        • Cada hilo accede a su porción de memoria compartida

          unsignedintsMemPos = threadIdx.x * tamaño; //tamaño 256

        • Ahora la forma de direccionar S cambia

          • Puntero al comienzo de S

            • Antes: S

            • Ahora: &(S[sMemPos])

          • Acceso a una posición de S

            • Antes: S[i]

            • Ahora: S[sMemPos+i]


    03 rc4 sm1

    03-RC4 SM

    • Análisis

      • Hilos por bloque máximo

        • 256*número hilo < 16384 (en teoría) -> 64

          • Hay que compilar para saber cuanto espacio de memoria compartida va a usar el driver de forma transparente

        • 256*número hilo < 49152 -> 192

        • OccupancyCalculator

      • No va a funcionar en una tarjeta no dedicada!!!!


    04 rc4 sm sin bloqueos

    04-RC4 SM sin bloqueos

    • Ejemplo

      • Ajuste de los datos por columnas en el kernel principal

        • Primer elemento: S[threadIdx.x]

        • Puntero al primer elemento: &(S[threadIdx.x])

        • Acceso al elemento i: S[threadIdx.x + numero de hilos por bloque * i]

      • Acceso al elemento i en rc4_init

        • S[número de hilos por bloque * i]


    04 rc4 sm sin bloqueos1

    04-RC4 SM sin bloqueos

    • Análisis

      • Aumento de rendimiento

      • Cada bloque de memoria compartida proporciona un entero de 32 bits

      • En nuestro caso se leen char, se bloquean 4 threads de cada vez

      • Solución

        • Desperdiciar el espacio (no hay suficiente SM)

        • Tipos de acceso más sofisticados por wraps de 8 hilos


    05 rc4 con memoria de constantes

    05-RC4 con memoria de constantes

    • Ejemplo

      • Ideal para meter datos pequeños (hasta 64k)

        • A los que acceden todos los hilos a la misma posición a la vez

        • Y sólo pueden leer

      • Se coloca la cadena a cifrar

      • Se declara de forma estática y global en el fichero donde se define el kernel:

        __device____constant__unsigned char d_lEnt[100];

      • La cadena de entrada se copia con otro tipo de llamada y no hace falta reservar espacio

        cutilSafeCall(cudaMemcpyToSymbol(d_lEnt, text, lText * sizeof(unsigned char)));


    05 rc4 con memoria de constantes1

    05-RC4 con memoria de constantes

    • Ejemplo

      • Quitar el parámetro text de la entrada del kernel

      • Sustituir la variable text por d_lEnt


    05 rc4 con memoria de constantes2

    05-RC4 con memoria de constantes

    • Análisis

      • Aumento de rendimiento

      • Se reduce un parámetro de entrada


    06 rc4 con coalescencia

    06-RC4 con coalescencia

    • Ejercicio

      • Organizar los datos de entrada para permitir la lectura/escritura simultánea de 16 hilos a datos consecutivos (32, 64, 128)

      • Se supone que los datos ya están ordenados

      • Se reserva la memoria garantizando la alineación de los datos

        • pitch: número de bytes por fila

        • Indicar el tamaño de la fila

        • Indicar el número de filas

          cutilSafeCall(cudaMallocPitch((void**)&d_keyP, &pitch, nKey * sizeof(unsignedchar), lKey));

          cutilSafeCall(cudaMallocPitch((void**)&d_cypherTp, &pitchS, nKey * sizeof(unsignedchar), lText));


    06 rc4 con coalescencia1

    06-RC4 con coalescencia

    • Ejercicio

      • Se copian los valores de entrada y salida de forma distinta

        cutilSafeCall(cudaMemcpy2D(d_keyP, pitch, key, nKey * sizeof(unsignedchar), nKey * sizeof(unsignedchar), lKey, cudaMemcpyHostToDevice));

        cutilSafeCall(cudaMemcpy2D(cypherT, nKey * sizeof(unsignedchar), d_cypherTp, pitchS, nKey * sizeof(unsignedchar), lText, cudaMemcpyDeviceToHost));


    06 rc4 con coalescencia2

    06-RC4 con coalescencia

    • Ejercicio

      • Añadir el pitch como variable de entrada al kernel

        __global__ void rc4_kernel_SMsin_const_coa(unsignedchar *key, unsignedintlKey, unsignedintlText, unsignedintnKey, unsignedchar *cypherT,unsigned int pitch, unsigned intpitchS)

      • En el kernel los datos se direccionan de forma distinta

        unsignedintdimXxIdxX = blockDim.x * blockIdx.x;

        unsigned int pos = dimXxIdxX + threadIdx.x;

      • Cambiar la forma de acceso a las variables key y cypherT

        • Acceso al elemento k-ésimo: cypherT[pos + k * pitchS]


    06 rc4 con coalescencia3

    06-RC4 con coalescencia

    • Análisis

      • Sólo se traen 8 bits a la vez por lectura

        • Se podrían intentar traer hasta 128 si se conoce bien el tamaño de entrada

      • Aumenta el rendimiento significativamente


    07 08 rc4 texturas

    07|08-RC4 texturas

    • Ejemplo

      • Para datos de sólo lectura (en Fermi se puede escribir también)

      • En el fichero del kernel se declara una variable de tipo texturas

        texture<unsigned char, 1, cudaReadModeElementType> textKey;

      • En el fichero que llama al kernel se crea un descriptor de textura

        cudaChannelFormatDescchannelDesc = cudaCreateChannelDesc<unsignedchar>();

      • Después de subir los datos se enlaza el descriptor con la textura

        cudaBindTexture(NULL, &textKey, d_keyP, &channelDesc, lKey * pitch * sizeof(unsignedchar));

      • Al finalizar se desenlaza la textura

        cudaUnbindTexture(textKey);


    07 08 rc4 texturas1

    07|08-RC4 texturas

    • Ejemplo

      • El kernel no necesita que se pasen las claves como entrada

        __global__ void rc4_kernel_SMsin_const_coa_text(unsignedintlKey, unsignedintlText, unsignedintnKey, unsignedchar *cypherT, unsigned int pitch, unsigned intpitchS)

      • Cambiar los accesos a “key” por una búsqueda en textura:

        tex1Dfetch(textKey, pos + (i & (key_length - 1) * pitch))

      • Las texturas no se pueden pasar como parámetro

        • Modificar RC4_Init para que reciba la posición de acceso

        • Cambiar los accesos a “key”


    07 08 rc4 texturas2

    07|08-RC4 texturas

    • Análisis

      • Mejora de rendimiento

      • Fermi


    09 rc4 m ltiples elementos por hilo

    09-RC4 Múltiples elementos por hilo

    • Cada hilo procesa más de una palabra

    • En la distribución de bloques es necesario tener en cuenta el número de claves que procesa cada hilo

      blocks= nkeys / (blockDim.x * n)

      gridDim.x = (blocks * blockDim.x* n <nkeys) ? blocks + 1:blocks;

    • Añadir al kernel el número de elementos que debe procesar cada hilo

      __global__ void rc4_kernel_SMsin_const_coa_mul(unsignedchar *key, unsignedintlKey, unsignedintlText, unsignedintnKey, unsignedchar *cypherT, unsigned int pitch, unsigned intpitchS, unsigned int n);


    09 rc4 m ltiples elementos por hilo1

    09-RC4 Múltiples elementos por hilo

    • Se debe adaptar el kernel para procesar varias claves

      • Posición de inicio de cada bloque

        unsignedintdimXxIdxX = (blockDim.x * n) * blockIdx.x;

      • Se añade un bucle que procesa todas las claves asignadas

        for (int l = 0; l < n; ++l)

      • Cambiar el acceso a cada clave

        • Primera posición de cada clave: &(key[pos + l * blockDim.x])

        • Elemento i-ésimo del kernel rc4_init: key[i*pitch]


    10 rc4 m ltiples tarjetas

    10-RC4 Múltiples tarjetas

    • Ejemplo

      • Forma de acceso

        • IP: 193.147.62.16

        • Usuario: gpu_user

        • Clave: CNI_UsEr_gpu

        • Se accede a través de PUTTY o WINSCP (en la carpeta material)

        • Existe una carpeta por usuario: NVIDIA_GPU_Computing_SDK_??


    10 rc4 m ltiples tarjetas1

    10-RC4 Múltiples tarjetas

    • Ejemplo

      • Compilación en Linux

        • La compilación se lleva a cabo desde un terminal

          • Se hace a través de un fichero MAKEFILE

            • make: compila el proyecto

            • makeclean: limpia un proyecto compilado

        • El SDK tiene un fichero de compilación (common.mk)

          • Copiar ese fichero en la carpeta de trabajo (C/common)

          • La versión que se adjunta puede compilar librerías dinámicas

          • Se necesita ubicar los proyectos en la carpeta C/src del directorio de trabajo

          • Los binarios son generados en la carpeta C/bin del directorio de trabajo

      • Explicar código


    10 rc4 m ltiples tarjetas2

    10-RC4 Múltiples tarjetas

    • Ejemplo

      • Código

        • Se genera un hilo en CPU por cada tarjeta gráfica

          • Se añade la librería multithreading del SDK

          • Se crea la estructura con la información que se le pasa a cada hilo

        • El kernel no cambia

    • Análisis

      • Es rentable para datos o computación masiva

      • CUDA 4 mejora el soporte


    11 rc4 y md5

    11-RC4 y MD5

    • Ejemplo

      • Código

        • El código calcula la clave haciendo sucesivas llamadas a MD5

        • La clave resultante se utiliza para cifrar una cadena con RC4

    • Tareas

      • Evitar las copias a CPU entre diferentes llamadas al kernel

      • Incluir el bucle en el kernel del MD5


    12 kernel concurrente

    12-Kernel concurrente

    • Ejemplo

      • Código


    Situaciones no tratadas

    Situaciones no tratadas

    • Operaciones de módulo potencia de 2

      • I % j == i & (j - 1)

    • Control de flujo

      • Claves con tamaños diferentes. Cada Wrap debería acceder a claves del mismo tamaño

    • Caché

      • Configuración del tamaño de la caché

      • Desactivación de la caché

    • Multitarjeta 4.0

      • Paso de parámetros entre tarjetas

      • Espacio unificado

    • Operaciones atómicas en cache

    • Desenrollado de bucles


  • Login