bloque iv
Download
Skip this Video
Download Presentation
Bloque IV

Loading in 2 Seconds...

play fullscreen
1 / 48

Bloque IV - PowerPoint PPT Presentation


  • 119 Views
  • Uploaded on

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

loader
I am the owner, or an agent authorized to act on behalf of the owner, of the copyrighted work described.
capcha
Download Presentation

PowerPoint Slideshow about ' Bloque IV' - ariane


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
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
ad