1 / 48

Bloque IV

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

ariane
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. Content is provided to you AS IS for your information and personal use only. Download presentation by click this link. While downloading, if for some reason you are not able to download a presentation, the publisher may have deleted the file from their server. During download, if you can't get a presentation, the file might be deleted by the publisher.

E N D

Presentation Transcript

  1. Bloque IV Prácticas de programación en CUDA David Miraut Marcos García Ricardo Suárez

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

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

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

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

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

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

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

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

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

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

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

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

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

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

  16. 000-RC4 CPU • Contenido • Implementación del RC4 en CPU • El fichero “main.c” contiene múltiples llamadas al RC4 que deberán paralelizarse

  17. 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));

  18. 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;

  19. 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();

  20. 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));

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

  22. 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);

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

  24. 02-RC4 SIN SM • Análisis • Fermi (256 – 100%-441ms)

  25. 02-RC4 SIN SM • Análisis • Fermi (128 – 66% - 416ms)

  26. 02-RC4 SIN SM • Análisis • Fermi (50% - 96 – 397MS)

  27. 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]

  28. 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!!!!

  29. 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]

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

  31. 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)));

  32. 05-RC4 con memoria de constantes • Ejemplo • Quitar el parámetro text de la entrada del kernel • Sustituir la variable text por d_lEnt

  33. 05-RC4 con memoria de constantes • Análisis • Aumento de rendimiento • Se reduce un parámetro de entrada

  34. 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));

  35. 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));

  36. 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]

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

  38. 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);

  39. 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”

  40. 07|08-RC4 texturas • Análisis • Mejora de rendimiento • Fermi

  41. 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);

  42. 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]

  43. 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_??

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

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

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

  47. 12-Kernel concurrente • Ejemplo • Código

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

More Related