Programación eficiente de Intel Xeon Phi: Caso de estudio con patrones stencil

Índice general

Resumen

Índice de figuras

Índice de tablas

I Motivación y fundamentos
1 Motivación
1.1 Introducción
1.2 Contexto del problema
1.2.1 Arquitectura HPC (High Performance Computing)
1.2.2 El acelerador Intel Xeon Phi
1.2.3 La programación paralela
1.2.4 Patrones de cálculo Stencil
1.3 Objetivos
1.4 Estructura del documento
2 Fundamentos
2.1 Introducción
2.2 Arquitectura del coprocesador Intel Xeon Phi
2.2.1 Visión general
2.2.2 Arquitectura de Intel Xeon Phi
2.2.3 Organización de la caché y accesos a memoria
2.2.4 Interfaz PCIe y DMA
2.3 Programación en Intel Xeon Phi
2.3.1 Modelos de programación en Intel Xeon Phi
2.3.2 Modelo de programación offload
2.3.3 Herramientas de desarrollo y nivel de aplicación
2.4 Arquitectura CUDA de Nvidia
2.5 Evolución de Xeon Phi: Knights Landing
2.6 Fundamentos de los códigos Stencil
2.6.1 Descripción
2.6.2 Modelos matemáticos de los códigos utilizados
2.6.3 Algoritmos Stencil
2.7 Herramientas para evaluar el rendimiento y la energía
2.7.1 La interfaz GUI micsmc
2.7.2 Intel Vtune Amplifier XE
2.7.3 Librería Papi
2.8 Conclusiones

II Acelerando patrones Stencil 3-D en arquitecturas basadas en Xeon Phi
3 Directrices para programar códigos Stencil 3-D en Intel Xeon Phi
3.1 Introducción
3.1.1 Objetivos
3.1.2 Trabajo relacionado
3.2 Migración de patrones Stencil a Intel Xeon Phi
3.2.1 Ejecución del código base
3.2.2 Optimizaciones escalares
3.2.3 Manejo de los hilos en Xeon Phi
3.2.4 Vectorización del código
3.3 Otras técnicas sencillas para mejor el rendimiento
3.4 Influencia de los datos de entrada
3.5 Evaluación
3.5.1 Entorno de evaluación
3.5.2 Evaluación del código base
3.5.3 Evaluación de otras optimizaciones y la escalabilidad
3.5.4 Evaluación de la influencia de los datos de entrada
3.6 Conclusiones
4 Optimizaciones avanzadas. Energía y GPUs
4.1 Introducción
4.1.1 Objetivos
4.1.2 Trabajo relacionado
4.2 Mejora de la paralelización
4.2.1 Políticas de planificación de bucles OpenMP
4.2.2 Políticas de afinidad
4.3 Optimizando el acceso a memoria
4.4 Eficiencia energética
4.5 Implementación del modelo de la ecuación de difusión acústica 3D en GPU
4.6 Evaluación
4.6.1 Entorno de evaluación
4.6.2 Evaluación de las políticas de planificación
4.6.3 Evaluación de la afinidad
4.6.4 Evaluación de la técnica del blocking
4.6.5 Evaluación del rendimiento energético
4.6.6 Evaluación comparativa Xeon Phi con CUDA
4.7 Conclusiones
5 Evaluación del modo de descarga de Xeon Phi
5.1 Introducción
5.1.1 Objetivos
5.1.2 Trabajo relacionado
5.2 Manejo de la descarga en un único coprocesador
5.2.1 Pasos para descargar nuestro código en Intel Xeon Phi
5.2.2 Evaluación del tiempo de la descarga
5.2.3 Comparación del rendimiento nativo vs offload
5.2.4 Otros aspectos evaluados
5.3 Manejo de la descarga en dos coprocesadores
5.3.1 Esquema básico para dos tarjetas
5.3.2 Reducción del tamaño de la transferencia de datos
5.3.3 Manejo de la descarga por parte del procesador anfitrión
5.4 Evaluación
5.4.1 Entorno de evaluación
5.4.2 Evaluación de la inicialización
5.4.3 Evaluación del uso de 2 hilos dentro del bucle global de cálculo
5.4.4 Evaluación del uso de 2 hilos con región paralela y barrera explícita
5.4.5 Evaluación con un hilo y tarjetas de forma asíncrona
5.4.6 Comparativa de las tres estrategias evaluadas en la descarga del Stencil 3-D en dos MICs
5.4.7 Evaluación de las tres estrategias con tamaños grandes de problema
5.5 Conclusiones

III Conclusiones
6 Conclusiones
6.1 Conclusiones
6.2 Contribuciones principales
6.3 Publicaciones científicas y colaboraciones
6.3.1 Revistas Internacionales
6.3.2 Actas de congresos

Apéndices
Apéndice A Manejando el Intel LEO
A.1 Modelo de memoria no compartida mediante LEO para Offload
A.1.1 Offload a varios coprocesadores
A.1.2 Personalización de código para la CPU y Xeon Phi
A.1.3 Declaración de variables y funciones
A.1.4 Movimiento de datos
A.1.5 Modo asíncrono de transferencia y ejecución
A.2 Usando el compilador de Intel para offload
A.2.1 Variables de entorno para el control de descarga
A.2.2 Opciones del compilador Intel
A.2.3 Archivos de cabecera
A.2.4 Opción vec-report utilizada con offload
A.2.5 Generación de un informe de offload
Apéndice B Códigos para el Stencil de difusión de calor
Apéndice C Lista de siglas
Apéndice D Fuentes de las Figuras

Referencias

Índice de figuras

1.1 Desarrollo de rendimiento proyectado (Top500)

2.1 Coprocesador Intel Xeon Phi

2.2 Procesador y coprocesador juntos en una plataforma

2.3 Microarquitectura de todo el coprocesador

2.4 Estructura individual de un núcleo

2.5 Modos de aplicación de la MIC

2.6 Modelo offload entre un procesador e Intel Xeon Phi

2.7 Modos de paralelización y vectorización en MIC

2.8 Herramientas de desarrollo y componentes de aplicaciones y de la interfaz

2.9 Hilos agrupados en bloques

2.10 Bloques que componen un Grid

2.11 Modelo de ejecución en CUDA

2.12 Procesador Knights Landing

2.13 Stencil 3-D de 11 puntos

2.14 Bordes o fronteras de los códigos Stencil

2.15 Patrón Stencil para 3D-FD del modelo de la ecuación de difusión acústica

2.16 Utilización de los cores en mic1 y mic

2.17 GUI de VTune Amplifier XE

2.18 Utilización de VTune en nuestras evaluaciones

3.1 Informe de vectorización de nuestro código Stencil

3.2 Código de la declaración de la función Stencil de 11 puntos

3.3 Informe de vectorización de nuestro código Stencil compilado con -O restrict

3.4 Alineación de los datos con _mm_malloc

3.5 Plataforma hardware

3.6 Resultados de rendimiento (GFLOPs) de Xeon y Xeon Phi

3.7 Comparación de rendimiento (GFLOPs) y aceleración entre Xeon y Xeon Phi

3.8 Evaluación de otras optimizaciones

3.9 Evaluación de la característica Turbo+NoECC

3.10 Evaluación de la escalabilidad en Xeon Phi para la versión (Sstore)

3.11 Influencia del tamaño de los datos de entrada

3.12 Influencia de la forma de los datos de entrada

4.1 Distribuciones de tipo de afinidad

4.2 Diseño de bloques CUDA 2D

4.3 Datos accesados por un bloque CUDA 2D

4.4 Políticas de planificación del bucle OpenMP

4.5 Planificación de la afinidad

4.6 Blocking aplicado a la solución de nuestros códigos Stencil variando los ejes Y y Z

4.7 Blocking aplicado a la solución de nuestros códigos Stencil variando el eje X

4.8 Energía y Potencia promedio para nuestros kernels Stencil

4.9 EDP y Potencia promedio para nuestros kernels Stencil

4.10 Tiempo de ejecución y aceleración para diferentes resoluciones de salas

4.11 Energía y Potencia promedio para diferentes resoluciones de salas

4.12 EDP y Potencia promedio para diferentes resoluciones de salas

5.1 Forma de asignar espacio en la MIC sin transferir datos de entrada

5.2 Forma de asignar espacio en la MIC transfiriendo datos de forma implícita

5.3 Formas de transferencia de datos a la MIC

5.4 Función para copiar el contorno de la matriz fin a la matriz fout

5.5 Descarga y ejecución del Stencil 3-D en la MIC

5.6 Copiado de datos de Xeon Phi al host

5.7 Liberación de datos en Xeon Phi

5.8 Actividades de offload del código Stencil 3-D de 11 puntos

5.9 Compilación en C del código Stencil 3-D 11 puntos

5.10 Informe de diagnóstico de offload del Stencil a mic

5.11 Vista de utilización de los cores en mic

5.12 Transferencia de datos a dos mics

5.13 Transfiriendo una mitad de la matriz a mic1 y la otra mitad a mic2

5.14 Comunicación entre dos MICs para el cálculo del Stencil 3-D

5.15 Paralelización del host Xeon con 2 hilos el bucle global

5.16 Paralelización con 2 hilos y una barrera explícita

5.17 Paralelización con un hilo de las dos tarjetas de forma asíncrona

5.18 Informe de diagnóstico de offload del Stencil a dos tarjetas

5.19 Monitorización de los cores al paralelizar con 2 hilos dentro de un

ciclo global de cálculo

5.20 Monitorización de los cores al paralelizar con 2 hilos y una barrera explícita

5.21 Monitorización de los núcleos usando un hilos y dos tarjetas de forma asíncrona

5.22 Rendimiento en GFLOPs para cada estrategia en función del error

5.23 Rendimiento en GFLOPs para cada estrategia en función de la activi dad analizada

5.24 Comparación de la estrategia(c) con la ejecución en modo nativo

A.1 Coprocesador offload

A.2 Código versión MIC o CPU

A.3 Declaración de funciones y variables globales

A.4 Múltiples declaraciones de funciones y variables globales

A.5 Ciclo de vida offload

A.6 Transfiriendo datos con offload_transfer

A.7 Ejemplo de asignación de memoria a partes de una matriz

A.8 Ejemplo de utilización del modificador into

A.9 Código para asignar datos dinámicamente

A.10 Macros para gestionar la memoria

A.11 Descarga en la mic continuando la actividad concurrente

A.12 Ejemplo de transferencia asíncrona del procesador al coprocesador

A.13 Ejemplo de transferencia asíncrona del procesador al coprocesador en diferentes momentos

A.14 Variables de entorno de MIC

B.1 Código base, no vectorizado

B.2 Código con optimizaciones escalares

B.3 Código vectorizado utilizando restrict y collapse

B.4 Código utilizando tamaño de paginas grandes (huge)

B.5 Código con (Streaming stores)

B.6 Código con políticas de planificación del bucle OpenMP (Schedule)B.7 Código con Optimización del acceso a memoria, mediante la técnica de Blocking

B.8 Código que corre en el host en modo offload en una sola tarjeta (mic2) B.9 Código que utiliza 2 hilos dentro del bucle global de cálculo, descar gando en mic1 y mic

B.10 Código que utiliza 2 hilos con región paralela y barrera explicita

B.11 Código usando un solo hilo y manejando las tarjetas asíncronamente

Índice de tablas

2.1 Parámetros clave de la caché del coprocesador Intel Xeon Phi

2.2 Parámetros TLB del coprocesador Intel Xeon Phi

3.1 Especificaciones del coprocesador Intel Xeon Phi 7120P

3.2 Comparativa entre Intel Xeon y Xeon Phi

3.3 Escalabilidad de nuestro código Stencil base en Xeon y Xeon Phi

3.4 Tiempos de ejecución de las tres primeras versiones de nuestros códi- gos Stencil

4.1 Tiempos de ejecución de tres versiones de nuestros códigos Stencil

4.2 Comparativa de Xeon Phi y CUDA en tiempo, energía y EDP

5.1 Tiempo de ejecución del código Stencil 3-D descargado en MIC

5.2 Tiempo del código Stencil 3-D nativo y descargado en MIC

5.3 Rendimiento del código Stencil 3-D nativo y descargado en MIC

5.4 Evaluación de afinidad del cálculo del Stencil 3-D descargado en MIC

5.5 Evaluación del número de iteraciones del cálculo del Stencil 3-D des cargado en MIC

5.6 Evaluación del tamaño de matriz del Stencil 3-D descargado en MIC

5.7 Evaluación del error al ejecutar el código en dos MICs sin comunica ción para diferentes valores del contorno

5.8 Evaluación del error al ejecutar el código en dos MICs actualizando parcialmente para diferentes valores del contorno

5.9 Tiempo utilizado por la inicialización

5.10 Evaluación del rendimiento con el uso de 2 hilos dentro del bucle

global de cálculo

5.11 Evaluación del tiempo con el uso de 2 hilos dentro del bucle global de cálculo

5.12 Evaluación del rendimiento con el uso de dos hilos y una barrera explícita

5.13 Evaluación del tiempo con el uso de 2 hilos y una barrera explícita

5.14 Evaluación del rendimiento usando un hilo y tarjetas de forma asíncrona

5.15 Evaluación del tiempo usando un hilo y tarjetas de forma asíncrona

5.16 Comparativa de rendimiento de las tres estrategias de descarga

5.17 Comparativa de tiempo de las tres estrategias de descarga

5.18 Evaluación del rendimiento del cálculo del Stencil 3-D descargado en dos MICs, para matrices que no caben en una MIC

A.1 Modificadores y cláusulas explícitas de descarga

A.2 Descripción de in/out/nocopy y uso de la cláusula longitud

Resumen

El auge de los campos de investigación basados en la simulación y el modelado, junto con el aumento de las necesidades del sector servicios (web y bases de da- tos), están poniendo a prueba las capacidades de cómputo de las arquitecturas de alto rendimiento. La tendencia actual es que los supercomputadores basen su di- seño en arquitecturas heterogéneas, donde se unen los tradicionales procesadores de altas prestaciones con un número importante de aceleradores/coprocesadores.Entre los patrones de cómputo más utilizados por los grandes centros de datos se encuentran los métodos de diferencias finitas, también conocidos como patrones Stencil, comúnmente utilizados para resolver ecuaciones en derivadas parciales.

Este documento se centra en el diseño, análisis y evaluación de aplicaciones científicas sobre plataformas de computación de altas prestaciones basadas en arquitecturas heterogéneas. En concreto, abordamos el desarrollo de códigos científicos basados en patrones Stencil sobre una arquitectura heterogénea x86, donde conviven procesadores orientados a latencia y procesadores orientados a rendimiento (throughput), es decir, procesadores Intel Xeon multinúcleo con tarjetas coprocesadoras Intel Xeon Phi. Hemos evaluado los dos modos de trabajo que tiene Xeon Phi, nativo y descarga (offload). Para el modo nativo ofrecemos una guías que ayudan a optimizar el rendimiento y la eficiencia energética de códigos basados en patrones Stencil. Utilizamos tres kernels a modo de ejemplo (difusión acústica, sísmica y de calor), mostrando cómo vectorizar y paralelizar el código, además de otras técnicas sencillas que mejoran el rendimiento. Para el modo offload mostramos cómo ejecutar el código usando dos tarjetas Xeon Phi, lo cual permite mejorar las prestaciones del código así como manejar tamaños de problema que superan la capacidad de memoria de una sola tarjeta. Para paliar la sobrecarga de las comunicaciones entre las tarjetas y el procesador Xeon proponemos una estrategia de intercambio de datos relajada, que como contrapartida reduce la precisión de los resultados obtenidos.

A lo largo de todo el documento hemos tenido como objetivo comparar, cuando ha sido posible, nuestros resultados con los obtenidos en otras dos arqui- tecturas de altas prestaciones: Xeon multinúcleo y GPUs de Nvidia (arquitectura CUDA). Dicha comparación tiene en cuenta tanto el tiempo de ejecución como la eficiencia energética. El camino para dominar la arquitectura Xeon Phi comienza en Knights Corner y su uso como acelerador, lo cual sentará las bases de conoci- miento para futuras generaciones como Knights Landing (KNL) o Knights Hill (KNH).

I Motivación y fundamentos

Capítulo 1 Motivación

1.1 Introducción

Una arquitectura de computadores heterogénea viene definida por la combinación de elementos de cómputo de distintas características, los cuales no tienen por qué compartir ni el espacio de direccionamiento de memoria ni el conjunto de instrucciones (ISA). El ejemplo más común es aquel en el que se combinan CPUs multinúcleo de propósito general con aceleradores especializados. Aunque este tipo de arquitecturas aparecieron hace años en el ámbito de la computación de alto rendimiento, ha sido en la última década cuando se han popularizado. Esto se debe principalmente a sus beneficios en términos de rendimiento y eficiencia energética.

En el año 2006 se presentó la arquitectura CUDA de Nvidia [20], con la nove- dosa propuesta de usar las tarjetas de procesamiento gráfico, o Graphics Processing Units (GPUs), para computación de propósito general. Esto produjo un profundo cambio tecnológico en el hardware y software que ha facilitado el uso y el acceso a arquitecturas heterogéneas masivamente paralelas. Dichas arquitecturas están presentes en todos los ámbitos de nuestra vida cotidiana: ordenadores de sobre- mesa, teléfonos móviles, entretenimiento, vehículos, medicina, etc. Más tarde, en el año 2012, se presentó la arquitectura Intel Many Integrated Core (MIC) [28] [52], más enfocada al ámbito empresarial y de computación de alto rendimiento.

Esta combinación CPU y aceleradores se ha convertido rápidamente en la solución más viable para construir supercomputadores de alto rendimiento, como se demuestra en los computadores que copan la lista Top500 [2]. Esta lista recoge los 500 supercomputadores más rápidos del mundo¹. A fecha de presentación de este documento, los dos supercomputadores más potentes están basados en aceleradores:

- Tianhe-2. Instalado en el NSCC y basado en aceleradores Intel Xeon Phi.
- Titan. Utiliza aceleradores (GPGPUs) de Nvidia y se encuentra emplazado en el ORNL.

Lamentablemente, desarrollar aplicaciones para estos sistemas no es una tarea sencilla, dada la enorme variabilidad del hardware subyacente. Los desarrollado- res tienen a su disposición nuevos modelos de programación que les ayudan a producir código para este tipo de arquitecturas heterogéneas. Estos modelos ga- rantizan la portabilidad del código fuente, pero, por regla general, no garantizan una ejecución óptima en términos de energía o tiempo. Con el fin de maximizar el rendimiento y la eficiencia energética de dichos sistemas, los desarrolladores de software necesitan hacer uso de los últimos avances tanto a nivel hardware, como en el campo específico de interés del problema. Este enfoque vertical, conocido como co-diseño hardware-software, permite a los desarrolles de aplicaciones sacar el máximo partido del hardware, y a los desarrolles hardware localizar los cuellos de botella de sus arquitecturas y mejorarlos en las siguientes revisiones del hardware.

La principal desventaja de este modelo de desarrollo es que conlleva un sobre-coste doble, tanto en la formación de los desarrolladores como en el tiempo de desarrollo. Por suerte, la mayoría de las aplicaciones utilizan métodos algorítmicos con patrones de cómputo y comunicación bien definidos (conocidos en el ámbito académico como Dwarfs [5]). Estos patrones se repiten en diferentes aplicaciones y, por tanto, las soluciones de hardware y software se pueden extrapolar entre áreas científicas. Un caso concreto de patrón de cómputo muy común a una amplia variedad de aplicaciones científicas y de ingeniería es el kernel computacional Stencil. Se trata de un kernel de cálculo que presenta unas propiedades muy adecuadas para mostrar las bondades y limitaciones de una arquitectura basada en aceleradores. El objetivo de este documento es la implementación, optimización y eva- luación de distintas aplicaciones basadas en dicho kernel sobre una plataforma heterogénea basada en el coprocesador Intel Xeon Phi, muy utilizada por los supercomputadores actuales. Dicha evaluación no sólo nos permitirá conocer el rendimiento y la eficiencia energética de dicha arquitectura heterogénea, sino que nos permitirá sintetizar una guía que ayudará a otros desarrolladores a optimizar sus aplicaciones para esta plataforma. El camino para dominar la arquitectura Xeon Phi comienza en Knights Corner y su uso como acelerador, lo cual sentará las bases de conocimiento para futuras generaciones como Knights Landing (KNL) o Knights Hill (KNH).

1.2 Contexto del problema

A continuación vamos a describir más ampliamente el contexto en el que se sitúa el desarrollo de este documento. Como ejes principales a desarrollar, mostraremos los siguientes: a) La importancia actual de las arquitecturas de alto rendimiento (HPC) y su necesidad de mejorar la eficiencia energética, b) la arquitectura masivamente paralela de Intel (Intel Xeon Phi), c) la programación paralela y d) los códigos Stencil aplicados a problemas de ciencia e ingeniería.

1.2.1 Arquitectura HPC (High Performance Computing)

El auge de los campos de investigación basados en la simulación y el modelado, junto a la teoría y experimentación, están poniendo a prueba las capacidades de cómputo de las arquitecturas de alto rendimiento. Campos tales como la mecánica cuántica, la predicción de fenómenos físicos, la investigación del clima, el análisis de ADN, etc, se suman a las necesidades del sector servicios (web y bases de datos) en un aumento constante e imparable de las necesidades de cómputo de los centros de datos [61, 33, 39, 5].

Un claro ejemplo de las elevadas necesidades de cómputo en la investigación actual es el proyecto Human Brain, que tiene como objetivo la simulación de la funcionalidad completa del cerebro humano. Los estudios preliminares sugieren que se necesitarán al menos 100 petabytes de memoria y 1 EFLOP (Exaflop) de potencia de cálculo para lograr este objetivo. El objetivo actual de la EU y del DARPA en EEUU es cruzar la frontera Exascale en el año 2020 6. Recordemos que el supercomputador más potente disponible en la actualidad apenas alcanza los 0,033 Exaflops, con unas necesidades totales de disipación de potencia eléctrica en el rango de los 20 MWatts. La tendencia actual es que los avances en la tecnología de procesador, así como en el subsistema de memoria y las redes de interconexión nos permitan doblar la potencia pico de cálculo cada 18 meses de una manera regular (Figura 1.1). Por lo tanto, podría parecer que, siguiendo esta tendencia, alcanzaríamos el hito de un pico de cálculo de 1 EFLOP dentro de los próximos 5-8 años.

Figura 1.1: Rendimiento previsto de los supercomputadores en los próximos cinco años, lista Top500.

Abbildung in dieser Leseprobe nicht enthalten

Sin embargo, las mejoras en la eficiencia energética de los sistemas actuales está llegando a su límite, complicando el desafío Exascale. Mientras los equipos más eficientes (green), de acuerdo a la lista Green500 - www.green500.org 3, alcanzan la cifra de 4 GFLOPs por vatio consumido, los ordenadores que ocupan las primeras posiciones en el Top500 sitúan su eficiencia energética a la mitad de dicho valor (es decir, en el orden de 2 GFLOPs por vatio). Incluso si pudiéramos escalar linealmente la eficiencia energética del equipo más eficiente de hoy en día, en 2020 un sistema de este tipo consumiría más de 200 MW. Nos enfrentamos pues a una mejora de la eficiencia energética de todos los componentes del sistema en un factor de 25x para los próximos 4 años (hasta alcanzar los 50 GFLOPS por vatio).

Los gobiernos y las agencias nacionales de investigación de todo el mundo han colocado a la investigación en arquitecturas HPC como objetivo estratégico prioritario para la próxima década (2015-2025) 34. La tendencia actual es que los supercomputadores basen su diseño en arquitecturas heterogéneas, donde se unen los tradicionales procesadores de altas prestaciones con un número impor- tante de aceleradores/coprocesadores. Los procesadores actúan como anfitriones o coordinadores, ejecutando tareas que requieren baja latencia, mientras que los aceleradores proporcionan un alto rendimiento computacional (throughput) y una elevada eficiencia energética. De hecho, la mayoría de supercomputadores en las últimas ediciones del Top500 están basados en arquitecturas heterogéneas.

1.2.2 El acelerador Intel Xeon Phi

El acelerador (también llamado coprocesador) Intel® Xeon Phi™ extiende el concepto de arquitectura multinúcleo de Intel a una arquitectura con muchos núcleos (hasta 61 en KNC y 68 en KNL). Esta arquitectura consigue un elevado rendimiento (throughput) y una alta eficiencia energética en aplicaciones masi- vamente paralelas. Al mismo tiempo, se beneficia de herramientas comunes de desarrollo de software con la arquitectura Intel multinúcleo tradicional, en lo que se refiere al análisis y depurado de las aplicaciones. Entre las características más importantes que presenta Intel Xeon Phi, tenemos las siguientes [54]:

- La migración de aplicaciones paralelas desarrolladas para entornos multi- núcleo (por ejemplo, Intel Xeon) es prácticamente directa, minimizando el esfuerzo de desarrollo para Xeon Phi. Aunque Xeon Phi es un coprocesador conectado al bus PCI, éste se puede comportar como si fuera un procesador Xeon con un número mucho más alto de núcleos (cuando empleamos el modo nativo de ejecución). Esto será particularmente cierto con la segunda generación de Intel Xeon Phi, Knights Landing, que puede trabajar de mane- ra totalmente autónoma, sin necesidad de que otro procesador actúe como host o anfitrión.

- Su compatibilidad con la arquitectura x86 también significa flexibilidad. Incluso cuando Xeon Phi trabaja como coprocesador, la gran diferencia con la forma tradicional de usar un acelerador es que no necesita ejecutar el código de forma lock-step. En otras palabras, permite una programación más cómoda usando un lenguaje de alto nivel, como OpenMP o Cilk+, sin tener que estar manejando las peculiaridades particular de cada hilo.

- Xeon Phi también se beneficia de las mismas técnicas y herramientas para optimizar las aplicaciones que se utilizan en los procesadores multinúcleo tradicionales de Intel. Podemos usar toda la suite de herramientas de Intel para analizar una aplicación en cuanto a su escalabilidad, uso de datos vectoriales, acceso a memoria, búsqueda de cuellos de botella, etc. Esto implica una curva de aprendizaje mucho más liviana que en el caso de otras arquitecturas (por ejemplo, CUDA) y, por tanto, una mayor productividad.

- Adicionalmente, y debido a que Xeon Phi tiene un ancho de banda de memoria significativamente mayor que una CPU, hay muchos códigos limitados por memoria (memory-bound) que se ejecutan más rápido tan sólo portándolos a Xeon Phi, sin tener en cuenta ninguna de las características específicas de la propia arquitectura.

- Por último, Intel Xeon Phi es la arquitectura de altas prestaciones más eficiente energéticamente de Intel. Es capaz de alcanzar una potencia de cálculo pico de algo más de 1 TFLOP usando datos en coma flotante de doble precisión con un TDP (Thermal Design Point) de 225-250 Watts. Esta arquitectura tiene alrededor de 5 mil millones de transistores (tipo 3D tri-gate) en un único chip (die) usando la tecnología de 22 nm.

Vamos a concluir insistiendo en que la gran ventaja que tiene la arquitectura Intel Xeon Phi es que añadiendo únicamente un parámetro al compilador, en nuestro caso -mmic, podemos obtener un código objeto compatible con la ar- quitectura Xeon Phi. Posteriormente, en un segundo proceso de optimización, podemos tener un cuenta las características específicas de la nueva arquitectura para mejorar el rendimiento. El camino para dominar la arquitectura Xeon Phi comienza en KNC y su uso como acelerador, para continuar en KNL y futuras generaciones como Knights Hill (KNH).

1.2.3 La programación paralela

Para obtener el máximo rendimiento posible, tanto de una arquitectura mul- tinúcleo como de una arquitectura masivamente paralela (tipo Xeon Phi), es necesario programar en paralelo. En la actualidad existe una gran diferencia en la calidad del código producido por un programador experto y un programador corriente debido a su conocimiento y experiencia. En el artículo aparecido en la conferencia ISCA del año 2012 59, denominan a este salto cualitativo “ninja gap”. Los autores se cuestionaban si existe un “ninja gap” insalvable a la hora de programar eficientemente en paralelo, y si sólo unos pocos elegidos podrían sacar partido a las nuevas arquitecturas paralelas que estaban apareciendo por entonces. En dicho artículo, los autores mostraban cómo gracias a la arquitectura Xeon Phi y su modelo de programación, los desafíos quedaban acotados y un programador normal podría desarrollar rápidamente programas para las nuevas plataformas paralelas.

Dicha tarea puede implicar un rediseño del programa e incluso un replantea- miento del algoritmo utilizado. J. Reinders y J. Jeffers hablan más de modernizar el código que de portar el código para el caso de Xeon Phi, por su peculiar carac- terística de que se puede ejecutar el mismo código que se ejecutaría en cualquier arquitectura Intel multinúcleo [56, 57]. El primer paso del proceso de moderni- zación del código es aprender a pensar en paralelo, es decir, a plantearse cómo el código debe ser estructurado para que se pueda ejecutar de forma concurrente por varios procesadores. También debemos plantearnos qué comunicaciones son necesarias entre los hilos, y cómo declarar y mapear los datos del programa para que puedan ser usados por los procesadores de forma local. Finalmente, nos interesa poder cargar y almacenar los datos en memoria por bloques que puedan ser fácilmente usados por las unidades vectoriales.

Por tanto, de cara a explotar adecuadamente una arquitectura paralela, hay tres factores clave que hay que tener en cuenta:

1. Paralelizar adecuadamente el código para que escale y pueda usar todos los núcleos (hilos) ofrecidos por la arquitectura.
2. Vectorizar adecuadamente dicho código para que pueda usar todas las unidades vectoriales ofrecidas por la arquitectura.
3. Manejar adecuadamente la localidad de los datos (tanto espacial como temporal) para que facilite una adecuada vectorización y paralelización.

Tanto en el campo de los procesadores multinúcleo como en Xeon Phi, pro- gramar en paralelo suele suponer extender un programa científico ya escrito en Fortran o C/C++ con unas directivas que permitan expresar el paralelismo necesario. Dichas directivas se encuentran en lenguajes como OpenMP, MPI, o OpenCL. Intel también ha desarrollado extensiones contenidas en lenguajes como Intel Threading Building Blocks o Intel Cilk+. De cara a la vectorización del código, el compilador de Intel presenta la característica de la autovectorización, que puede alcanzar buenos resultados sin un gran esfuerzo por parte del pro- gramador. En caso necesario, la versión 4.0 de OpenMP incorpora la directiva #pragma SIMD que fuerza la vectorización del código si el compilador no ha sido capaz de hacerlo. Existen también una serie de directivas y consejos relativos al alineamiento y conversión de las estructuras de datos que ayudan a obtener mejores resultados.

Por último, comentar que Intel Xeon Phi proporciona dos modos de progra- mación que le permiten ser utilizado como un coprocesador (utilizando el modo descarga), o como un nodo independiente (utilizando el modo nativo). Elegir un modo u otro dependerá de las necesidades de cómputo y de los requerimientos de memoria de la aplicación.

1.2.4 Patrones de cálculo Stencil

Existen numerosas aplicaciones en una amplia variedad de campos de la ciencia e ingeniería que requieren para su funcionamiento de la resolución de ecuaciones en derivadas parciales. Este tipo de problemas pueden ser resueltos mediante el método de las diferencias finitas, que calcula una solución de forma aproximada. Frecuentemente, este tipo de operaciones implican el acceso a grandes volúmenes de datos siguiendo patrones regulares. Estos patrones son los que se conocen como operaciones Stencil, y se llevan a cabo accediendo a la matriz donde se guardan los datos y actualizando cada uno de los elementos con valores correctamente ponderados de los elementos vecinos.

Los coeficientes usados para ponderar los valores de los diferentes elementos vecinos son diferentes según el problema. Esta operación sobre la matriz se aplica durante un número finito de veces o unidades de tiempo, ya sea buscando la convergencia de los datos, o con la intención de analizar el estado de la matriz tras un cierto número de iteraciones.

Una característica muy interesante del patrón de cálculo Stencil es que trata problemas que necesitan acceder a una gran cantidad de datos para realizar su cálculo, estando, por tanto, limitados por el ancho de banda de memoria que tenga la plataforma donde se ejecuten dichos códigos. Esta característica se pone más o menos de manifiesto dependiendo del número de vecinos que considera el kernel en cada cálculo, y la dimensionalidad del mismo.

1.3 Objetivos

Este documento se centra en el diseño, análisis y evaluación de aplicaciones científicas sobre plataformas de computación de altas prestaciones basadas en arquitecturas heterogéneas. En concreto, nos centraremos en el desarrollo de códigos científicos basados en patrones Stencil sobre la arquitectura heterogénea x86, donde conviven procesadores orientados a latencia y procesadores orientados a rendimiento (throughput), es decir, el procesador Intel Xeon multinúcleo y el coprocesador Intel Xeon Phi.

Los coprocesadores Intel Xeon Phi analizados están basados en la arquitectura (Knights Corner). Este coprocesador presenta dos modos principales de funciona- miento, el modo acelerador (también denominado modo descarga u offload), y el modo nativo. El modo descarga es la forma tradicional de usar los aceleradores, donde el procesador inicia la ejecución de la aplicación y ejecuta remotamente partes del código en el acelerador, normalmente secciones críticas masivamente paralelas. Por el contrario, el modo nativo permite a Xeon Phi comportarse como un procesador autónomo que puede ejecutar por sí mismo una aplicación de prin- cipio a fin. Este modo nativo va a tener continuación en la segunda generación de Xeon Phi (Knights Landing).

Comenzaremos centrándonos en el modo nativo, evaluándolo con diversos códigos Stencil, dado el carácter novedoso de dicho modo de trabajo para una tarjeta aceleradora. Describiremos en detalle los pasos necesarios para conseguir un rendimiento óptimo en Xeon Phi, comparando su potencia de cálculo y con- sumo energético contra la arquitectura CUDA de Nvidia y contra la arquitectura Intel Xeon multinúcleo. En la parte final del documento evaluaremos experimen- talmente el modo offload, centrándonos en el Stencil de difusión de calor 3-D de 11 puntos. Además de experimentar con el coprocesador en modo acelerador, usaremos esta característica para mejorar el rendimiento utilizando múltiples aceleradores, aumentando la memoria total disponible para la resolución del problema, lo cual nos permite la ejecución del código con tamaños de problema mayores.

Concretando, los principales objetivos que hemos procurado cubrir en este documento han sido los siguientes:

a) Desarrollar el estado del arte de nuestro problema. Este documento se enmar- ca dentro del ámbito de las investigaciones multidisciplinares. Por tanto, resulta necesario estudiar los trabajos realizados en:

- Arquitecturas heterogéneas para HPC: Intel Xeon Phi y Nvidia CUDA
- Programación paralela en Xeon Phi y herramientas relacionadas
- Códigos Stencil aplicados a problemas de ciencia e ingeniería
- Consumo de energía y su medición en los aceleradores

b) Ofrecer guías básicas para programar códigos Stencil. Ya que una de las carac- terísticas más atractivas de Xeon Phi es su facilidad para desarrollar programas paralelos, nuestro primer objetivo ha sido evaluar la veracidad de esta afir- mación, y ofrecer al programador inexperto unas sencillas directrices que le permitan usar esta arquitectura de manera eficiente. Este objetivo se desarro- lla utilizando el modo nativo de ejecución. Tomando como caso de uso tres códigos Stencil 3-D con diferente números de vecinos, hemos estudiado:

- Cómo vectorizar y paralelizar el código
- La influencia de los datos de entrada
- Técnicas sencillas para mejorar el rendimiento

Como punto de referencia para establecer la bondad de esta nueva arquitectura se comparan nuestros resultados con los obtenidos en Intel Xeon multinúcleo.

c) Estudio de la eficiencia energética y comparativa con CUDA. Nuestro tercer objetivo ha sido estudiar la eficiencia energética de las distintas platafor- mas evaluadas en este documento. En primer lugar, y por medio del uso de optimizaciones avanzadas (planificación de hilos y mejora del acceso a memoria), hemos mejorado el rendimiento de nuestros códigos Stencil 3-D. Posteriormente realizamos una comparativa con Xeon multinúcleo utilizando dos métricas importantes: a) Energía (Julios) y potencia (vatios) consumida al ejecutar nuestros códigos de prueba y b) eficiencia energética, utilizando la métrica EDP (Energy Delay Product), que nos muestra el producto del tiempo de ejecución por la energía consumida en dicha ejecución. Dado que Xeon Phi no es el único acelerador disponible en el mercado, también hemos comparado nuestros resultados con los obtenidos en diversas GPUs de Nvidia con soporte CUDA, utilizando un patrón Stencil de difusión acústica de 7 puntos. Esta comparación se ha realizado en términos de tiempo de ejecución así como de eficiencia energética.

d) Evaluación de Xeon Phi como coprocesador (acelerador). En los objetivos ante- riores hemos utilizado Xeon Phi en modo nativo, es decir, ejecutando los diferentes códigos en Xeon Phi de forma totalmente independiente de la má- quina anfitrión en el que está instalada la tarjeta aceleradora. Nuestro último objetivo, de cara a cerrar el trabajo del documento, se centra en evaluar el funcionamiento de Xeon Phi cuando trabaja en el modo de descarga (offload). Este modo de funcionamiento nos permite escalar el código a varias tarjetas, así como ampliar el tamaño de los datos de entrada al de la suma de las memorias individuales de cada tarjeta. En este modo se nos presenta un problema muy interesante, la comunicación de los datos obtenidos después de cada iteración del algoritmo entre las diversas tarjetas que están resolviendo el problema. Para reducir este intercambio de datos, hemos propuesto una solución basada en buscar un compromiso entre la precisión de los resultados y la velocidad de cómputo de los mismos. El manejo Xeon Phi en modo offload lo hemos abordado haciendo uso de las extensiones LEO (Language Extension Offload) de Intel.

1.4 Estructura del documento

Este documento se estructura en tres secciones principales y una Sección de apéndices. La Parte I está dedicada a la motivación y fundamentos. Esta Sección la integran dos Capítulos donde se introduce y motiva el trabajo desarrollado, mencionando los antecedentes necesarios para entender perfectamente el resto del documento. El primer Capítulo de esta parte presenta un breve resumen del documento, describiendo la motivación, objetivos principales y la estructura del documento.

El segundo Capítulo introduce al lector en el contexto de la computación de alto rendimiento, las arquitecturas heterogéneas basadas en aceleradores (con especial énfasis en la arquitectura Intel Xeon Phi), las herramientas para programar estas arquitecturas, y los fundamentos teóricos de los códigos Stencil 3-D. También presenta los algoritmos utilizados para la generación de los códigos objeto utilizados en las Secciones de evaluación de este documento, incluyendo: i) Stencil 3-D de difusión acústica de 7 puntos, ii) Stencil 3-D isotrópico de onda sísmica de 25 puntos, y iii) Stencil 3-D de difusión de calor de 11 puntos.

La Parte II muestra el manejo y la evaluación de la arquitectura Intel Xeon Phi a través de los casos de uso basados en patrones Stencil 3-D, tanto en el modo nativo de programación como en modo descarga. El tercer Capítulo presenta los resultados de la evaluación de Intel Xeon Phi en modo nativo contra un proce- sador Xeon multinúcleo ejecutando los tres códigos Stencil 3-D anteriormente mencionados. Esta evaluación nos permite proponer una serie de directrices para mejorar los códigos Stencil 3-D en el coprocesador Intel Xeon Phi.

El cuarto Capítulo presenta una comparativa en términos de eficiencia energé- tica entre Xeon multinúcleo y Xeon Phi. En primer lugar, mejoramos los códigos que habíamos desarrollado en el Capítulo anterior por medio del uso de optimiza- ciones avanzadas: a) planificar la distribución de las iteraciones que corresponden a cada hilo, b) elegir la mejor política de afinidad entre hilos y núcleos y, c) mejorar el acceso a memoria por medio de la reutilización de datos utilizando mecanismos de blocking. Posteriormente pasamos a comparar la eficiencia energé- tica entre Xeon y Xeon Phi en términos de EDP (Energy Delay Product). Viendo el buen comportamiento obtenido por la arquitectura Xeon Phi, nos planteamos extender dicha comparación con la arquitectura CUDA de Nvidia (usando las GPUs: Tesla k40c y GeForce 980 GTX). Esta última comparativa la realizamos sólo para el Stencil 3-D de difusión acústica.

En el quinto Capítulo analizamos el modelo de programación de descarga heterogénea (offload) para arquitecturas basadas en Xeon e Intel Xeon Phi. También mostramos los beneficios que conseguimos distribuyendo el código entre varias tarjetas, tanto en términos de rendimiento como en la capacidad de trabajar con entradas más grandes. También discutimos la relación que se establece entre rendimiento y precisión en los resultados obtenidos al ejecutar nuestros kernels Stencil.

Por ultimo, la Parte III resume el trabajo presentado en este documento, mos- trando las conclusiones generales de esta investigación, así como el trabajo futuro que se vislumbra en las arquitecturas masivamente paralelas. Además, se incluye en este Capítulo las principales contribuciones realizadas y las publicaciones relacionadas obtenidas con esta investigación.

El documento finaliza con una Sección de Apéndices, tres, más concretamente. En el Apéndice A damos toda la información necesaria para usar adecuadamente la extensiones LEO de Intel, a fin de manejar Xeon Phi en modo offload. En el Apéndice B ofrecemos el código fuente del kernel de difusión de calor de 11 puntos en cada una de las versiones evaluadas en los Capítulos 3, 4 y 5. Finalmente, en el Apéndice C mostramos una relación de las siglas utilizadas a lo largo del documento.

Capítulo 2 Fundamentos

2.1 Introducción

En este Capítulo se presentan los fundamentos de patrones Stencil 3-D así como las características más representativas de las arquitecturas Intel Xeon e Intel Xeon Phi (Knights Corner), que serán nuestras principales plataformas de evaluación en este documento. También se introducen brevemente la arquitectura de las GPUs de Nvidia, así como su lenguaje de programación CUDA. No obstante, el marco de trabajo principal de este documento es el paradigma propuesto por Intel y la programación para la arquitectura Intel Xeon Phi. Las tarjetas Nvidia se han utilizado con propósitos comparativos como se detalla en el Capítulo 4.

En este documento se han evaluado diversas soluciones Stencil 3-D utilizadas en múltiples campos científicos. Cabe destacar que las distintas implementaciones del patrón Stencil 3-D para las peculiaridades de un entorno concreto presentan características computacionales distintas, que son descritas con detalle en los fundamentos de los códigos Stencil de este Capítulo para los códigos de difusión de calor, acústica y onda sísmica.

Considerando que los lenguajes, herramientas y aplicaciones son compatibles entre los procesadores Intel Xeon y los coprocesadores Intel Xeon Phi, podríamos pensar en “reutilizar” en lugar de “recodificar” nuestras aplicaciones. Así que no estamos trabajando sobre un nuevo modelo de programación, más bien adaptamos, recompilamos y optimizamos nuestros códigos Stencil 3-D utilizando herramientas, librerías y programas. Al disponer de un único código fuente compartido entre los procesadores Intel Xeon y los coprocesadores Xeon Phi, sólo necesitamos optimizar una vez para extraer paralelismo y hacer uso de las unidades vectoriales, maximizando el rendimiento tanto en el procesador como en el coprocesador.

Finalmente, se concluye este Capítulo con una introducción a la segunda revisión de la arquitectura de Intel Xeon Phi, con nombre clave Knights Landing, ya que representa un paradigma emergente para la computación científica, solu- cionando muchos de los problemas de la primera generación. Además, será el primer procesador que soporta en su totalidad las extensiones AVX512 e incluirá memoria de alto ancho de banda (High Bandwidth Memory) en el propio die.

2.2 Arquitectura del coprocesador Intel Xeon Phi

2.2.1 Visión general

En noviembre de 2012 Intel lanzó al mercado el primer producto basado en la arquitectura Intel MIC (del ingles Intel Many Integrated Core) con nombre clave "Knights Corner"(ver Figura 2.1). La arquitectura Intel MIC se dirige a usuarios y aplicaciones capaces de sacar partido de un sistema procesamiento altamente paralelo, poniendo a disposición de los desarrolladores un gran número de núcleos de procesamiento de bajo consumo. Estos procesadores están fabricados con una tecnología de 22 nanómetros y un chip puede contener hasta 61 núcleos de procesamiento 52. Dichos núcleos disponen de unidades de procesamiento vectorial (en inglés Vector Processing Units - VPU) de ancho 512-bit. La VPU se puede utilizar para procesar 16 elementos de simple precisión u 8 de doble precisión por instrucción. Para mantener la disipación de energía por unidad de área bajo control, estos núcleos ejecutan las instrucciones en orden y funcionan a baja frecuencia (<1.3 Ghz). La arquitectura está respaldada por grandes cachés y un gran ancho de banda de memoria. Intel Xeon Phi se basa en la ISA x86, lo que le permite un cierto grado de compatibilidad con los procesadores x86 convencionales. No obstante, como veremos más adelante, Knights Corner no puede ejecutar los mismos ficheros binarios que otras plataformas Intel, y necesita que el código se compile de manera explícita para esta arquitectura. La segunda revisión de dicha arquitectura, Knights Landing, si es capaz de ejecutar binarios compilados para un procesador x86, aunque Intel recomienda recompilar para dicha arquitectura con el fin de obtener el mejor rendimiento de las unidades vectoriales. Knights Landing, a su vez, no es compatible con binarios de la anterior generación, Knights Corner.

Figura 2.1: Knights Corner - El primer coprocesador Intel Xeon Phi.

Abbildung in dieser Leseprobe nicht enthalten

La arquitectura MIC proporciona a los desarrolladores la ventaja de poder ser programada utilizando lenguajes de programación como C/C++ o Fortran, per- mitiendo mantener las herramientas y métodos de programación estándar, como son compiladores, librerías multi-thread y matemáticas para HPC, herramientas de depurado, etc. Intel Xeon Phi utiliza una distribución completa del sistema operativo Linux, de tal manera que el mismo código fuente de programas escritos para Intel MIC pueden recompilarse y ejecutarse en un procesador Intel Xeon estándar.

Figura 2.2: Procesador Intel Xeon y coprocesador Intel Xeon Phi juntos en una plataforma.

Abbildung in dieser Leseprobe nicht enthalten

La Figura 2.2 presenta un esquema típico del montaje de un nodo de cómputo de la plataforma Intel Xeon Phi. En cada nodo de cómputo coexisten uno o varios procesadores Intel Xeon (CPUs), denominados “host”, y de 1 a 8 coprocesadores Intel Xeon Phi, conectados a través de un bus PCI Express. Dado que el coprocesa- dor ejecuta de forma autónoma un sistema operativo Linux, es posible virtualizar una comunicación TCP/IP entre éste y los procesadores, permitiendo al usuario acceder al coprocesador como si se tratase de un nodo más de la red. Por tanto, cualquier usuario puede conectarse a un coprocesador Intel Xeon Phi a través de una sesión ssh (secure shell) y ejecutar sus aplicaciones. Este modo de ejecución se conoce como modo “nativo”. Otra forma de ejecutar aplicaciones en Xeon Phi es el modo “Offload”, en el cual una parte de la aplicación se ejecutaría en el host y otra en el coprocesador (ejecución heterogénea). Cuando un nodo cuenta con varios Intel Xeon Phi la comunicación entre ellos se puede establecer sin intervención del host.

Intel Xeon Phi se utiliza en la actualidad en muchos sistemas de alto ren- dimiento, entre ellos el supercomputador Tianhe-2 de la Universidad de Sun Yat-Sen en China, listado en el índice Top 500 como el supercomputador más rá- pido del mundo desde junio de 2013. Los 16.000 nodos de este supercomputador contienen dos procesadores Xeon (Ivy Bridge) y tres coprocesadores Xeon Phi cada uno, contando con un total de 3.120.000 núcleos que ofrecen un pico total de cálculo teórico de 54.9 PetaFLOPS.

2.2.2 Arquitectura de Intel Xeon Phi

Como hemos mencionado anteriormente, el coprocesador Intel Xeon Phi 52 28 está formado por hasta 61 núcleos de procesamiento basados en la arquitectura x86, memorias caché, controladores de memoria, lógica de cliente PCIe y un anillo de interconexión bidireccional que proporciona un elevado ancho de banda al sistema (ODI, on-die interconnect). La comunicación sobre el ODI es transparente para el desarrollador, dado que las transacciones se gestionan únicamente por el hardware. Los controladores de memoria y la lógica de cliente PCIe proporcionan un interfaz directa con la memoria GDDR5 del coprocesador y el bus PCIe respectivamente (Figura 2.3).

La Figura 2.4 muestra un diagrama de alto nivel que describe cada núcleo de procesamiento. Dichos núcleos están basados en una arquitectura superescalar en orden derivada del diseño del procesador Intel Pentium P54c con importantes mejoras, incluyendo soporte de instrucciones de 64 bits, unidades vectoriales, ejecución de múltiples hilos por núcleo, administración avanzada de energía, etc. El código se transfiere de la memoria del coprocesador a la caché de ins- trucciones del núcleo seleccionado por el planificador para ser decodificada y

Figura 2.3: Esquema general de la arquitectura de Xeon Phi.

Abbildung in dieser Leseprobe nicht enthalten

posteriormente enviada a ejecución. La arquitectura está diseñada para poder ejecutar hasta cuatro hilos por núcleo de forma independiente, donde cada núcleo (en orden) puede decodificar hasta dos instrucciones por ciclo. Esto permite una velocidad de ciclo más alta y facilita la gestión de los hilos. A diferencia de ar- quitecturas orientadas a latencia, la arquitectura MIC asume que las aplicaciones que se ejecutan en el sistema serán altamente paralelas y escalables. Con el fin de ocultar la latencia caché-memoria, la política de planificación intercambia hilos en cada ciclo. Cuando una aplicación se ejecuta en un solo hilo por núcleo, el planificador cambia a un hilo nulo especial antes de volver al hilo de la aplicación. Por tanto, Intel recomienda tener en ejecución al menos dos hilos por núcleo, mientras el óptimo puede variar entre 2 y 4. La ejecución de un solo hilo por núcleo reducirá la capacidad máxima del sistema a la mitad. Existen dos unida- des de procesamiento, una unidad escalar que ejecuta código x86 y x87, y una unidad de procesamiento vectorial (VPU - Vector Processing Unit).

La unidad de procesamiento vectorial cuenta con un repertorio de instruccio- nes SIMD (Single Instruction Multiple Data) de 512 bits conocido como Intel IMCI (Intel Initial Many Core Instructions). Esto se traduce en la ejecución simultánea de 16 operaciones de simple precisión (SP) u 8 de doble precisión (DP) por ciclo de reloj. Estas VPU también soportan instrucciones FMA (Fused Multiply-Add), que permite multiplicar y sumar en una misma instrucción, aumentando la precisión y mejorando el rendimiento. Asimismo, la arquitectura dispone de 32 registros

Abbildung in dieser Leseprobe nicht enthalten

Figura 2.4: Estructura individual de un núcleo.

vectoriales más 8 registros de máscara que facilitan la vectorización de bucles con flujos de ejecución condicionales. La VPU también soporta instrucciones de tipo gather y scatter por hardware. Este tipo de instrucciones facilitan enormemente la vectorización de aplicaciones con patrones de accesos a memoria irregula- res. Por último, la VPU también cuenta con una unidad matemática extendida (EMU - Extended Math Unit), que puede ejecutar funciones trascendentes (recípro- co, raíces cuadradas, funciones logarítmicas, etc). La EMU funciona calculando aproximaciones polinómicas de estas funciones.

Las unidades vectoriales proporcionan una mejora energética considerable, ya que una única operación codifica una gran cantidad de trabajo eliminando el coste energético que supondrían todas las etapas del pipeline para la ejecución de múltiples instrucciones.

2.2.3 Organización de la caché y accesos a memoria

Cada núcleo tiene una caché de segundo nivel de 512-KB privada inclusiva (L2), es decir, el contenido de la caché de primer nivel (L1) está replicado en la L2. Las cachés de primer nivel tienen una capacidad de 32-KB para datos y 32-KB para instrucciones. Las cachés L2 son coherentes y pueden intercambiar datos con otras cachés a través del anillo de interconexión en el chip. La coherencia se mantiene gracias a un directorio de etiquetas distribuido denominado TD (Tag Directory). Los parámetros clave de la estructura de la caché del coprocesador se enumeran en la Tabla 2.1.

Tabla 2.1: Parámetros clave de la caché del coprocesador Intel Xeon Phi

Abbildung in dieser Leseprobe nicht enthalten

Xeon Phi soporta la gestión de memoria virtual con tamaños de página de 4 KB (estándar), 64 KB (no estándar) y 2 MB (tamaños de pagina grandes). También incluye un TLB (Translation Lookaside Buffer) para acelerar la búsqueda de traducción de direcciones física a virtual. Las características del TLB se muestran en la Tabla 2.2.

Tabla 2.2: Parámetros TLB del coprocesador Intel Xeon Phi

Abbildung in dieser Leseprobe nicht enthalten

2.2.4 Interfaz PCIe y DMA

La versión comercial del coprocesador Intel Xeon Phi (en formato tarjeta) cumple con la especificación PCI Express Gen2 x16, permitiendo la comunicación direc- ta con otros dispositivos PCI express en la plataforma sin requerimientos de la memoria del host. Todos los intercambios de datos de entrada y salida del coproce- sador Xeon Phi utilizan el bus PCI Express. Los dos medios que se proporcionan para intercambiar datos entre el host y Xeon Phi son: a) el direccionamiento virtual asignado en memoria, recomendado para las transferencias de datos y mensajes de tamaño pequeño, y b) DMA (direct memory access), recomendado para tamaños grandes de datos.

2.3 Programación en Intel Xeon Phi

2.3.1 Modelos de programación en Intel Xeon Phi

La Figura 2.5 ilustra el amplio espectro de aplicaciones que se pueden abarcar utilizando procesadores y coprocesadores. Dependiendo de las necesidades de computo de la aplicación, la ejecución puede ser iniciada en un procesador central o en uno o más coprocesadores. En cualquier caso, el rendimiento óptimo de la aplicación se puede alcanzar en muchos casos utilizando una combinación de cálculo entre el procesador y el coprocesador.

Para ello, Intel Xeon Phi proporciona diversos modos de programación extre- madamente flexibles, pudiendo ser utilizado como un coprocesador o como un nodo independiente. Podemos distinguir los siguientes cuatro modos de uso, los mismos que se muestran en la Figura 2.5 y que describimos a continuación:

Abbildung in dieser Leseprobe nicht enthalten

Figura 2.5: Modos de aplicación de MIC.

1. Ejecución centrada en CPU. Este modo se usa en dos circunstancias: i) computación CPU pura, que comúnmente se aplica a los programas secuen- ciales y/o con bajo grado de paralelismo, y ii) computación centrada en CPU, con cooperación de MIC, que se aplica a los programas de cómputo secuencial que contienen segmentos de computación altamente paralela. La CPU inicia la función principal del programa. Cuando se llega a la parte de computación altamente paralela, ésta se ejecuta en la MIC.
2. Ejecución descargada (Offload) en MIC. Aplicado a programas de compu- tación altamente paralelos con un segmento parcial de computación secuen- cial (habitualmente al inicio y fin de la aplicación para las operaciones de E/S). Dichos programas son iniciados por la CPU, descargados a la MIC para que ejecute la computación paralela, y finalmente terminados en la CPU. La CPU se encarga principalmente de las operaciones de control y tareas de transmisión de datos (raramente realiza tareas de computación).
3. Ejecución CPU y MIC de igual a igual. Se aplica a múltiples programas de computación en paralelo. Dichos programas son iniciados por CPU y MIC al mismo tiempo.
4. Ejecución en MIC (nativo). Este modo se aplica a programas de compu- tación altamente paralelos que se ejecutan nativamente en MIC.

2.3.2 Modelo de programación offload

Esta sección describe el modelo de programación de descarga (offload) para Intel Xeon Phi, ya que va a ser objeto de estudio y evaluación en el Capítulo 5.

El modelo de programación offload ejecuta un programa en el procesador anfitrión (host), que puede opcionalmente lanzar o “descargar” porciones de código a uno o más coprocesadores Intel Xeon Phi. Los desarrolladores identifican líneas o secciones de código que consideran adecuadas para un procesamiento masivamente paralelo e insertan directivas similares a las utilizadas en OpenMP para marcarlas. Host y Xeon Phi no poseen una memoria unificada, por lo que es necesario mover los datos de ida y vuelta entre ambos. La Figura 2.6 muestra una visión conceptual del modo de uso offload.

El código es compilado para el host por un compilador compatible con la arquitectura Xeon Phi, que cuando encuentra las directivas de offload transfiere de manera transparente al programador tanto el código como los datos entre el host y Xeon Phi, los ejecuta, y recupera los datos de salida (también de

Abbildung in dieser Leseprobe nicht enthalten

Figura 2.6: Modelo offload sencillo entre un CPU (host) y un coprocesador Intel

manera transparente). Si Xeon Phi en la misma plataforma. manera transparente). Si Xeon Phi no está disponible, la línea o bloque de código se ejecutan en el host. Las palabras clave utilizadas para marcar una zona de código para offload son las directivas #pragma offload target(mic). Es importante mencionar que el rendimiento puede verse limitado por la falta de concurrencia entre el host y Xeon Phi, así como el movimiento implícito de datos antes y después de descargar el código.

El compilador soporta dos modelos de programación diferentes, que difieren de su enfoque para hacer frente a la falta de memoria compartida entre CPU y aceleradores: el modelo de memoria no compartida (non-shared memory model) y el modelo de memoria virtual-compartida (virtual-shared memory model).

El modo offload es un buen enfoque de programación siempre y cuando se cumplan una serie de condiciones: a) el código pase gran parte del tiempo ha- ciendo cálculos sin entrada/salida, b) la parte intensiva computacionalmente del código y los datos sobre los que trabaja sean relativamente fáciles de identificar y encapsular, c) el tiempo de cálculo sea sustancialmente mayor que el tiempo de transferencia de datos (al menos, N2cálculos para N datos), y d) los datos se ajustan a la memoria del coprocesador.

El coprocesador Intel Xeon Phi ofrece soporte para OpenMP 4.0 y el Lenguaje de Extensión para offload (Intel LEO). Este último provee a los desarrolladores de un modelo compartido implícito que va mas allá de lo que OpenMP 4.0 soporta. LEO nos permite establecer un modelo de memoria de programación compartida (llamado Intel MYO) entre el procesadores Intel Xeon y los coprocesadores Intel Xeon Phi. En el Apéndice A damos un extenso resumen de cómo manejar las extensiones Intel LEO.

2.3.3 Herramientas de desarrollo y nivel de aplicación

Abbildung in dieser Leseprobe nicht enthalten

Figura 2.7: Modos de paralelización y vectorización en MIC.

Podemos afirmar que no existe una diferencia significativa entre el modelo de programación MIC y el modelo de programación de una CPU. Si se quiere portar un programa secuencial existente para su ejecución en la tarjeta MIC, será necesario reescribir el programa secuencial en paralelo. Para facilitar esta tarea, los compiladores compatibles con MIC soportan lenguajes de programación que permiten expresar de manera explícita el paralelismo (OpenMP, OpenCL, pthreads o Cilk). La Figura 2.7 muestra los diferentes modelos de paralelización y vectorización soportados por los compiladores compatibles con la arquitectura MIC, donde la facilidad de uso suele ser inversamente proporcional al grado de control disponible por el programador de la arquitectura.

Al contrario que las GPUs, la arquitectura MIC no tiene un ISA independiente o un modelo de programación específico. La programación para MIC se basa, fundamentalmente, en una extensión del lenguaje C/C++/Fortran al que se agregan una serie de directivas del compilador (pragmas) parecidos a OpenMP que sirven para marcar regiones de código que se deben ejecutar en la MIC.

Figura 2.8: Herramientas de desarrollo y componentes de aplicaciones y compo- nentes de la interfaz.

Abbildung in dieser Leseprobe nicht enthalten

La Figura 2.8 representa una parte de la arquitectura de software, incluyendo los bloques marcados “Herramientas y Aplicaciones”. Las herramientas de desa- rrollo incluyen compiladores, librerías funcionales (como el Math Kernel Library (MKL) y Message Passing Interface (MPI) de Intel, depuradores, herramientas de optimización de programas (VTune), etc. Las aplicaciones de procesador y coprocesador se construyen para incluir los componentes de ejecución y librerías, que a su vez invocan el sistema operativo y a los componentes de middleware correspondientes.

La CPU y la MIC trabajan en espacios de direcciones de memoria indepen- dientes: la memoria del anfitrión (host) y la memoria del dispositivo (device). Los intercambios de datos entre el anfitrión y el dispositivo se realizan a través de una función API o de instrucciones específicas insertadas por del compilador. La memoria de dispositivo se utiliza casi de la misma forma que la del anfitrión, normalmente de manera transparente al usuario. El MPSS (Intel MIC Platform Software Stack) ofrece un API para que los usuarios de niveles más avanzados puedan hacer optimizaciones de rendimiento específicas más complejas.

2.4 Arquitectura CUDA de Nvidia

Este documento se centra en la programación y evaluación de la arquitectura Intel Xeon Phi para los patrones de códigos Stencil 3-D, tanto en modo nativo como en modo descarga. En el Capítulo 4, una vez optimizados dichos códigos, se presenta una comparativa de rendimiento y consumo de energía con la versión del código para GPUs basadas en la arquitectura CUDA, utilizando un código Stencil de 7 puntos (difusión acústica del sonido). Esta sección describe brevemente el modelo de programación CUDA y las principales características de esta arquitectura.

En noviembre de 2006 Nvidia lanzó CUDA 48 (Compute Unified Device Architecture), permitiendo a los usuarios programar sus tarjetas gráficas (GPUs) para la computación de propósito general. En pocos años, la programación de las GPUs para problemas científicos de diversa índole (GPGPU General Purpose for GPUs) amplió su campo y se convirtió en una de las mejores maneras de conseguir altas prestaciones en un gran conjunto de aplicaciones, logrando un alto rendimiento dentro de los procesadores disponibles en el mercado. Combinar el uso de la GPU y la CPU para acelerar cálculos ha propiciado un gran avance en cuanto a rendimiento de las aplicaciones en los últimos años a través de la programación en CUDA. CUDA hace referencia tanto a un compilador como a un conjunto de herramientas de desarrollo creadas por Nvidia. CUDA está basado en C con algunas extensiones, pero da soporte para C++, Fortran, .NET, Python, Java y otros lenguajes de programación ampliamente utilizados.

CUDA se basa en una jerarquía de capas de abstracción; el hilo (thread) es la unidad de ejecución básica; Los hilos se agrupan en bloques (blocks), cada uno de los cuales se ejecuta en un sólo multiprocesador, donde pueden compartir datos a través de la memoria compartida, una memoria pequeña pero con un gran ancho de banda 2.9.

Un grid se compone de una serie de bloques, que se distribuyen y planifican por igual entre todos los multiprocesadores. Las secciones paralelas de una aplicación que se desean ejecutar en la GPU se definen como kernels. Estos kernels se ejecutan siguiendo una paradigma de programación conocido como SPMD (Single Program Multiple Data), es decir, todos los hilos ejecutan el mismo código pero sobre distintos datos. Los hilos de un mismo bloque se dividen en grupos de 32 (en las arquitecturas Nvidia actuales), denominados Warps. El Warp es la

Figura 2.9: Hilos agrupados en bloques.

Abbildung in dieser Leseprobe nicht enthalten

unidad de planificación que se mapea a los cores de un multiprocesador, y se ejecuta en modo SIMD.

Figura 2.10: Bloques que componen un Grid.

Abbildung in dieser Leseprobe nicht enthalten

El modelo de programación CUDA asume que el kernel CUDA se ejecuta en un dispositivo físicamente separado que opera como un coprocesador corriendo un programa C en el host, tal y como se muestra en la Figura 2.11.

Abbildung in dieser Leseprobe nicht enthalten

Figura 2.11: Modelo de ejecución en CUDA.

El modelo de programación CUDA también asume que tanto el host como el device mantienen sus propios espacios de memoria por separado que, además, están implementados con diferentes tecnologías (DDR y GDDR), ofreciendo así diferentes anchos de banda.

Kepler and Maxwell son las últimas generaciones de la arquitectura CUDA de Nvidia. En comparación con diseños anteriores (Fermi), éstas amplían el número de núcleos dentro de un multiprocesador de 32 a 192 (128 SMM en Maxwell), y las unidades de programación de 2 a un máximo de 8 warps simultáneos. Además, la caché L2 duplica su tamaño. A pesar del aumento de los recursos (que resulta en una elevada cantidad de transistores por unidad de superficie), las GPU modernas son tres (seis) veces más eficientes energéticamente que las generaciones anteriores. Esto se consigue principalmente manteniendo la frecuencia por debajo de 1 GHz y utilizando un proceso de fabricación de 28 nm.

2.5 Evolución de Xeon Phi: Knights Landing

Anunciado para finales de 2015, y retrasado hasta mediados de 2016, en breve estará disponible el nuevo Intel Xeon Phi, conocido como Knights Landing. Aunque no va a llegar a tiempo para evaluar las implementaciones realizadas en este documento, nos gustaría aprovechar esta sección de fundamentos para dar una primera descripción de las características de esta nueva arquitectura.

Knights Landing utiliza núcleos de CPU x86 compatibles plenamente con extensiones AVX512 para ofrecer un rendimiento vectorial similar a los obtenidos por las GPUs: 3 TFLOPs trabajando con datos en doble precisión, y 6 TFLOPs en simple precisión, operando a una frecuencia de reloj de 1.31 GHz. Knights Landing utiliza una microarquitectura de CPU basada en el diseño Silvermont de Intel optimizada para HPC, eliminando la necesidad de un procesador anfitrión separado (host). Se espera que obtengan un rendimiento aproximado de un 80 % del de un servidor dual-socket consumiendo sólo una cuarta parte de la energía, o lo que es lo mismo, una mejora de la eficiencia energética de 3.5x FLOPs/Watt.

Knights Landing estará disponible en tres familias de productos, los cuales tendrán hasta 16 GB de MCDRAM. La diferencia entre estas familias será cómo interactúa el procesador con el resto del sistema. Dos de ellas podrán acceder a memoria (DDR4), mientras que la tercera es una tarjeta coprocesadora concebida como una actualización para los clientes Knights Corner que carecerá de soporte DDR4 y sólo utilizará MCDRAM.

En todas ellas la memoria MCDRAM se podrá configurar de tres maneras en el momento del arranque. Como memoria caché tradicional gestionada por el hardware, como memoria caché gestionada de forma explícita (tipo scratchpad), o bien como una mezcla de ambas soluciones (4 GB u 8 GB podrían funcionar como memoria caché, y el resto como memoria administrada de forma explícita).

El nuevo Knights Landing será el primer procesador que soporte las extensiones del conjunto de instrucciones AVX512, que aparecerán en los procesadores del servidor Skylake el próximo año. Las instrucciones AVX512 son un superconjunto las extensiones Knights Corner existentes (IMCI), aunque se codifican de distinta forma, por lo que el código binario compilado para la actual generación de coprocesadores Intel MIC necesitará ser recompilado.

Su arquitectura interna está organizada en una malla de 54 bloques (tiles), con ocho controladores MCDRAM colocados a lo largo de los bordes norte y sur de la malla, y dos controladores de memoria DDR4-2400, cada uno de ellos con tres canales, tal y como se muestra en la Figura 2.12.

Figura 2.12: Procesador Knights Landing.

Abbildung in dieser Leseprobe nicht enthalten

Los 38 bloques restantes contienen cada uno dos núcleos de CPU Knights Landing y 1 MB de caché L2, aunque para optimizar el rendimiento y el consumo de energía, los primeros modelos anunciados tendrán desactivados dos bloques, por lo que el procesador tendrá 72 núcleos de procesamiento y 36 MB de caché L2.

Cada núcleo Knights Landing puede emitir dos instrucciones por ciclo de reloj, tiene cuatro vías multithreading simultáneas, y posee dos unidades vectoriales de 512 bits. También puede ejecutar seis uops por ciclo con capacidad fuera de orden, ofreciendo hasta 32 FLOP de doble precisión por ciclo de reloj usando las dos unidades vectoriales AVX512.

La conexión de datos entre agentes de malla es de 32 bytes de ancho. Estima- ciones realizadas con el benchmark de memoria stream, sitúan el ancho de banda de cada EDC en más de 50GB/s, lo que significa que la malla debe operar a más de 1.05 GHz. Dentro de una columna vertical, cada salto en la malla tarda un único ciclo de reloj, sin embargo moverse a través de una fila horizontal tiene dos ciclos de latencia. Debido a que la mayoría de los agentes que sientan el tráfico de malla de inyección en la parte superior e inferior, la comunicación es enrutada y-x moviéndose primero verticalmente y luego horizontalmente.

2.6 Fundamentos de los códigos Stencil

Los códigos Stencil [67, 37, 11, 66, 30] son un tipo de kernels iterativos los cuales actualizan los datos de entrada de acuerdo con algún conjunto fijo predetermi- nado o patrón. Estos algoritmos se puede utilizar, por ejemplo, para calcular el valor de varios elementos de una matriz en un determinado punto de la ejecución en función de los valores de sus vecinos (y de ellos mismos) en instantes de tiempo anteriores. Los patrones Stencil son la base de los algoritmos de diferen- cias finitas utilizados para resolver ecuaciones diferenciales parciales (PDE) a gran escala y de gran dimensión. Las PDE proporcionan aproximaciones numé- ricas a problemas computacionales complejos, siendo ampliamente utilizadas en muchos campos científicos y de ingeniería [68, 12, 19, 76]. Esto permite a los científicos modelar con precisión fenómenos tales como la propagación de ondas escalares [35], conducción de calor, difusión acústica, etc.

2.6.1 Descripción

Existe una gran diversidad de códigos Stencil. Para diferenciarlos entre sí se pueden utilizar las siguientes características:

- Dimensión: Se refiere al número de ejes que tiene la estructura de datos sobre la que se realiza el cálculo del Stencil. La dimensión de un Stencil tiene un fuerte impacto en las optimizaciones que se puedan aplicar a éste. En un Stencil unidimensional, la reutilización de datos es muy sencilla, debido a que los datos están almacenados de forma adyacente en memoria y a que la distancia entre el primer y último elemento usado para cada cálculo es corta. Cuanta mayor sea la distancia entre elementos peor será su localidad espacial, perjudicando la reutilización de los datos. Éste es el caso de los códigos Stencil de 3 o mas dimensiones.

- Orden: Este parámetro especifica la distancia a la que se encuentran el vecino más alejado que se necesita para el cómputo del elemento actual en cualquiera de los ejes dimensionales. A mayor orden, mayor distancia entre los elementos de un mismo cálculo y mayor posibilidad de existencia de problemas de localidad espacial.

- Tamaño: Esta característica define el número total de elementos de la matriz que están involucrados en el cálculo. Alguno de los tamaños mas conocidos son de 5 puntos, 7 puntos, 9 puntos o 25 puntos. En la Figura 2.13 se muestra un Stencil 3-D de 11 puntos donde se considera un vecino en cada una de las tres dimensiones, más los vecinos que están en la diagonal del plano X-Y.

Abbildung in dieser Leseprobe nicht enthalten

Figura 2.13: Stencil 3-D de 11 puntos.

- Patrones iterativos: Esta característica define dónde y en qué orden se actualizan cada uno de los elementos. Los tres tipos básicos son:

- Jacobi. Este patrón iterativo es ampliamente utilizado por su sencillez y capacidad para ser paralelizado. La base del patrón consiste en realizar las iteraciones sobre una copia los datos que se utiliza sólo para lectura mientras se escriben los resultados en otra copia de los datos. Al final de cada iteración se intercambian las copias de los datos, la nueva entrada es la salida de la iteración anterior y la antigua copia pasa a ser la salida de la nueva iteración. Al ser necesario mantener dos copias del volumen de datos, necesitamos el doble de recursos de almacenamiento, consumiendo también el doble de ancho de banda a memoria. Esta duplicación de los datos es a la vez su mayor ventaja. Este patrón se puede paralelizar como convenga al programador al no existir dependencias de datos dentro de cada iteración, ya que los volúmenes de datos para lecturas y escrituras están totalmente aislados.

- Gauss-Seidel. En este patrón iterativo se llevan a cabo las escrituras en la misma copia que se usa para leer los datos. En consecuencia, los accesos memoria se reducen drásticamente. Este patrón dificulta el proceso de paralelización, ya que los accesos a los datos no son completamente independientes. Es decir, algunos elementos serán actualizados antes de que puedan ser consultados por los elementos vecinos para sus cálculos. Se hace necesario establecer un patrón de actualización para que el código sea determinista y se pueda replicar el mismo resultado final.

- Gauss-Seidel Red-Black (GSRB). Al igual que en Gauss-Seidel, este patrón escribe el resultado del cálculo en la misma copia de donde se leen los datos. La diferencia radica en que el patrón de actualización viene dado desde el comienzo y permite que la ejecución se realice en dos pasos independientes. Para esto se dividen los datos en “puntos” rojos y negros de forma que los puntos adyacentes tengan un color diferente. El patrón actualizará primero los puntos de un color y después los del otro. Como los puntos de un mismo color no forman parte del cálculo del resto de puntos de ese color, cada elemento de un color es independiente del resto. Sin embargo, aunque los requisitos de memoria son los mismos que en Gauss-Seidel, no es así con el ancho de banda, ya que al dividir la iteración de forma simple calculando primero los puntos de un color y después los del otro, aumenta la tasa de fallos de caché y, en consecuencia, el tráfico de datos a memoria.

- Bordes o fronteras. Para determinados algoritmos Stencil es necesario rea- lizar un trato especial sobre las partes limítrofes de los datos (bordes o fronteras). Destacando los siguientes:

- Borde constante. Se definen así los elementos limítrofes de la entrada de datos que no se actualizan. Suelen estar situados al comienzo y final de cada una de las dimensiones de los datos de entrada. Esta situación se puede simular creando un halo alrededor de los datos de entrada cuyos elementos no se actualicen con el paso del tiempo y mantengan valores fijos (Figura 2.14a).

- Borde periódico. Este tipo de fronteras se utiliza cuando los datos limítrofes de la matriz de datos se solapan, es decir, en una entrada de 0 a N, el elemento N+1 vuelve a ser 0 (Figura 2.14b).

Abbildung in dieser Leseprobe nicht enthalten

Figura 2.14: Bordes o fronteras de los códigos Stencil.

- Coeficientes. Los coeficientes son los valores mediante los cuales se pondera el peso de cada uno de los elementos y pueden ser constantes o variables.

2.6.2 Modelos matemáticos de los códigos utilizados

En esta sección vamos a describir los diversos modelos matemáticos en que se basan los códigos Stencil 3-D que hemos utilizado en el desarrollo de este documento. Más concretamente definiremos los modelos de difusión de calor, difusión acústica, y por último la ecuación isotrópica de onda sísmica. Todos ellos resuelven las ecuaciones diferenciales que representan los problemas matemáticos por medio del uso de métodos de diferencias finitas.

El método de diferencias finitas es una técnica numérica utilizada para resolver una ecuación diferencial en una región determinada, sujeta a condiciones de contorno especificadas. Cuando se utiliza un método de diferencias finitas, el dominio del problema se discretiza de manera que los valores de la variable dependiente desconocidos se consideran únicamente en un número finito de puntos nodales o celdas en lugar de en cada punto sobre la región. Una función discretizada se define como (Ecuación 2.1):

Abbildung in dieser Leseprobe nicht enthalten

El índice temporal n y los índices espaciales i, j y k se han introducido junto con la discretización temporal Δt y discretizaciones espaciales en el eje cartesiano Δx, Δy y Δz. Estos últimos se definen como la inversa de las resoluciones temporales y espaciales respectivamente. Por ejemplo, la baja resolución implica tamaños de celdas grandes y de alta resolución implica tamaños de celdas pequeñas.

2.6.2.1 Ecuación de difusión de calor

Consideremos un espacio de tres dimensiones representado por una matriz 3D XxYxZ, donde ut(x,y,z) es la temperatura de un punto (x,y,z) en el instante t y α es la difusión térmica. En dicho entorno, la difusión de calor se representa como (Ecuación 2.2):

Abbildung in dieser Leseprobe nicht enthalten

Esta ecuación diferencial parcial se puede resolver por aproximación mediante el uso de un cálculo del Stencil sobre el espacio discretizado y el tiempo, mediante el uso de la siguiente ecuación de actualización estilo Jacobi (Ecuación 2.3):

Abbildung in dieser Leseprobe nicht enthalten

2.6.2.2 Ecuación de difusión acústica

La ecuación de difusión para la densidad de energía del sonido w(r,t) en la posición r definida en un dominio V e instante de tiempo t y que incluye un término de fuente de sonido P(t) localizado en la posición rs, consiste de una ecuación diferencial parcial con condiciones de contorno mixtas 69 que se muestra en la ecuación (2.4).

Abbildung in dieser Leseprobe nicht enthalten

Dicha ecuación (2.4) es una ecuación diferencial homogénea parabólica parcial, donde ∇2 es el operador de Laplace y D = λc/3 es el coeficiente de difusión con c siendo la velocidad del sonido. Este coeficiente de difusión tiene en cuenta la geometría de la sala a través de su camino libre medio λ, lo que indica la distancia media que una partícula de sonido viaja entre dos colisiones consecutivas 43. En la teoría acústica clásica, el recorrido libre medio de una sala viene dado por λ = 4V/St, con volumen V y el área interior total St. El termino cmw(r,t) representa la atenuación atmosférica dentro de la sala, donde m es el coeficiente de absorción del aire [7].

La ecuación (2.5) es una condición de frontera mixta que los modelos de los efectos locales en el campo del sonido inducen por diferentes grados de absorción en las superficies. El termino n representa el vector de la unidad normal a la superficie limite. Esta ecuación permite expresar la distribución de las propiedades de absorción de la superficie a través del factor de absorción AX = AX(r,ff), donde α es el coeficiente de absorción. Diferentes definiciones de AX se han presentado en la literatura técnica, cada una en función de la teoría física asumida. En esta sección, el factor de absorción modificado [46, 29] se adopta para realizar las simulaciones en la ecuación (2.6):

Abbildung in dieser Leseprobe nicht enthalten

En este documento, la solución 3D-FD elegida para la ecuación de difusión acústica sigue el esquema de Dufort-Frankel 14, publicado en 45, por su idoneidad y características atractivas (por ejemplo, es incondicionalmente estable). Para simplificar, vamos a definir β0ν = (2DΔt)/(Δν)2, donde ν = [x, y, z], y β0 = [Abbildung in dieser Leseprobe nicht enthalten] por lo que el esquema de diferencias finitas para la ecuación (2.4) es la mostrada en la ecuación (2.7),

Abbildung in dieser Leseprobe nicht enthalten

El término fuente, incluido como una fuente suave, se añade en la posición adecuada como [illustration not visible in this excerpt]

Se necesitan relaciones adicionales para hacer que número de ecuaciones sea igual al número de variables desconocidas. Estas variables se obtienen a partir de las condiciones de contorno y se explican a continuación.

Una diferencia exacta de segundo orden de las condiciones de contorno de la ecuación (2.5) se utiliza para asegurar la exactitud de la aproximación. Por simplicidad, presentamos una aproximación de diferencias finitas de la superficie límite orientada en el eje x en ambas posiciones x = 0 y x = lx.

Abbildung in dieser Leseprobe nicht enthalten

La derivación de la dimensión y y la dimensión z es sencilla.

Por lo tanto, la ecuación (2.8) y la ecuación (2.9), junto con la ecuación (2.7) son las aproximaciones en diferencias finitas completas de la ecuación de difusión acústica 3D sometidas a condiciones de contorno mixtas.

2.6.2.3 Ecuación isotrópica de onda sísmica

La ecuación de onda es una ecuación diferencial parcial hiperbólica. La ecuación de onda isotrópica que consideramos en este documento es la (2.10):

Abbildung in dieser Leseprobe nicht enthalten

Donde ∇2p es el operador de Laplace, p los campos de presión y ∇ el campo de velocidad de onda-p.

Las soluciones de esta ecuación describen la propagación de alteraciones fuera de la región a una velocidad fija en una o en todas las direcciones del espacio, al igual que las ondas físicas del plano o fuentes localizadas; la constante c se identifica con la velocidad de propagación de la onda. Esta ecuación es lineal, por lo que la suma de las dos soluciones es de nuevo una solución (principio de superposición).

Desarrollando esta ecuación obtenemos:

Abbildung in dieser Leseprobe nicht enthalten

Luego, utilizando diferencias finitas:

Abbildung in dieser Leseprobe nicht enthalten

Donde:

Abbildung in dieser Leseprobe nicht enthalten

Aquí Co a Cn son los coeficientes de las diferencias finitas. Entonces, la presión en el siguiente instante de tiempo puede expresarse a partir de la presión en los instantes de tiempo actuales y anteriores:

Abbildung in dieser Leseprobe nicht enthalten

2.6.3 Algoritmos Stencil

Los algoritmos Stencil 3D se implementan como un triple bucle anidado que recorre la estructura de datos completa, mientras se actualiza cada punto del grid. El cálculo de cada elemento de salida por lo general requiere: a) la contribución ponderada de algunos vecinos cercanos en cada dirección definida por la física del problema, b) el valor anterior de ese elemento en el instante de tiempo t-1 y, c) un único punto correspondiente de otras matrices de entrada. Dependiendo del orden del código en el tiempo, el algoritmo utiliza dos o tres copias del grid espacial (en pasos de tiempo t y t+1, y a veces t-1), intercambiando sus roles como fuente y destino en pasos de tiempo alternos. El Algoritmo 1 muestra el pseudocódigo de una solución genérica del kernel Stencil 3-D para bucles espaciales con orden 2 en el tiempo.

El bucle más externo del Algoritmo 1 simula el tiempo de ejecución. Tras cada iteración, se comprueba la condición de convergencia para los valores de la matriz solución, y cuando se satisface dicha condición se finaliza la ejecución

Abbildung in dieser Leseprobe nicht enthalten

del algoritmo. De cara a la versión paralela de este algoritmo, es común evitar en cada iteración temporal la comprobación de dicha condición de convergencia pues impide obtener las aceleraciones en la ejecución del código prometidas por la arquitectura paralela destino. Por contra, se suele comprobar a priori cuantas iteraciones temporales han sido necesarias realizar para converger a la solución deseada y, posteriormente, el algoritmo siempre se ejecuta poniendo un bucle con un número fijo de iteraciones temporales igual al obtenido de forma off-line. Este es el procedimiento que hemos seguido a la hora de evaluar los algoritmos presentados en este documento.

Una característica importante de estos algoritmos es que los kernels Stencil 3-D generalmente sufren de una alta tasa de fallos de caché y pobre localidad de datos. La razón es que, para los tamaños de entrada que exceden la capacidad de la memoria caché, cuando reutilizamos una entrada del conjunto de datos ésta ya ha sido reemplazada en caché. Además, el patrón de acceso no lineal a memoria de las implementaciones 3-D aumenta la presión sobre el sistema de memoria. Como resultado, las implementaciones estándar de Stencils 3-D sólo alcanzan una pequeña fracción del rendimiento máximo del hardware [55].

2.6.3.1 Algoritmo Stencil 3-D de difusión acústica (7-puntos)

El Algoritmo 2 muestra la implementación 3D-FD del kernel del modelo de la ecuación de difusión acústica presentado anteriormente (7-puntos). Se implemen- ta como un triple bucle anidado atravesando el dominio computacional completo para actualizar cada punto del grid. Utilizamos tres matrices tridimensionales en el bucle del cómputo principal (wn+1, wn−1 y wn), pues es de segundo orden en

Abbildung in dieser Leseprobe nicht enthalten

el tiempo (para los instantes t-1, t y t+1). La Figura 2.15 muestra gráficamente el método multigrid empleado para implementar dicho cálculo Stencil.

Abbildung in dieser Leseprobe nicht enthalten

Figura 2.15: Patrón Stencil para 3D-FD del modelo de la ecuación de difusión acústica. Hemos considerado como variables el coeficiente de absorción del medio (en este caso el aire dentro de la sala (m)), la velocidad del sonido (c) y la discretización de tiempo (dt), además de la ecuación de la constante dentro de la sala (β0).

2.6.3.2 Algoritmo Stencil 3-D isotrópico de onda sísmica (25-puntos)

La Figura 3 muestra el algoritmo Stencil 3-D isotrópico de onda sísmica de 25- puntos. Está implementado como un cuádruple bucle anidado atravesando el dominio computacional completo para actualizar cada punto del grid.

Abbildung in dieser Leseprobe nicht enthalten

También aquí utilizamos tres matrices tridimensionales en el bucle del cómpu- to principal (prev[z,y,x], current[z,y,x], next[z,y,x]) pues es de segundo orden en el tiempo (para los instantes t-1, t y t+1). Además, utilizamos otra matriz tridimen- sional (vel[z,y,x]) para las constantes que modulan cada punto, y que para este algoritmo son dependientes de la posición de dicho punto.

2.6.3.3 Algoritmo Stencil 3-D de difusión de calor (11-puntos)

El Algoritmo 4 muestra el código fuente del kernel del Stencil 3-D de jacobi de 11- puntos. Se implementa como un tripe bucle anidado que atraviesa completamente el dominio computacional del problema mediante la actualización de cada punto del grid.

La actualización de cada elemento requiere del valor de entrada anterior, junto con los vecinos más próximos en cada dirección, y un único punto de la

Abbildung in dieser Leseprobe nicht enthalten

segunda matriz de entrada correspondiente. El código utiliza dos copias de la matriz espacial, intercambiando sus papeles en pasos de tiempo alternativos.

2.7 Herramientas para evaluar el rendimiento y la energía

2.7.1 La interfaz GUI micsmc

Dentro del paquete de herramientas Intel MPSS existen varias utilidades para gestionar y recuperar información de la tarjeta del coprocesador. Una de las más interesantes, y que más hemos utilizado a lo largo del documento, es la interfaz micsmc. Se trata de una aplicación para Intel MIC que permite la configuración y la gestión del sistema (desde el sistema anfitrión) de los coprocesadores Intel Xeon Phi conectados a dicho sistema. También cuenta con una versión de línea de comandos que ofrece la misma funcionalidad.

Como se puede apreciar en la Figura 2.16, la aplicación micsmc permite monitorizar la actividad del coprocesador utilizando un panel de control donde se muestra gráficamente información sobre todos los coprocesadores del sistema, incluyendo la temperatura del coprocesador, la utilización de los núcleos, el uso de la memoria, etc.

Abbildung in dieser Leseprobe nicht enthalten

Figura 2.16: Utilización de los cores en mic1 y mic2.

Una de las principales utilidades es el monitoreo de la utilización de los núcleos. En la Figura anterior se puede ver cómo el coprocesador mic1 está utilizando los diversos núcleos que tiene al 50 %, mientras que en mic2 se utilizan al 100%. De esta manera tan gráfica y sencilla podemos apreciar si nuestros códigos están utilizando de manera eficiente los recursos del coprocesador.

2.7.2 Intel Vtune Amplifier XE

Intel VTune Amplifier XE es una herramienta de línea de comandos y una interfaz gráfica de usuario que proporciona metainformación sobre la ejecución de una aplicación, es decir, contadores de rendimiento tales como tasa de fallos de caché, instrucciones ejecutadas, etc. Esta información ayuda a detectar las secciones de código críticas, también conocidos como hot spots, y optimizar el código para un mejor rendimiento en procesadores Intel. También puede proporcionar información crítica, como el ancho de banda utilizado por la aplicación durante su tiempo de ejecución, el comportamiento multihilo, o el posible desequilibrio de la carga, ayudando a detectar y eliminar posibles cuellos de botella que afectan al rendimiento. Ver Figura 2.17

Figura 2.17: GUI de VTune Amplifier XE 2015.

Abbildung in dieser Leseprobe nicht enthalten

El Vtune Amplificador XE viene con un controlador de muestreo que se instala en el sistema operativo del coprocesador Intel MIC. Este controlador se encarga de la comunicación con la unidad de supervisión del rendimiento del hardware Intel MIC y de la recolección de eventos del hardware correspondientes a las aplicaciones ejecutadas en el coprocesador. Se puede iniciar el proceso de recopilación de eventos desde el host mediante la interfaz GUI o línea de comandos proporcionada por la herramienta Vtune Amplificador XE.

Es importante establecer un proyecto de forma correcta sobre el VTune Am- plifier XE que sea responsable de la ejecución de la aplicación que se ejecutará en el coprocesador Intel MIC. Debido a que la aplicación se ejecuta desde el host, es necesario asegurarse de que el script de la configuración de especificación de destino realiza la carga necesaria del ejecutable con los archivos dependientes necesarios y ejecuta la aplicación en Xeon Phi. También es necesario construir su binario con los símbolos de depuración adecuados (-g -debug inline-debug-info), así como especificar los directorios de búsqueda usando la pantalla de configuración para que los archivos de origen sean encontrados por la herramienta. Esta ha sido una de las herramientas que hemos utilizado para explorar con más detalle las características de rendimiento de bajo nivel de las implementaciones de nuestros códigos Stencil 3-D (ver Figura 2.18).

Figura 2.18: Utilización de VTune en nuestras evaluaciones.

Abbildung in dieser Leseprobe nicht enthalten

2.7.3 Librería Papi 5.0

PAPI (Performance Application Programming Interface) 44 ofrece a los diseñadores y desarrolladores de aplicaciones un interfaz consistente y una metodología para el uso de contadores de rendimiento del hardware disponibles en la mayoría de microprocesadores comerciales. PAPI permite a los ingenieros de software ver, en tiempo real, la relación entre el rendimiento del software y los eventos del procesador.

En concreto en este documento utilizamos el módulo “rapl” de la librería PAPI para realizar mediciones de consumo (energía). Los procesadores Xeon estudiados incluyen tres registros de energía, uno para todo el paquete (PKG), uno para los núcleos (Power Plane 0 - PP0), y otro para la memoria (DRAM). Intel Xeon Phi también está soportado por PAPI, apoyándose en un componente externo (micpower). Este componente permite a los desarrolladores extraer las lecturas de potencia instantáneas (no de energía) para diferentes partes de la placa del Intel Xeon Phi, incluyendo la potencia disipada por los núcleos, el total de la tarjeta, la potencia del PCIe, etc.

2.8 Conclusiones

En la primera parte de este Capítulo se describe en detalle los fundamentos de la arquitectura Intel Xeon Phi, desde detalles hardware hasta particularidades de su programación. Entre estos detalles de programación destacamos los diversos modos de trabajo que ofrece. En los siguientes Capítulos del documento evalua- mos los dos más importantes, el modo nativo y el modo de Offload. Asimismo, describimos brevemente la arquitectura CUDA de Nvidia que utilizaremos para comparar las prestaciones con Intel Xeon Phi en el Capítulo 4. Para terminar, hemos presentado una pequeña introducción al nuevo Xeon Phi (Knights Landing) que se comercializará próximamente.

Desde un punto de vista software, hemos descrito los fundamentos de los patrones de cómputo Stencil, mostrando con todo detalle su fundamento matemá- tico, el proceso de cómputo y los principales algoritmos que vamos a implementar. En este documento se evalúan tres kernels basados en Stencils 3-D, concretamente el modelo de la ecuación de difusión de calor, el modelo de la ecuación de difu- sión acústica, y por último, el modelo de la ecuación isotrópica de onda sísmica. Dichos algoritmos Stencil van a ser nuestros códigos de prueba para evaluar el rendimiento de la arquitectura Intel Xeon Phi (Knights Corner).

Parte II Acelerando patrones Stencil 3-D en arquitecturas basadas en Xeon Phi Capítulo

3 Directrices para programar códigos Stencil 3-D en Xeon Phi

3.1 Introducción

Los patrones de cómputo Stencil comprenden una amplia familia de aplicaciones para resolver problemas científicos y de ingeniería de gran calado. Estas aplicacio- nes se basan en kernels iterativos, que operan a través de una estructura de datos N-dimensional. El diseño eficiente de este tipo de aplicaciones ha demostrado ser una tarea difícil, siendo un problema ampliamente estudiado por la industria y la academia.

Tradicionalmente, la mayoría de aplicaciones basadas en este tipo de patro- nes Stencil se ha desarrollado en lenguajes de programación de alto nivel como Matlab 63. La principal razón son los campos de aplicación de dichos patrones, presentes principalmente en dominios de simulaciones científicas y de ingeniería. Matlab proporciona un punto de entrada accesible para programadores inexper- tos (por ejemplo, Ingenieros o Físicos) para crear un modelo de aproximación computacional de su problema de interés. Esta abstracción de la programación de bajo nivel se logra a través de un gran número de librerías matemáticas complejas. Sin embargo, esta facilidad de programación por lo general compro- mete el rendimiento. Además, los lenguajes de programación inherentemente secuenciales conducen a una explotación limitada de los recursos disponibles en los procesadores y aceleradores modernos 68.

En este Capítulo evaluamos el rendimiento de la arquitectura Intel Xeon Phi trabajando en modo nativo utilizando tres aplicaciones cuyo patrón de cómputo se centra en kernels Stencil: difusión de calor (tamaño de 11 puntos), (2) difusión acústica (tamaño de 7 puntos), y (3) isótropico de onda sísmica (tamaño de 25 puntos). El rendimiento obtenido se compara con el conseguido en una arquitectura Intel Xeon multinúcleo. Además, los cálculos de estos kernels se desarrollan con números de coma flotante en doble precisión, poniendo al límite las necesidades de ancho de banda de nuestras aplicaciones. Partimos de una versión paralela simple (naïve) en C/C++ usando OpenMP (obtenida de forma inmediata de la versión secuencial que implementa en Matlab el Algoritmo 1 descrito en la Sección 2.6.3). Seguidamente mostraremos, paso a paso, el proceso de paralelización y vectorización, la escalabilidad del código, así como la influencia del tamaño y la forma de los datos de entrada.

3.1.1 Objetivos

El objetivo de este Capítulo es guiar a los desarrolladores de software sin expe- riencia en el desarrollo de códigos Stencil para la arquitectura Intel Xeon Phi. Por ello, consideraremos sólo aquellos aspectos cuya dificultad de programación sea media-baja, es decir, evitamos un desarrollo de bajo nivel que requiera co- nocimientos exhaustivos de la arquitectura como son las operaciones intrinsics o la programación en ensamblador. A pesar de ello, perseguimos desarrollar un código lo más eficiente posible, apoyándonos en la capacidad que tiene el compilador de Intel para vectorizar y paralelizar, e intentando explotar al máximo posible la reutilización del código desarrollado, tanto para la arquitectura Intel Xeón multinúcleo como para el Intel Xeon Phi, realizando una comparativa con el Intel Xeon multinúcleo en términos de rendimiento (expresado en GFLOPs). Los principales aspectos a evaluar son los siguientes:

- Facilidad de programación: El primer punto importante de la arquitectura Xeon Phi es su facilidad de programación. Según Intel, los programas desarrollados para sus arquitecturas multinúcleo pueden ser compilados y ejecutados directamente sobre la arquitectura Intel Xeon Phi.

- Escalabilidad: La arquitectura Intel Xeon Phi está diseñada para optimizar el rendimiento (throughput) de la aplicación. Esto se consigue gracias a un elevado número de núcleos, y la posibilidad de usar hasta 4 hilos por núcleo (es decir, un total de 244 hilos). Al tener nuestros códigos Stencil base un patrón limitado principalmente por sus accesos a memoria, nos planteamos conocer la escalabilidad que presentan dichos códigos al ir aumentando el número de hilos, tanto en el caso del Xeon multinúcleo como para la arquitectura Xeon Phi.

- Vectorización: El otro gran beneficio del Intel Xeon Phi es la gran anchura que presentan sus unidades vectoriales (512 bits). Este objetivo evalúa la facilidad de conseguir códigos vectorizados manejando adecuadamente la capacidad que tiene el compilador de Intel para auto-vectorizar. Para ayudar al compilador, exploraremos el uso de algunas directivas SIMD específicas de la arquitectura, así como el uso de técnicas sencillas para el manejo de los datos (alineamiento y padding). En cualquier caso, como hemos comentado, queremos evitar el uso de construcciones de muy bajo nivel como instrucciones en ensamblador o los intrinsics de Intel.

- Comparativa con Xeon multinúcleo: De cara a evaluar adecuadamente esta nueva arquitectura, queremos comparar los resultados que obtenemos con aquellos que se obtienen ejecutando los mismos códigos en el procesador anfitrión. Esta comprativa la vamos a realizar en términos del tiempo de ejecución o de los GFLOPs obtenidos indistintamente.

- Influencia de los datos de entrada: Finalmente, queremos evaluar cómo los datos de entrada (su tamaño y su forma) influyen en el rendimiento de estos kernels. Analizamos tres tamaños para las matrices de datos (pequeña, mediana y grande), así como cuatro formas diferentes para la disposición (layout) de los datos (una cuboide perfecta y tres cubos rectangulares).

3.1.2 Trabajo relacionado

A continuación se ofrece una breve reseña de los principales trabajos relacionados respecto a la codificación de códigos Stencil para arquitecturas Intel Xeon e Intel Xeon Phi.

El tema de modelado de cálculos Stencil ha sido muy estudiado en los últimos años. Los sistemas multinúcleo 51 proporcionan buenas oportunidades de paralelización de aplicaciones basadas en patrones Stencil. Kamil et al. 32 presentaron modelos de coste para capturar el rendimiento de los Stencil de 7 puntos teniendo en cuenta tres tipos de accesos a memoria en una memoria de jerarquía plana. Los autores en 13 presentan una metodología completa para evaluar y predecir el rendimiento del código Stencil en arquitecturas HPC complejas. En 53 se desarrolla un conjunto de fórmulas a través de análisis de regresión para modelar el rendimiento general de cálculos en Stencil de 7 y 27 puntos para Jacobi y Gauss-Seidel. La técnica de tiempo de sesgo (time-skewing) también se ha modelado por Strzodka et al. 65, que propusieron un modelo de rendimiento de su algoritmo de tiempo de sesgo de caché exacta (CATS), donde el sistema y el ancho de banda de memoria caché fueron estimados usando un análisis de regresión lineal.

Los autores en 49 introducen una metodología que dirige los esfuerzos del programador hacia las regiones de código con más probabilidades de benefi- ciarse de la migración a Intel Xeon Phi, así como proporcionar estimaciones de aceleración (speedup). Otros investigadores 15 han estudiado la portabilidad y optimización de la prueba del problema básico de simulación N-body para el coprocesador Intel Xeon Phi 71, que es también la base de una serie de aplicaciones en astrofísica computacional y biofísica.

Por último, hay varios libros recientes que muestran las altas capacidades de computación del coprocesador Intel Xeon Phi para muchas aplicaciones HPC, como [40, 28, 57, 56, 52, 70]. También es interesante la contribución de 16. Específicamente, sobre códigos Stencil tenemos las referencias [26, 41], y 4 trata sobre el código de difusión sísmica.

3.2 Migración de patrones Stencil a Intel Xeon Phi

En esta Sección se muestra el proceso de ejecutar en la arquitectura Intel Xeon Phi los tres códigos Stencil que usamos a lo largo del documento, describiendo en detalle el proceso de paralelización y vectorización, así como el manejo de la escalabilidad del código.

3.2.1 Ejecución del código base

Nuestro punto de partida es una implementación directa basada en la tupla C/C++ y OpenMP que implementa el Algoritmo 1 (descrito en la Sección 2.6.3) para cada uno de nuestros tres kernel Stencil. Esta paralelización base es una implementación sencilla (naïve) del código secuencial. Dicho código base secuen- cial ha sido desarrollado usando el lenguaje C y paralelizado con extensiones OpenMP. OpenMP está basado en declaraciones #pragma que son capturadas por el compilador, validadas y traducidas a llamadas a funciones apropiadas de la librería OpenMP en tiempo de ejecución del sistema. Las extensiones proporcio- nadas por OpenMP permiten crear y manejar los diferentes hilos de ejecución en un programa paralelo usando el modelo de memoria compartida para todos los hilos. Adicionalmente, para que el compilador reconozca las extensiones OpenMP hay que incluir el archivo de cabecera <omp.h> en el código fuente de dichos programas, así como añadir la opción de compilación -openmp para enlazar con las librerías OpenMP.

Para ejecutar estos códigos en la arquitectura Intel Xeon Phi en modo nativo, tan sólo fue necesario compilar con el compilador icc de Intel agregando la directiva de compilación -mmic. Los códigos objeto resultantes pudieron ser ejecutados sin problemas en la arquitectura Xeon Phi, obteniendo una mejora de rendimiento con respecto a la ejecución en la arquitectura Xeon multinúcleo gracias al uso del mayor número de núcleos que tiene la arquitectura Xeon Phi.

Sin embargo, los rendimientos obtenidos fueron muy pobres debido a que estos códigos escalan muy mal por estar limitados por el acceso a memoria. Además, los códigos no estaban optimizados para su ejecución escalar, ni estaban bien paralelizados, ni se estaba aplicando la vectorización¹ en ambas arquitecturas. Por tanto, nuestro siguiente paso fue la mejora de estas características.

3.2.2 Optimizaciones escalares

Antes de aplicar algunas estrategias de vectorización, consideramos necesario revisar nuestro código para aplicar algunas optimizaciones escalares que pudieran mejorar el rendimiento de cada tarea paralela. En este sentido, se aplican las siguientes optimizaciones:

- Optimización de expresiones aritméticas: En esta optimización se trabajó en la reorganización de las operaciones del cálculo del Stencil para aprove- char mejor la localidad temporal y espacial, de forma que afecten lo mínimo a los fallos de caché. Asimismo, se modificó también el orden de acceso a los datos en el cálculo del Stencil, accediendo en el siguiente orden (se muestra para el caso del código de la difusión de calor, en los otros dos códigos se aplica de forma similar): Noroeste (nw), norte(n), noreste(ne), oeste(w), centro(c), este(e), suroeste(sw), sur(s), sureste(se), diagonal norte (nl) y diagonal sur (ns). Por último, y debido a que teníamos expresiones que se reutilizaban varias veces sin modificarse dentro del bucle más in- terno de cálculo (bucle en la dimensión x), estas expresiones las movimos al siguiente bucle exterior (el bucle en la dimensión y).

- Reducción de precisión de las operaciones aritméticas. En los códigos Sten- cil hemos procurado sustituir operaciones más costosas por otras que sean de menor coste. Por ejemplo, sustitución de divisiones por la multiplicación de su inversa, etc. Asimismo, hemos compilado con las opciones -fimf-precision=low -fast-transcendentals -no-prec-sqrt -no-prec-div, para tener una precisión baja en nuestras operaciones en coma flotante, y con la opción -fp-model fast=2 para minimizar el impac- to de tratar con números denormalizados del estándar IEEE 754, pues no influían en el valor final de la matriz de salida de nuestros códigos.

- Eliminación del uso del operador ternario. El operador ternario : ?, expre- sión abreviada para la instrucción if ... else que permite escribir un código de forma muy sencilla y compacta, puede causar cálculos redundan- tes en el código y bajar el rendimiento de los mismos. Hemos eliminado dicho operador de nuestro código.

- Uso del calificador const. En todas aquellas variables locales o argumentos de función que no cambian de valor en el código, hemos procedido a declararlas con el calificador const. Este calificador permite al compilador generar un código más eficiente.

- Uso de registros. Hemos alojado en registros aquellas variables que se usan con mucha frecuencia en la ejecución de un bucle. Para ello, hemos utilizado el calificador register en la declaración de dichas variables. Esta optimización ayuda a reducir las necesidades de memoria del código y a acelerar su ejecución.

- Uso de las matrices de datos. Las matrices de datos para nuestro Stencil las hemos declarado como una matriz lineal de tamaño ancho*alto*profundidad (width, height, depth en nuestro código), evitando usar matrices tridimensio- nales que empeoran el rendimiento del mismo. Por tanto, hemos asignado todas las filas de las matrices de nuestros códigos Stencil consecutivamente en memoria. Asimismo, accedemos a cada uno de los elementos de la matriz utilizando un desplazamiento a partir de la dirección base de la matriz, y evitando acceder a dichos elementos a base de utilizar punteros relativos a cada fila o plano de datos.

3.2.3 Manejo de los hilos en Xeon Phi

En esta Sección queremos mostrar como mejorar el rendimiento de los códigos base desarrollados por medio de un mejor manejo de los hilos en Intel Xeon Phi. Como ya hemos comentado, los códigos base de que partimos han sido paralelizados por medio de añadir un pragma en el código antes del triple bucle anidado que recorre los datos de entrada y realiza la actualización de la matriz de salida (#pragma omp parallel for).

De cara a la mejor paralelización de nuestros códigos ha sido muy adecuado añadir el modificador collapse(2) a dicho pragma. Dicho modificador combina (colapsa) los dos bucles más externos de nuestro kernel en un mismo bucle y entonces se paraleliza, con lo que aumenta el número de unidades de trabajo que se pueden repartir por cada hilo produciéndose un mejor balanceo de la carga entre los hilos. Por tanto, se aumenta el rendimiento de la paralelización, especialmente en aquellos escenarios en que el número de hilos sea elevado, como es el caso de la arquitectura Xeon Phi.

Finalmente, estamos también interesados en conocer y evaluar la escalabi- lidad de nuestros códigos Stencil en función del número de hilos utilizados al ejecutar dichos códigos. Para ello, en los experimentos de escalabilidad que realicemos, vamos a variar el número de hilos de cada prueba a 1, 2, 3 ó 4 hilos por núcleo de ejecución. Esto lo podemos hacer de dos formas, o bien usando la variable de entorno OMP_NUM_THREADS, o por medio del uso de la función omp_set_num_threads() insertada en el código fuente. Como nuestra arquitectu- ra Xeon Phi tiene 61 cores, vamos a experimentar escalando nuestro código a 61, 122, 183 y 244 hilos.

3.2.4 Vectorización del código

Hay varias estrategias para vectorizar un código para las arquitecturas Intel. En nuestro caso queremos evitar el uso de construcciones de muy bajo nivel como instrucciones en ensamblador o los intrinsics de Intel, por su elevada complejidad en el desarrollo y por su falta de portabilidad entre diferentes versiones de la arquitectura. En cambio, proponemos el uso de la capacidad de autovectorización que tiene el compilador de Intel. Esta opción se basa en añadir algunas de las directivas SIMD específicas que se ofrecen, así como el uso de técnicas sencillas para el manejo de los datos como el alineamiento o el padding.

Para observar cómo ha ido el proceso de vectorización por parte del compi- lador, éste proporciona una opción de informe de vectorización con dos clases importantes de información. En primer lugar, el informe de vectorización nos informa acerca de qué bucles de nuestro código han sido vectorizados (es decir, el compilador ha generado para dichos bucles vectorizados un paquete de instruc- ciones SIMD). En segundo lugar, el informe nos muestra los detalles acerca de por qué el compilador no ha podido vectorizar un bucle. Esta información es muy útil ya que permite identificar las barreras que el compilador encuentra para la vectorización. La directiva -vec-report=<n> le indica al compilador que genere los informes de vectorización con diferentes niveles de información, donde el argumento <n> se usa para especificar la información presentada (<n> puede variar de 0 a 6).

Por tanto, lo primero que hicimos fue compilar nuestro código con -openmp -O3 -vec-report=6. El nivel de optimización -O3 que utilizamos es más agresivo que el -O2, y además de incluir las características de -O2 (la vectorización auto- mática viene incluida en ambos niveles), proporciona funciones de optimización para procesos en línea, eliminación de código muerto y desenrollado de bucles. La Figura 3.1 muestra un extracto del informe de vectorización en el que se muestra que el compilador asume dependencias, motivo por el cual, nuestro código no es vectorizado. Salvo que digamos expresamente lo contrario, todos los ejemplos de este Capítulo son mostrados para el caso del código de difusión de calor de 11 puntos.

A continuación se muestra como hemos conseguido vectorizar el código y obtener un buen rendimiento en el mismo:

1. Evitar dependencias: Para resolver las dependencias de datos tenemos que ver el código que presenta problemas, pues en principio y tal y como se mostró en el Algoritmo 1 base (descrito en la Sección 2.6.3) del código, no debía haber ninguna dependencia de datos. Si nos fijamos en la cabecera de la función que calcula el Stencil (mostrada en la Figura 3.2), nos damos cuenta de que dicha función tiene como parámetros de entrada los punteros fin y fout que contiene las matrices de entrada y salida del algoritmo. Al estar la declaración de dichas matrices en el programa principal, el compilador no conoce las direcciones exactas de dichas variables, por lo que al ir a compilar el código asume que pudiera haber alguna superposición (overlapping) en la zona de memoria de dichas variables, y por seguridad supone la existencia de dichas dependencias entre los diversos elementos de dichas variables.

Para el kernel acústico y de calor, la forma en que le podemos ayudar al compilador a vectorizar el bucle es utilizando la palabra clave restrict,

Abbildung in dieser Leseprobe nicht enthalten

Figura 3.1: Informe de vectorización de nuestro código Stencil compilado con -O3

Abbildung in dieser Leseprobe nicht enthalten

Figura 3.2: Código de la declaración de la función Stencil de 11 puntos.

con la que informamos al compilador de que la referencia a memoria pasada por un puntero no es un alias, es decir, que no se accede a ella de ninguna otra forma. De esta forma, el compilador ya puede comprobar si hay o no dependencias y, si no las hay como es nuestro caso, generar un código vectorizado. Para ello, modificamos ligeramente la declaración de las variables de la entrada de la función utilizando (REAL* restrict fin, REAL* restrict fout) y, además, en el momento de compilar, tenemos que agregar la directiva -restrict para que el compilador de Intel conozca esta palabra como palabra reservada. Como se puede ver en el informe de la Figura 3.3, nuestro código fue vectorizado, pero los accesos a los datos están desalineados. Esto último se debe a que por defecto el compilador no puede saber ni asumir que los datos están alineados dentro de un bucle.

Abbildung in dieser Leseprobe nicht enthalten

Figura 3.3: Informe de vectorización de nuestro código Stencil compilado con -O3 -restrict.

En el caso del kernel sísmico, debido a la complejidad de las operaciones que realiza, la anterior directiva no le ayuda a resolver las dependencias entre datos. Otra forma de forzar al compilador a vectorizar y que no tenga en cuenta posibles dependencias es añadiendo la directiva #pragma ivdep antes del bucle interno que va a ser vectorizado, con lo que se le notifica al compilador que éste puede asumir que las direcciones de memoria de las distintas matrices apuntan a posiciones disjuntas. En el caso del kernel sísmico, ésta ha sido la aproximación realizada, y con esto se han resuelto las dependencias. Comentar también que esta directiva es muy útil, ya que además es muy genérica y se puede colocar antes de cualquier bucle que sepamos que se puede vectorizar y que el compilador no lo hace debido a que pudiera existir algún tipo de dependencia entre los datos. Sin estos consejos, el compilador podría no ser capaz de identificar correctamente el bucle interno como vectorizable, y podría fracasar en la vectorización del código como habíamos mostrado anteriormente.

2. Mejorar la vectorización: Para mejorar el rendimiento de la vectorización, debemos manejar adecuadamente la asignación de memoria y su uso, y para ello es necesaria la ayuda por parte del programador en los siguientes aspectos:

- La alineación de datos: Aspecto clave para la vectorización en la ar- quitectura Intel Xeon Phi. La alineación de datos es un método para forzar al compilador a crear objetos de datos en la memoria en limites específicos de bytes, para aumentar la eficiencia del almacenamiento y cargas de los datos a y desde el procesador. Para el coprocesador Intel Xeon Phi, el movimiento de memoria es óptimo cuando los datos se sitúan en direcciones de inicio múltiplos de 64 bytes.

Para ayudar al compilador, es necesario llevar a cabo dos pasos: 1) alinear los datos y 2) utilizar pragmas/directivas y cláusulas para decirle al compilador que los accesos a memoria están alineados. Para realizar la alineación de datos correctamente, la asignación de memoria dinámica con llamadas a funciones en C/C++ (malloc() y free()) han de ser reemplazadas por implementaciones alterna- tivas que soportan alineación de datos (es decir, _mm_malloc() y _mm_free()), pasándole como parámetro a dichas llamadas el factor de alineación. En el caso de la arquitectura Xeon Phi hay que seleccionar un factor de alineación de 64 bytes. En la Figura 3.4 se muestra la declaración que utilizamos para nuestras matrices de datos alineadas a 64 bytes.

Abbildung in dieser Leseprobe nicht enthalten

Figura 3.4: Alineación de los datos con _mm_malloc.

Por otra parte, es conveniente informar al compilador de que nuestros datos están alineados en la zona de código donde vayan a ser usados.

Para ello debemos utilizar la cláusula __assume_aligned(Data, 64) para proveer al compilador con la información respecto a la alineación de Data. Sin esta información, el compilador podría no ser capaz de identificar correctamente la alineación utilizada por la estructura de datos. En el caso de que las estructuras de datos alineadas sean de tipo vector, podemos usar el pragma #pragma vector aligned colocándolo justamente antes del bucle más interno que vaya a ser vectorizado.

En nuestro caso, y aunque los datos de entrada los alineamos a 64 bytes, debido a los bordes (halo) del Stencil, el acceso a los puntos internos para el cálculo fuerza un acceso no alineado a dichos valores. Aunque probamos a usar el pragma anterior, el informe del compilador mostró que no era posible un acceso alineado a los datos.

- Padding: El padding o “relleno” es una técnica interesante que utilizamos para asegurarnos que el primer elemento de cada fila está alineado en la dirección deseada (64 bytes para Intel Xeon Phi). Para ello, a cada una de las filas de la matriz de datos en memoria se les añade un espacio adicional (que no es utilizado en realidad para los cálculos del algoritmo). Aunque utilizamos un mayor espacio en memoria, dicho espacio adicional “ficticio” nos ayuda en el manejo de los bloques de memoria al traerlos a la caché. En nuestro caso, el padding ha sido una técnica muy conveniente en nuestros códigos Stencil por dos razones principales:

a) Evita la desalineación entre filas. La estructura del conjunto de da- tos se asigna en la memoria como un todo, es decir, usando una so- la instrucción _mm_malloc(). Por tanto, dependiendo de la anchura de la dimensión X (de tamaño width), puede haber una desalinea- ción de los datos entre cada una de las filas. Para solucionarlo usa- mos una nueva anchura agregando algunos elementos (si es necesa- rio) para asegurar de que el primer elemento de cada fila está ali- neado en la dirección deseada (64 bytes en Xeon Phi). La nueva an- chura con padding la calculamos como width_PADD = ((((width ∗ sizeof(REAL)) + 63)/64) ∗ (64/sizeof(REAL))). REAL define el tipo de datos en coma flotante utilizado en el kernel (doble preci- sión, en nuestro caso).

b) Evita fallos de conflicto patológicos. Los fallos de caché por con- flicto son fallos causados por líneas de caché que mapean al mismo conjunto de caché, por lo que causan una baja tasa de aciertos degradando el rendimiento de un algoritmo. A pesar de que la aso- ciatividad de las cachés L1 y L2 en el Intel Xeon Phi es alta (8 vías), cuando utilizamos datos de tres dimensiones en los códigos Stencil, pueden aparecer fallos por conflicto bajo ciertas combinaciones del tamaño de la matriz de entrada (width x height x depth). Más concretamente cuando un kernel accede a datos con una distancia de 4KB (L1) o de 64KB (L2).

3.3 Otras técnicas sencillas para mejor el rendimiento

En la Sección anterior se muestra cómo portar el código base al Xeon Phi, optimi- zar su comportamiento secuencial por medio de optimizaciones escalares, y cómo mejorar su rendimiento gracias al uso adecuado de la capacidad de vectorizar y paralelizar que tiene esta arquitectura. A continuación, se describen otras técnicas sencillas para Xeon Phi, y que son fáciles de usar por los programadores sin experiencia y que ayudan a mejorar el rendimiento de los códigos Stencil.

1. Tamaño de paginas grandes. La arquitectura Intel Xeon Phi puede ser con- figurada para utilizar varios tamaños de página: 4 KB ó 2 MB (también conocido como páginas grandes). Esta configuración permite al TLB del Intel Xeon Phi mapear 128 MB de memoria en comparación con los 256 KB mapeados por defecto. Esta configuración podría llegar a reducir los fallos de pagina del TLB significativamente (hasta un 15 %) en ciertas aplicaciones. La asignación de paginas grandes se realiza mediante la sustitución de las llamadas malloc() con la función mmap().

2. Streaming stores. La caché del Xeon Phi tiene por defecto una política de localización en escritura (write-allocate) para los fallos en escritura (común en cachés de postescritura). Esta política implica que cuando una aplicación escribe su salida en una ubicación de memoria y se produce un fallo de escritura en caché, el bloque de datos destino es cargado de la memoria y se mueve a través de la jerarquía de memoria hasta que alcanza la L1 y entonces se realiza la escritura. Pero esta política se puede modificar, y se puede cambiar a una política de no localización en escritura (write- non-allocate) para fallo de escritura en caché, que Intel denomina Streaming stores. Esta política de no localización en escritura puede ser interesante en el caso de que los datos de salida no se requieran para cualesquiera otros cálculos, evitando entonces llenar la caché y el resto de la jerarquía de memoria con datos que no presentan ningún uso adicional. La técnica de Streaming stores almacena los datos usando un almacenamiento (buffer) no temporal, mejorando por tanto la utilización del ancho de banda de me- moria, y potencialmente mejorando el rendimiento global de la aplicación. Para aplicar esta técnica, hay que colocar la siguiente directiva #pragma vector non temporal () antes del uso de los datos de salida, poniendo entre paréntesis las variables que se desean que tengan el comportamiento de no localización en escritura en el caso de un fallo en escritura en ca- ché. En nuestros códigos, hemos colocado la directiva #pragma vector non temporal (fout) antes del bucle tridimensional más interno (dimensión en X). Alternativamente, el programador puede utilizar la directiva de compilación -opt-streaming-stores para solicitar al compilador la gene- ración de Streaming stores para aquellas variables que el compilador estime conveniente.

3. Memoria ECC y turbo. La arquitectura Intel Xeon Phi ha sido diseñada para proporcionar un entorno de trabajo de alta robustez (reliability). Por ello, utiliza el Código de Corrección de Errores (ECC) para proporcionar corrección y detección de errores en el acceso a la memoria RAM. La plataforma actual de Intel Xeon Phi incorpora la memoria gráfica GDDR5 (Graphics Double Data Rate), la cual no soporta directamente en hardware la característica ECC. Por ello, dicha característica es implementada en Xeon Phi por software. Ello implica que parte de la memoria se usa para almacenar dichos códigos de corrección de error (Intel cifra en un 12,5 % el impacto en la reducción de memoria por el ECC), así como en una ligera disminución en el rendimiento de las aplicaciones (no cuantificado por Intel). Para aquellos usuarios que lo deseen, se puede desactivar esta característica para obtener mayor espacio en memoria y algo más de rendimiento. ECC puede ser deshabilitada/habilitada utilizando el comando micsmc ecc disable/enable micn (donde n es el número de la tarjeta mic).

Intel incorpora también la tecnología Intel Turbo Boost a la arquitectura Xeon Phi para permitir forzar la capacidad de aumento de frecuencia de reloj (overclocking), en función del uso de energía actual y la temperatura. Esta capacidad puede ofrecer una mejora del rendimiento de la aplicación de hasta un 10%. La capacidad turbo puede ser habilitada/deshabilitada utilizando el comando micsmc turbo enable/disable mic0 (en este caso, se aplica sólo a la tarjeta mic0).

3.4 Influencia de los datos de entrada

La última Sección de este Capítulo 3 del documento se dedica a evaluar la influencia que tienen los datos de entrada en el rendimiento que obtenemos al ejecutar los códigos Stencil. Nos parece que dicha influencia la podemos evaluar teniendo en cuenta dos características distintas: el tamaño que presentan los datos de entrada así como la forma de los mismos.

El tamaño de entrada podría influir en los resultados obtenidos al estar traba- jando con algoritmos limitados por memoria. Por ello, vemos adecuado evaluar su influencia en relación al rendimiento obtenido en el código. En este Capítulo hemos considerado tres tamaños: uno pequeño (matriz de 400x400x400), uno mediano (800x400x600) y uno grande (800x900x900). Es importante mencionar que estos tamaños se refieren sólo al tamaño de una de las matrices de entrada del kernel (denominada en nuestros códigos como fin).

Con la finalidad de conocer si la forma de la matriz de entrada afecta el rendimiento de nuestros códigos Stencil, evaluamos formas diferentes para el tamaño de nuestra matriz. Dependiendo de la profundidad de la matriz (dimen- sión Z) y el paso de acceso (stride) a memoria, el compilador podría ser capaz de vectorizar mejor el bucle interno de nuestros algoritmos, lo que permitiría un mejor aprovechamiento del ancho de banda de acceso a memoria en el Intel Xeon Phi. Hemos evaluado cuatro formas diferentes para el tamaño de la matriz de entrada (con un número de elementos de doble precisión igual a ancho x alto x profundidad): una forma cuboide perfecta y tres formas cuboides rectangulares, en donde cada una de estas formas rectangulares tiene una dimensión mayor que las otras dos dimensiones (que son iguales en tamaño). Específicamente, las formas evaluadas para nuestro tamaño de matriz son: forma cuboide perfecta (400x400x400), tres formas cuboides rectangulares (800x400x600, 600x800x400 y 400x600x800).

Además, se ha tenido en cuenta que las aplicaciones basadas en cálculos Stencil necesitan discretizar el conjunto de datos en las matrices como parte de la representación numérica del problema. Con la evaluación de esta característica queremos conocer el impacto de la forma de la matriz en el rendimiento alcanza- do, de tal manera que el programador pueda seleccionar más adecuadamente la forma de la matriz para su simulación.

3.5 Evaluación

En esta Sección se evalúa todas las propuestas descritas anteriormente. Empeza- remos describiendo nuestro entorno de evaluación.

3.5.1 Entorno de evaluación

Antes de comenzar la evaluación, está Sección describe la plataforma hardware heterogénea utilizada en nuestros experimentos. Igualmente, detallamos los parámetros específicos que hemos considerado para cada uno de los tres códigos Stencil 3-D que hemos evaluado, el concepto de intensidad aritmética aplicado a nuestros entornos hardware y software, y la metodología en la realización de los experimentos que hemos seguido a lo largo de todo el documento.

3.5.1.1 Plataforma hardware

La plataforma de evaluación donde se llevaron a cabo los experimentos, denomi- nada pacioli.inf.um.es, está equipada con dos Ivy Bridge-EP Intel Xeon E5-2650 CPUs (2x8 núcleos en total) y tres coprocesadores de la gama Intel Xeon Phi 7120P (61 cores en cada coprocesador). En la Figura 3.5 se muestra una fotografía de una plataforma similar a la utilizada en nuestros experimentos.

Cada chip Intel Xeon CPU E5-2650 v2 funciona a una frecuencia de reloj de 2.6 GHz, tiene 8 núcleos físicos y 16 hilos, con una caché L1 privada para instrucciones y otra para datos de 32 KB cada una, una caché L2 compartida de 20 MB, y soporta hasta 32 GBGB de memoria principal DDR3-1600. Proporciona un ancho de banda teórico de acceso a memoria de 59.7 GB/s por procesador (es decir, por chip). Esta máquina es utilizada como nuestra CPU evaluada (Xeon en nuestras Figuras).

El Xeon Phi 7120P tiene 61 cores trabajando a 1.238 GHz, y cada núcleo tiene una unidad vectorial de 512 bits, por lo que puede procesar 8 elementos de datos de doble precisión a la vez, con un máximo de 2 operaciones (multiplicación- suma) por ciclo en cada vía (es decir, un elemento de vector). Por lo tanto, el rendimiento pico teórico es de 1208 GFLOPs (doble precisión). Por otra parte, la característica más importante de los coprocesadores Xeon Phi es el ancho de banda de memoria. El Xeon Phi evaluado tiene 16 canales de memoria (ancho de 32 bits). Con hasta 6 GT/s de velocidad de transferencia, el 7120P ofrece un ancho de banda teórico de 384 GB/s. En la Tabla 3.1 se mencionan con más detalle las características de este sistema.

Abbildung in dieser Leseprobe nicht enthalten

Figura 3.5: Plataforma hardware.

Tabla 3.1: Especificaciones del coprocesador Intel Xeon Phi 7120P.

Abbildung in dieser Leseprobe nicht enthalten

Además, se ha utilizado el sistema operativo Linux CentOS 6.5 con kernel 2.6.32, el compilador Intel icc (versión 14.0.2), y la versión 3.4.3 de Intel MPSS.

3.5.1.2 Aplicaciones software

En este documento hemos evaluado tres códigos Stencil de diferentes campos científicos. Estas soluciones cubren una área amplia de investigación y tienen características computacionales distintas.

El código Stencil mas común utilizado en problemas de Ingeniería es repre- sentado por nuestro código Stencil 3-D de difusión acústica, que utiliza un Stencil de 7 puntos vecinos espaciales y segundo orden en tiempo. Utiliza tres matrices diferentes del mismo tamaño para el cálculo del kernel. Nuestro siguiente kernel es el Stencil 3-D isotrópico de onda sísmica de 25 puntos vecinos espaciales y también de segundo orden en el tiempo. Por ultimo, hemos evaluado también un código kernel más simple, el Stencil 3-D de difusión de calor de 11 puntos vecinos espaciales y de primer orden en el tiempo, el cual sólo utiliza dos matrices para los cálculos del Stencil. En el caso del Stencil sísmico, al utilizar una matriz espa- cial diferente para el almacenamiento de características físicas necesita utilizar cuatro matrices de datos en total. Finalmente, comentar que para los casos de los códigos acústico y sísmico, al ser de segundo orden temporal, aplicamos la optimización de encontrar el valor del punto en el instante t-1 en la matriz de salida fout, ya que al haber hecho el cambio de matrices al acabar cada iteración, en la matriz de entrada fin se tienen los datos del instante t, mientras que en la fout se tienen los del instante t-1, y se almacenan ahí los del instante t+1. Aplicando esta optimización, nos evitamos el uso de una de las matrices de entrada y se reducen las necesidades de ancho de banda con memoria.

En relación a la forma de la matriz, hemos elegido para los experimentos formas cuboides rectangulares de ancho x alto x profundo (width x height x depth). Esta es una característica importante ya que dicha forma dicta los patrones de acceso de memoria. Por último, comentar que cuando se habla del tamaño de la matriz de entrada de los kernels se refiere sólo al tamaño de una de las matrices de entrada en cada uno de los kernels utilizados.

Para las simulaciones de los tres kernels Stencil hemos elegido un valor de 1.000 iteraciones para la variable IterCount que nos garantice suficientemente la convergencia del problema. Tal y como viene recomendado en 4, hemos realizado una ejecución previa de 2 iteraciones como calentamiento que no hemos incluido en los datos de rendimiento recogidos de la aplicación.

Como parámetros específicos para cada kernel, hemos empleado los siguien- tes: En el kernel acústico, tenemos un valor de 0,01 para el parámetro ant, un valor de 3,5 para el parámetro cons, un valor de 0,13 para el parámetro ctrl, un valor de 0,135 para el parámetro next, y un valor de 0,09 para el parámetro length. En el caso del kernel sísmico, hemos considerado los siguientes valores: un valor de 0,002 para el parámetro DT, un valor de 50,0 para el parámetro DXYZ, y un valor de 4 para el parámetro HALF_LENGTH. Los valores de los coeficientes que se aplican a cada punto para el caso del kernel de 25 puntos son {-9.164532312924e-1, +1.777777777777e-1, -3.111111111111e-1, +7.542087542087e-2, -1.767676767676e-2}. Finalmente, en el caso del kernel de calor hemos usado un valor de 0,1 para el parámetro ctrl, un valor de 0,045 para los parámetros diag, un valor de 0,135 para los parámetros de next, y un valor de 0,09 para los parámetros length. La temperatura inicial de los puntos de la matriz de entrada ha sido fijada a los siguientes valores: 10 º C para el interior del cuerpo, 150 º C para los bordes izquierdo y derecho, 70 º C para los bordes superior e inferior, y también 70 º C para los bordes anterior y posterior.

Todas nuestras evaluaciones han sido realizadas utilizando un tipo de datos de coma flotante de doble precisión (64 bits), para de esta forma presionar más sobre el ancho de banda de acceso a memoria.

3.5.1.3 Intensidad aritmética (AI)

Para cualquier plataforma de cálculo, las especificaciones hardware definen una capacidad máxima para el cálculo de FLOPs y de transferencia de datos hacia y desde la memoria (ancho de banda de memoria). Dichos valores están relacionados por el concepto de intensidad aritmética (AI por sus siglas en inglés), que nos muestra el número de operaciones de punto flotante que dicha plataforma (arquitectura) puede realizar por operando recibido de memoria55. La expresión para AI es la siguiente:

Abbildung in dieser Leseprobe nicht enthalten

Es decir, debido a que el ancho de banda a memoria para acceder a los datos de una aplicación está limitado, para alcanzar el máximo rendimiento se deben realizar un determinado número de operaciones en coma flotante por cada dato traído de memoria. En caso contrario, el algoritmo estará limitado por el ancho de banda de memoria y será imposible obtener el rendimiento máximo de la arquitectura de que se trate. El concepto de AI fue propuesto dentro del modelo Roofline 74 y permite determinar, para una arquitectura dada, el número de operaciones que hace que un algoritmo sea limitado por memoria o limitado por cálculo. Esto además nos permite estimar, de forma teórica, cual va a ser el rendimiento real alcanzable por un algoritmo dado.

En el caso de Xeon Phi, como acabamos de decir en la Sección 3.5.1.1, su rendimiento pico teórico es de 1208 GFLOPs (doble precisión), con un ancho de banda teórico de 384 GB/s. Por tanto obtenemos un valor de AI de AI =[1208] ≈ 25, es decir, por cada operando (en doble precisión) leído necesitamos ejecutar 25 operaciones en coma flotante para poder obtener el rendimiento máximo del Intel Xeon Phi. Pero si tenemos en cuenta que el ancho de banda real de memoria alcanzable está limitado al 50-60% del ancho de banda de memoria máxima (es decir, en torno a los 230 GB/s), para explotar el rendimiento total del Xeon Phi obtenemos un valor de AI de 40 FLOPs/Operando buscado (en DP). Esto significa que para Xeon Phi podemos caracterizar un kernel dado como limitado por cálculo si su AI es mayor que 40 FLOPs/Operando (DP), o limitado por memoria en caso contrario.

Vamos ahora a mostrar la intensidad aritmética que presentan los kernels evaluados en este documento. Esté cálculo está influido por el número de puntos que tiene cada uno de los kernels así como el número de operaciones que se realizan para el cálculo del nuevo punto en cada iteración. El Stencil de difusión acústica tiene una intensidad aritmética muy baja (alrededor de 6.5). El Stencil sísmico tiene mayor intensidad aritmética que el anterior, en torno a 9,7. Finalmente, la intensidad aritmética del kernel de calor esta cercana a 10.5. Como podemos ver, todos ellos son kernels limitados por memoria y, por tanto, sólo podrían alcanzar una fracción del rendimiento máximo de Xeon Phi.

3.5.1.4 Metodología experimental

Para las diversas evaluaciones que hemos realizado a lo largo de todo el do- cumento, hemos tenido en cuenta que la plataforma donde hemos realizado los experimentos (pacioli.inf.um.es) esté en estado de reposo, sin ningún otro usuario conectado, y desactivados la mayoría de los daemons que pudieran afectar significativamente a los tiempos de ejecución de los experimentos realizados.

Aún así, y para tener una mayor precisión en la medición del tiempo de ejecución de cada una de las versiones de nuestros códigos, hemos ejecutado cada uno de ellos 10 veces. De los tiempos obtenidos en estas ejecuciones, hemos eliminado el mayor y el menor, y hemos computado la media de los 8 restantes. Dicha media es la que mostramos en nuestras tablas y gráficas como el tiempo obtenido. No ofrecemos la desviación típica pues ésta es muy pequeña.

Como valores por defecto, y salvo que digamos lo contrario, ejecutamos nuestros códigos con 32 hilos en el caso de Xeon y 244 hilos para Xeon Phi. La afinidad de los hilos por defecto es compact para Xeon y balanced para Xeon Phi. La política de planificación tomada por defecto ha sido static.

Tabla 3.2: Comparativa entre Intel Xeon y Xeon Phi.

Abbildung in dieser Leseprobe nicht enthalten

Por último, y ya que parte de nuestra evaluación experimental es la compara- tiva de las prestaciones obtenidas por nuestros códigos entre Xeon multinúcleo y Xeon Phi, vamos a concluir esta Sección con la Tabla 3.2 que nos muestra una comparativa de las características principales de cada plataforma.

3.5.2 Evaluación del código base

En esta Sección vamos a empezar a evaluar el rendimiento que hemos obtenido al ejecutar los tres códigos Stencil en la arquitectura Intel Xeon Phi. Como comen- tamos en la Sección 3.2, los códigos utilizados son los mismos que ya teníamos codificados para Xeon, y tan sólo han sido compilados con la opción -mmic. En la Tabla 3.3 podemos ver los tiempos obtenidos en función del número de hilos para los diversos códigos, tanto en el procesador anfitrión como en Xeon Phi.

Debemos resaltar de esta primera evaluación:

- Como ya se había anunciado, los códigos Stencil no escalan debido a su alta necesidad de datos, por lo que están limitados por la memoria.
- En el caso de Xeon, la mayor reducción de tiempos se observa con 32 hilos (aunque a partir de 8 la ganancia obtenida es mínima), consiguiendo una aceleración entre 2 y 3 dependiendo del código elegido.

Tabla 3.3: Escalabilidad de nuestro código Stencil base en Xeon y Xeon Phi (tiempo en segs.).

Abbildung in dieser Leseprobe nicht enthalten

- En el caso de Xeon Phi, la mayor reducción de tiempos se observa con 122 hilos (183 y 244 empeoran los resultados), consiguiendo una aceleración superior a 50 sobre el tiempo de ejecución con un solo hilo.
- La aceleración del mejor tiempo obtenido para Xeon Phi contra el mejor tiempo de Xeon va de un 10 % para el caso del código de difusión sísmica, a 2.5 veces para los códigos de difusión acústica y de calor.
- Como ya se había dicho en la Sección 2.2.2, la arquitectura Xeon Phi no está pensada para usar sólo un hilo, pues cuando una aplicación se ejecuta así, el planificador cambia a un hilo nulo especial antes de volver al hilo de la aplicación. Esto hace que se obtengan tiempos de ejecución malos, requiriendo, al menos, de 2 hilos por núcleo de ejecución.

Esperamos que al mejorar los códigos mejore algo su escalabilidad en Xeon Phi y podamos sacar provecho del mayor número de hilos que tiene la arquitectura. Pasamos por tanto a aplicar las optimizaciones escalares que habíamos descrito en la Sección 3.2.2. Compilamos con la opción -O3 de máxima optimización de código, y obtuvimos una nueva versión del código que ofrece ligeras mejoras de rendimiento (entre el 1 % y el 5 %, dependiendo del kernel que se trate).

Pasamos pues a evaluar la vectorización de nuestros códigos donde esperamos obtener mejoras significativas. En la Tabla 3.4 podemos ver los tiempos obtenidos por las tres versiones desarrolladas hasta ahora para los diversos códigos ejecuta- das tanto en el procesador anfitrión como en Xeon Phi: base (códigos originales), escalar (código base más las optimizaciones escalares descritas en la Sección 3.2.2, y la versión vectorizada más optimizada, como se describimos en la Sección 3.2.4. Recordar que la versión paralelizada también incluye el mejor manejo de los hilos con la directiva collapse(2) añadida al #pragma parallel for, por lo que hemos lanzado las ejecuciones con 32 hilos en Xeon multinúcleo, y con 244 hilos en Xeon Phi.

Tabla 3.4: Tiempos de ejecución de las tres primeras versiones de nuestros códigos Stencil (seg.).

Abbildung in dieser Leseprobe nicht enthalten

A partir de estos resultados, podemos indicar lo siguiente:

- El proceso de vectorización en Xeon Phi es muy efectivo, obteniendo una reducción del tiempo de ejecución en torno a 4x.

- Sorprendentemente, la versión vectorizada en Xeon no mejora sobre las otras versiones ejecutadas en el procesador anfitrión. Hemos investigado el porqué de este comportamiento, llegando a la conclusión de que el sistema de memoria del procesador anfitrión no tiene capacidad para proporcionar los datos que ayuden a mejorar su velocidad de ejecución. Para comprobarlo, y usando sólo el kernel de difusión de calor, hemos utilizado la herramienta PAPI descrita en la Sección 2.7.3 para comprobar cómo el número total de instrucciones ejecutadas en el código base ha pasado de 2.566.555.437.707 a 1.452.678.072.994 (un 1.76x menos) en el código vectorizado, y el núme- ro total de accesos a caché ha pasado de 1.335.534.090.796 en el base a 676.933.483.773 (un 1.97x menos) en el vectorizado, pero se ha mantenido el número total de fallos de caché (78.432.746.558 en el base y 76.853.201.057 en el vectorizado). Por lo que vemos que el problema viene en que la tasa de fallos de caché (L1) es de 5,9% en el caso base, y sube al 11,4% para la versión vectorizada. Para Xeon Phi el proceso de vectorización funciona mucho mejor, ya que tiene un mayor número de cachés L1D, el sistema de memoria tiene un ancho de banda mayor, y al usar 4 hilos se oculta mejor la latencia de memoria y se mejora el flujo de datos a las unidades vectoriales.

En la Figura 3.6 podemos ver los mismos resultados para los tres kernels y para las tres versiones de cada código, mostrando en esta ocasión el número de GFLOPs alcanzado.

Abbildung in dieser Leseprobe nicht enthalten

Figura 3.6: Resultados de rendimiento (GFLOPs) de Xeon y Xeon Phi.

Como podemos observar, el rendimiento en el Intel Xeon Phi en el caso del código vectorizado es mejor para todos los kernels y todas las versiones del código. La Figura 3.7 muestra los mismos valores anteriores, pero en esta ocasión agrupando para cada kernel las versiones Xeon y Xeon Phi, y dando la aceleración obtenida por cada una de ellas en relación a la versión base de cada código ejecutada en Xeon multinúcleo. Podemos observar como casi sin esfuerzo hemos obtenido unas aceleraciones de un factor cercano a 7 para todos los códigos entre la versión vectorizada para Xeon Phi y la versión base de Xeon. Si comparamos con la versión base ejecutada en el propio Xeon Phi, dicha mejora es alrededor de 4.

3.5.3 Evaluación de otras optimizaciones y la escalabilidad

A continuación vamos a evaluar las otras optimizaciones que habíamos propuesto en la Sección 3.3 para la arquitectura Xeon Phi. En la Figura 3.8 podemos ver los resultados de rendimiento y aceleración (cada uno de ellos medido sobre su base) para los tres kernels y para las tres versiones de cada código: la etiquetada como Sstore añade al código vectorizado la optimización de la no localización en escritura por fallo de escritura (Streaming store), la denominada Huge que incorpora, sobre el código vectorizado, el manejo de páginas grandes de memoria y, por último, la etiquetada como Turbo+NoECC, que evalúa el impacto de tener activada la tecnología Turbo de Intel junto con la desactivación de ECC de la memoria. Para tener una visión de conjunto, incluimos también las versiones anteriores.

Abbildung in dieser Leseprobe nicht enthalten

Figura 3.7: Comparación de rendimiento (GFLOPs) y aceleración entre Xeon y Xeon Phi.

Como podemos ver, el manejo de páginas grandes no nos ofrece ninguna ganancia en rendimiento, pues el acceso a los datos en estos kernels 3D no consigue aprovechar esta característica, ya que en la actualización de cada punto tenemos que acceder a los vecinos en tres dimensiones, los cuales están separados en la memoria por una distancia grande que impide aprovechar el manejo de las páginas grandes. Es decir, para este código el usar las páginas grandes (opción Huge) no mejora la tasa de fallos del TLB.

En cambio, la opción Sstore sí que es muy beneficiosa para el caso del kernel de la difusión de calor, consiguiendo hasta casi un 15% de ganancia adicio- nal (recuérdese que tan sólo hay que colocar la directiva #pragma vector non

Figura 3.8: Evaluación de otras optimizaciones.

Abbildung in dieser Leseprobe nicht enthalten

temporal (fout) antes del bucle tridimensional más interno en la dimensión en X). En los otros dos kernels de difusión acústica y sísmica, no se obtiene ninguna mejora significativa, pues como ya comentamos en la Sección 2.6.3, estos dos kernel son de segundo orden en el tiempo. Para evitar tener una matriz más, hemos realizado la optimización descrita en la Sección 3.5.1.2, pero esto implica que en la ecuación de actualización de un punto, su salida depende del valor de su entrada y ello inhabilita la optimización de no localización en escritura (write-non-allocate) para fallo de escritura en caché (Sstore).

Por último, sobre la versión con Sstore hemos evaluado el uso de (overclocking) y el no manejo de ECC por software, añadiendo otro beneficio adicional entre 10-20% al rendimiento. Este resultado está en línea con lo que anuncia Intel acerca de su tecnología Turbo así como del beneficio de no usar la capacidad de corrección de errores en la memoria (aunque la mayoría de la ventaja de rendimiento proviene de la desactivación de ECC). El problema del uso de esta optimización es el elevado aumento de temperatura que sufren los núcleos de la tarjeta a la que se aplica dicha optimización, como se puede apreciar en la Figura 3.9. En dicha Figura mostramos la ejecución en dos tarjetas mic de dos versiones del código, la versión Sstore (corriendo en la mic1) y la versión Turbo+NoECC (corriendo en la mic2). Como podíamos esperar, ambas versiones muestran una

Abbildung in dieser Leseprobe nicht enthalten

Figura 3.9: Evaluación de la característica Turbo+NoECC.

utilización al 100 % de los núcleos de Xeon Phi de cada tarjeta, pero en el caso de la versión Turbo+NoECC la temperatura media de cada núcleo ha aumentado 6 grados centígrados (pudiendo llegar en ocasiones hasta los 10º).

De los resultados obtenidos hasta ahora, podemos decir que el código Sstore va a ser el código sobre el que vamos a realizar las pruebas subsecuentes. Hemos descartado el Huge por no aportar ninguna ganancia, y también la opción Turbo por no ser aconsejable su utilización de forma continuada.

Por último, vamos a volver a evaluar la escalabilidad de nuestros códigos corriendo en Xeon Phi. Como vimos para la versión base, al estar los códigos limitados por su acceso a memoria, no escalaban a un elevado número de hilos, obteniendo un mejor tiempo de ejecución para 122 hilos (2 hilos por núcleo). En la Figura 3.10 se muestra la evaluación de los tres códigos Stencil para la versión de código optimizada con código Sstore para 122, 183 y 244 hilos. Como podemos observar, para esta versión del código optimizado la escalabilidad obtenida es muy pequeña y no se aprecian diferencias importantes en el rendimiento obtenido, lo que nos confirma que dichas versiones tienen prácticamente ya saturado el ancho de banda de memoria a partir de 2 hilos por núcleo.

Puesto que no obtenemos peores rendimientos, que para el caso del código de difusión de calor si que hay un cierto beneficio, y que pensamos que las versiones

Abbildung in dieser Leseprobe nicht enthalten

Figura 3.10: Evaluación de la escalabilidad en Xeon Phi para la versión (Sstore).

más optimizadas que mostraremos en el Capítulo 4 se van a beneficiar del mayor número de hilos, vamos a tomar a partir de ahora el número de 244 hilos para todas las evaluaciones que realicemos en el resto del documento.

3.5.4 Evaluación de la influencia de los datos de entrada

Para concluir este Capítulo 3 acerca de las directrices básicas para programar códigos Stencil 3-D en una arquitectura Intel Xeon Phi, nos resta evaluar la influencia que tienen los datos de entrada sobre los resultados obtenidos.

Vamos a estudiar dicha influencia bajo dos aspectos: el tamaño de los datos de entrada, así como la forma de dichos datos. Ambos aspectos tienen que ver con la influencia que la memoria está jugando en estos kernels, por lo que vemos necesario evaluar hasta qué punto dichos aspectos puedan influir en los resultados y las conclusiones obtenidas.

Vamos a empezar con la influencia del tamaño de la matriz de entrada. Como ya dijimos en la Sección 3.4, hemos considerado tres tamaños de datos, pequeño (matriz de 400x400x400 elementos), mediano (800x400x600) y grande (800x900x900). El tamaño mayor de los datos está limitado al tamaño de problema máximo posible para poder ejecutar cada uno de los tres kernels en Xeon Phi, y hemos elegido un factor de 3 entre cada una de las tres clases de tamaños. Es importante mencionar que estos tamaños se refieren sólo al tamaño de una de las matrices de entrada del kernel (denominada en nuestros códigos como fin).

Abbildung in dieser Leseprobe nicht enthalten

Figura 3.11: Influencia del tamaño de los datos de entrada.

En la Figura 3.11 podemos ver la evaluación para dichos tamaños. Como se puede apreciar, los resultados obtenidos son similares, lo que nos confirma que todos los tamaños elegidos saturan el ancho de banda a memoria (incluido el tamaño pequeño), por lo que no tiene ninguna influencia significativa el tamaño del problema. Por ello, y salvo que se diga lo contrario, para el resto de nuestras evaluaciones se tomará un tamaño mediano de problema (con un número de elementos en doble precisión de 800x400x600).

Por último, vamos a evaluar la forma de la matriz de entrada. Elegimos una forma cuboide perfecta (400x400x400) y tres formas cuboides rectangula- res (800x400x600, 600x800x400 y 400x600x800). La Figura 3.12 nos muestra los resultados obtenidos para nuestros tres kernels estudiados.

De nuevo, podemos apreciar que las diferencias no son significativas, y que el rendimiento obtenido en la ejecución de nuestros códigos es bastante independiente de la forma que puedan tener las matrices de los datos de entrada. De nuevo, tenemos que insistir en que lo significativo es la saturación del ancho de banda de memoria, y que el detalle de la forma de entrada tiene una muy pequeña influencia. Por tanto, y salvo que se diga lo contrario, para el resto de

Abbildung in dieser Leseprobe nicht enthalten

Figura 3.12: Influencia de la forma de los datos de entrada.

nuestras evaluaciones se tomará una forma cuboide rectangular de 800x400x600 elementos con la dimensión mayor situada sobre el eje X.

3.6 Conclusiones

En este Capítulo hemos presentado un conjunto de directrices para mejorar el ren- dimiento de las aplicaciones basadas en patrones Stencil 3-D para la arquitectura Intel Xeon Phi (Knights Corner). La programabilidad en arquitecturas masiva- mente paralelas puede ser un desafío para los desarrolladores sin experiencia, pero las directrices propuestas facilitarán la portabilidad de cualquier aplicación basada en Stencil a estas arquitecturas. Hemos utilizado el lenguaje C/C++ con extensiones OpenMP para codificar los tres kernels de los Stencil evaluados.

Los resultados experimentales obtenidos nos llevan a las siguientes conclusiones:

- La portabilidad de los códigos Stencil 3D para su ejecución en una arquitec- tura Intel Xeon Phi es inmediata. A partir de los códigos base que teníamos para una arquitectura Intel multinúcleo, su ejecución en una arquitectura Intel Xeon Phi se reduce a volver a compilar dichos códigos añadiendo una opción de compilación (-mmic) que le indique al compilador Intel que debe generar código para la arquitectura Xeon Phi.

- De cara a la mejor paralelización de nuestros códigos ha sido muy adecuado añadir el modificador collapse(2) al pragma utilizado para paralelizar el triple bucle anidado que recorre los datos de entrada (#pragma omp parallel for). De esta forma facilitamos el uso de un mayor número de hilos (hasta 244) de la arquitectura Intel Xeon Phi, combinando los dos bucles mas externos de nuestro kernel, ofreciendo a los distintos hilos una mayor cantidad de trabajo que se puede realizar en paralelo.

- El mayor beneficio obtenido en la ejecución sobre Xeon Phi se logra por medio de facilitar que el compilador vectorice adecuadamente el código. Para ello, ha sido de gran utilidad tener una alineación de los datos de entrada a 64 bytes (por medio de usar la función _mm_malloc), así como utilizar la técnica de padding para rellenar adecuadamente la estructura de datos de entrada y lograr que todas las filas de la matriz estén alineadas. Para evitar que el compilador sea incapaz de vectorizar adecuadamente el bucle interno de cálculo, ha sido necesario añadir para nuestros kernels acústico y de calor la palabra reservada restrict en la declaración de la función que realiza el cálculo Stencil, así como añadir la opción -restrict al compilar dichos códigos. En el caso del kernel sísmico, hemos utilizado la directiva #pragma ivdep para forzar al compilador a que ignore las posibles dependencias que encuentra y vectorice el código.

- Aunque siempre es muy adecuado optimizar el código escalar todo lo posible, hemos comprobado que la ganancia obtenida es muy pequeña (menos de miles de millones de FLOPs) para que merezca la pena hacer un gran esfuerzo en dicha optimización. Simplemente compilando con el nivel de optimizaciones más elevado del compilador (opción -O3) se obtiene un nivel de prestaciones adecuado.

- Con respecto a la escalabilidad de estos códigos en Xeon Phi, hemos com- probado que dicha escalabilidad es pésima debido a la gran cantidad de datos que necesitan y a la baja intensidad aritmética que tienen. Para las versiones de los códigos menos optimizadas, obtenemos mejores rendimientos utilizando únicamente 122 hilos (2 hilos por núcleo). Para la mejor versión obtenida en este Capítulo (con la optimización de Streaming Stores), hemos comprobado que no hay diferencia significativa entre usar 122, 183 o 244 hilos, aunque hay una ligera diferencia a favor de tener 244 hilos.

- Con respecto a la comparativa entre Xeon Phi y Xeon multinúcleo, hemos comprobado que las versiones no optimizadas obtienen una aceleración en tiempo de ejecución en Xeon Phi alrededor del 30%, mientras que en las versiones optimizadas hemos llegado a un factor cercano a 7 para todos los códigos entre la versión de Xeon Phi y la versión de Xeon.

- La opción de usar la caché sin localización de bloques en caso de fallo de escritura (denominada por Intel Streaming Stores) no afecta en el caso del código acústico y sísmico (debido al tipo de dependencia temporal que tiene los datos), siendo muy conveniente en el caso del kernel de difusión de calor. Esta opción tiene la gran ventaja de ser muy sencilla de usar (basta añadir la directiva #pragma vector non temporal (fout) antes del bucle tridimensional más interno en la dimensión en X).

- La opción de usar páginas grandes de memoria no ha tenido ningún efecto significativo. Si que hemos comprobado, por otra parte, que la tecnología de overclocking Intel Boost, junto a deshabilitar el manejo de ECC por software, producen una ganancia significativa. Sin embargo, y por problemas de calentamiento del procesador, así como de posibles errores de cálculo, no aconsejamos su uso de forma habitual ni continuada.

- Finalmente, hemos comprobado que el tamaño y la forma de los datos de entrada no tienen una influencia significativa, al tratarse de unos códigos que están limitados por el ancho de banda de acceso a memoria. Por lo tanto, y para el resto del documento, hemos elegido usar el tamaño de datos mediano con forma rectangular cuboide con la dimensión mayor en el eje X (matriz de entrada de 800x400x600 elementos de doble precisión).

Concluir diciendo que las aplicaciones del mundo real con base en cálculos Stencil pueden beneficiarse enormemente de las optimizaciones propuestas, ya que no sólo proporcionan resultados más rápidos, sino que también mejoran la precisión al permitir simulaciones más detalladas de los diferentes fenómenos.

Capítulo 4 Optimizaciones avanzadas. Comparativa en energía y con GPUs

4.1 Introducción

En el Capítulo 3 se ha evaluado el rendimiento de la arquitectura Intel Xeon Phi para tres kernels Stencil 3-D: difusión de calor (11 puntos), difusión acústica (7 puntos), e isotrópico de onda sísmica (25 puntos). Para todos ellos se ha usado una representación de los datos de entrada con valores de punto flotante de doble precisión (8 bytes). Partiendo de una versión secuencial sencilla escrita en C/C++ para una arquitectura Intel Xeon, mostramos cómo el código se puede ejecutar de manera sencilla en Xeon Phi, compilando con el modificador -mmic. Una vez comprobada lo sencillo que es portar del código a Xeon Phi se han ofrecido diversas guías que muestran paso a paso el proceso básico de paralelización con OpenMP y vectorización mediante el compilador para Xeon Phi.

En este Capítulo queremos ir un paso más allá, y estudiar otras técnicas de optimización que pueden incrementar el rendimiento de nuestras aplicaciones. En concreto, nos centramos en dos áreas de trabajo: (1) la optimización en el proceso de paralelización del código, teniendo en cuenta aspectos tales como la afinidad y la planificación de los bucles, y (2) la optimización en el acceso a memoria, estudiando la técnica de tiling de datos.

Una vez presentadas las técnicas para incrementar el rendimiento de nuestros códigos en términos de FLOPS (número de operaciones en coma flotante por segundo), en este Capítulo nos planteamos otra de las métricas fundamentales en el campo del HPC: el consumo energético. Para ello se evalúa el consumo energético de nuestras propuestas algorítmicas en las distintas plataformas objeto de estudio. Entre estas plataformas, además de la plataforma paralela y orientada a latencia Intel Xeon, añadimos una de las arquitecturas más avanzadas del mercado; la arquitectura CUDA de Nvidia. Aprovechando que en nuestro grupo de investigación teníamos portado el código del kernel acústico de 7 puntos a la arquitectura CUDA, hemos comparado nuestros desarrollos en Intel Xeon Phi con los obtenidos en CUDA, tanto desde el punto de vista del rendimiento en GFLOPs como del consumo energético.

4.1.1 Objetivos

El objetivo de este Capítulo es profundizar en los aspectos más novedosos de la arquitectura Xeon Phi: su elevado número de hilos, sus unidades vectoriales de 512 bits y su gran ancho de banda a memoria). A través de la aplicación de optimizaciones avanzadas podemos mejorar el rendimiento que hemos obtenido en nuestros códigos en el Capítulo 3. Junto a eso, también queremos estudiar la eficiencia energética de esta arquitectura midiendo el consumo energético de cada una de las versiones que hemos desarrollado para los diferentes códigos. Finalmente, nos parece interesante realizar una comparativa con la arquitectura CUDA de Nvidia. Los principales objetivos que queremos cubrir en este Capítulo son los siguientes:

- Planificación de los hilos: El lenguaje OpenMP permite realizar tres tipos de planificación cuando se ejecutan en paralelo diversos hilos en el procesador: static, guided, y dynamic. Al tener la arquitectura Xeon Phi un elevado número de hilos, se recomienda evaluar esta característica que podría influir en el rendimiento del código a ejecutar.
- Afinidad: La arquitectura Xeon Phi presenta tres tipos de afinidad (balanced, scatter y compact) para asignar los diversos hilos lógicos que se estén ejecutando a los núcleos físicos que tiene la arquitectura.
- Optimizar el acceso a memoria: La arquitectura Intel Xeon Phi requiere un uso adecuado de las cachés para lograr un buen rendimiento cuando ejecutamos aplicaciones Stencil. Una de las técnicas más conocidas para mejorar la localidad de datos es la de acceder a los datos de memoria por bloques, a fin de reutilizar todo lo que podamos los datos a los que hemos accedido. Dicha técnica es conocida como tiling o blocking y su implementación no es nada trivial siendo muy dependiente del kernel que estemos ejecutando. Queremos, por tanto, implementar esta técnica en nuestros códigos, y evaluar el tipo de blocking más adecuado (en la dimensión X, o Y, o Z, o combinaciones de ellas) así como ajustar el tamaño de los mismos.
- Eficiencia energética: En comparación con las arquitecturas multinúcleo tradicionales, que cuentan con unidades agresivas capaces de realizar un ejecución especulativa y desarrollar cálculos fuera de orden, el acelerador Xeon Phi (y también las GPUs de Nvidia) está diseñado mediante la in- clusión de muchos núcleos de computación sencillos, cuya ejecución es en orden. Uno de los principales beneficios de este diseño es que Xeon Phi presenta una mejor eficiencia energética. Queremos, por tanto, evaluar el consumo de energía realizado por las diversas versiones de nuestros códigos, y ver cuál de ellas es la mejor si tomamos como parámetro de medida la eficiencia energética.
- Comparativa con CUDA: Finalmente, queremos evaluar cómo los datos obtenidos para la arquitectura Xeon Phi se relacionan con los obtenidos en otra arquitectura habitual utilizada para acelerar códigos científicos, CUDA de Nvidia. Realizaremos dicha comparativa en términos de rendimiento y consumo energético para el kernel Stencil de difusión acústica de 7 puntos.

4.1.2 Trabajo relacionado

A continuación se ofrece un resumen de los principales trabajos relacionados con los desarrollos presentados en este Capítulo.

Muchas propuestas se han centrado en mejorar la reutilización de caché por medio del tiling o blocking. En [58, 64], los autores se centran en la explotación de localidad de datos mediante la aplicación de técnicas de tiling. En [31, 36, 38, 77] consideran problemas de localidad y paralelismo. Los autores en 36 desarrollaron un enfoque para la paralelización automática de los códigos Stencil que aborda explícitamente la cuestión de la ejecución de carga equilibrada de los tiles. Kamil et al. 31 examina varias optimizaciones específicas para mejorar la reutilización de caché a través de barridos del Stencil. Esto permite que múltiples iteraciones del Stencil reutilicen la porción de datos residentes en la caché. J. Fang et al. 17 evalúan Xeon Phi por medio de microbenchmarks así como diversos kernels útiles, uno de ellos el código stencil, encontrando para el kernel Stencil que no llega al 10% del rendimiento pico de Xeon Phi por problemas de saturación del ancho de banda de la memoria (cuando la técnica de blocking es utilizada).

Los cálculos Stencil con grids regulares han sido ampliamente estudiados en GPUs, debido a que los cálculos se pueden mejorar mediante el aprovechamiento del ancho de banda de la DRAM de la GPU. Micikevicius reportó la implemen- tación de un método de aplicación Stencil de diferencias finitas en CUDA con (blocking) de memoria compartida 42. Datta et al. 12 estudiaron el rendimiento y optimización Stencil en varias arquitecturas multinúcleo y aceleradores. Las op- timizaciones más comunes utilizados en estos estudios anteriores son el blocking en los registros y memoria compartida.

Para Intel Xeon Phi, en60podemos encontrar diversas optimizaciones para mejorar el código vectorizado, y en75se describe una herramienta para medir el consumo energético en esta plataforma.

La técnica de blocking temporal también ha sido ampliamente estudiada para diversos cálculos científicos. En 67 se estudia para el patrón Stencil. Este también ha sido evaluado en el entorno de las GPUs [25, 47, 78]. En [11, 30] los autores examinan múltiples técnicas de optimización para códigos stencil tanto para arquitecturas CPU como GPU. Por último, referimos al lector a 42 y 10 para encontrar optimizaciones desarrolladas para los patrones Stencil en arquitecturas CUDA de Nvidia.

En 50 se describe el modelo que se utiliza para predecir el campo de sonido en salas de forma arbitraria y una distribución no homogénea de la absorción del sonido.

4.2 Mejora de la paralelización

Esta Sección muestra cómo mejorar el rendimiento de la paralelización de los tres códigos Stencil que utilizamos en este documento. Para ello evaluamos dos características fundamentales en la paralelización de un código. En primer lugar, la planificación de la carga de trabajo asignada a cada uno de los hilos ejecutados; esto es, cómo distribuir las diferentes iteraciones del bucle del cálculo del kernel Stencil entre los hilos de Xeon Phi. En segundo lugar comprobamos la afinidad para ver cómo asignar los diversos hilos lógicos a los núcleos físicos de la arquitectura. Debido al número de hilos y de núcleos que presenta Xeon Phi, se recomienda evaluar ambas características para obtener un mayor rendimiento en la ejecución de cualquier código paralelo.

4.2. Mejora de la paralelización

4.2.1 Políticas de planificación de bucles OPENMP

Uno de los atributos más importantes para el rendimiento de las aplicaciones para- lelas que se están ejecutando usando varios hilos es asegurar el balanceo de carga, es decir, que la división del trabajo entre los hilos es equitativa. El balanceo de la carga de trabajo en los diversos hilos es extremadamente importante, ya que garantiza que los procesadores físicos (los núcleos) están ocupados de forma similar, obteniendo por tanto el mejor rendimiento posible. Sin una carga equili- brada, algunos hilos pueden terminar significativamente antes que los demás, dejando inactivos algunos de los núcleos del procesador, y desperdiciando así oportunidades para mejorar el rendimiento.

Por lo general, mediante el examen del código fuente es fácil determinar la variabilidad en el tiempo de cálculo de las diversas iteraciones de un bucle. En muchos casos, las diversas iteraciones del bucle consumen una cantidad similar de tiempo, por lo que no hace falta realizar ninguna planificación para distribuir la carga de trabajo entre los hilos que van a ejecutar dicho bucle (por defecto, se realiza una distribución uniforme). Cuando esto no es cierto, es necesario ajustar el balanceo de la carga dando información adicional en la programación del bucle en OpenMP para distribuir mejor las iteraciones del bucle entre los hilos (y por tanto entre los núcleos).

Esta información adicional es lo que se conoce como política de planificación de un bucle, y consiste en dar la información necesaria para que en la ejecución del código se puedan distribuir las iteraciones del bucle entre los hilos en cantidades aproximadamente iguales. Además, las diversas políticas de planificación tratan también de reducir al mínimo los posibles conflictos de memoria que pueden surgir debido al problema de la falsa compartición (false sharing) de bloques de memoria. A veces, puede suceder que la mejor opción para resolver los problemas de memoria puede ser mala para el balanceo de carga entre los hilos, y al contrario. Por ello, se deben evaluar las diversas políticas de planificación para encontrar un equilibrio entre el uso óptimo de la memoria y el balanceo de la carga óptima mediante la medición del rendimiento, a fin de determinar qué política produce los mejores resultados.

En el lenguaje OpenMP hay que agregar la clausula schedule al pragma que paraleliza un bucle para indicarle a OpenMP qué política se quiere aplicar a la hora de asignar trabajo a cada hilo de ejecución. Por tanto, la directiva queda de la siguiente forma: #pragma omp parallel for schedule(tipo [,tamaño¹ ]). Se pueden proporcionar cuatro tipos de planificación más uno adicional que permite en tiempo de ejecución elegir entre uno de estos cuatro tipos. A continuación describimos cada una de estas políticas de planificación:

1. static: Divide el bucle en trozos de igual tamaño o tan igual como sea posible en el caso de que el número de iteraciones del bucle no sea divisible por el número de hilos multiplicado por el tamaño especificado. Por defecto, el tamaño de iteraciones asignado a cada hilo es igual a: Número de iteraciones / Número de hilos. Si se pone un tamaño de 1, se intercalan las iteraciones entre los hilos disponibles.
2. dynamic: Usa una cola de trabajo interna para dar un bloque de iteraciones del bucle a cada hilo en tiempo de ejecución. Cuando un hilo termina, recupera el siguiente bloque de iteraciones del bucle de la cola de trabajo. Por defecto, el tamaño es 1 (se va asignando iteración a iteración), aunque se puede modificar dicho valor cambiando el tamaño en la directiva. Hay que tener cuidado al usar este tipo de planificación debido a la sobrecarga adicional que tiene en tiempo de ejecución.
3. guided: Similar a la planificación dinámica, pero con un tamaño de carga de trabajo (número de iteraciones) que comienza grande y va disminuyendo para lograr un mejor equilibrio entre la carga de los hilos. El parámetro opcional de tamaño especifica el tamaño de carga de trabajo más pequeña a considerar. Por defecto, dicho tamaño mínimo es igual a: Número de iteraciones / Número de hilos.
4. auto: Cuando se especifica una planificación automática, la decisión con respecto a la planificación se delega en el compilador. El programador cede al compilador la libertad de elegir cualquier posible asignación en el número de iteraciones para cada hilo.
5. runtime: En este caso, se utiliza la variable de entorno OMP_SCHEDULE para especificar cuál de los tres primeros tipos de planificación debe ser utilizado. La variable OMP_SCHEDULE es una cadena con el mismo formato que aparece en la directiva de asignación de la planificación.

En la evaluación de dichas políticas hemos probado con tamaños de 1, 2 y 4, para las diferentes políticas de planificación, más el tamaño por defecto que tiene cada una de dichas estrategias.

4.2.2 Políticas de afinidad

La librería de ejecución o (runtime) que tiene OpenMP posee la capacidad de asignar los hilos lógicos a diferentes núcleos físicos de nuestra arquitectura. De- pendiendo de la topología de la conexión entre los núcleos, de las características de comunicación de la aplicación, y del sistema operativo, la afinidad puede producir un cierto beneficio en el rendimiento de una aplicación. Para activar un tipo u otro de afinidad, se usa la variable de entorno KMP_AFFINITY. En Intel Xeon Phi tenemos 3 tipos de afinidad:

- compact: Con este tipo de afinidad el runtime de OpenMP asigna el hilo <n>+1 al núcleo más cercano (o el mismo núcleo si es posible) donde se asignó el hilo <n>. Por ejemplo, con este tipo de afinidad tendríamos que el núcleo físico 0 estaría ejecutando los hilos 0..3, en el núcleo físico 1 estarían los hilos 4..7, y así sucesivamente.
- scatter: Con este tipo de afinidad la librería runtime de OpenMP asigna los hilos de la manera más dispersa posible. Por ejemplo, con este tipo de afinidad tendríamos que en el núcleo físico 0 estaría el hilo 0, en el núcleo físico 1 estaría el hilo 1, y así sucesivamente. La afinidad scatter es la contraria a la afinidad compact.
- balanced: Este tipo de afinidad es exclusivo de la arquitectura Intel Xeon Phi. Con este tipo de afinidad el runtime de OpenMP asigna los hilos a diferentes núcleos hasta que todos los núcleos tengan al menos 1 hilo. En esto es similar a cómo funciona la afinidad scatter. Sin embargo, cuando el runtime tiene que asignar varios hilos a un núcleo, la afinidad balanced asegura que los hilos con numeración más parecida estarán en el mismo núcleo, cosa que no lo asegura la afinidad scatter.

En la Figura 4.1 se muestra un ejemplo de cada tipo de afinidad para una configuración de 8 hilos y 4 núcleos. Para probar el efecto de la afinidad, en nuestros experimentos hemos modificado la variable de entorno KMP_AFFINITY a compact, balanced, o scatter.

4.3 Optimizando el acceso a memoria

Esta Sección realiza un estudio de una de las técnicas más utilizadas en arquitec- turas con jerarquías de caché multinivel. Esta técnica es conocida como blocking o

Figura 4.1: Distribuciones de tipo de afinidad (KMP_AFFINITY).

Abbildung in dieser Leseprobe nicht enthalten

tiling. Esta técnica promueve la reutilización de datos en los niveles más altos de la jerarquía de memoria con el objetivo de evitar costosos accesos a memoria y/o disco, cuyas latencias de acceso son órdenes de magnitud superiores a las de las memorias cachés. Está demostrado que esta técnica mejora el rendimiento de la ejecución de una aplicación paralela limitada por su ancho de banda a memoria, siempre y cuando identifiquemos el mejor parámetro de tamaño de bloque, que depende principalmente del problema y la arquitectura donde se está ejecutando dicha aplicación.

La mayoría de los Stencil exhiben un alto grado de localidad espacial, ya que cada operación de actualización accede a los valores vecinos del grid. En un Stencil 3-D, hay reutilización de datos a lo largo de las tres dimensiones (dimensión X, dimensión Y, dimensión Z). Los códigos Stencil con un tamaño de entrada que no caben en los niveles más altos de la memoria caché del procesador, como nuestro caso, experimentarán una degradación significativa del rendimiento, debido a los fallos de capacidad de la caché ya que los accesos a los datos están por lo general demasiado separados. Las transformaciones del código que mejoren la localidad de datos pueden ser útiles para ocultar la complejidad de los accesos a memoria, mejorando el rendimiento general de los códigos Stencil. Estas transformaciones básicas incluyen transformaciones de bucles² y transformaciones de datos.³ Con la técnica de blocking pretendemos minimizar los fallos de caché, y por tanto reducir los requisitos de ancho de banda de la aplicación.

El blocking es una transformación que agrupa iteraciones de bucles en sub- conjuntos de tamaño N (o tiles). El tamaño de los tiles debe ser ajustado para que quepa en la memoria caché con el fin de obtener el máximo de ganancia de rendimiento en la explotación de localidad de datos. De esta manera, los fallos

Abbildung in dieser Leseprobe nicht enthalten

de caché se pueden minimizar trayendo bloques de datos a caché una vez para todos los accesos necesarios.

En nuestros códigos Stencil explotamos la localidad de datos, centrándose en el aumento de la reutilización de los elementos del plano (X-Y) a lo largo del eje Z. El primer paso que consideramos fue crear tiles de tamaños width_Tblock, height_Tblock y depth_Tblock (para la dimensión X, Y y Z, respectivamente). A continuación, se crean tres bucles adicionales sobre los tres bucles ya existentes para recorrer el conjunto de datos en los tiles de los tamaños seleccionados. El Algoritmo 5 muestra una versión genérica de blocking para nuestros códigos Stencil 3-D.

Para lograr el mejor tamaño del bloque, hemos llevado a cabo un análisis de diferentes tamaños de bloque de la siguiente manera:

a) En primer lugar, buscamos los mejores tamaños de bloque para el plano XY, manteniendo el tamaño de blocking en el eje X (width_Tblock) igual al tamaño de la dimensión del eje X (width_Tblock=width). Para ello, variamos los valores sobre el eje Y (height_Tblock= 2, 4, 8 y 16) y sobre el eje Z (depth_Tblock= 1, 2, 4, 8, 16 y 32). Esta evaluación permite obtener mejoras en el rendimiento con determinados tamaños de blocking, pero también encontramos que con otros tamaños puede disminuir.
b) Posteriormente buscamos el mejor tamaño de blocking en el eje X (width_Tblock), variando los valores en dicho eje (width_Tblock= 150, 200, 300, 400), para un tamaño de la dimensión X, que en nuestro caso es de 800. Para esto, tomamos los mejores tamaños de blocking del eje Y y eje Z, identificados en el punto anterior.
c) Finalmente evaluamos los mejores tamaños de bloque identificados en los dos puntos anteriores combinado con la planificación del bucle OpenMP (por ejemplo, schedule(dynamic,2)) en la búsqueda de un mejor rendimiento.

4.4 Eficiencia energética

Uno de los aspectos interesantes que presenta la arquitectura Xeon Phi (y los aceleradores en general) es su alta eficiencia energética. Es decir, además de pro- porcionar un elevado rendimiento en potencia de cálculo (tiempo de ejecución), el consumo energético asociado a dicho rendimiento es mucho mejor que el que se obtiene en las tradicionales arquitecturas multinúcleo con ejecución agresiva de instrucciones fuera de orden. Esto implica que los aceleradores proporcionan un ratio FLOPs / Vatio consumido mucho mayor que las arquitecturas habituales en el campo de la computación de altas prestaciones.

Es por ello que en la actualidad la mayoría de propuestas de plataformas de servidores y máquinas de HPC poseen una arquitectura heterogénea que integra un buen número de tarjetas aceleradoras, ya sean Intel Xeon Phi o GPUs basadas en la arquitectura CUDA de Nvidia.

Por ello, en este documento, una vez mejorados códigos Stencil 3-D al máximo, nos planteamos la medición de su consumo energético para evaluar su eficiencia energética. Para ello vamos a utilizar la figura de mérito EDP (Energy Delay Product o producto de tiempo de ejecución por energía consumida en dicha ejecución), ya que este parámetro tiene en cuenta tanto el rendimiento en potencia de cálculo de una aplicación (medido en segundos) como su consumo energético. Dicha evaluación la queremos hacer en dos planos:

- Por una parte, comparando cada una de los versiones de nuestros códigos en Xeon Phi desde el punto de vista de la eficiencia energética. Hasta ahora hemos buscado obtener la mejor versión en Xeon Phi tomando como medida la potencia de cálculo desarrollada (GFLOPs). En este apartado queremos tener en cuenta también la energía consumida por cada versión, y ver cuál de todas las versiones desarrolladas es mejor en términos de su eficiencia energética.
- Por otra parte, comparando los resultados obtenidos en Xeon Phi con los obtenidos por las mismas versiones en Xeon multinúcleo. Esto nos va a permitir visualizar cuan bueno es esta nueva plataforma en el terreno de la eficiencia energética en relación a las plataformas más tradicionales de HPC.

En relación a cómo se realiza la medición de la energía consumida por una aplicación, durante décadas las estimaciones de potencia y energía se basaron en modelos complejos o dispositivos externos caros. Hoy en día, muchos desarro- lladores de hardware proporcionan medios para dar información de potencia y energía a desarrolladores de aplicaciones y centros de cómputo. Para las arqui- tecturas Intel, la segunda generación de procesadores de núcleo i5 e i7 extendió su MSRs⁴ con información de energía, permitiendo a las aplicaciones examinar su eficiencia energética sin necesidad de utilizar dispositivos externos.

El procesador Xeon incluye tres registros de energía, uno para todo el paquete (PKG), uno para los núcleos (Power Plane 0 - PP0), y otro para la memoria (DRAM). Las últimas versiones de las librerías PAPI (5.0 o superior) [44] inclu- yen soporte para la lectura de registros de energía de Intel. Plataformas como Intel Xeon Phi también son soportadas por PAPI apoyándose en componentes externos (micpower). Este componente permite a los desarrolladores extraer las lecturas de potencia instantáneas para diferentes partes de la placa de Xeon Phi, incluyendo la potencia del núcleo, potencia total de la tarjeta, potencia de la PCIe, etc. Por lo tanto, en la Sección de evaluación de este apartado presentaremos para Xeon multinúcleo la energía consumida total obtenida a partir del registro PKG, mientras que para la tarjeta Intel Xeon Phi mostraremos la energía total consumida por la misma.

Un último comentario importante a tener en cuenta cuando se mide la energía consumida. En el caso de Xeon Phi, donde los registros proporcionan información “instantánea” de la potencia disipada, no energía antes y después de acabar de ejecutar el kernel, crearemos un proceso que irá muestreando la potencia disipada cada segundo. Debido a que este tipo de instrumentación introduce una sobrecarga adicional en el tiempo de ejecución proporcional a la velocidad de muestreo, no es posible medir la energía y el rendimiento en tiempo al mismo tiempo de una aplicación. Por ello, la información de rendimiento y del consumo de energía se obtuvieron en ensayos separados para minimizar la sobrecarga mencionada. Adicionalmente, hay que tener en cuenta que la temperatura que tenga el procesador afecta a las lecturas de energía hasta en un 5% cada 8-10 grados centígrados [9]. En nuestro caso, este problema lo hemos solucionado porque como ya dijimos en la Sección 3.5.1.4 de la metodología de nuestros experimentos, ejecutamos cada aplicación 10 veces consecutivas, con lo que hemos comprobado que hay tiempo suficiente para que el procesador se caliente y alcance una temperatura estable. Lógicamente, la medición energética la realizamos a partir de las pruebas en que ya se ha estabilizado la temperatura.

4.5 Implementación del modelo de la ecuación de difusión acústica 3D en GPU

Aprovechando que dentro del grupo de investigación se disponía de una imple- mentación del kernel del modelo de la ecuación de difusión acústica 3D-FD para la arquitectura CUDA de Nvidia, nos ha parecido muy interesante comparar los resultados obtenidos en la arquitectura Xeon Phi con el obtenido en CUDA, tanto en cuanto a rendimiento como a consumo energético.

En esta Sección vamos a describir brevemente el proceso de codificación de dicho kernel de difusión acústica para CUDA (basado en el paralelismo de datos desarrollado para la arquitectura CUDA).

En primer lugar, comentar que el algoritmo propuesto en esta Sección se aplica a la ecuación (2.7), es decir, a aquellas celdas dentro del volumen del recinto simulado, que en realidad representa cerca del 95 % del tiempo total de ejecución.

Abbildung in dieser Leseprobe nicht enthalten

Figura 4.2: Diseño de bloques CUDA 2D. Cada hilo CUDA necesita un bucle sobre el eje Z.

Existen varios enfoques que hemos desarrollado para explotar el paralelismo de datos en la arquitectura CUDA, pero aquí tan sólo mostraremos el mejor diseño de los realizados en nuestro grupo. La Figura 4.2 muestra nuestro diseño CUDA. Se basa en el diseño de bloques de dos dimensiones al dividir el cubo en bloques verticales, uno para cada coordenada (x, y) de tamaño z. Un hilo CUDA se asigna a cada punto del grid (x, y) ∈ [1, NZ]. Esos hilos iteran sobre el eje z.

La Figura 4.3 muestra los datos de acceso por cada bloque 2D de CUDA. Mostramos los datos acostados en el eje X. De esta manera, los hilos dentro del mismo bloque pueden compartir información a través de la memoria compartida en el chip y los registros, tomando ventaja de la localidad de datos que está disponible en este patrón computacional.

Abbildung in dieser Leseprobe nicht enthalten

Figura 4.3: Datos accesados por un bloque CUDA 2D.

4.6 Evaluación

En esta Sección presentamos la evaluación de las estrategias de optimización avanzadas descritas en las secciones anteriores, la comparativa entre Xeon Phi y Xeon multinúcleo en términos de eficiencia energética, así como el estudio comparativo entre Xeon Phi y la arquitectura CUDA de Nvidia. En primer lugar, presentamos la descripción de nuestro entorno de evaluación.

4.6.1 Entorno de evaluación

En el Capítulo 3, en las secciones 3.5.1.1 y 3.5.1.2 ya hemos mencionado el entorno de evaluación hardware y software para Intel Xeon Phi, que también utilizamos en este Capítulo. En esta Sección, por tanto, vamos a describir únicamente el entorno de hardware y software, los conjuntos de datos de entrada y la medición de energía, para todo lo referente a la implementación en CUDA.

1. Plataformas Hardware

La arquitectura destino de nuestras pruebas está basada en el diseño Kepler (la tarjeta Tesla K40c) y en el diseño Maxwell (la GPU 980GTX). La Kepler K40c tiene 2.880 núcleos corriendo con un reloj de 0,88 GHz, ofreciendo un procesamiento pico de cálculo hasta de 5.068 GFLOPs. La velocidad de la memoria es 3,0 GHz con un ancho de bus de memoria de 384 bits que ofrece un ancho de banda de 288 GB/seg. El tamaño de la memoria es de 12 GB de GDDR5 con capacidades ECC. Por otra parte, la GPU 980GTX tiene 2.048 núcleos funcionando a 1,1 GHz, 4 GB de memoria GDDR5 y un ancho de banda de memoria máximo de 224 GB/seg.

2. Entorno Software Hemos utilizado el modelo de programación CUDA para programar las GPUs de Nvidia (Tesla K40c y GeForce 980GTX). Más precisamente, se ha utilizado el SDK de CUDA 6.5. La plataforma evaluada utiliza el sistema operativo Ubuntu 14.04.1. Hemos usado la media acotada (0.3) para reducir los efectos de los valores extremos en las mediciones.

3. Conjunto de datos de entrada

A continuación describimos los datos de los diferentes parámetros de entrada utilizados para la experimentación con CUDA, pues estos datos son diferentes en este estudio al conjunto de datos que hemos utilizado a través de todo el documento.

El recinto discretizado consiste en una sala de forma cúbica con dimensiones de 8x8x8m3. La discretización espacial y temporal fueron elegidas para cum- plir con las conclusiones enunciadas en 45. En consecuencia, (Δt)2/(Δν)2 debe ser del orden de 10−8 para asegurar que la implementación 3D-FD converge a un valor fijo con un error bajo. Aplicando este criterio de con- vergencia, cuatro resoluciones diferentes del grid fueron creadas para las simulaciones, con un número de celdas igual a 128x128x128 (32 MBytes), 256x256x256 (256 MBytes), 384x384x384 (864 MBytes) y 512x512x512 (2048 MBytes). Estos tamaños de grid, finalmente, determinan los requisitos de memoria de las implementaciones. El uso de memoria proviene principal- mente de cuatro matrices necesarias para los cálculos. Cada una utiliza ( INPUT_SIZE3 ) x sizeof(float) bytes de almacenamiento. Destacamos que para la evaluación con CUDA se ha utilizado un dato en coma flotante de tipo float (con el tipo double no teníamos espacio suficiente en memo- ria). Por tanto, estas evaluaciones también las hemos ejecutado en Xeon Phi con dicho tipo de dato.

4. Medición de energía en las GPUs

A continuación vamos a describir el proceso seguido para la evaluación del consumo de energía en las GPUs.

Con respecto a la medición de energía en la GPU, la medición de potencia en tiempo real de los componentes individuales de la GPU también se puede hacer en las modernas GPUs de Nvidia a través de la lectura de ciertos registros específicos que tiene la arquitectura. Esto se hace mediante el uso de la librería NVML (Nvidia Management Library) 1, que informa del uso de energía de la GPU en tiempo real. Afortunadamente, PAPI 44 también proporciona un componente para interactuar con NVML y recuperar la información de la energía desde la propia aplicación.

Hemos usado 2 plataformas CUDA de Nvidia distintas para la comparativa con Xeon Phi, la tarjeta K40c y la tarjeta 980GTX. En el caso de la K40c, el procedimiento seguido para medir la energía ha sido por medio de PAPI. Por desgracia, en el caso de la 980GTX , no hemos podido acceder a los contadores de energía con nuestra configuración actual (driver + versión NVML). Para resolver este problema, hemos utilizado un medidor externo que mide la potencia disipada por todo el sistema (at-the-wall power). A dicho valor le restamos la energía utilizada por la plataforma en que se encuentra dicha tarjeta gráfica ejecutando una aplicación de un solo hilo en la CPU (sin usar dicha GPU).

4.6.2 Evaluación de las políticas de planificación

En esta Sección vamos a evaluar el rendimiento que hemos obtenido al aplicar las políticas de planificación a nuestros tres códigos Stencil en la arquitectura Xeon Phi. Partimos de los tres códigos Stencil (Acústico, Sísmico y Calor) identificados en el Capítulo 3 con la etiqueta (Sstore), es decir, código vectorizado optimizado con streaming store. Como comentamos en la Sección 4.2.1, han sido tres las políticas de planificación evaluadas (dynamic, static y guide), con un tamaño de 2, 4 y el valor por defecto.

En la Figura 4.4 se muestran los resultados de las optimizaciones realizadas mediante las políticas de planificación del bucle de cálculo OpenMP. Como se puede apreciar, static es la opción que menos rendimiento presenta en los tres kernels, aunque en el caso de (static,2) el rendimiento sube un 15,23% en los tres kernels, y con (static,4) sólo sube un 8,77 % para el caso del sísmico y el de calor.

Asimismo, podemos ver que usando la política guided de planificación, con cualquier parámetro, el rendimiento alcanzado es el mismo, pero un 5,51 % más que la obtenida con static. La mejor opción ha sido la política de planificación dynamic, concretamente (dynamic,4) para el caso de los kernels acústico y sísmico, y (dynamic,2) para el kernel de calor.

Abbildung in dieser Leseprobe nicht enthalten

Figura 4.4: Políticas de planificación del bucle OpenMP.

Si comparamos nuestros resultados con los obtenidos sin especificar ninguna política de planificación (que es lo mismo que usar shedule(static), ya que se toma por defecto), logramos un 34 % de aceleración para el kernel acústico, y un 42 % para el kernel sísmico, ambos usando la política shedule(dynamic,4), y para el kernel de calor obtenemos un 30 % usando una política (dynamic,2).

4.6.3 Evaluación de la afinidad

Esta Sección evalúa los resultados obtenidos por nuestros tres kernels al variar el parámetro de la afinidad. Hemos evaluado los tres tipos de afinidad (compact, scatter, balanced) junto a no usar ninguna afinidad (none). La Figura 4.5 muestra los resultados obtenidos, donde podemos observar que la afinidad tiene poca influencia en la ejecución de nuestros tres kernels.

Aunque no logramos una mejora significativa con ninguno de los tres tipos de afinidad, entendemos que para otros algoritmos este parámetro debería ser

Abbildung in dieser Leseprobe nicht enthalten

Figura 4.5: Planificación de la afinidad.

tenido en cuenta y evaluado. Elegimos la distribución balanced como base para ser utilizada en subsecuentes optimizaciones, a menos que se diga lo contrario.

4.6.4 Evaluación de la técnica del blocking

En esta Sección vamos a evaluar la técnica del blocking para optimizar el acceso a memoria.

En primer lugar, vamos a determinar el mejor tamaño de los bloques para los ejes YZ, para posteriormente buscar el tamaño del bloque en X. En la Figura 4.6 se muestra el rendimiento alcanzado por nuestros tres kernels Stencil, identificados con la etiqueta Acústico, Sísmico y Calor. En esta gráfica hemos establecido el tamaño del bloque en eje X igual al tamaño de la dimensión X (width_Tblock=800). Por lo que el rendimiento alcanzado sólo está en función del blocking utilizado en los ejes Y y Z.

Nuestros mejores resultados muestran que para el Stencil acústico, las mejores parejas de tamaños de bloque tienen un valor de 2 en el eje Y, mientras que para el eje Z sirven los tamaños de 8, 16 y 32 por igual. En el caso del Stencil sísmico, las mejores parejas de tamaños de bloque tienen un valor de 2 en el eje Y, mientras que para el eje Z sirven los tamaños de 2 y 4 por igual. Por último, para el Stencil de calor, las mejores parejas de tamaños de bloque tienen un valor de 2 y 4 en el eje Y, mientras que para el eje Z sirven los tamaños de 2 y 4 por igual.

Abbildung in dieser Leseprobe nicht enthalten

Figura 4.6: Blocking aplicado a la solución de nuestros códigos Stencil variando los ejes Y y Z.

Una vez obtenidos los mejores tamaños de bloque para las dimensiones Y y Z, vamos a pasar a evaluar el tamaño de bloque en la dimensión X. En la Figura 4.7 se muestra los tamaños de bloque que nos generan un mayor rendimiento variando el eje X. Hemos evaluado el tamaño del bloque en el eje X usando los valores 150, 200, 250, 400 y 600. Teniendo en cuenta los resultados de la Figura 4.6 anterior, para el eje Y se establecieron los valores 2 y 4, mientras que hemos elegido para el eje Z el valor de 4. Como podemos ver en esta Figura, el añadir la estrategia de blocking también a la dimensión X nos permite una ganancia adicional sobre el blocking en dos dimensiones de hasta un 15,19 % para el kernel sísmico, y hasta un 3,20% para el kernel de calor, mientras que para el kernel acústico no obtiene ningún beneficio.

Finalmente, comentar que en relación a las mejores versiones de los códi- gos optimizados del Capítulo 3 (usando la optimización de Streaming stores), que no tienen implementada la técnica de blocking ni una adecuada política de planificación, hemos obtenido una mejora sustancial en el rendimiento cuando hemos implementado una política adecuada de planificación (concretamente, schedule(dynamic)) junto a la técnica de mejora en el acceso a memoria (blocking)

Abbildung in dieser Leseprobe nicht enthalten

Figura 4.7: Blocking aplicado a la solución de nuestros códigos Stencil variando el eje X.

con unos valores correctos para los tamaños de bloque en cada dimensión. Concre- tamente, y usando para todos los códigos una planificación schedule(dynamic), hemos alcanzado una aceleración de 38,97% para el Kernel acústico para un tamaño de blocking en la dimensión X de 400, y 4 para las dimensiones Y y Z. El kernel sísmico alcanza una aceleración de 75,03 % para un tamaño de blocking en la dimensión X de 400, 2 en la dimensión Y, y 4 en la dimensión Z. Por último, el kernel de calor alcanza una aceleración de 46,23 %, para un tamaño de blocking en la dimensión X de 400, y 4 para las dimensiones Y y Z.

Finalmente, tanto elegir la política adecuada de planificación como establecer el tamaño adecuado de bloque en una dimensión determinada es un proceso un tanto artesanal que depende del kernel que se esté utilizando. Por ello, son unas técnicas muy útiles que hay que manejar con cuidado, porque en caso contrario podríamos no obtener ninguna mejora significativa o incluso empeorar el rendimiento del código.

4.6.5 Evaluación del rendimiento energético

En esta Sección vamos a mostrar la evaluación del rendimiento energético en Intel Xeon Phi, y lo vamos a comparar con lo obtenido en Xeon multinúcleo. Este es 4. Optimizaciones avanzadas. Energía y GPUs un punto importante del documento, pues queremos comprobar si efectivamente la arquitectura Intel Xeon Phi es tan eficiente energéticamente como dice el fabricante, pues reducir el consumo de energía en una plataforma HPC es hoy en día el parámetro clave en su diseño.

En la Figura 4.8 se muestra la energía consumida y la disipación de potencia media (eje Y secundario) para los tres códigos Stencil evaluados. La energía consumida en Xeon multinúcleo es alrededor de 4 veces más elevada que en Xeon Phi, llegando a valores de 50 o 60 KJoules, mientras que en Xeon Phi como máximo alcanza los 12 KJoules. En cuanto a la potencia media, para Xeon multinúcleo se mantiene bastante constante y está cerca de los 145 Watts. Para Xeon Phi muestra una mayor variabilidad dependiendo del código ejecutado, pero se todos los valores están alrededor de 200 Watts.

Abbildung in dieser Leseprobe nicht enthalten

Figura 4.8: Energía (prim.Y) y Potencia promedio (sec.Y) para nuestros kernels Stencil.

Antes de mostrar la métrica EDP, vamos a dar en una Tabla los tiempos de ejecución que hemos obtenido para cada una de las versiones que vamos a evaluar con la Figura EDP en ambas plataformas (Tabla 4.1).

A partir de la Tabla anterior, y de los datos mostrados en la Figura 4.8, obtenemos la Figura 4.9 que nos muestra la métrica EDP (eje Y primario) y el factor de ganancia en EDP entre Xeon multinúcleo y Xeon Phi para la misma versión de los códigos (eje Y secundario). Los resultados muestran una eficiencia

Abbildung in dieser Leseprobe nicht enthalten

Tabla 4.1: Tiempos de ejecución de tres versiones de nuestros códigos Stencil (seg.).

energética global superior en Xeon Phi, mientras que proporciona una menor eficiencia en la arquitectura Xeon.

Abbildung in dieser Leseprobe nicht enthalten

Figura 4.9: EDP (prim.Y) y Potencia promedio (sec.Y) para nuestros kernels Stencil.

De nuestra evaluación de rendimiento energético destacamos:

- A medida que mejora el rendimiento en GFLOPs para cada una de nuestras versiones, también consumen menos energía, teniendo una menor EDP. Así que nuestros tres kernels muestran mejoras en la eficiencia energética obtenida para ambas plataformas cuando son optimizados.
- Hemos obtenido unas mejoras de eficiencia energética de 28X para el kernel acústico, de 10X para el kernel sísmico, y de 21X para el kernel de calor, entre la mejor versión corriendo en Xeon Phi y esa misma mejor versión ejecutándose en Xeon multinúcleo.

4.6.6 Evaluación comparativa Xeon Phi con CUDA

Esta Sección compara las prestaciones de la tarjeta Intel Xeon Phi con las ob- tenidas en dos tarjetas GPUs de Nvidia. Para la comparación, hemos tomado un código CUDA que ya se tenía dentro del grupo de investigación. La mejor configuración de hilos para cada plataforma es: 32 hilos para Xeon multinúcleo, 244 hilos para Xeon Phi, y 256 hilos por bloque para las GPUs respectivamente.

A continuación, se muestran información del rendimiento, energía y EDP (Energy Delay Product) para las diferentes resoluciones de sala cúbica (128x128x128 (32 MBytes), 256x256x256 (256 MBytes), 384x384x384 (864 MBytes) y 512x512x512 (2048 MBytes) celdas) explicado anteriormente en la Sección 3.

Abbildung in dieser Leseprobe nicht enthalten

Figura 4.10: Tiempo de ejecución (prim. Y) y aceleración (sec. Y) para diferentes resoluciones de salas.

La Figura 4.10 muestra la información de tiempo de ejecución (escala loga- rítmica), así como la aceleración alcanzada⁵ (Eje secundario Y) para todas las plataformas evaluadas y número de hilos óptimo. Para la sala de tamaño pequeño, el tiempo de ejecución de Xeon Phi es mucho mejor que el obtenido en la 980GTX, y es casi dos veces menor que en la K40. Podemos notar que, a medida que aumenta el tamaño del problema, Xeon Phi pierde rendimiento comparado con las GPUs 980GTX y la K40, mostrando unos tiempos de ejecución bastante similares (véase también los resultados numéricos mostrados en la Tabla 4.2).

Queremos hacer notar que estos resultados los hemos obtenido en Xeon Phi usando una versión del código de difusión acústica muy optimizado. De hecho, en una primera versión de la comparativa que tenemos publicada en 24, los resultados eran más favorables a la implementación en CUDA. En aquella ocasión utilizamos la versión del código resultante del Capítulo 3. Cuando a dicha versión le hemos añadido la planificación dinámica y la mejora en el acceso a memoria utilizando el blocking, obtenemos que para todas las resoluciones existe una ligera ventaja de Xeon Phi sobre ambas implementaciones de las GPUs (980GTX y K40). Sin embargo, es posible que con una mejor implementación de la versión en CUDA se obtuvieran mejores resultados. Pero estas mejoras están fuera del alcance de este documento.

Por otro lado, las Figuras 4.11 y 4.12 muestran tanto energía y EDP, al igual que la disipación de potencia media (eje Y secundario) dado el número de hilos óptimo para cada plataforma. Es importante señalar que las GPUs requieren una CPU para trabajar, pero la energía de la CPU no se contabilizó en los cálculos de GPU. También vale la pena mencionar que la segunda generación de aceleradores Xeon Phi puede funcionar sin una CPU principal. Los resultados muestran un consumo de energía global similar entre Xeon Phi y las GPUs, con un pequeño margen superior para Xeon Phi. La disipación de energía promedio de Xeon Phi no cambia mucho con los tamaños de los problemas, y está cerca de 206 Watts, mientras que para las GPU está cerca de 150 Watts. La potencia media del CPU Xeon se mantiene alrededor de 144 Watts.

La métrica EDP da importancia adicional al rendimiento frente a la disipación de potencia. Para esta métrica (menos es mejor), Xeon Phi se encuentra cerca de las GPU, con una ligera ventaja a favor de Xeon Phi.

La Tabla 4.2 proporciona un breve resumen de los resultados de tiempo, energía y EDP obtenidos en nuestra evaluación para los diferentes tamaños de sala.

Abbildung in dieser Leseprobe nicht enthalten

Figura 4.11: Energía (prim. Y) y Potencia promedio (sec. Y) para diferentes resoluciones de salas.

Tabla 4.2: Comparativa de Xeon Phi y CUDA en tiempo, energía y EDP. Plataformas

Abbildung in dieser Leseprobe nicht enthalten

Tiempo (seg.), Energía (Kilo Joules) y EDP (KiloJoules*Seg.) para diferentes resoluciones de salas.

Específicamente, hemos obtenido una reducción de tiempo de ejecución superior a 7 entre Xeon Phi y Xeon multinúcleo (entre 5 y 6 cuando comparamos las GPUs con Xeon) para los diversos tamaños de salas considerados, mientras que en cuanto a su eficiencia energética, Xeon Phi ha obtenido un factor de mejora energética entre 35 y 40 con respecto al Xeon, por un factor de mejora entre 30 y 35 cuando comparamos las GPUs con Xeon.

Por lo tanto, concluimos que tanto Xeon Phi como CUDA (GPUs K40c y 980GTX) ofrecen beneficios claros para este kernel, tanto desde el punto de vista

Abbildung in dieser Leseprobe nicht enthalten

Figura 4.12: EDP (prim. Y) y Potencia promedio (sec. Y) para diferentes resolu- ciones de salas.

de la aceleración del tiempo de ejecución como especialmente de su eficiencia energética.

4.7 Conclusiones

En este Capítulo hemos profundizado en los aspectos novedosos y avanzados de la arquitectura del Intel Xeon Phi. Hemos tomado como punto de partida nuestros códigos optimizados en el Capítulo 3 y a través de la aplicación de optimizaciones avanzadas alcanzamos entre un 25 % y 40 % más de rendimiento. Otro aspecto de gran importancia hoy en día es el consumo energético, motivo por el cual incluimos en este Capítulo un estudio de la eficiencia energética de esta arquitectura que arroja que las optimizaciones a nuestros códigos consumen menos energía y tiene un EDP mejor. Así también, hemos llevado a cabo una comparativa con la arquitectura CUDA de Nvidia en cuanto a rendimiento y consumo energético. Específicamente, en este Capitulo hemos cubierto los siguientes aspectos:

- Hemos evaluado la planificación de hilos utilizando el lenguaje OpenMP, mediante las cláusulas static, guided, y dynamic. Nuestro estudio arroja una aceleración de 34% para el kernel acústico, y de 42% para el kernel sísmico, ambos usando la política shedule(dynamic,4), y de 30% para el kernel de calor usando una política (dynamic,2).
- Hemos evaluado los tres tipos de afinidad (balanced, scatter y compact) para esta arquitectura. En nuestra evaluación se mantuvo el rendimiento alcanzado mencionado en el punto anterior, mostrando una diferencia mínima en rendimiento entre los tres tipos de afinidad evaluados, eligiendo como la mejor opción la afinidad balanced.
- Hemos evaluado la optimización de los accesos a memoria, implementado la técnica del blocking en nuestros códigos, que se basa en acceder a los datos de memoria por bloques para reutilizarlos en los últimos niveles de la jerarquía de memoria. En nuestro estudio, hemos buscado el tamaño de blocking más adecuado (en la dimensión X, o Y, o Z). Tomando como base nuestros mejores códigos optimizados del Capítulo 3, que no tienen implementada esta técnica. Nuestros resultados arrojan una aceleración del 38,97% para el kernel acústico, del 75,03% para el kernel sísmico, y del 46,23% para el kernel de calor cuando utilizamos un tamaño de bloque de 400 en la dimensión X, de 4 en la dimensión Y (2 en el caso del sísmico), y también de 4 en la dimensión Z, usando en este caso para todos los códigos una planificación schedule(dynamic).
- Hemos evaluado la eficiencia energética de Xeon Phi y la hemos comparado con la obtenida en Xeon multinúcleo. Para ello, hemos usado la figura de mérito EDP (Energy Delay Product). Hemos evaluado el consumo de energía de diversas versiones de nuestros códigos, observando que a medida que optimizamos nuestros códigos, estos también consumen menos energía y tienen un mejor EDP. Hemos obtenido unas mejoras de eficiencia energética de 28X para el kernel acústico, de 10X para el kernel sísmico, y de 21X para el kernel de calor, entre la mejor versión corriendo en Xeon Phi y esa misma mejor versión ejecutándose en Xeon multinúcleo.
- Finalmente evaluamos los datos obtenidos para la arquitectura Xeon Phi con otra arquitectura utilizada para acelerar códigos científicos, en este caso la arquitectura CUDA. Para este estudio, sólo realizamos la comparativa para el kernel Stencil de difusión acústica de 7 puntos, en términos rendimiento y consumo energético. Específicamente, hemos obtenido una reducción de tiempo de ejecución superior a 7 entre Xeon Phi y Xeon multinúcleo (entre 5 y 6 cuando comparamos las GPUs con Xeon) para los diversos tamaños de salas considerados, mientras que en cuanto a su eficiencia energética, Xeon Phi ha obtenido un factor de mejora energética entre 35 y 40⁶ con respecto al Xeon, por un factor de mejora entre 30 y 35 cuando comparamos las GPUs con Xeon. Por lo tanto, concluimos que tanto Xeon Phi como CUDA (GPUs K40c y 980GTX) ofrecen beneficios claros para este kernel con limitaciones en los accesos a memoria, tanto desde el punto de vista de la aceleración del tiempo de ejecución como de su eficiencia energética.

Capítulo 5 Evaluación del modo offload en Xeon Phi para problemas de gran tamaño

5.1 Introducción

Hasta ahora nos hemos centrado en el modo nativo de ejecución del Intel Xeon Phi, desarrollando estrategías de paralelización y compilación para la ejecución directa de nuestros kernels Stencil en el coprocesador. Esta es la manera más versatil e inmediata de usar el Intel Xeon Phi, y propone una diferencia competitiva contra las GPUs discretas de Nvidia; líderes en el mercado de los aceleradores para desarrollo de propósito general. Es por ello que hemos centrado una parte importante de este documento a entender cómo mejorar el rendimiento en este modo, dedicando dos de los tres Capítulos principales a este propósito.

Sin embargo, hay situaciones en donde podría ser adecuado usar el copro- cesador como un acelerador. Por una parte, cuando necesitamos más potencia de cálculo para resolver un problema. Si usando un coprocesador Xeon Phi no tenemos bastante rendimiento, podemos intentar dividir el problema entre dos o más tarjetas aceleradoras para obtener una mayor aceleración (speedup). El éxito de esta estrategia dependerá de la cantidad de datos que se necesiten intercambiar entre las diversas tarjetas que colaboran en la resolución del problema. Al realizarse la comunicación a través del bus PCI-Express, una elevada transferencia de datos entre las tarjetas puede implicar un bajo (o incluso nulo) speedup. Por otra parte, el manejo de aplicaciones reales, en muchos casos, implica la necesidad de tratar con una ingente cantidad de datos. Para poder manejar esta gran cantidad de datos es fundamental escalar a varios dispositivos, a fin de disponer de más memoria y recursos computacionales.

Para poder manejar el coprocesador como un acelerador, Intel Xeon Phi presenta el modelo de programación de descarga (offload), donde la ejecución de una aplicación comienza en el procesador (CPU) anfitrión, encargado de lanzar las partes de código indicadas por el programador (normalmente intensivas en cómputo y basadas en paralelismo de datos) al acelerador Intel Xeon Phi para obtener una mejora del rendimiento (véase la Sección 2.3.1 para un mayor detalle). En este Capítulo se muestra cómo manejar Xeon Phi en el modo offload, tanto para mejorar el rendimiento de los códigos Stencil utilizados, como para poder abordar problemas de un tamaño mayor de lo que cabe en la memoria de una única tarjeta aceleradora.

5.1.1 Objetivos

Como hemos visto hasta ahora, cuando nuestros códigos Stencil se ejecutan en modo nativo en un única tarjeta aceleradora, obtenemos un rendimiento en torno a 80 GFLOPs para los códigos acústico y sísmico, y en torno a 140 GFLOPs para el código de calor. Nuestro primer objetivo de este Capítulo es intentar mejorar dicho rendimiento, usando varias tarjetas aceleradoras. Esta tarea no es trivial ya que, debido al patrón de cómputo que estamos utilizando, existen dependencias de datos que hacen difícil la división de los mismo en distintos espacios de memoria, forzando así una elevada comunicación de datos entre las tarjetas. En este Capítulo realizamos diversas pruebas utilizando únicamente nuestro Stencil 3-D de difusión de calor de 11 puntos. Específicamente, perseguimos los siguientes objetivos:

- Describir el funcionamiento del modo offload para el Intel Xeon Phi y apli- carlo a nuestro código Stencil.
- Evaluar el rendimiento de nuestro código descargándolo tan sólo a una tarjeta. De esta forma, mediremos el impacto de las comunicaciones de datos entre el procesador Xeon multinúcleo anfitrión y la tarjeta aceleradora.
- Estudiar la forma de dividir los datos entre varias tarjetas (en nuestro caso, hemos considerado dos tarjetas). Evaluar la repercusión que tienen las comunicaciones en el rendimiento obtenido y estudiar cómo mejorarlo retrasando las comunicaciones de los datos compartidos, lo cual impactará en la precisión de la solución encontrada.
- Estudiar diversas formas de manejo de las tarjetas aceleradoras por parte del procesador anfitrión, evaluando las ventajas e inconvenien- tes de cada una de ellas.
- Evaluar el rendimiento que se obtiene cuando se usan varias tarjetas (dos en nuestro caso), tanto para acelerar un tamaño de problema determinado, como para permitir la simulación de tamaños de problemas grandes.

5.1.2 Trabajo relacionado

El modo de programación offload para la arquitectura Intel Xeon Phi no ha sido muy estudiado en la comunidad científica. La mayoría de referencias son libros relacionados con el uso de Xeon Phi donde viene explicado este modo de programación [27, 73, 18]. En 52 el autor presenta en términos genéricos algunos pragmas de offload que están incluidos en OpenMP 4.0. Se apoya en C++/Fortran para mostrar pequeñas aplicaciones de descarga al coprocesador.

En 72 el autor presenta aplicaciones muy simples que utilizan el modo offload, poniendo el énfasis en la transferencia de datos hacia Xeon Phi. Presenta además un resumen de la gramática de offload.

En [57, 56] los autores muestran muchos ejemplos de aplicaciones utilizando Xeon Phi, y algunas de ellas ejecutándose en modo descarga. Como ejemplos, se muestra cómo ejecutar simulaciones de química cuántica usando NWChem en modo descarga, el algoritmo de Black-Scholes para mercados financieros, o cómo usar el módulo pyMIC para descargar al Xeon Phi aplicaciones escritas en el lenguaje Python.

Con respecto a otros trabajos de investigación, en 62 se presentan algunos métodos de Monte Carlo para el coprocesador Intel Xeon Phi, para calcular las integrales de Feynman en la física de alta energía. En este trabajo se presentan las aplicaciones ejecutadas en el coprocesador Intel Xeon Phi tanto en modo de offload como en modo nativo. Finalmente, en 8 se describen las modificaciones al código de dinámica molecular LAMMPS que permiten su ejecución simultánea en la CPU y Xeon Phi por medio de directivas offload, obteniendo una aceleración de hasta 4.7X.

5.2 Manejo de la descarga en un único coprocesador

Esta Sección describe cómo usar la arquitectura Intel Xeon Phi para que tra- baje con el modelo de “offload”, a fin de maximizar el rendimiento utilizando eficientemente todos los recursos disponibles. Todos los experimentos que se muestran en este Capítulo han sido realizados sobre nuestra plataforma habitual ya descrita en la Sección 3.5.1.1, que tiene un Intel Xeon multinúcleo junto con tres tarjetas coprocesadoras Intel Xeon Phi 7120P conectadas a través del bus PCIe. Es importante mencionar que los dispositivos MIC están numerados del 0 al número de dispositivos en el sistema menos 1. En nuestro caso, contamos con 3 dispositivos MIC instalados (mic0, mic1 y mic2). En esta Sección, para los experimentos realizados utilizamos únicamente un coprocesador, el mic2.

Como caso de uso, hemos utilizado el kernel de cálculo stencil 3-D de difusión de calor (11 puntos) con valores de punto flotante de doble precisión (8 bytes). Antes de empezar, resaltar dos detalles importantes. En primer lugar, en el modo descarga el sistema operativo que corre en Intel Xeon Phi utiliza exclusivamente un núcleo para manejar adecuadamente la comunicación con el procesador anfitrión, por lo que en todas nuestras evaluaciones usaremos un número de hilos igual a 240 (4 hilos por cada uno de los otros 60 núcleos restantes de la tarjeta). En segundo lugar, el manejo de las operaciones de descarga entre Xeon y Xeon Phi se puede realizar mediante el lenguaje MPI o por medio de las extensiones Intel LEO. En nuestro caso, este Capítulo se basa en la utilización de las primitivas Intel LEO. Una explicación detallada de estas primitivas se puede encontrar en el Apéndice A.

5.2.1 Pasos para descargar nuestro código en Intel Xeon Phi

Hemos identificado seis actividades principales necesarias para la descarga de nuestro código en el coprocesador Intel Xeon Phi, la Figura 5.8 muestra las actividades de offload para nuestro código Stencil. Las seis actividades se detallan a continuación:

1. Inicialización. En todas las pruebas de descarga habilitamos la variable de entorno OFFLOAD_INIT mediante “export OFFLOAD_INIT=on_start” para inicializar los dispositivos MIC del sistema, lo que permite ahorrar unos segundos en la ejecución de la aplicación. Los dispositivos MIC se inicializan al entrar al programa principal (main), antes de utilizarlos en las descargas.

En esta primera actividad de inicialización el sistema evalúa las condiciones para la descarga. El sistema operativo verifica, en tiempo de ejecución, la presencia de una librería de descarga a nivel de usuario llamada “infra- estructura Intel del coprocesador de descarga” (Intel Coprocessor Offload Infrastructure (COI)) y la disponibilidad de los coprocesadores Intel Xeon Phi en el sistema. Si se satisfacen todas las condiciones de descarga, el sis- tema crea los procesos necesarios en el coprocesador que se corresponden con el proceso del anfitrión (host).

2. Asignación de espacio de datos en la MIC. Dado que nuestro código Sten- cil 3-D utiliza dos matrices de tres dimensiones, una matriz principal (para lectura) y una matriz secundaria (para escritura), asignamos espacio en la MIC para ambas. Intel LEO proporciona diferentes mecanismos para asignar espacio de datos en la MIC, en nuestro caso experimentamos con dos estrategias:

a) Sólo asignar espacio en la MIC, sin transferir los datos de entrada. Para asignar sólo espacio sin transferir los datos (tanto de entrada como de salida) utilizamos la cláusula nocopy, a través del pragma mostrado en la Figura 5.1.

Abbildung in dieser Leseprobe nicht enthalten

Figura 5.1: Forma de asignar espacio en la MIC sin transferir datos de entrada.

En nuestra evaluación experimental utilizamos este mecanismo para asignar espacio en la MIC, ya que es más eficiente (no implica transfe- rencia de datos). Manejaremos la transferencia de datos posteriormente, de acuerdo a nuestras necesidades.

b) Asignar espacio en la MIC, transfiriendo los datos de manera implícita. Es importante mencionar que también se podría asignar espacio en la MIC transfiriendo los datos de entrada y/o salida de manera implícita, utilizando las cláusulas in y out (copia de entrada y copia de salida). El pragma para llevar a cabo dicha tarea sería el que se muestra en la Figura 5.2

Abbildung in dieser Leseprobe nicht enthalten

Figura 5.2: Forma de asignar espacio en la MIC transfiriendo datos de forma implícita.

3. Transferencia de datos (copia de datos de entrada a la memoria de la MIC). Cuando los datos son descargados al coprocesador, las transferencias de datos entre el coprocesador y el procesador tiene lugar, al menos al principio y al final del cálculo, pero por lo general también durante la ejecución. El comportamiento de las descargas de datos hacia y desde el coprocesador, en el inicio y el final de la descarga puede ser modificado mediante la inclusión de las cláusulas in, out y nocopy. Por lo tanto, hemos identificado tres formas de realizar esta actividad, detalladas a continuación:

a) Transferir la matriz principal y la matriz secundaria a la MIC.
b) Transferir sólo la matriz principal a la MIC, y en MIC copiar los datos de la matriz principal a la secundaria.
c) Transferir sólo la matriz principal a la MIC, y en MIC copiar sólo el contorno de la matriz principal a la secundaria.

De estas tres posibilidades, la más eficiente es la segunda: Copiar la matriz principal (fin) a MIC, y a la matriz secundaria (fout) sólo copiar el contorno de la matriz principal (fin) con memcpy. De aquí en adelante nos referiremos a la matriz principal o de lectura con el término "fin", y a la matriz secun- daria o de escritura como "fout". La Figura 5.3 muestra un ejemplo de la transferencia de datos a la MIC.

Una vez que los datos de la matriz fin están en la MIC, podremos copiar los datos de la matriz fin a la matriz fout, o bien copiar sólo el contorno de fin a fout. Una de las maneras que hemos experimentado para el copiado del contorno de la matriz fin a la matriz fout, es mediante una función ejecutada en MIC que se muestra en la Figura 5.4.

4. Ejecución del código descargado en la MIC. La ejecución del cálculo Stencil 3-D en Intel Xeon Phi en modo descarga la realizamos mediante la llamada a una función denominada stencil11pt_base. En la Figura 5.5 se muestra el

Abbildung in dieser Leseprobe nicht enthalten

Figura 5.3: Formas de transferencia de datos a la MIC.

Abbildung in dieser Leseprobe nicht enthalten

Figura 5.4: Función para copiar el contorno de la matriz fin a la matriz fout.

código que se ejecuta en el procesador anfitrión (Xeon multinúcleo) para realizar dicho cálculo en Intel Xeon Phi.

Abbildung in dieser Leseprobe nicht enthalten

Figura 5.5: Descarga y ejecución del Stencil 3-D en la MIC.

Una característica de nuestro código Stencil 3-D es que los datos en las matrices fin y fout contienen un contorno en la primera y última posición de cada dimensión que permanece constante durante todo el cálculo.

5. Copiado de los resultados de Xeon Phi a la memoria del host. Una vez rea- lizado el cálculo del Stencil 3-D en Intel Xeon Phi, necesitamos copiar los resultados de la matriz de salida (fout) de Xeon Phi al procesador anfitrión (host) mediante la cláusula out, tal como aparece en la Figura 5.6.

Abbildung in dieser Leseprobe nicht enthalten

Figura 5.6: Copiado de datos de Xeon Phi al host.

6. Liberación del espacio de datos asignados en Intel Xeon Phi. Para liberar la memoria de Intel Xeon Phi, utilizamos la macro FREE, que no asig- na memoria (alloc_if(0)) y si libera memoria (free_if(1)), tal como aparece en la Figura 5.7.

Abbildung in dieser Leseprobe nicht enthalten

Figura 5.7: Liberación de datos en Xeon Phi.

5.2.2 Evaluación del tiempo de la descarga

Para la evaluación experimental tomamos como base nuestro código nativo Stencil 3-D de difusión de calor (11 puntos) más optimizado (el que desarrollamos con la técnica de blocking (ver Sección 4.3 en el Capítulo 4). En cuanto a la versión con soporte para offload, la única diferencia con el código nativo es que utilizamos solamente 60 cores (240 hilos). Para ambos códigos utilizaremos una matriz de tamaño 800x400x600 elementos, estableciendo 4 hilos por núcleo y afinidad balanced, con tamaño de bloque en el eje X de 400 (width_TBlock=400), 4 (height_TBlock=4) en el eje Y y 4 (length_Tblock=4) en el eje Z.

La Figura 5.9 muestra los parámetros de compilación utilizados en la ver- sión con soporte para descarga, con las directivas para paralelizar, vectorizar y optimizar.

Una vez codificado el algoritmo llevamos a cabo un diagnóstico del proceso de descarga activando la variable de entorno OFFLOAD_REPORT, mediante “export OFFLOAD_REPORT=2”. El informe del diagnóstico de descarga de nuestro

Abbildung in dieser Leseprobe nicht enthalten

Figura 5.8: Actividades de offload del código Stencil 3-D de 11 puntos.

Abbildung in dieser Leseprobe nicht enthalten

Figura 5.9: Compilación en C del código Stencil 3-D 11 puntos.

código Stencil 3-D a la mic2 se muestra en la Figura 5.10. La información más destacable del informe es la siguiente:

1. Inicialización y asignación (Tag 0). El tiempo utilizado por el host es de 9.165445 segundos y el tiempo MIC es de 0.000169 segundos. Sin transfe- rencia de datos.
2. Transferencia de la matriz principal del host a Intel Xeon Phi (Tag 1). El tiempo invertido por el host es de 0.230212 segundos, mientras que el tiempo MIC es de 0.000023 segundos. Los datos transferidos de CPU a MIC son 1536000008 (bytes). La matriz transferida es de 800x400x600 valores de doble precisión (8 bytes).
3. Copiar matriz secundaria (Tag 2). El tiempo del host es de 0.024667 segundos y el tiempo MIC es de 0.022581 segundos. No hay transferencia de datos.
4. Cálculo del Stencil (Tag 3). El tiempo CPU es de 31.716863 segundos y el tiempo MIC es de 29.322435 segundos. No hay transferencia de datos.
5. Liberación de memoria de la MIC (Tag 4). El tiempo del host es de 0.224737 segundos y el tiempo MIC es de 0.000138 segundos. 1536000000 (bytes) de datos transferidos de MIC al host. El tamaño de la matriz es de 800x400x600 valores de doble precisión (8 bytes).

5.2.3 Comparación del rendimiento nativo vs offload

Para comparar el rendimiento de los códigos Stencil descargado en MIC y nativo necesitamos disponer en la MIC de las matrices principal (fin) y secundaria (fout). La matriz principal (fin) debe estar inicializada con los valores correspondientes al contorno y los valores centrales (matriz de lectura). La matriz secundaria (fout) debe contener los valores correspondientes al contorno y los valores centrales serán calculados en base a la lectura de la matriz principal (fin) y el cálculo propio del Stencil (matriz de escritura). La Tabla 5.1 resume los tiempos (en seg.) de las diversas actividades mencionadas en la sección anterior.

La Tabla 5.2 muestra el tiempo de ejecución (seg.) para el código nativo y el código descargado en el coprocesador Intel Xeon Phi. Como se puede apreciar, hemos agrupado las diversas actividades mencionadas en la Tabla 5.1.

De manera similar, la Tabla 5.3 muestra el rendimiento en GFLOPs para el código nativo y el código descargado en el coprocesador Xeon Phi. Podemos apreciar una diferencia mínima entre los GFLOPs obtenidos entre las dos versio- nes en relación al cálculo. Este resultado se debe, en parte, a que en modo nativo utilizamos 61 cores (244 hilos) y en modo offload utilizamos 60 cores (240 hilos).

Abbildung in dieser Leseprobe nicht enthalten

Figura 5.10: Informe de diagnóstico de offload del Stencil a mic2.

Como mencionamos anteriormente, el core 61 se reserva para manejar el sistema operativo. Además debemos contar con el tiempo adicional que conllevan las actividades de descarga en Xeon Phi. Así que en base a estos resultados podemos

Tabla 5.1: Tiempo de ejecución (seg.) del código Stencil 3-D descargado en MIC. Actividad

Abbildung in dieser Leseprobe nicht enthalten

Tabla 5.2: Tiempo (seg.) del código Stencil 3-D nativo y descargado en MIC.

Abbildung in dieser Leseprobe nicht enthalten

afirmar que la parte de cálculo apenas pierde rendimiento, pasando de 133.27 GFLOPs en nativo a 127.15 GFLOPs en modo offload (aproximadamente un 4.6% menos de rendimiento).

Tabla 5.3: Rendimiento (Gflops) del código Stencil 3-D nativo y descargado en MIC.

Abbildung in dieser Leseprobe nicht enthalten

Incluso si no tenemos en cuenta la fase de inicialización (asumiendo que el host puede realizar trabajo útil en modo asíncrono de ejecución (cláusula signal), siempre obtendremos mejores resultados en modo nativo. Sin embargo, el modo offload nos permite aumentar el rendimiento al utilizar múltiples tarjetas, o bien poder ejecutar entradas de tamaños que no caben en la memoria de una única tarjeta. Es importante destacar que se han utilizado de manera eficiente los cores de la tarjeta MIC. Hemos monitorizado la utilización de los cores a través de la interfaz GUI micsmc, lo que nos permite observar un histograma de utilización de los cores de la tarjeta MIC (Figura 5.11, 100 % de uso).

Figura 5.11: Vista del porcentaje de utilización de los cores en la tarjeta mic2.

5.2.4 Otros aspectos evaluados

Al igual que ya hicimos cuando trabajamos con Xeon Phi en modo nativo, también ahora queremos evaluar la influencia que algunos parámetros pudieran tener en el rendimiento que ofrece nuestro código en el modo descarga. Nos hemos centrado en la influencia de la afinidad, del número de iteraciones del algoritmo, y del tamaño de la matriz. Como mostramos a continuación, ninguno de estos aspectos tiene una influencia significativa.

- Afinidad. Hemos probado el efecto de la afinidad de los hilos para nues- tro código Stencil descargado en MIC, cambiando la variable de entorno KMP_AFFINITY a compact, balanced, o scatter. Esta variable se encarga de controlar la forma en que se asignan los hilos a los cores. A continuación mostramos como configuramos dicha variable de entorno:

Para afinidad scatter:

export MIC_ENV_PREFIX=MIC

export MIC_KMP_AFFINITY="granularity=thread,scatter"

Para afinidad balanced:

export MIC_ENV_PREFIX=MIC

export MIC_KMP_AFFINITY="granularity=thread,balanced"

Para afinidad compact:

export MIC_ENV_PREFIX=MIC

export MIC_KMP_AFFINITY="granularity=thread,compact"

La Tabla 5.4 muestra los resultados de rendimiento obtenidos cuando variamos la política de afinidad entre las tres opciones. Podemos apreciar una mínima mejora en cuanto al rendimiento con balanced, motivo por el cual la utilizaremos por defecto en posteriores evaluaciones.

Tabla 5.4: Evaluación de afinidad del cálculo del Stencil 3-D descargado en MIC.

Abbildung in dieser Leseprobe nicht enthalten

- Número de iteraciones. En cuanto al número de iteraciones totales de eje- cución del kernel del Stencil 3-D, evaluamos nuestro código con 500, 1000, 2000 y 5000 iteraciones.

La Tabla 5.5 muestra los resultados obtenidos del cálculo del Stencil 3-D cuando variamos el número total de iteraciones. En base a estos resultados utilizaremos 1000 iteraciones para el resto de evaluaciones de rendimiento, como compromiso entre rendimiento y tiempo total de ejecución.

- Tamaño de matriz. Otro aspecto importante en la evaluación de algoritmos tipo Stencil es la variabilidad de rendimiento respecto a la entrada. La tabla 5.6 muestra los resultados obtenidos al variar el tamaño de la matriz de entrada, mostrando poca variación en el rendimiento. Para el resto de nuestra evaluaciones utilizaremos un tamaño de matriz de 800x400x600.

Tabla 5.5: Evaluación del número de iteraciones del cálculo del Stencil 3-D descar- gado en MIC.

Abbildung in dieser Leseprobe nicht enthalten

Tabla 5.6: Evaluación del tamaño de matriz del Stencil 3-D descargado en MIC.

Abbildung in dieser Leseprobe nicht enthalten

5.3 Manejo de la descarga en dos coprocesadores

Esta Sección evalúa diferentes estrategias para descargar nuestro código Stencil 3-D a Xeon Phi en dos tarjetas (mic1 y mic2), para una entrada de tamaño inferior a la que cabe en la memoria de una ellas.

Lo primero que tenemos que decir es que al partir los datos de entrada del algoritmo Stencil 3D en varias tarjetas, se nos crean unas necesidades de comunicación entre las mismas por cada iteración realizada. Debido a que el algoritmo utiliza los elementos vecinos en las tres dimensiones (11 puntos) para el cálculo de los elementos de una iteración (iteración i), cuando el algoritmo se ejecuta en una única tarjeta todos los datos residen en su memoria, por lo que no necesitamos realizar ningún trato especial de los mismos.

Decidimos dividir la entrada en distintas submatrices partiéndolas por el eje Z, y transfiriendo cada submatriz a una tarjeta. Al partir los datos entre diferentes tarjetas, cada tarjeta necesitará para el cálculo de los planos extremos, en cada iteración, los valores de la iteración anterior que tenían dichos planos que se habían calculado en la otra tarjeta (ver Figura 5.12). Dichos planos extremos o planos frontera tendrán que ser intercomunicados entre las tarjetas a través del procesador anfitrión. En una división del trabajo entre N tarjetas, cada tarjeta necesita enviar y recibir dos planos X-Y a otras dos tarjetas.

Un adecuado manejo de esta comunicación es de vital importancia, pues si el

Abbildung in dieser Leseprobe nicht enthalten

Figura 5.12: Paralelización del Stencil 3-D sobre N tarjetas MIC.

tiempo en comunicar las fronteras es muy alto en relación al tiempo de cómputo del algoritmo, el rendimiento de la implementación que usa varias tarjetas Xeon Phi puede verse comprometido, y podría no beneficiarse en nada la descarga de nuestro algoritmo a varias tarjetas.

En esta sección vamos a ver el esquema básico que hemos seguido para comunicar estos planos frontera entre dos tarjetas, y cómo podríamos reducir el tamaño de los datos que comunicamos entre las dos tarjetas a costa de perder un poco de precisión en los resultados obtenidos por el algoritmo.

5.3.1 Esquema básico para dos tarjetas

Partimos de un tamaño de la matriz de entrada de 800x400x600, con la intención de repartir los datos para el cálculo del Stencil 3-D en las dos tarjetas. Dividimos la entrada en dos submatrices de tamaño 800x400x300, partiéndolas por el eje Z, y transfiriendo cada submatriz a una tarjeta. La Figura 5.13 muestra la transferencia de datos a las dos tarjetas.

Abbildung in dieser Leseprobe nicht enthalten

Figura 5.13: Transfiriendo una mitad de la matriz a mic1 y la otra mitad a mic2.

Es importante recordar que nuestros datos de entrada para el cálculo del Stencil 3-D están en una matriz de tres dimensiones, y que el inicio y el final de cada dimensión tiene un valor considerado como contorno (borde). Dicho contorno global permanece constante durante todo el proceso de cálculo (es decir, el contorno no se modifica).

Otra cuestión distinta es el contorno que le corresponde a cada tarjeta. Al partir los datos por el eje Z, los planos X-Y extremos (o bordes) de cada tarjeta si que son modificados como se puede ver en la Figura 5.12. En este caso de tan sólo dos tarjetas, uno de los bordes coincide con el borde global de la matriz (constante entre iteraciones), por lo que únicamente hay que enviar y recibir el plano inmediatamente superior e inferior a la mitad del eje Z.

El tener una mitad de los datos en la mic1 y la otra mitad en la mic2 nos obliga a comunicar los planos (length/2)-1 y (length/2)+1 entre ambas tarjetas para realizar el cálculo del Stencil 3-D en cada iteración del algoritmo. Dicha comunicación de datos no se puede realizar directamente y se debe realizar por medio del procesador anfitrión.

La Figura 5.14 muestra la forma en que se comunican las dos tarjetas paraactualizar los planos (length/2)-1 y (length/2)+1, los cuales se establecen como contornos de la dimensión del eje Z de las submatrices.

Abbildung in dieser Leseprobe nicht enthalten

Figura 5.14: Comunicación entre dos MICs para el cálculo del Stencil 3-D.

Vamos a explicar con más detalle la forma como se utiliza el anfitrión (host) y las dos MICs para realizar la transferencia de datos necesaria para realizar los cálculos del código Stencil 3-D:

- El anfitrión reparte los datos (las dos submatrices) para el cálculo del Stencil 3-D en dos tarjetas (mic1 y mic2), estableciendo un mecanismo de actualización de los planos length/2-1 y length/2+1 en cada iteración del algoritmo, recibiendo de la tarjeta mic1 el nuevo valor del plano length/2-1 y enviándolo a la tarjeta mic2, y lo mismo con el plano length/2+1 entre la tarjeta mic2 y la mic1.

- En mic1, una vez recibida la mitad de los datos totales (en este caso 800x400x300), se actualiza el plano (length/2)-1 (sobre el eje Z), copian- do dichos valores de una variable auxiliar almacenada en new_border2. La primera vez obtendremos los valores iniciales, pero para subsecuentes actua- lizaciones esta variable contendrá los valores recibidos de mic2 (enviados a través del anfitrión). A continuación se lleva a cabo el cálculo del código del Stencil, resultando actualizado el plano (length/2)-1. Dicho plano se copia en memoria a la variable auxiliar almacenada en (new_border1), la cual se transferirá a la tarjeta mic2 a través del procesador anfitrión.

- En mic2 se realiza un proceso similar al de mic1, salvo que tomamos como entrada para el plano (length/2)+1 los valores de new_border1 y copiamos a la salida los nuevos valores del borde a new_border2.

5.3.2 Reducción del tamaño de la transferencia de datos

El algoritmo Stencil 3-D analizado presenta una cantidad de transferencia de datos que podría ser muy elevada, dependiendo del número de tarjetas utilizadas y del tamaño de los datos de entrada del código.

En este apartado vamos a estudiar cómo podemos reducir el tamaño de la transferencia de datos. Para ello, proponemos reducir el número de veces que se transfieren dichos datos. Aunque el algoritmo obliga a comunicar a cada iteración que se realiza, nos planteamos reducir la comunicación de los bordes intermedios al realizar los cómputos de manera independiente durante algunas iteraciones en los dos coprocesadores. Está claro que los resultados no serán todo lo correctos que sería deseable, pero podría ser posible que para un determinado problema este modo de trabajar fuera asumible. La primera pregunta que nos hacemos es: ¿Es admisible esta implementación?

En la Tabla 5.7 se muestra el margen de error en los datos de salida de nuestro kernel Stencil 3D de difusión de calor de 11 puntos cuando no comunicamos nunca dichos planos X-Y, es decir, hemos comparado la salida del código ejecutado en una única tarjeta contra la salida del código ejecutado en dos tarjetas (cada una con la mitad de los datos) sin ninguna comunicación. Los resultados obtenidos los hemos evaluado para diferentes valores del contorno.

Como se puede apreciar en la Tabla 5.7, el margen de error ( % error) para los casos que probamos va desde un 44.41 % a un 100 %, dependiendo de los valores del contorno (Centro, Bordes izquierdo, derecho, superior, inferior, atrás y frente). Cabe destacar que cuando el valor del centro (por ejemplo C=10) es más cercano a los valores de los bordes, el margen de error es menor. Para el caso del centro con valor de 5 (C=5) y todos los bordes con valor de 150, encontramos un 100.0 % de error. En base a los resultados obtenidos, consideramos que lógicamente esta estrategia no es admisible, ya que el margen de error es muy alto.

Por tanto, ahora nos planteamos encontrar el margen de error al establecer comunicación entre las dos MICs, actualizando cada 1, 2, 5 y 10 iteraciones de

Tabla 5.7: Evaluación del error al ejecutar el código en dos MICs sin comunicación para diferentes valores del contorno.

Abbildung in dieser Leseprobe nicht enthalten

(*) C=Centro, I=Borde Izq., D=Borde Der., S=Borde Sup., In=Borde Inf., A=Borde Atr., F=Borde Fre.

la ejecución del algoritmo. Esto nos va a permitir encontrar un margen de error aceptable.

Tabla 5.8: Evaluación del error al ejecutar el código en dos MICs actualizando parcialmente para diferentes valores del contorno.

Abbildung in dieser Leseprobe nicht enthalten

(*) C=Centro, I=Borde Izq., D=Borde Der., S=Borde Sup., In=Borde Inf., A=Borde Atr., F=Borde Fre.

La tabla 5.8 muestra los resultados de la evaluación del error que tendría- mos con diferentes valores de contorno (bordes) en el cálculo del Stencil 3-D descargado en dos mics, actualizando cada 1, 2, 5 y 10 iteraciones (en base a 1000 iteraciones) los planos (length/2)-1 y (length/2)+1. Destacamos que el error que depende en gran medida de los valores del contorno. Con los valores de contorno que experimentamos encontramos un margen de error de 0.0% si ac- tualizamos cada 1 iteración (caso base, actualizar en cada iteración), de 0.15% a 0.71% de error si actualizamos cada 2 iteraciones, de 0.60% a 2.91% de error si actualizamos cada 5 iteraciones y de 1.43% a 6.73% de error si actualizamos cada 10 iteraciones. Estos resultados son muy satisfactorios, pues nos muestran que relajando un poco las comunicaciones entre las tarjetas, el error obtenido es mínimo y pensamos que puede ser asumido cuando estemos ejecutando una aplicación real.

5.3.3 Manejo de la descarga por parte del procesador anfitrión

Hasta ahora hemos visto cómo cada una de las dos tarjetas recibe los datos que le corresponden, ejecuta el código Stencil, y en función de la frecuencia de comunicación, la precisión que obtenemos en los resultados de salida puede variar. En esta Sección nos vamos a fijar en cómo se maneja por parte del procesador anfitrión la descarga de los datos y el código a las dos tarjetas aceleradoras. Hemos identificado 3 opciones, dependiendo de si tenemos en el anfitrión 2 hilos de ejecución para manejar a cada tarjeta, o bien un único hilo que las controle de forma asíncrona.

5.3.3.1 Uso de 2 hilos dentro del bucle global de cálculo

Una primera estrategia para la versión offload de nuestro código Stencil 3-D con- siste en colocar dentro del bucle global de cálculo (el bucle for correspondiente al número de iteraciones a realizar para que el algoritmo Stencil converja) una región paralela con dos hilos (#pragma omp parallel), de tal manera que el hilo 0 lleve a cabo el manejo del cálculo en una tarjeta (la mic1), y el hilo 1 lleve a cabo el manejo del cálculo en la otra (la mic2).

En la Figura 5.15) se puede ver el esquema básico de esta opción, y a conti- nuación describimos el manejo de cada hilo con cada tarjeta:

- Hilo 0 con mic1:

- Descarga los datos (la mitad de los datos totales) a la mic1 (#pragma offload target (mic1)).
- Actualiza el plano (length/2)-1 (sobre el eje Z) copiando los valores almacenados en new_border2, que habían sido actualizados con los cálculos parciales recibidos de mic2.
- Lanza la ejecución en mic1 del cálculo del Stencil con la actualización del plano (length/2)-1.
- Recibe y copia en memoria en la (new_border1) los datos del plano que se utilizarán como entrada en mic2.

Abbildung in dieser Leseprobe nicht enthalten

Figura 5.15: Paralelización del host Xeon con 2 hilos en el bucle global.

- Hilo 1 con mic2:

- Descarga los datos (la mitad de los datos totales) a la mic2 (#pragma
offload target (mic2)).
- Actualiza el plano (length/2)+1 (sobre el eje Z) copiando los valores almacenados en new_border1, que habían sido actualizados con los cálculos parciales recibidos de mic1.
- Lanza la ejecución en mic2 del cálculo del Stencil con la actualización del plano (length/2)+1.
- Recibe y copia en memoria en la (new_border2) los datos del plano que se utilizarán como entrada en mic1.

5.3.3.2 Uso de 2 hilos con región paralela y barrera explícita

Una segunda estrategia para descargar y manejar nuestro código Stencil 3-D en las dos tarjetas (mic1 y mic2) consiste en definir una región paralela con dos hilos (#pragma omp parallel) y una barrera explícita, de tal forma que cada hilo maneje las iteraciones de la tarjeta que le corresponda, y usemos la barrera para sincronizar los hilos en cada iteración del bucle (o en el número de iteraciones que nos parezca oportuno). En la Figura 5.16 se detalla de forma esquemática la estrategia descrita, y a continuación la describimos:

Abbildung in dieser Leseprobe nicht enthalten

Figura 5.16: Paralelización con 2 hilos y una barrera explícita en el host Xeon.

- Hilo 0 con mic1:

- Inicia un bucle de N iteraciones para realizar los cálculos del Stencil (for loop), descargando lo primero los datos (la mitad de los datos totales) en mic1 por medio de la directiva (#pragma offload target (mic1)).
- Actualiza el plano (length/2)-1 (sobre el eje Z) copiando los valores almacenados en new_border2, que habían sido actualizados con los cálculos parciales recibidos de mic2.
- Lanza la ejecución en mic1 del cálculo del Stencil con la actualización del plano (length/2)-1.
- Recibe y copia en memoria en la (new_border1) los datos del plano que se utilizarán como entrada en mic2.
- Sincroniza con el hilo 1 por medio de una barrera explícita para sincronizar las iteraciones parciales de mic1 con las iteraciones parciales de mic2, continuando posteriormente el bucle de N iteraciones.

- Hilo 1 con mic2:

- Inicia un bucle de N iteraciones para realizar los cálculos del Stencil (for loop), descargando lo primero los datos (la mitad de los datos totales) en mic2 por medio de la directiva (#pragma offload target (mic2)).
- Actualiza el plano (length/2)+1 (sobre el eje Z) copiando los valores almacenados en new_border1, que habían sido actualizados con los cálculos parciales recibidos de mic1.
- Lanza la ejecución en mic1 del cálculo del Stencil con la actualización del plano (length/2)+1.
- Recibe y copia en memoria en la (new_border2) los datos del plano que se utilizarán como entrada en mic1.
- Sincroniza con el hilo 0 por medio de una barrera explícita para sincronizar las iteraciones parciales de mic2 con las iteraciones parciales de mic1, continuando posteriormente el bucle de N iteraciones.

5.3.3.3 Uso de un hilo y manejo de las tarjetas de forma asíncrona

La tercera y última estrategia que hemos desarrollado para descargar y manejar nuestro código Stencil 3-D en las dos tarjetas (mic1 y mic2) consiste en definir un bucle global de cálculo (for loop) dentro del cual se lleve a cabo la descarga y el cálculo en cada tarjeta de manera asíncrona.

Para implementar esta estrategia hemos utilizado la descarga asíncrona me- diante el uso de las cláusulas signal y wait. Cuando añadimos la cláusula signal a una directiva de descarga, la sección de descarga de nuestro código Stencil 3-D ocurre de forma asíncrona. Esto nos permite el uso simultáneo (concurrente) del procesador y del coprocesador. Hay que añadir la cláusula wait para esperar a la finalización del cálculo antes de iniciar otra iteración del bucle. En la Figura 5.17 se muestra esta estrategia. A continuación se explica con más detalle los pasos que ejecuta el hilo para el manejo de las dos tarjetas:

- Se descargan los datos (la mitad de los datos totales) a la tarjeta mic1 de for- ma asíncrona (#pragma offload target (mic1)), con la cláusula signal.
- Al ser la transferencia asíncrona, el hilo continúa con la actualización del plano (length/2)-1 (sobre el eje Z) copiando los valores almacenados en new_border2, que habían sido actualizados con los cálculos parciales recibidos de mic2.
- Continúa lanzando la ejecución en mic1 del cálculo del Stencil con la actua- lización del plano (length/2)-1.
- Recibe y copia en memoria en la (new_border1) los datos del plano que se utilizarán como entrada en mic2.
- Se descargan los datos (la mitad de los datos totales) a la tarjeta mic2 de for- ma asíncrona (#pragma offload target (mic2)), con la cláusula signal.
- Al ser la transferencia asíncrona, el hilo continúa con la actualización del plano (length/2)+1 (sobre el eje Z) copiando los valores almacenados en new_border1, que habían sido actualizados con los cálculos parciales recibidos de mic1.
- Continúa lanzando la ejecución en mic2 del cálculo del Stencil con la actua- lización del plano (length/2)+1.
- Recibe y copia en memoria en la (new_border2) los datos del plano que se utilizarán como entrada en mic1.
- La cláusula signal nos permite continuar la actividad entre el hilo y el procesador anfitrión de forma concurrente, y al llegar a dicha cláusula nos sirve de punto de sincronización entre las dos tarjetas.

Es importante destacar que esta estrategia se puede aplicar con entradas de datos más grandes, y haciendo uso de más tarjetas MIC, si nuestro sistema contara con ellas. También se podría asignar parte del cálculo al procesador anfitrión, siempre y cuando tengamos en cuenta que su rendimiento será menor al logrado por las tarjetas MIC.

Abbildung in dieser Leseprobe nicht enthalten

Figura 5.17: Paralelización con un hilo de las dos tarjetas de forma asíncrona por parte del host Xeon.

5.4 Evaluación

5.4.1 Entorno de evaluación

Los resultados que mostramos en esta Sección corresponden a la evaluación experimental de las estrategias mencionadas en las Secciones 5.3.3.1, 5.3.3.2 y 5.3.3.3.

El entorno de evaluación hardware para Intel Xeon Phi es el mismo que fue descrito en la sección 3.5.1.1 del Capítulo 3, mientras que el entorno software fue descrito en la sección 3.5.1.2 de dicho Capítulo (aquí únicamente usamos el kernel de difusión de calor). En los resultados que se muestran a continuación tan sólo subrayar que hemos establecido el contorno (bordes) para todos nuestros experimentos a los siguientes valores (en grados centígrados): centro en 10, borde izquierdo en 150, borde derecho en 150, borde superior en 70, borde inferior en 70, borde anterior en 70 y borde posterior en 70.

Por último, comentar que los resultados mostrados en esta Sección siguen como referencia las 5 actividades mencionadas en la Sección 5.2.2, pero evaluando por separado la inicialización (actividad (1)).

5.4.2 Evaluación de la inicialización

En primer lugar vamos a evaluar la inicialización de las tarjetas, ya que ésta es una tarea común a todas las estrategias. En este primer paso de inicialización, el sistema evalúa las condiciones para la descarga (descrito en la Sección 5.2.1). Para reducir esta sobrecarga al máximo, hemos modificado la variable de entorno OFFLOAD_INIT para que nada más comience la ejecución en el anfitrión se inicia- licen las tarjetas y se reduzca dicho tiempo de activación (hemos ejecutado el comando export OFFLOAD_INIT=on_start en el anfitrión.

Además de la inicialización, hay otros dos aspectos que hemos considerado: el tiempo consumido en la transferencia de los datos de entrada del anfitrión a la mic1 y la mic2, así como el tiempo de copia de los datos de salida de los coprocesadores al anfitrión.

La Tabla 5.9 muestra los resultados que corresponden al tiempo de inicializa- ción y transferencia para todas nuestras estrategias de evaluación, tanto por el anfitrión como por el coprocesador MIC.

Tabla 5.9: Tiempo utilizado por la inicialización.

Abbildung in dieser Leseprobe nicht enthalten

(*) 768000008 (bytes) a cada MIC.

Los valores de tiempo mostrados en la Tabla 5.9 se obtuvieron mediante un informe de diagnóstico de la descarga de nuestro código Stencil 3-D, y corres- ponden a la suma de los tiempos consumidos por las dos tarjetas (mic1 y mic2). Los datos transferidos (768000008 bytes) corresponden a dos submatrices de 800x400x300. El informe de diagnóstico lo obtuvimos activando la variable de entorno “export OFFLOAD_REPORT=2”, y se muestra en la Figura 5.18.

Abbildung in dieser Leseprobe nicht enthalten

Figura 5.18: Informe de diagnóstico de offload del Stencil a dos tarjetas (mic1 y mic2.

5.4.3 Evaluación del uso de 2 hilos dentro del bucle global de cálculo

Esta Sección evalúa la codificación del Stencil 3-D al colocar dentro de un bucle global de cálculo una región paralela con dos hilos, tal como se detalla en la Sección 5.3.3.1. La Tabla 5.10 y la Tabla 5.11 muestran los resultados de nuestra evaluación experimental en rendimiento (GFLOPs) y tiempo (segundos) respec- tivamente, para diversos escenarios de comunicación de los planos extremos o bordes (véase la Sección 5.3.2).

Tabla 5.10: Evaluación del rendimiento con el uso de 2 hilos dentro del bucle global de cálculo.

Abbildung in dieser Leseprobe nicht enthalten

(*) Porcentaje de error

Tabla 5.11: Evaluación del tiempo con el uso de 2 hilos dentro del bucle global de cálculo.

Abbildung in dieser Leseprobe nicht enthalten

(*) Porcentaje de error

Si consideramos que el cálculo del rendimiento del Stencil 3-D implica trans- ferir la matriz principal, copiado de la matriz secundaria, cálculo del Stencil 3-D y liberación de memoria, con un 0.0% de error, obtenemos 180.44 GFLOPs. Si permitimos un error de 6.73 %, entonces el rendimiento crece a 216.90 GFLOPs. Es decir, si nos podemos permitir un margen de error del 6.73 % nuestra aplicación se ejecutará un 16 % más rápido. Este mismo porcentaje de aceleración (16 %) se obtiene en cada una de las demás actividades evaluadas.

En la Figura 5.19 se puede apreciar la utilización de los núcleos de las tarjetas, tanto en mic1 como en mic2, usando la interfaz GUI micsmc. Cabe destacar que en mic2 la utilización no llega al 100 %, debido a que ambas tarjetas (mic1 y mic2) deben sincronizarse para transferir los bordes, motivo por el cual los hilos de la tarjeta mic2 esperan en cierto momento a que terminen los hilos de la tarjeta mic1.

Figura 5.19: Monitorización de la utilización de los cores al paralelizar con 2 hilos dentro de un ciclo global de cálculo.

Abbildung in dieser Leseprobe nicht enthalten

5.4.4 Evaluación del uso de 2 hilos con región paralela y barrera explícita

En esta implementación del Stencil 3-D definimos una región paralela con dos hilos y una barrera explícita detallada en la Sección 5.3.3.2. La Tabla 5.12 y la Tabla 5.13 muestran los resultados de nuestra evaluación experimental en rendi- miento (GFLOPs) y en tiempo (seg.) respectivamente, para diversos escenarios de comunicación de los planos extremos o bordes (véase la Sección 5.3.2).

Tabla 5.12: Evaluación del rendimiento con el uso de dos hilos y una barrera explícita.

Abbildung in dieser Leseprobe nicht enthalten

Tabla 5.13: Evaluación del tiempo con el uso de 2 hilos y una barrera explícita. Actividad Tiempo (seg.)

Abbildung in dieser Leseprobe nicht enthalten

(*) Porcentaje de error

Abbildung in dieser Leseprobe nicht enthalten

Figura 5.20: Monitorización de la utilización de los cores al paralelizar con 2 hilos y una barrera explícita.

Consideramos también en esta evaluación las actividades que agrupa el cálcu- lo del Stencil: copiado de matriz secundaria, transferencia de datos y liberación de memoria (Actividades Cálculo+MatrizSec.+Tranf.+Lib.(4,3,2 y 5) para el cálculo del rendimiento del Stencil 3-D.

Si no queremos asumir ningún error, obtenemos un rendimiento aproximado de 182.16 GFLOPs. Asumiendo un error del 6.73% obtenemos un rendimiento de 218.87 GFLOPs. Si miramos la utilización de los núcleos, en la Figura 5.21 se nos muestra de nuevo que la mic2 está infrautilizada. El motivo para esta infrautilización vuelve a ser la sincronización entre ambas tarjetas.

5.4.5 Evaluación al utilizar un hilo y tarjetas de forma asíncrona

Finalmente mostramos resultados del cálculo del Stencil 3-D al utilizar un hilo con las dos tarjetas (mic1 y mic2) de forma asíncrona, tal como se detalla en la Sección 5.3.3.3. La Tabla 5.14 y la Tabla 5.15 muestran los resultados de nuestra evaluación experimental en rendimiento (GFLOPs) y en tiempo (segundos) respectivamente, para diversos escenarios de comunicación de los planos extremos o bordes (véase la Sección 5.3.2).

Tabla 5.14: Evaluación del rendimiento usando un hilo y tarjetas de forma asín- crona.

Abbildung in dieser Leseprobe nicht enthalten

(*) Porcentaje de error

Tabla 5.15: Evaluación del tiempo usando un hilo y tarjetas de forma asíncrona.

Abbildung in dieser Leseprobe nicht enthalten

(*) Porcentaje de error

El rendimiento medio para esta estrategia ronda los 191.02 GFLOPs para un error del 0.0 %. Si permitimos un error del 6.73 % obtenemos un rendimiento de 235.06 GFLOPs (mejora del 18.7 %). Al igual que en las estrategias anteriores, la interfaz GUI micsmc nos permite monitorear los núcleos, tanto en mic1 como en mic2. En la Figura 5.21 se puede apreciar que la utilización llega al 100 %. Esto se ha logrado al utilizar ambas tarjetas en forma asíncrona.

Esta tercera estrategia ha resultado ser la más eficiente con un 6% mas de rendimiento que las otras estrategias anteriores.

Abbildung in dieser Leseprobe nicht enthalten

Figura 5.21: Monitorización de los núcleos usando un hilos y dos tarjetas de forma asíncrona.

5.4.6 Comparativa de las tres estrategias evaluadas en la descarga del Stencil 3-D en dos MICs.

Esta Sección agrupa las tres estrategias evaluadas en las Secciones 5.4.3, 5.4.4 y 5.4.5, identificadas como: estrategia (a), paralelizando con 2 hilos dentro del bucle global de cálculo; estrategia (b), paralelizando con 2 hilos y una barrera explícita; estrategia (c), usando un hilo con dos tarjetas de forma asíncrona.

La Tabla 5.16 y la Tabla 5.17 muestran los resultados comparativos de las tres estrategias de nuestra evaluación experimental en rendimiento (GFLOPs) y en tiempo (seg.) respectivamente, para diversos escenarios de comunicación de los planos extremos o bordes (véase la Sección 5.3.2). Hemos comparado las tres estrategias sobre la actividad que agrupa el cálculo del Stencil 3-D, copiado de matriz secundaria, transferencia de datos y liberación (Actividades Cálculo+MatrizSec.+Tranf.+Lib.(4,3,2 y 5) de nuestra clasificación en la Sección 5.2.1).

La estrategia (c) es la que mayor rendimiento logra, mejorando en un 6% a las demás. En base a estos resultados, consideramos que la estrategia (c) es la forma más adecuada de realizar trabajo de cómputo de manera concurrente en Xeon Phi, en nuestro caso utilizando dos tarjetas (mic1 y mic2) para la evaluación experimental de nuestro kernel de Stencil 3-D.

Tabla 5.16: Comparativa de rendimiento (GFLOPs) de las tres estrategias de descarga.

Abbildung in dieser Leseprobe nicht enthalten

(*) Porcentaje de error

Tabla 5.17: Comparativa de tiempo (seg.) de las tres estrategias de descarga.

Abbildung in dieser Leseprobe nicht enthalten

(*) Porcentaje de error

En la Figura 5.22 se muestra gráficamente el rendimiento en GFLOPs para cada una de las estrategias evaluadas en función del porcentaje de error. Se puede apreciar que con un margen de error de 6.73 %, logramos acelerar el rendimiento en un 16% en relación al rendimiento obtenido con un margen de error del 0%. También vemos el comportamiento del rendimiento en relación a la actividad eva- lúe (Calc.(4), Calc.+MatrizSec.+Tranf.(4,3 y2) ó Calc.+MatrizSec.+Tranf.+Lib.(4,3,2 y 5), con una aceleración mayor en la actividad que implica sólo el cálculo.

En la Figura 5.23 se muestra gráficamente el rendimiento en GFLOPs para cada una de las estrategias evaluadas en función de la actividad, al variar el margen de error para cada actividad (0.00 %, 0.71 %, 2.91 % o 6.73 %). Como se puede apreciar, si tenemos en cuenta todas las actividades (Calc.+MatrizSec.+Tranf.+Lib.(4,3,2 y 5)) el rendimiento alcanzado decrece levemente en relación al rendimiento obtenido por la actividad de cálculo (Calc.)

La Figura 5.24 muestra la comparación del rendimiento alcanzado en el modo nativo con la estrategia (c). Dicha estrategia muestra el rendimiento alcanzado con un margen del 0.00% y del 6.73% de error. La mejora alcanzada en la descarga en dos MICs, con un margen de error del 0.00% es de un 50%, mientras que con un margen de error del 6.73 % es de 80 %. También se puede apreciar que si

Abbildung in dieser Leseprobe nicht enthalten

Figura 5.22: Rendimiento en GFLOPs para cada estrategia en función del error (el % de error se presenta en el eje principal (eje X) de la gráfica).

utilizamos una única tarjeta, el rendimiento decrece en un 4.5% en relación al modo nativo.

5.4.7 Evaluación de las tres estrategias con tamaños grandes de problema.

Una de las ventajas del uso del modelo de offload es que podemos dividir el problema entre varias tarjetas cuando una de ellas no tiene memoria suficiente para almacenar los datos de entrada. Para analizar esta hipótesis, realizamos experimentos con matrices grandes (1000x800x1000), con los resultados que se muestran en la Tabla 5.18. Como se puede apreciar el comportamiento del rendimiento es parecido al mostrado en la sección 5.4.6.

5.5 Conclusiones

En este Capítulo hemos explorado y experimentado con el modelo de programa- ción offload (que ejecuta un programa en el host y descarga parte del trabajo al

Abbildung in dieser Leseprobe nicht enthalten

Figura 5.23: Rendimiento en GFLOPs para cada estrategia en función de la actividad analizada.

Tabla 5.18: Evaluación del rendimiento del cálculo del Stencil 3-D descargado en dos MICs, para matrices que no caben en una MIC.

Abbildung in dieser Leseprobe nicht enthalten

(*) Porcentaje de error

coprocesador) para el coprocesador Intel Xeon Phi, tomando como punto de par- tida el modelo de programación nativo (las aplicaciones se ejecutan directamente en el coprocesador).

Específicamente, en este Capítulo hemos cubierto los siguientes aspectos:

- Uno de los puntos que consideramos importantes es comprender el modelo de programación de offload, una vez comprendido lo aplicamos a nuestro código Stencil. Hemos tomado en cuenta la sobrecarga asociada al movi- miento de datos hacia y desde el coprocesador Intel Xeon Phi, con todo lo

Abbildung in dieser Leseprobe nicht enthalten

Figura 5.24: Comparación de la estrategia c) con la ejecución en modo nativo.

que esto implica: el movimiento de datos, la transferencia y la forma de gestionar la descarga de los datos.

- Tomando como punto de partida nuestro código nativo, en nuestras pri- meras evaluaciones experimentales al descargar nuestro kernel Stencil 3-D utilizamos una sola tarjeta. Para el manejo de las operaciones de descarga utilizamos las extensiones de Intel LEO. Esto nos permitió identificar las actividades principales necesarias para la descarga de nuestro código en el coprocesador Intel Xeon Phi. Apreciamos una degradación del rendimiento respecto a la ejecución de nuestro kernel en modo nativo. Esto se debe principalmente al sobrecoste asociado que tiene el modelo de offload. - Recomendamos que siempre que sea posible se utilice el modelo nativo. Por el contrario, si disponemos de varios coprocesadores o si nuestros datos no caben en una tarjeta MIC, recomendamos utilizar el modelo de offload.

- Estudiamos la forma de dividir los datos en dos tarjetas y evaluamos la repercusión en el rendimiento obtenido cuando se comunican nuestras dos tarjetas. Mostramos que reduciendo la exactitud entre un 0.71 % a un 6.73 %, alcanzamos una aceleración de un 20% mas, que si no nos permitiéramos ningún margen de error.

- En nuestras evaluaciones experimentales mediante la utilización de dos tarjetas (mic1 y mic2) evaluamos tres estrategias de solución para nuestro kernel Stencil 3-D y mostramos de manera detallada en que consiste cada es- trategia, así como los resultados de la evaluación experimental, presentando en este Capítulo los resultados más relevantes.

Consideramos que el modelo de offload plantea nuevos retos que, aunque incrementan la complejidad de la programación, permiten acceder a más recursos, tanto computaciones como de memoria. Estamos seguros que este modelo puede presentar amplios beneficios en la resolución de modelos complejos que manejen ingentes cantidades de datos, y que estos beneficios se pueden ver elevados si incrementamos el número de dispositivos. Además, sabemos que en el modelo heterogéneo presentado, la CPU sólo se utiliza como árbitro de los coprocesadores, repartiendo y sincronizando las distintas ejecuciones. Creemos que la CPU puede jugar un papel más activo, interviniendo en la computación para acelerar más el cómputo global. No obstante, esto requeriría una estrategia avanzada de balanceo de carga dinámico para asignar datos a cada procesador que intervenga, y que podría incrementar, todavía más si cabe, la complejidad de la programación del sistema. Es por ello, que pensamos que debe ser el compilador y/o runtime el que maneje los recursos en sistemas heterogéneos.

Parte III Conclusiones

Capítulo 6 Conclusiones

6.1 Conclusiones

En este documento analizamos una de las tendencias de cómputo de alto rendi- miento más atractivas de la actualidad que ofrece, desde nuestro punto de vista, grandes oportunidades para el desarrollo eficiente y sostenible de aplicaciones con una gran demanda de requisitos computacionales. Esta tendencia está guiada por una visión de computación heterogénea, donde el procesador tradicional, orientado a ejecutar hilos o tareas donde premia la latencia, cohabita con una serie de aceleradores, diseñados para ejecutar un elevado número de hilos de ejecución, explotando el paralelismo (throughput). En concreto, nos centramos en la arquitectura propuesta por Intel, donde los aceleradores están basados en la arquitectura Xeon Phi. Esta arquitectura posee, desde nuestro punto de vista, una ventaja competitiva fundamental con otros aceleradores del mercado. Su ISA basado en el tradicional x86, ofrece una curva de aprendizaje baja y un esfuerzo para migrar las aplicaciones existentes relativamente pequeño. Además, el modo nativo de programación de Intel Xeon Phi, consolidado en las nuevas generaciones de esta arquitectura, ofrece una programación sencilla evitando manejar el dispositivo como si fuese un periférico. No obstante, Intel Xeon Phi también permite el manejo de varias tarjetas aceleradoras para abordar problemas con mayor demanda de memoria de la que un único dispositivo es capaz de proveer, logrando una mayor aceleración para problemas que así lo necesiten.

En este documento evaluamos los dos modos de programación más destaca- dos de la arquitectura Xeon Phi, el modo nativo y el modo offload. En los capítulos 3 y 4 mostramos los principales pasos necesarios para la optimización de códigos Stencil, a fin de conseguir un rendimiento óptimo en esta arquitectura. En el proceso comprobamos la facilidad de portar nuestros códigos base a dicha arquitectura con dos acciones muy sencillas, paralelizar el código con un pragma de OpenMP, y compilar el código para Xeon Phi agregando la directiva -mmic. Una vez portado nuestro código a Xeon Phi, llevamos a cabo algunas optimizaciones escalares, y también vectorizamos los códigos utilizando las clausulas restrict y collapse, logrando con esto una mejora media del rendimiento por un factor cercano a 7 veces. La vectorización fue posible tras realizar un alineado de los datos mediante técnicas del padding. En cuanto a los datos de entrada, evaluamos tanto su tamaño como su forma, mostrando que no tienen ninguna influencia significativa y que el rendimiento obtenido es independiente de estos aspectos.

A fin de mejorar la paralelización de nuestros códigos también evaluamos la política de no traer datos a caché en caso de fallo de escritura en caché (denomi- nada por Intel Streaming non-temporal store instructions). La técnica Streaming stores nos permitió reducir la polución de las cachés, logrando una mejora significativa del rendimiento para el kernel de difusión de calor. Otra modificación importante es la inclusión del modificador collapse al pragma que utilizamos para paraleli- zar el triple bucle anidado que recorre los datos de entrada. Este modificador combina varios niveles de bucles for, aumentando la cantidad de trabajo que el planificador OpenMP tiene disponible. Además, las políticas de planificación del bucle OpenMP nos permitieron mejorar aún más el rendimiento, específicamente cuando agregamos la cláusula shedule al pragma que paraleliza el bucle del cálculo de nuestros Stencil, indicándole a OpenMP qué política queremos aplicar a la hora de asignar trabajo a cada hilo de ejecución. Nuestros resultados experimentales ofrecieron una mejora del rendimiento promedio de 34 % para el kernel acústico, de 42% para el kernel sísmico, ambos usando la política shedule(dynamic,4), y de 30% para el kernel de calor usando una política (dynamic,2). En cuanto a la política de afinidad de los hilos, nuestra evaluación muestra un rendimiento similar para todas ellas, con una mejora mínima con la afinidad balanced.

Otra de las optimizaciones que ofrecen una mejora de rendimiento significati- va es la técnica de blocking. Esta técnica mejora el uso del sistema de memoria. Hemos logrado identificar el mejor tamaño de bloque para nuestros tres kernels (en las dimensiones X, Y ó Z). Nuestros resultados muestran una mejora del rendimiento del 38,97 % para el kernel acústico, 75,03 % para el kernel sísmico, y 46,23% para el kernel de calor, cuando consideramos unos tamaños de bloque de 400 en la dimensión X, y 4 para las dimensiones Y y Z (en el caso del código sísmico tiene un bloque de 2 en la dimensión Y), usando siempre para todos los códigos una planificación schedule(dynamic).

En cuanto a rendimiento energético, cabe destacar que, conforme mejoramos el rendimiento en GFLOPs en las distintas versiones de nuestros códigos, observamos una mejora en la eficiencia energética equivalente. Si comparamos el EDP de Xeon Phi contra el de Xeon obtenemos un factor de mejora energética entre 35 y 40 veces para los diversos kernels. En relación a los resultados obtenidos al comparar con GPUs de Nvidia, podemos apreciar unos resultados similares para las GPUs a los obtenidos con Xeon Phi para el kernel que modela la ecuación de difusión acústica, pues ambas plataformas son muy eficientes energéticamente.

Si comparamos los modos de programación offload y nativo de Xeon Phi, hay una degradación en rendimiento de un 4.5 % entre el modo descarga y el modo nativo (bastante pequeño teniendo en cuanta la sobrecarga asociada al movimiento de datos hacia/desde Xeon Phi). El uso del modo offload con dos tarjetas, utilizando la estrategia de utilizar un sólo hilo y las tarjetas de forma asíncrona por parte del procesador anfitrión, nos ofrece una mejora del 50 % para una precisión en los resultados del 100%. Si podemos relajar dicha precisión a tan sólo el 6.73 % de margen de error, la mejora ya sube al 80 % (es decir, estamos prácticamente doblando el rendimiento al usar dos tarjetas). Consideramos que aún podemos explotar más si cabe el modo offload, utilizando más tarjetas y distribuyendo trabajo entre procesador y coprocesador.

6.2 Contribuciones principales

Los principales resultados obtenidos en este trabajo para la arquitectura Xeon Phi son los siguientes:

- Hemos comprobado que la portabilidad de los códigos Stencil 3-D para su ejecución en una arquitectura Xeon Phi es inmediata cuando partimos de un código base paralelizado de forma sencilla (naïve) para una arquitectura Intel Xeon multinúcleo. Únicamente hay que añadir la opción de compilación (-mmic) al compilador de Intel. El problema que hemos tenido es la falta de escalabilidad de estos códigos con el número de hilos usados (tanto para Xeon multinúcleo como para Xeon Phi), pues los códigos Stencil están muy limitados por la gran cantidad de datos que necesitan traer de memoria y las pocas operaciones que realizan con dichos datos.

- Para mejorar el rendimiento de nuestros códigos, hemos tenido que mejorar la paralelización y la vectorización del código

- Asimismo, hemos necesitado mejorar el uso del sistema de memoria por medio reutilizando los datos lo máximo posible, mediante la técnica de blocking en cada una de las dimensiones de las matrices de datos. La opción de usar la caché sin localización de bloques en caso de fallo de escritura (Streaming Stores) produce un beneficio para el código de calor, y no per- judica en el caso de los códigos de difusión acústica y sísmica. También hemos comprobado que la planificación dinámica de los hilos mejora el tiempo de ejecución de todos nuestros códigos.

- Hemos comparado Xeon Phi contra Xeon multinúcleo para todos nuestros códigos en términos de potencia de cálculo y de eficiencia energética. La mejora de rendimiento en tiempo de ejecución oscila entre 5x y 7x, y la mejora en la eficiencia energética entre 10x y 28x.

- Una vez certificado que el Xeon Phi es mejor en ambas métricas (especial- mente en la eficiencia energética), hemos evaluado también la arquitectura CUDA de Nvidia (usando las GPUs K40c y 980GTX) para el kernel Stencil de difusión acústica de 7 puntos. Hemos comparado los resultados obtenidos en Xeon Phi y en CUDA con los obtenidos en el Xeon multinúcleo, obtenien- do resultados similares en ambas plataformas, con mejoras en el tiempo de ejecución superiores a 6, y mejoras en la eficiencia energética mayores de 30 en ambas plataformas (ligeramente superiores en el caso del Xeon Phi). Por lo tanto, concluimos que tanto Xeon Phi como CUDA (GPUs K40c y 980GTX) ofrecen beneficios claros para este kernel, tanto desde el punto de vista de la aceleración del tiempo de ejecución como especialmente de su eficiencia energética.

- Finalmente, hemos evaluado el Xeon Phi actuando como un acelerador, manejándolo en el modo descarga (offload). Para el caso de nuestros códigos, se nos ha planteado el compromiso de elegir entre precisión y aceleración cuando ejecutamos el código en varias tarjetas aceleradoras, debido a las comunicaciones de datos que se deben realizar al finalizar cada iteración del bucle Stencil. Hemos encontrado que, con tan sólo un 6 % de pérdida de precisión en los valores finales de salida de nuestros códigos, conseguimos unas aceleraciones en dos tarjetas alrededor del 80%, mientras que si queremos mantener la precisión en el 100% el rendimiento baja al 50%.

El modo descarga del Xeon Phi es atractivo porque nos permite abordar mayores tamaños de problema, así como reducir el tiempo de ejecución usando varias tarjetas aceleradoras en paralelo. Hemos comprobado que utilizando las extensiones LEO proporcionadas por Intel el manejo de dicho modo es sencillo, y el usuario tiene control en cada momento de qué datos se transfieren entre la memoria del anfitrión y la tarjeta aceleradora, así como qué partes del código se van a ejecutar en el acelerador.

6.3 Publicaciones científicas y colaboraciones

La investigación llevada cabo durante el desarrollo de este documento ha sido publicada en una revista internacional recogida en el listado JCR, y varias confe- rencias internacionales y nacionales. En esta sección se mencionan las publicacio- nes científicas directamente relacionadas con nuestro trabajo aquí presentado, y que apoyan nuestras futuras lineas de investigación.

6.3.1 Revistas Internacionales

1. Mario Hernández, Baldomero Imbernón, Juan M. Navarro, José M. García, Juan M. Cebrián, and José M. Cecilia. Evaluation of the 3-D finite difference implementation of the acoustic diffusion equation model on massively parallel architectures. Computers & Electrical Engineering, Volumen 46, páginas 190 - 201, 2015.Doi:10.1016/j.compeleceng.2015.07.001.24

En este articulo mostramos la hoja de ruta para utilizar arquitecturas masi- vamente paralelas en una simulación 3D-FD. Describimos la aceleración de un kernel 3D-FD dentro de la simulación acústica de la sala para mostrar los beneficios de co-diseño hardware-software. Partimos de una implemen- tación Matlab reescrita en ANSI C. Portamos esta aplicación a CUDA y a la arquitectura Intel Xeon Phi. Nuestras versiones mejoran tanto en el tiempo de ejecución como en el consumo de energía respecto al uso de una CPU multinúcleo. Los resultados mostrados en este artículo se obtuvieron en el caso del Xeon Phi para con unos códigos optimizados siguiendo tan sólo los desarrollos del Capítulo 3.

6.3.2 Actas de congresos

1. Mario Hernández, José M. García y José M. Cecilia. Early experiences run- ning the 3D Stencil Jacobi method in Intel Xeon Phi. In H. Weghorn, editor, 11th International Conference Applied Computing 2014, pages 85-92, Por- to (Portugal), Oct. 2014. International association for development of the information society.23

En este articulo presentamos nuestras primeras experiencias con el pro- cesador Intel Xeon Phi. Nuestro objetivo es explorar si (y cuándo) hay ventajas utilizando el procesador Intel Xeon Phi para el procesamiento de los cálculos Stencil 3D de Jacobi. Nos encontramos con que, en general, el software basado en código C/C++ puede ser ejecutado nativamente en Xeon Phi sin recodificación extra. Nuestras optimizaciones revelan que la arquitectura Intel Xeon Phi obtiene hasta un factor de aceleración de 29.4x en el rendimiento en comparación con su versión secuencial ejecutada en un procesador Intel Xeon.

2. Mario Hernández, Juan M. Cebrián, José M. Cecilia y José M. García. Eva- luating 3-D Stencil codes on Intel Xeon Phi: Limitations and Trade-offs. In XXVI Jornadas de Paralelismo, pages 568-573, Córdoba (Spain), Sept. 2015. Uni- versidad de Córdoba.21

En este articulo evaluamos tres núcleos basados en el Stencil 3-D con cálcu- los de valores de punto flotante de doble precisión (8 bytes) seleccionados de diferentes campos de investigación: 1) Stencil de difusión de calor 3-D (11 puntos), 2) Stencil de difusión acústica 3-D (7 puntos) y 3) Stencil 3-D de onda sísmica isotrópica (25 puntos). Nos centramos en cuestiones clave que deben tenerse en cuenta a fin de lograr un rendimiento óptimo en la arquitectura Intel MIC. Mostramos cómo lograr el máximo rendimiento afi- nando algunos parámetros inherentes específicos en relación tanto al núcleo como a la arquitectura de hardware. Entre ellos, destacan la escalabilidad, la afinidad, el blocking y el efecto de la forma del grid. Nuestros resultados experimentales producen pequeñas mejoras de rendimiento usando la afini- dad scatter, lo que demuestra que el tamaño de blocking afecta fuertemente el rendimiento del kernel. Además revela que la forma del grid tiene un impacto mínimo en el rendimiento, siempre y cuando se seleccione el mejor tamaño de bloque.

3. Mario Hernández, Juan M. Cebrián, José M. Cecilia y José M. García. Eva-luation of 3-D Stencil Codes on the Intel Xeon Phi Coprocessor. In Parallel Computing (ParCo 2015). Edimburgo (Reino Unido), Sept. 2015.22

En este articulo presentamos la evaluación de Intel Xeon Phi (Knights Corner) para códigos Stencil 3-D utilizando diferentes estrategias de optimización. Nuestra evaluación se basa en tres núcleos que son ampliamente aplicados para simular el calor, la difusión acústica, así como ecuaciones de onda sísmica isotrópica. Nuestros resultados experimentales producen mejoras de rendimiento entorno a 25x en comparación con implementaciones se- cuenciales de alto nivel. Las mediciones de energía muestran una tendencia similar a la del rendimiento. Además, proponemos una serie de consejos para optimizar los códigos del Stencil para estar arquitectura, basándonos en una implementación OpenMP C/C++.

Apéndices

Apéndice A Manejando el Intel LEO

A.1 Modelo de memoria no compartida mediante Lenguaje de Extensiones (LEO)

El modelo de memoria no compartida Intel LEO usa las directivas o pragmas con el prefijo offload_. Este modelo es apropiado para hacer frente a estructuras de datos planas como escalares, matrices y estructuras que son transferibles bit a bit (por ejemplo, que no contienen punteros y no invocan a constructores o destructores). Los datos de este modelo se copian de ida y vuelta entre el procesador y el coprocesador en torno a las regiones de código descargado. Estos datos son una combinación de las variables transferidas implícitamente porque están léxicamente referenciadas dentro del código descargado, y las variables que figuran explícitamente en las cláusulas del pragma.

A.1.1 Offload a varios coprocesadores

Un sistema puede tener varios coprocesadores Intel Xeon Phi. El sistema de asignación de código a cada coprocesador puede ser programado mediante un bucle paralelo en el anfitrión que distribuye el trabajo de forma explícita como se muestra en la Figura A.1. Otra alternativa es hacerlo sin hilos (quitando la directiva OpenMP), mediante un bucle secuencial utilizando la cláusula signal(& var) (en lugar de simplemente offload). El uso de signal convierte las descargas en asíncronas, es decir, lanzan trabajo al coprocesador y devuelven el control al anfitrión, sin esperar a que éste se complete. El anfitrión puede usar una cláusula wait(& var) para esperar la finalización de las tareas descargadas.

Abbildung in dieser Leseprobe nicht enthalten

Figura A.1: Código sencillo de descarga a múltiples coprocesadores.

A.1.2 Personalización de código para la CPU y Xeon Phi

En ocasiones pueden ser necesario optimizar el código genérico para el coprocesa- dor, pero éste puede dejar de ser compatible con la CPU del host (por ejemplo, si utilizamos intrínsecos que no están disponibles en ambos procesadores). Cuando sea posible, se recomienda utilizar una verificación dinámica para seleccionar el código a ejecutar en el coprocesador.

No es recomendable dejar que sea el preprocesador de C quien seleccione el código a ejecutar (Figura A.2), ya que tiene el potencial de crear un desajuste entre las variables del coprocesador y las de la CPU.

Abbildung in dieser Leseprobe nicht enthalten

Figura A.2: Código versión MIC o CPU.

A.1.3 Declaración de variables y funciones

Cuando diseñamos una aplicación siguiendo el modelo de offload tenemos que considerar que las funciones declaradas en el código principal (fuera de la sección descargada) no se pueden utilizar como tal, por lo que es necesario emplear una gramática específica que asegure que dichas funciones están disponibles en el coprocesador. Para este tipo de declaración única, tenemos que añadir un modificador al código existente. Para C/C++, __declspec(target (mic)) o bien __atribute__ (( target (mic))) (Figura A.3).

Notar que hay dos niveles de paréntesis en el atributo, si escribimos un único nivel causará un error de compilación. Esta declaración se puede utilizar tanto para funciones como para variables. Si se utiliza en una variable, ésta debe ser una variable global. Una vez declarados, funciones y variables se pueden utilizar tanto por la CPU como por Xeon Phi. Compilar sólo funciones y datos marcados explícitamente con el atributo target(mic) en el binario del coprocesador asegura que el código en el coprocesador es tan pequeño como sea posible y que las variables y funciones están disponibles en el coprocesador.

Cuando se tiene varios datos y-o declaraciones de funciones, en lugar de especificar cada declaración con su propio atributo target, se puede incluir un grupo de datos y declaraciones de funciones entre las variantes push y pop del pragma o directiva offload_attribute, como el que se muestra en la Figura A.4.

Abbildung in dieser Leseprobe nicht enthalten

Figura A.3: Declaración de funciones y variables globales.

A.1.4 Movimiento de datos

Como se ha mencionado anteriormente, la memoria del procesador y de los coprocesadores están separadas tanto física como virtualmente. Conceptualmente,

Abbildung in dieser Leseprobe nicht enthalten

Figura A.4: Múltiples declaraciones de funciones y variables globales.

el ciclo de vida de una descarga de datos dinámica, mostrado en la Figura A.5, consiste de cinco actividades:

1. Asignación de espacio de datos en Xeon Phi.
2. Copias de datos de entrada a la memoria de Xeon Phi.
3. Ejecución de descarga en la MIC.
4. Copiado de resultados de regreso a la memoria del anfitrión.
5. Liberación del espacio de datos asignado en Xeon Phi.

Abbildung in dieser Leseprobe nicht enthalten

Figura A.5: Ciclo de vida del offload.

El compilador proporciona mecanismos para controlar cada una de estas actividades de forma individual, o bien estas cinco actividades se pueden englo- bar en un sólo bloque (etiquetado con el pragma de descarga). Esto permite la computación y comunicación solapada, así como la posibilidad de permitir datos persistentes en la MIC.

La ventaja utilizar Intel LEO es que nos permite controlar el movimiento de datos de manera muy precisa, pero con cierta pérdida de generalidad, ya que debemos especificar los datos a ser transferidos. Los datos no necesitan ser declarados o asignados de manera especial. En este enfoque, las estructuras de datos nos permite especificar precisamente que bloques de datos deben ser transferidos desde y hacia Xeon Phi. Se puede colocar un pragma de descarga antes de cualquier sentencia, incluyendo una sentencia compuesta. La sentencia prefijada con el pragma de descarga puede también tener un pragma paralelo OpenMP.

Las cláusulas para el movimiento de datos son:

- in(var1 [,...]): Copia del anfitrión a Xeon Phi. Define una variable como estrictamente de entrada a Xeon Phi.
- out(var1 [,...]): Copia de Xeon Phi al anfitrión. Define una variable como estrictamente de salida de Xeon Phi.
- inout(var1 [,...]): Copia del anfitrión a Xeon Phi y de nuevo de Xeon Phi al anfitrión. Define una variable que es a la vez copiada del anfitrión a Xeon Phi y regresada desde Xeon Phi al anfitrión.
- nocopy(var1 [,...]): No copia las variables seleccionadas.

La cláusula #pragma offload target(mic) in(a), out(c,d), inout(b) es un ejemplo que muestra de manera clara el movimiento de datos de manera explícita.

Gracias al uso de estas API asíncronas podemos ocultar en parte la latencia PCIe, mediante la superposición de la comunicación con el cálculo tanto de el anfitrión como en el coprocesador. El componente clave para hacer este trabajo es implementar un mecanismo para detectar la finalización de las copias de datos asíncronos a-y-desde el coprocesador. Para ello, todos los datos descargados y computados de forma asíncrona deben estar asociados con una cláusula de señal (signal) que describiremos en detalle en las siguientes secciones. Detectar la terminación de esta operación se logra con llamadas explícitas a la API.

Consideramos necesario remarcar que se debe intentar reducir al mínimo la transferencia de datos a Xeon Phi, realizando cálculos simples en el anfitrión. Además, si los valores de los datos al final de una descarga son necesarios para una descarga posterior, se deben mantener en Xeon Phi de manera persistente. Cuando trabajamos en un entorno con varios coprocesadores, la descarga debe realizarse en el mismo coprocesador, mediante el uso de un número de MIC explícito en la cláusula de destino. Los datos sólo se pueden copiar entre los coprocesadores y el anfitrión usando directivas offload; no existe un método para copiar de una MIC a otra MIC. Se debe utilizar la cláusula NOCOPY cuando se reutilicen los datos anteriores. La Tabla A.1 muestra los modificadores/cláusulas utilizadas en las descargas explícitas.

Tabla A.1: Modificadores y cláusulas explícitas de descarga.

Abbildung in dieser Leseprobe nicht enthalten

A.1.4.1 Transferir datos con offload_transfer

También existe la posibilidad de realizar una descarga sólo de datos, que realiza la transferencia pero no ejecuta código en el coprocesador. Para ello se puede utilizar el pragma offload_transfer. En C/C++: #pragma offload_transfer target(mic)[: Numero_coprocesador])[,cláusula...]. Es importante destacar que todas las cláusulas del pragma de descarga también se aplican a offload_transfer. El código de la Figura A.6 muestra el uso del pragma offload_transfer.

Abbildung in dieser Leseprobe nicht enthalten

Figura A.6: Transfiriendo datos con offload_transfer.

A.1.4.2 Asignación de memoria para subconjuntos de matrices en C/C++

En el caso de trabajar con grandes conjuntos de datos podemos utilizar el modificador alloc junto a los pragmas de descarga para referenciar un subconjunto de dichos datos. Sólo podemos utilizar una variable con este modificador en esa sección. Cuando trabajamos con una matriz multidimensional, la segunda y posteriores expresiones índice deben usar como base los elementos de la primera dimensión (más un offset). La reserva de memoria de la matriz debe ser contigua. Los datos se transfieren en la porción de la matriz especificada por la expresión in o out.

Cuando el limite inferior de la primera dimensión de una sección utilizada en el modificador alloc es distinto de cero, entonces la asignación de memoria comienza en ese elemento. La memoria anterior al limite inferior es liberada y no debe ser referenciada por el programa en el coprocesador. Al no hacer referencia a ella se habilita una sección más pequeña de la matriz para ser transferida al coprocesador. En la Figura A.7 se muestra un ejemplo que considera lo siguiente:

- El modificador alloc(p[5:1000]) reserva memoria para 1000 elementos en el coprocesador.
- El primer elemento utilizable tiene índice de 5 y el último tiene índice de 1004. Por lo tanto sólo los elementos 5 a 1004 tienen memoria reservada en el coprocesador.
- La transferencia de datos se especifica mediante la cláusula in(p[10:100]). Se transfieren 100 elementos a la memoria asignada en el rango de p10 hasta p109.

A.1.4.3 Moviendo datos entre variables

El modificador into permite transferir datos de una variable en el procesador a otra en el coprocesador (cláusula in), y viceversa (cláusula out). Sólo está

Abbildung in dieser Leseprobe nicht enthalten

Figura A.7: Ejemplo de asignación de memoria a partes de una matriz.

permitida la transferencia de un item en la lista de variables-ref cuando se utiliza el modificador into. De este modo se establece una correspondencia uno a uno entre el anfitrión y un coprocesador. Hay que tener en cuenta que esto puede limitar la portabilidad del código, ya que los nombres de variables dentro de la región de descarga van a tener un nombre diferente a las del anfitrión debido a la cláusula into.

Los modificadores alloc_if, free_if, y alloc se aplican a la expresión into, pero no está permitido con las cláusulas inout y nocopy.

Cuando se utiliza el modificador into, la expresión fuente genera un stream de los elementos que van a copiarse en los rangos de memoria especificados por la expresión into. Algunos ejemplos en C/C++, se muestran en la Figura A.8.

Abbildung in dieser Leseprobe nicht enthalten

Figura A.8: Ejemplo de utilización del modificador into.

A.1.4.4 Gestionando explícitamente datos asignados al heap

El código que se ejecuta en el coprocesador puede llamar a malloc/free para asignar/liberar memoria dinámica. Las variables puntero que apuntan a la memoria dinámica asignada de esta manera, son escalares y están sujetos a las normas de persistencia de datos, dependiendo del alcance de su definición (estática o función local).

Para evitar la interferencia entre la asignación dinámica gestionada por el compilador y la asignación dinámica explícita, se utiliza la cláusula NOCOPY para las variables puntero referenciadas dentro de las regiones de offload que están siendo gestionadas explícitamente.

Tabla A.2: Descripción de in/out/nocopy y uso de la cláusula longitud.

Abbildung in dieser Leseprobe nicht enthalten

A.1.4.5 Punteros locales frente a punteros usados a través de Offload

Los punteros utilizados dentro de las regiones de offload son por defecto inout, es decir, los datos asociados a ellos se transfieren al comienzo y final de la ejecución del código descargado. A veces, un puntero puede ser utilizado estrictamente a nivel local, es decir, se asigna y se utiliza sólo en el coprocesador. La cláusula NOCOPY es útil en este caso, ya que deja el puntero sin modificar por la cláusula offload, y permitiendo que el programador gestione explícitamente el valor del puntero. En otros casos, los datos son transferidos al apuntador de la CPU, y una subsecuente descarga puede querer o bien, a) usar la misma memoria asignada y transferir nuevos datos en el, o b) mantener la misma memoria y reusar los mismos datos. Para el caso a), necesitamos una cláusula in con longitud igual al número de elementos a transferir. Para el caso b) usaremos una cláusula con longitud 0 para “refrescar” el puntero pero evita cualquier transferencia de datos.

La descripción completa de in/out/NOCOPY y el uso de la cláusula de longitud de datos se muestra en la Tabla A.2. La Figura A.9 muestra un ejemplo de datos asignados dinámicamente.

Abbildung in dieser Leseprobe nicht enthalten

Figura A.9: Código para asignar datos dinámicamente.

A.1.4.6 Persistencia: Datos asignados al heap

El heap del coprocesador es persistente entre descargas. Hay dos formas de utilizar la memoria heap en Xeon Phi. Dejar que el compilador gestione la memoria dinámica usando el #pragma o gestionando este usando malloc/free.

Cuando los datos son gestionados por el compilador, la asignación de memoria es controlada por las cláusulas alloc_if y free_if y la transferencia de datos es controlada por in/out/inout/nocopy. Los dos son independientes, pero los datos sólo pueden ser transferidos in y out de la memoria asignada.

Dada la importancia de gestionar adecuadamente la memoria en Xeon Phi, especialmente en lo referente a la asignación, reuso y liberación de la misma, utilizaremos las macros que se muestran en la Figura A.10 y que detallamos a continuación:

Macro ALLOC. Lo utilizamos para asignar memoria (alloc_if(1)) y no liberarla (free_if(0)).

Abbildung in dieser Leseprobe nicht enthalten

Figura A.10: Macros para gestionar la memoria.

Macro FREE. Lo utilizamos para liberar memoria, no asigna memoria (alloc_if(0)), pero si la libera (free_if(1)).

Macro REUSE. Lo utilizamos para reutilizar los datos que estén en memoria, no asigna memoria (alloc_if(0)) y no la libera (free_if(0)).

A.1.5 Modo asíncrono de transferencia y ejecución

Como ya hemos comentado anteriormente desarrollador puede definir el com- portamiento del anfitrión durante la descarga al coprocesador:

1. El anfitrión espera a que el coprocesador finalice la tarea.
2. El anfitrión trabaja en una tarea diferente.
3. El anfitrión trabaja en una parte de la misma tarea.

Por defecto, el pragma o directiva de descarga hace que el hilo del procesador que se encuentra con el pragma o directiva de descarga espere a la finalización de la descarga antes de continuar con su ejecución.

Para especificar una descarga asíncrona, se debe especificar una cláusula signal en el pragma o directiva de descarga para iniciar el cálculo, y posteriormente utilizar el pragma offload_wait o la directiva OFFLOAD_WAIT con el modificador target(mic[: número-de-tarjeta]) que espera a la finalización del cálculo descar- gado al coprocesador indicado. El ejemplo mostrado en la Figura A.11 permite al procesador emitir descargas de cálculos y continuar la actividad concurrente sin usar hilos adicionales del procesador. Opcionalmente podemos reemplazar la cláusula “do” por un “if” y utilizar una variable para comprobar en tiempo de ejecución si se debe intentar iniciar la descarga o no.

También es posible seguir estos comportamientos dentro una tarea MPI con múltiples coprocesadores.

Abbildung in dieser Leseprobe nicht enthalten

Figura A.11: Descarga en la mic continuando la actividad concurrente.

A.1.5.1 Transferencia de datos asíncrona

Usando offload_transfer con una cláusula signal podemos realizar la transferencia de datos asíncrona. La etiqueta especificada en la cláusula signal es una expresión de dirección asociada con ese conjunto de datos. La transferencia de datos se inicia y el procesador puede continuar más allá del pragma de descarga.

Un pragma escrita adelante con una cláusula wait hace que la actividad especi- ficada en el pragma comience sólo después de que se hayan recibido todos los datos asociados con la etiqueta. Los datos se colocan en las variables especifica- das cuando se inició la transferencia de datos. Estas variables todavía deben ser accesibles. Para transferir datos de forma asíncrona desde el procesador hasta el coprocesador, se utiliza una cláusula de señal en un offload_transfer con cláusulas in. Si, durante una descarga asíncrona, una señal es creada en un hilo, hilo A, y esperó en un hilo diferente, hilo B, se debe asegurar que el hilo B no consulta la señal antes de que el hilo A ha iniciado la descarga asíncrona para crear la señal.

En el ejemplo C/C++ mostrado en la Figura A.12, la transferencia de datos de las matrices de punto flotante f1 y f2 se inicia en la línea 11. La descarga no inicia un cálculo. Su único propósito es iniciar la transferencia de f1 y f2 al coprocesador. En la línea 22 el procesador inicia el cálculo de la función foo en el coprocesador. La función utiliza los datos de f1 y f2, cuya transferencia se inició anteriormente. La ejecución de la región de descarga en el coprocesador comienza sólo después de que la transferencia de f1 y f2 se haya completado.

Podemos realizar múltiples transferencias de datos asíncronos independien- tes en cualquier momento. El ejemplo mostrado en la Figura A.13 utiliza of- fload_transfer para enviar f1 y f2 al coprocesador en diferentes momentos, prime- ro f1 en la línea 11, y luego f2 en la línea 21.

Figura A.12: Ejemplo de transferencia asíncrona del procesador al coprocesador.

Abbildung in dieser Leseprobe nicht enthalten

A.2 Usando el compilador de Intel para offload

A.2.1 Variables de entorno para el control de descarga

El sistema MPSS de MIC utiliza varias variables de entorno para facilitar la ejecución o depuración del programa objetivo. Identificamos dos categorías de variables de entorno:

1. Aquellas que afectan a la forma en que la librería en tiempo de ejecución opera la descarga.

2. Aquellas que se transmiten a través del entorno de ejecución del coprocesa- dor por la librería de descarga.

En primer lugar describiremos las variables de entorno en la categoría 1. Estas utilizan los prefijos, “MIC_” o “OFFLOAD_”. El prefijo es fijo, como es el nombre de la variable de entorno. La variable de entorno especial MIC_ENV_PREFIX se utiliza para distinguir las variables en la categoría 2. Presentamos a continuación algunas variables de entorno de uso general; para una versión más completa y actualizada, se remite al lector a consultar el manual de MPSS.

MIC_STACKSIZE. Define el tamaño de la pila del proceso de descarga en la MIC. Su valor por defecto es de 2 MB. Las unidades que soporta son B, K,

Abbildung in dieser Leseprobe nicht enthalten

Figura A.13: Ejemplo de transferencia asíncrona del procesador al coprocesador en diferentes momentos.

M, G y T. Por ejemplo: export MIC_STACKSIZE = 5M, que se corresponde con 5MB. Se debe utilizar MIC_OMP_STACKSIZE para modificar el tamaño de cada hilo OpenMP.

MIC_LD_LIBRARY_PATH. Define la ruta de las librerías dinámicas para el pro- grama de MIC. La ruta especificada corresponde a la localización de las libre- rías dentro del sistema de ficheros del coprocesador. MIC_LD_LIBRARY_PATH normalmente apunta a las librerías dinámicas definidas por el usuario, las librerías estándar están localizadas por defecto en la carpeta “/lib64”.

MIC_ENV_PREFIX. Establece el prefijo de las variables de entorno que pertene- cen a MIC. Este prefijo ayuda a distinguir entre las variables de entorno del coprocesador y las del anfitrión. Si no se hace esta distinción, la variable de entorno del anfitrión también se utilizará en el coprocesador. Por ejemplo, si definimos OMP_NUM_THREADS = 8 en el anfitrión, OpenMP puede lanzar hasta 8 hilos de ejecución, tanto en el anfitrión como en el coprocesador, lo cual no es deseable. Por lo tanto, es importante distinguir las variables de entorno entre esos dos sistemas. La Figura A.14 muestra un ejemplo.

El ajuste de esta variable de entorno no tiene efecto sobre las variables de entorno MIC_* fijos, a saber MIC_USE_2MB_BUFFERS, MIC_STACKSIZE y MIC_LD_LIBRARY_PATH. Esos nombres son fijos.

Abbildung in dieser Leseprobe nicht enthalten

Figura A.14: Variables de entorno de MIC.

A.2.2 Opciones del compilador Intel

Existen multitud de opciones de compilación que se pueden seleccionar tanto vía parámetros como por variables de entorno. Estas opciones nos permiten, entre otras cosas:

1. Ignorar constructores del lenguaje de descarga (no-offload).
2. Construir una aplicación que se ejecute de manera nativa en la arquitectura Intel MIC (mmic).
3. Establecer un modificador para todas las rutinas globales y objetos de datos en el archivo fuente con el atributo de descarga target(mic)(offload-attribute- target).
4. Especificar las opciones que se utilizarán para un objetivo específico y herramientas (offload-option).
5. Especificar en que fase se genera el informe de optimización (opt-report- phase=offload).
6. Descargas específicas de argumentos al compilador Intel.
7. Advertencias.

Por defecto todo el conjunto de variables de entorno en el anfitrión se copian al coprocesador y durante la ejecución en modo offload, pero este comportamiento puede alterarse usando MIC_ENV_PREFIX para controlar las variables transferidas al coprocesador.

A.2.3 Archivos de cabecera

En C/C++, si se utiliza la API de la MIC, necesitamos que nuestro código contenga el archivo de cabecera offload.h. Por ejemplo para obtener el número de dispositivos.

A.2.4 Opción vec-report utilizada con offload

Por defecto, cuando compilamos un programa con opciones de descarga y usamos una opción de compilador vec-report, obtenemos un informe del compilador para el código compilado en el anfitrión y otro para el código compilado para el coprocesador. Los mensajes se etiquetan para indicar pertenecen al coprocesador (MIC) o no, pero puede ser confuso porque se entremezclan. Es especialmente confuso si un bucle está vectorizado sólo para una de las plataformas. Por ello puede ser interesante limitar los informes de vectorización para una de las plataformas.

Para obtener sólo los mensajes que afectan al coprocesador podemos compilar con:

-vec-report0 -offload-option,mic,compiler, -vec-report3

Del mismo modo, para mostrar sólo los mensajes que afectan al anfitrión podemos utilizar:

-vec-report3 -offload-option,mic,compiler, -vec-report0

A.2.5 Generación de un informe de offload

El entorno MPSS proporciona herramientas para generar informes sobre el proceso de descarga, que incluyen la siguiente información:

- Tiempo que se tarda en ejecutar una región de código descargado.
- Cantidad de datos transferidos entre el procesador y el coprocesador.
- Detalles adicionales, incluyendo inicialización de dispositivos y transferen- cias individuales de variables.

Los siguientes mecanismos habilitan y deshabilitan el informe de descarga:

- La variable de entorno OFFLOAD_REPORT.
- La API _offload_report.

El informe del compilador se inicia con una línea [Offload] para marcar claramente los eventos que corresponden a la descarga al coprocesador.

La actividad en el anfitrión está marcada con la palabra clave [HOST], mientras que la actividad en el coprocesador está marcada con [MIC n], donde n es el número lógico del coprocesador al que se envía la descarga. La parte superior del informe muestra el mapeo de los dispositivos lógicos a los dispositivos físicos.

Debido a que múltiples descargas pueden estar en curso simultáneamente, es necesario etiquetar toda la salida asociada con una directiva de descarga específica. Una etiqueta de la forma [Tag n] identifica de forma exclusiva las líneas en el informe de descarga que pertenecen a una descarga particular.

Para cada descarga, las primeras dos líneas del informe son el nombre de archivo fuente y el número de línea de la directiva de descarga. Después de eso muestran una línea que se le asigna una etiqueta a la que la descarga se imprime.

El informe contiene además, una línea para cada actividad principal. Estas lí- neas contienen una anotación de la actividad seguida de la etiqueta que identifica la actividad a la que pertenece la descarga. Las anotaciones son las siguientes:

[State] La actividad que se lleva a cabo como parte de la descarga.

[Var] El nombre de una variable transferida y la dirección(s) de la transfe- rencia.

[CPU Time] El tiempo total medido por dicha directiva de descarga en el procesador.

[MIC Time] El tiempo total medido para la ejecución de la descarga en el destino (Excluye el tiempo de transferencia de datos entre el procesador y el coprocesador, y sólo cuenta el tiempo de ejecución en el coprocesador).

[CPU->MIC Data] El número de bytes de datos transferidos desde el procesador al coprocesador.

[MIC->CPU Data] El número de bytes de datos transferidos desde el coprocesador al procesador.

La función _Offload_report controla la impresión de la descarga en tiempo de ejecución y la cantidad de datos transferidos. La sintaxis es: _offload_report(número entero). Para “número entero” los valores posibles son: 1, 2 ó 3. A continuación se detallan:

1 Imprime el tiempo de cálculo de descarga, en segundos.
2 Además de la información que se produce con el valor 1, agrega la cantidad de datos transferidos entre la CPU y el coprocesador, en bytes.
3 Además de la información producida con el valor 2, da más detalles so- bre la actividad de descarga, incluyendo inicialización de dispositivos, y transferencia de variables individuales.

Esta API habilita el tiempo en las regiones de descarga y mide la cantidad de datos transferidos durante su ejecución. El uso de esta API equivale a establecer la variable de entorno OFFLOAD_REPORT.

Apéndice B Códigos para el Stencil de difusión de calor

A continuación vamos a mostrar todas las versiones que hemos desarrollado para el caso del kernel de difusión de calor de 11 puntos.

Abbildung in dieser Leseprobe nicht enthalten

Figura B.1: Código base, no vectorizado.

Abbildung in dieser Leseprobe nicht enthalten

Figura B.2: Código con optimizaciones escalares.

Abbildung in dieser Leseprobe nicht enthalten

Figura B.3: Código vectorizado utilizando restrict y collapse.

Abbildung in dieser Leseprobe nicht enthalten

Figura B.4: Código utilizando tamaño de paginas grandes (huge).

Abbildung in dieser Leseprobe nicht enthalten

Figura B.5: Código con (Streaming stores).

Abbildung in dieser Leseprobe nicht enthalten

Figura B.6: Código con políticas de planificación del bucle OpenMP (Schedule).

Abbildung in dieser Leseprobe nicht enthalten

Figura B.7: Código con Optimización del acceso a memoria, mediante la técnica de Blocking.

Abbildung in dieser Leseprobe nicht enthalten

Figura B.8: Código que corre en el host en modo offload en una sola tarjeta (mic2).

Abbildung in dieser Leseprobe nicht enthalten

Figura B.9: Código que utiliza 2 hilos dentro del bucle global de cálculo, descar- gando en mic1 y mic2.

Abbildung in dieser Leseprobe nicht enthalten

Figura B.10: Código que utiliza 2 hilos con región paralela y barrera explicita.

Abbildung in dieser Leseprobe nicht enthalten

Figura B.11: Código usando un solo hilo y manejando las tarjetas asíncronamente.

Apéndice C Lista de siglas

3D-FD : 3-D Finite Difference, Diferencias Finitas 3-D.

AI : Arithmetic Intensity, Intensidad Aritmética.

API : Application Programming Interface, Interfaz de Programación de Aplicación .

AVX : Advanced Vector Extensions, Extensiones de Vector Avanzadas.

COI : Coprocessor Offload Infrastructure, Infraestructura de Descarga del Coproce- sador.

CPU : Central Processing Unit, Unidad Central de Procesamiento.

CUDA : Compute Unified Device Architecture, Arquitectura Unificada de Dispositi- vos de Cómputo.

DDR : Double Data Rate, Doble Transferencia de Datos.

DEM : Discrete Element Method, Método de Elementos Discretos.

DMA : Direct Memory Access, Acceso Directo a Memoria.

DP : Double Precision, Doble precisión.

DRAM : Dynamic Random Access Memory, Memoria de Acceso Dinámico Aleato- rio.

ECC : Error Check and Correct, Detección y Corrección de Errores.

EDP : Energy Delay Product, Producto de Retardo por Energía.

EMU : Extended Math Unit, Unidad Matemática Extendida.

FMA : Fused Multiply-Add, Suma-Multiplicación Fusionada.

GAP : Guided Auto-Parallelization, Auto-Paralelización Guiada.

GPGPUs : General Purpose Computing on GPUs, Computación de Propósito Gene- ral en GPUs.

GPUs : Graphics Processing Units, Unidades de Procesamiento Gráfico.

GSRB : Gauss-Seidel Red-Black, Gauss-Seidel Negro-Rojo.

GUI : Graphical User Interface, Interfaz Gráfica de Usuario.

HPC : High Performance Computing, Computación de Altas Prestaciones.

IMCI : Intel Initial Many Core Instructions, Instrucciones Intel para Muchos Nú- cleos.

ISA : Instruction Set Architecture, Arquitectura del Repertorio de Instrucciones.

KNC : Knights Corner, Nombre clave de la arquitectura Intel MIC actual.

KNL : Knights Landing, Nombre clave de la arquitectura Intel MIC nueva.

LEO : Language Extension Offload, Lenguaje de Extensiones para Descarga.

MAD : Multiply-Add, Suma-Multiplicación Fusionada (como FMA).

MIC : Many Integrated Core, Muchos Núcleos Integrados.

Intel MKL : Intel Math Kernel Library, Librera de Rutinas Matemáticas de Intel.

MLP : Memory Level Parallelism, Paralelismo de Nivel de Memoria.

MPI : Message Passing Interface, Interfaz de Paso de Mensajes.

MPSS : Manycore Platform Software Stack, Pila de Programa de la Plataforma de Muchos Núcleos.

MSRs : Model Specific Registers, Registros Específicos del modelo.

MYO : Mine-Yours-Ours.

NFS : Network File System, Sistema de Ficheros de Red.

NVML : Nvidia Management Library, Biblioteca de Administración de Nvidia.

ODI : on-die interconnect, Interconexión dentro del chip.

PAPI : Performance Application Programming Interface, Interfaz de Programación de Aplicaciones de Rendimiento.

PCIe : Peripheral Component Interconnect express, Interconector de componentes periféricos rápido.

PDEs : Partial Differential Equations, Ecuaciones Diferenciales Parciales.

RAM : Random Access Memory, Memoria de Acceso Aleatorio.

SSE : Streaming SIMD Extensions, Extensiones SIMD.

SIMD : Simple Instruction Multiple Data, Instrucciones Sencillas Múltiples Datos.

SP : Single Precision, Simple Precisión.

TCP/IP : Transmission Control Protocol/Internet Protocol, Protocolo de Control de Transmisión/Protocolo de Internet.

TD : Tag Directory, Directorio de Etiquetas.

TDP : Thermal Design Point, Punto de Diseño Térmico.

TLB : Translation Lookside Buffer, Búfer de Traducción de Direcciones.

VPU : Vector Processing Units, Unidades de Procesamiento Vectorial.

Apéndice D Fuentes de las Figuras

Figura 1.1: Top 500 website. Performance Development. 2016.

Figura 2.1: Intel®. Xeon PhiTM Coprocessor. 2015.

Figura 2.2: Intel®. Ecosistema Xeon PhiTM. 2015.

Figura 2.3: Intel®. Esquema interconexión en Xeon PhiTM. 2015.

Figura 2.4: Intel®. Esquema de un Core en Xeon PhiTM. 2015.

Figura 2.5: Creada por los autores. 2016.

Figura 2.6: Creada por los autores. 2016.

Figura 2.7: Creada por los autores. 2016.

Figura 2.8: Intel®. Herramientas de desarrollo. 2015.

Figura 2.9: Nvidia®. Grupo de hilos CUDATM. 2011.

Figura 2.10: Nvidia®. Grid CUDATM. 2011.

Figura 2.11: Nvidia®. Modelo CUDATM. 2011.

Figura 2.12: Intel®. Esquema Knights LandingTM. 2016.

Figura 2.13: Creada por los autores. 2016.

Figura 2.14: Creada por los autores. 2016.

Figura 2.15: Creada por los autores. 2016.

Figura 2.16: Creada por los autores. 2016.

Figura 2.17: Creada por los autores. 2016.

Figura 2.18: Creada por los autores. 2016.

Figura 3.1: Creada por los autores. 2016.

Figura 3.2: Creada por los autores. 2016.

Figura 3.3: Creada por los autores. 2016.

Figura 3.4: Creada por los autores. 2016.

Figura 3.5: Autor Desconocido.

Figura 3.6: Creada por los autores. 2016.

Figura 3.7: Creada por los autores.2016.

Figura 3.8: Creada por los autores.2016.

Figura 3.9: Creada por los autores.2016

Figura 3.10: Creada por los autores.2016.

Figura 3.11: Creada por los autores.2016.

Figura 3.12: Creada por los autores.2016.

Figura 4.1: Creada por los autores.2016.

Figura 4.2: Creada por los autores.2016.

Figura 4.3: Creada por los autores.2016.

Figura 4.4: Creada por los autores.2016.

Figura 4.5: Creada por los autores.2016.

Figura 4.6: Creada por los autores.2016.

Figura 4.7: Creada por los autores.2016.

Figura 4.8: Creada por los autores.2016.

Figura 4.9: Creada por los autores.2016.

Figura 4.10: Creada por los autores.2016.

Figura 4.11: Creada por los autores.2016.

Figura 4.12: Creada por los autores.2016.

Figura 5.1: Creada por los autores.2016.

Figura 5.2: Creada por los autores.2016.

Figura 5.3: Creada por los autores.2016.

Figura 5.4: Creada por los autores.2016.

Figura 5.5: Creada por los autores. 2016.

Figura 5.6: Creada por los autores. 2016.

Figura 5.7: Creada por los autores. 2016.

Figura 5.8: Creada por los autores. 2016.

Figura 5.9: Creada por los autores. 2016.

Figura 5.10: Creada por los autores. 2016.

Figura 5.11: Creada por los autores. 2016.

Figura 5.12: Creada por los autores. 2016.

Figura 5.13: Creada por los autores. 2016.

Figura 5.14: Creada por los autores. 2016.

Figura 5.15: Creada por los autores. 2016.

Figura 5.16: Creada por los autores. 2016.

Figura 5.17: Creada por los autores. 2016.

Figura 5.18: Creada por los autores. 2016.

Figura 5.19: Creada por los autores. 2016.

Figura 5.20: Creada por los autores. 2016.

Figura 5.21: Creada por los autores. 2016.

Figura 5.22: Creada por los autores. 2016.

Figura 5.23: Creada por los autores. 2016.

Figura 5.24: Creada por los autores. 2016.

Figura A.1: Creada por los autores. 2016.

Figura A.2: Creada por los autores. 2016.

Figura A.3: Creada por los autores. 2016.

Figura A.4: Creada por los autores. 2016.

Figura A.5: Creada por los autores. 2016.

Figura A.6: Creada por los autores. 2016.

Figura A.7: Creada por los autores. 2016.

Figura A.8: Creada por los autores. 2016.

Figura A.9: Creada por los autores. 2016.

Figura A.10: Creada por los autores. 2016.

Figura A.11: Creada por los autores. 2016.

Figura A.12: Creada por los autores. 2016.

Figura A.13: Creada por los autores. 2016.

Figura A.14: Creada por los autores. 2016.

Figura B.1: Creada por los autores. 2016.

Figura B.2: Creada por los autores. 2016.

Figura B.3: Creada por los autores. 2016.

Figura B.4: Creada por los autores. 2016.

Figura B.5: Creada por los autores. 2016.

Figura B.6: Creada por los autores. 2016.

Figura B.7: Creada por los autores. 2016.

Figura B.8: Creada por los autores. 2016.

Figura B.9: Creada por los autores. 2016.

Figura B.10: Creada por los autores. 2016.

Figura B.11: Creada por los autores. 2016.

Referencias

1 Nvidia Corporation. NVML API Reference, note = http://developer. download.nvidia.com/assets/cuda/files/cudadownloads/nvml/nvml. pdf, year =[last access 15 May 2015].

2 Top 500 supercomputer site. http://www.top500.org/[último acceso 15 Nov. 2015]. 1.1

3 Top Green500 List. http://www.green500.org/[último acceso 15 Nov 2015]. 1.2.1

4 C. Andreolli, P. Thierry, L. Borges, G. Skinner, and C. Yount. Chapter 23 - characterization and optimization methodology applied to stencil compu- tations. In J. Reinders and J. Jeffers, editors, High Performance Parallelism Pearls: Multicore and Many-core Programming Approaches, volume 1, pages 377 - 396. Morgan Kaufmann, Boston, MA, USA, 2015. 3.1.2, 3.5.1.2

5 K. Asanovic, R. Bodik, B. C. Catanzaro, J. J. Gebis, P. Husbands, K. Keutzer, D. A. Patterson, W. L. Plishker, J. Shalf, S. W. Williams, et al. The landscape of parallel computing research: A view from berkeley. Technical report, Technical Report UCB/EECS-2006-183, EECS Department, University of California, Berkeley, 2006. 1.1, 1.2.1

6 K. Bergman, S. Borkar, D. Campbell, W. Carlson, W. Dally, M. Denneau, P. Franzon, W. Harrod, K. Hill, J. Hiller, et al. Exascale computing study: Technology challenges in achieving exascale systems. Defense Advanced Research Projects Agency Information Processing Techniques Office (DARPA IPTO), Tech. Rep, 15, 2008. 1.2.1

7 A. Billon, J. Picaut, C. Foy, V. Valeau, and A. Sakout. Introducing atmospheric attenuation within a diffusion model for room-acoustic predictions. The Journal of the Acoustical Society of America, 123 (6):4040-4043, 2008. 2.6.2.2

8 W. M. Brown, J.-M. Y. Carrillo, N. Gavhane, F. M. Thakkar, and S. J. Plimpton. Optimizing legacy molecular dynamics software with directive-based offload. Computer Physics Communications, 195:95 - 101, 2015. 5.1.2

9 J. M. Cebrian and L. Natvig. Temperature effects on on-chip energy measu- rements. In Proceedings IGCC 2013, Arlington, USA, 2013, pages 1-6, 2013. 4.4

10 J. M. Cecilia, J. L. Abellán, J. Fernández, M. E. Acacio, J. M. García, and M. Ujaldón. Stencil computations on heterogeneous platforms for the jacobi method: GPUs versus Cell BE. The Journal of Supercomputing, 62(2):787-803, 2012. 4.1.2

11 M. Christen, O. Schenk, and H. Burkhart. Patus: A code generation and autotuning framework for parallel iterative stencil computations on modern microarchitectures. In Parallel & Distributed Processing Symposium (IPDPS), pages 676-687. IEEE, 2011. 2.6, 4.1.2

12 K. Datta, M. Murphy, V. Volkov, S. Williams, J. Carter, L. Oliker, D. Patterson, J. Shalf, and K. Yelick. Stencil Computation Optimization and Auto-tuning on State-of-the-art Multicore Architectures. In Proceedings of the ACM/IEEE Conference on Supercomputing, SC ’08, page 4, 2008. 2.6, 4.1.2

13 R. de la Cruz and M. Araya-Polo. Modeling stencil computations on modern HPC architectures. In 5th Int. Workshop (PMBS14) held as part of SC14. Springer, 2014. 3.1.2

14 E. C. Dufort and S. P. Frankel. Stability conditions in the numerical treatment of parabolic differential equations. Mathematical tables and others aids to computation, 7:135-152, 1953. 2.6.2.2

15 J. Fang, H. Sips, L. Zhang, C. Xu, Y. Che, and A. L. Varbanescu. Test- driving Intel Xeon Phi. In Proc. of the 5th ACM/SPEC Int. Conf. on Performance Engineering, pages 137-148. ACM, 2014. 3.1.2

16 J. Fang, A. L. Varbanescu, H. Sips, L. Zhang, Y. Che, and C. Xu. An Empirical Study of Intel Xeon Phi. arXiv preprint arXiv:1310.5842, 2013. 3.1.2

17 J. Fang, A. L. Varbanescu, H. Sips, L. Zhang, Y. Che, and C. Xu. Benchmar- king intel xeon phi to guide kernel design. Delft University of Technology Parallel and Distributed Systems Report Series, PDS-2013-005, 2013. 4.1.2

18 L. Feng. Chapter 20 - data transfer using the intel coi library. In J. Reinders and J. Jeffers, editors, High Performance Parallelism Pearls: Multicore and Many- core Programming Approaches, volume 1, pages 341 - 348. Morgan Kaufmann, Boston, MA, USA, 2015. 5.1.2

19 M. Frigo and V. Strumpen. Cache Oblivious Stencil Computations. In Proc. of the 19th Annual International Conference on Supercomputing, ICS ’05, pages 361-366, New York, USA, 2005. ACM. 2.6

20 M. Garland, S. Le Grand, J. Nickolls, J. Anderson, J. Hardwick, S. Morton, E. Phillips, Y. Zhang, and V. Volkov. Parallel Computing Experiences with CUDA. IEEE, Micro, 28(4):13-27, July 2008. 1.1

21 M. Hernández, J. M. Cebrian, J. M. Cecilia, and J. M. García. Evaluating 3-d stencil codes on intel xeon phi: Limitations and trade-offs. In XXVI Jornadas de Paralelismo, pages 568-573, Córdoba (Spain), Sept. 2015. Universidad de Córdoba. 2

22 M. Hernández, J. M. Cebrián, J. M. Cecilia, and J. M. García. Evaluation of 3-d stencil codes on the intel xeon phi coprocessor. In Parallel Computing (ParCo) 2015, Edimburgo, Reino Unido, Sept. 2015. 3

23 M. Hernández, J. M. García, and J. M. Cecilia. Early experiences running the 3d stencil jacobi method in intel xeon phi. In H. Weghorn, editor, 11th InternationalL Conference. Applied Computing 2014, pages 85-92, Porto (Portu- gal), Oct. 2014. International association for development of the information society. 1

24 M. Hernández, B. Imbernón, J. M. Navarro, J. M. García, J. M. Cebrián, and J. M. Cecilia. Evaluation of the 3-d finite difference implementation of the acoustic diffusion equation model on massively parallel architectures. Computers & Electrical Engineering, 46:190 - 201, 2015. 4.6.6, 1

25 J. Holewinski, L.-N. Pouchet, and P. Sadayappan. High-performance code generation for stencil computations on gpu architectures. In Proceedings of the 26th ACM international conference on Supercomputing, pages 311-320. ACM, 2012. 4.1.2

26 J. Jeffers and J. Reinders. Chapter 4 - driving around town: Optimizing a real-world code example. In Intel Xeon Phi Coprocessor High Performance Programming, pages 83 - 106. Morgan Kaufmann Publishers Inc., Boston, MA, USA, 2013. 3.1.2

27 J. Jeffers and J. Reinders. Chapter 7 - offload. In Intel Xeon Phi Coprocessor High Performance Programming, pages 189 - 241. Morgan Kaufmann Publishers Inc., Boston, MA, USA, 2013. 5.1.2

28 J. Jeffers and J. Reinders. Intel Xeon Phi Coprocessor High Performance Programming. Morgan Kaufmann Publishers Inc., Boston, MA, USA, 2013. 1.1, 2.2.2, 3.1.2

29 Y. Jing and N. Xiang. On boundary conditions for the diffusion equation in room acoustic predictions: Theory, simulations, and experiments. J. Acoust. Soc. Am., 123:145-153, 2008. 2.6.2.2

30 S. Kamil, C. Chan, L. Oliker, J. Shalf, and S. Williams. An auto-tuning framework for parallel multicore stencil computations. In Proc. of the Int. Symp. on Parallel & Distributed Processing (IPDPS), pages 1-12. IEEE, 2010. 2.6, 4.1.2

31 S. Kamil, K. Datta, S. Williams, L. Oliker, J. Shalf, and K. Yelick. Implicit and Explicit Optimizations for Stencil Computations. In Proc. of the Workshop on Memory System Performance and Correctness, MSPC ’06, pages 51-60, New York, USA, 2006. ACM. 4.1.2

32 S. Kamil, P. Husbands, L. Oliker, J. Shalf, and K. Yelick. Impact of modern memory subsystems on cache optimizations for stencil computations. In Proceedings of the 2005 Workshop on Memory System Performance, MSP ’05, pages 36-43, New York, NY, USA, 2005. ACM. 3.1.2

33 S. W. Keckler, W. J. Dally, B. Khailany, M. Garland, and D. Glasco. Gpus and the future of parallel computing. IEEE Micro, 31(5):7-17, 2011. 1.2.1

34 P. Kogge, K. Bergman, S. Borkar, D. Campbell, W. Carson, W. Dally, M. Den- neau, P. Franzon, W. Harrod, K. Hill, et al. Exascale computing study: Technology challenges in achieving exascale systems. 2008. 1.2.1

35 D. Komatitsch, G. Erlebacher, D. Göddeke, and D. Michéa. High-order finite-element seismic wave propagation modeling with MPI on a large GPU cluster. Journal of computational physics, 229(20):7692-7714, 2010. 2.6

36 S. Krishnamoorthy, M. Baskaran, U. Bondhugula, J. Ramanujam, A. Rountev, and P. Sadayappan. Effective Automatic Parallelization of Stencil Compu- tations. In Proc. of the ACM SIGPLAN Conf. on Programming Language Design and Implementation, PLDI ’07, pages 235-244, New York, USA, 2007. ACM. 4.1.2

37 S. Kronawitter and C. Lengauer. Optimization of two Jacobi Smoother Kernels by Domain-Specific Program Transformation. In Proceedings of the 1st International Workshop on High-Performance Stencil Computations, pages 75-80, 2014. 2.6

38 Z. Li and Y. Song. Automatic tiling of iterative stencil loops. ACM Transactions on Programming Languages and Systems (TOPLAS), 26(6):975-1028, 2004. 4.1.2

39 H. Markram, K. Meier, T. Lippert, S. Grillner, R. Frackowiak, S. Dehaene, A. Knoll, H. Sompolinsky, K. Verstreken, J. DeFelipe, et al. Introducing the human brain project. Procedia Computer Science, 7:39-42, 2011. 1.2.1

40 M. McCool, A. Robison, and J. Reinders. Structured Parallel Programming: Patterns for Efficient Computation. Morgan Kaufmann Publishers Inc., Boston, MA, USA, 2012. 3.1.2

41 M. McCool, A. D. Robison, and J. Reinders. Chapter 7 - stencil and recurrence. In Structured Parallel Programming: Patterns for Efficient Computation, pages 199 - 207. Morgan Kaufmann Publishers Inc., Boston, MA, USA, 2012. 3.1.2

42 P. Micikevicius. 3D finite difference computation on GPUs using CUDA. In Proceedings of 2nd Workshop on General Purpose Processing on Graphics Processing Units, pages 79-84. ACM, 2009. 4.1.2

43 P. M. Morse and H. Feshbach. Methods of Theoretical Physics. McGraw-Hill, New York, 1953. 2.6.2.2

44 P. J. Mucci, S. Browne, C. Deane, and G. Ho. PAPI: A portable interface to hardware performance counters. In Proc. of HPCMP Users Group Conf., 1999, pages 7-10, 1999. 2.7.3, 4.4, 4

45 J. M. Navarro, J. Escolano, and J. J. López. Implementation and evaluation of a diffusion equation model based on finite difference schemes for sound field prediction in rooms. Applied Acoustics, 73(6-7):659 - 665, 2012. 2.6.2.2, 3

46 J. M. Navarro, F. Jacobsen, J. Escolano, and J. J. López. A theoretical approach to room acoustic simulations based on a radiative transfer model. Acta Acustica United with Acustica, 96:1078-1089, 2010. 2.6.2.2

47 A. Nguyen, N. Satish, J. Chhugani, C. Kim, and P. Dubey. 3.5-D blocking optimization for stencil computations on modern CPUs and GPUs. In Proceedings of the 2010 ACM/IEEE International Conference for High Performance Computing, Networking, Storage and Analysis, pages 1-13. IEEE Computer Society, 2010. 4.1.2

48 NVIDIA. NVIDIA CUDA C Programming Guide 6.5. 2014. 2.4

49 J. Peraza, A. Tiwari, M. Laurenzano, L. Carrington, W. Ward, and R. Camp- bell. Understanding the performance of stencil computations on Intel’s Xeon Phi. In Int. Conf. on Cluster Computing (CLUSTER), pages 1-5, Sept 2013. 3.1.2

50 J. Picaut, L. Simon, and J.-D. Polack. A mathematical model of diffuse sound field based on a diffusion equation. Acta Acust. united Ac., 83(4):614-621, 1997. 4.1.2

51 S. Potluri, A. Venkatesh, D. Bureddy, K. Kandalla, and D. K. Panda. Efficient intra-node communication on intel-mic clusters. In 13th Int. Symp. on Cluster, Cloud and Grid Computing (CCGrid), pages 128-135. IEEE/ACM, 2013. 3.1.2

52 R. Rahman. Intel Xeon Phi Coprocessor Architecture and Tools: The Guide for Application Developers. Apress, Berkely, CA, USA, 1st edition, 2013. 1.1, 2.2.1, 2.2.2, 3.1.2, 5.1.2

53 S. M. F. Rahman, Q. Yi, and A. Qasem. Understanding Stencil Code Perfor- mance on Multicore Architectures. In Proceedings of the 8th ACM International Conference on Computing Frontiers, CF ’11, pages 30:1-30:10, New York, NY, USA, 2011. ACM. 3.1.2

54 J. Reinders. A conversation with james reinders. HPCwire: Global News and Information on High Performance Computer, january 2016. 1.2.2

55 J. Reinders and J. Jeffers. High Performance Parallelism Pearls, Multicore and Many-core Programming Approaches, chapter Characterization and Auto- tuning of 3DFD, pages 377-396. Morgan Kaufmann, 2014. 2.6.3, 3.5.1.3

56 J. Reinders and J. Jeffers, editors. High Performance Parallelism Pearls: Multicore and Many-core Programming Approaches, volume 2. Morgan Kaufmann Publishers Inc., Boston, MA, USA, 1st edition, 2015. 1.2.3, 3.1.2, 5.1.2

57 J. Reinders and J. Jeffers, editors. High Performance Parallelism Pearls: Multicore and Many-core Programming Approaches, volume 1. Morgan Kaufmann Publishers Inc., Boston, MA, USA, 1st edition, 2015. 1.2.3, 3.1.2, 5.1.2

58 G. Rivera and C.-W. Tseng. Tiling Optimizations for 3D Scientific Compu- tations. In Supercomputing, ACM/IEEE Conference, pages 32-32, Nov 2000. 4.1.2

59 N. Satish, C. Kim, J. Chhugani, H. Saito, R. Krishnaiyer, M. Smelyanskiy, M. Girkar, and P. Dubey. Can traditional programming bridge the ninja performance gap for parallel computing applications? In Proceedings of the 39th Annual International Symposium on Computer Architecture, ISCA ’12, pages 440-451, Washington, DC, USA, 2012. IEEE Computer Society. 1.2.3

60 M. Seaton, L. Mason, Z. A. Matveev, and S. Blair-Chappell. Chapter 23 - vectorization advice. In J. Reinders and J. Jeffers, editors, High Performance Parallelism Pearls Volume Two: Multicore and Many-core Programming Approa- ches, volume 2, pages 441 - 462. Morgan Kaufmann, Boston, MA, USA, 2015. 4.1.2

61 J. Shalf, S. Dosanjh, and J. Morrison. Exascale computing technology cha- llenges. In Proceedings of the 9th International Conference on High Performance Computing for Computational Science, VECPAR’10, pages 1-25, Berlin, Heidel- berg, 2011. Springer-Verlag. 1.2.1

62 B. Shareef, E. de Doncker, and J. Kapenga. Monte carlo simulations on intel xeon phi: Offload and native mode. In High Performance Extreme Computing Conference (HPEC), 2015 IEEE, pages 1-6, Sept 2015. 5.1.2

63 T. Siauw and A. M. Bayen. An introduction to MATLAB programming and numerical methods. Elsevier, 2015. 3.1

64 Y. Song, R. Xu, C. Wang, and Z. Li. Data locality enhancement by memory reduction. In Proc. of the 15th int. conf. on Supercomputing, pages 50-64. ACM, 2001. 4.1.2

65 R. Strzodka, M. Shaheen, D. Pajak, and W. Pomeranian. Impact of system and cache bandwidth on stencil computations across multiple processor generations. In Proceedings of the Workshop on Applications for Multi-and Many-Core Processors (A4MMC) at ISCA, 2011. 3.1.2

66 L. Szustak, K. Rojek, R. Wyrzykowski, and P. Gepner. Toward efficient distribution of MPDATA stencil computation on Intel MIC architecture. Proce. HiStencils, 14:51-56, 2014. 2.6

67 Y. Tang, R. A. Chowdhury, B. C. Kuszmaul, C.-K. Luk, and C. E. Leiserson. The pochoir stencil compiler. In Proceedings of the Twenty-third Annual ACM Symposium on Parallelism in Algorithms and Architectures, SPAA ’11, pages 117-128, New York, NY, USA, 2011. ACM. 2.6, 4.1.2

68 U. Trottenberg, C. W. Oosterlee, and A. Schuller. Multigrid. Academic press, 2000. 2.6, 3.1

69 V. Valeau, J. Picaut, and M. Hodgson. On the use of a diffusion equation for room-acoustic prediction. J. Acoust. Soc. Am., 119(3):1504-1513, 2006. 2.6.2.2

70 A. Vladimirov, R. Asai, and V. Karpusenko, editors. Parallel Programming and Optimization with Intel Xeon Phi Coprocessors, volume 1. Colfax International, CA, USA, 2da edition, 2015. 3.1.2

71 K. V. Vladimirov Andrey. Test-driving Intel Xeon Phi coprocessors with a basic N-body simulation. Coflax International, 2013. 3.1.2

72 E. Wang, Q. Zhang, B. Shen, G. Zhang, X. Lu, Q. Wu, and Y. Wang. High- Performance Computing on the Intel Xeon Phi. Springer, 2014. 5.1.2

73 F. Wende, M. Klemm, T. Steinke, and A. Reinefeld. Chapter 12 - concurrent kernel offloading. In J. Reinders and J. Jeffers, editors, High Performance Parallelism Pearls: Multicore and Many-core Programming Approaches, volume 1, pages 201 - 223. Morgan Kaufmann, Boston, MA, USA, 2015. 5.1.2

74 S. Williams, A. Waterman, and D. Patterson. Roofline: An insightful visual performance model for multicore architectures. Commun. ACM, 52(4):65-76, Apr. 2009. 3.5.1.3

75 C. J. Wright. Chapter 14 - power analysis on the intel xeon phi coprocessor. In J. Reinders and J. Jeffers, editors, High Performance Parallelism Pearls: Multicore

76 V. Zhukov, M. Krasnov, N. Novikova, and O. Feodoritova. Multigrid effec- tiveness on modern computing architectures. Programming and Computer Software, 41(1):14-22, 2015. 2.6

77 Y. Zou and S. Rajopadhye. Cache efficient parallelizations for uniform dependence computations. (TR/CS-14-101), May 2014. 4.1.2

78 G. Zumbusch. Vectorized higher order finite difference kernels. In Applied Parallel and Scientific Computing, pages 343-357. Springer, 2012. 4.1.2

[...]

---

¹ La lista es actualizada 2 veces por año, en mayo y noviembre.

¹ Para asegurar que el compilador no vectorizaba, se compilaron los códigos con el modificador-no-vec.

¹ Este parámetro opcional debe ser un entero positivo.

² Reordenar bucles.

³ Cambiando el diseño de datos.

⁴ Model Specific Registers.

⁵ Comparada con la implementación en Xeon con 32 hilos.

⁶ Utilizando la métrica EDP.