Optimización de aplicaciones de procesado de señales digitales empleando como plataforma hardware NVIDIA Jetson TK1

Departamento de Sistemas Inform´ aticos y Computaci´ on M´ ASTER EN COMPUTACI´ ON PARALELA Y DISTRIBUIDA TESIS DE M´ ASTER Optimizaci´ on de aplicaciones de procesado de se˜ nales digitales empleando como plataforma hardware NVIDIA Jetson TK1 VALENCIA,A21 DE SEPTIEMBRE DE 2015 TESIS DE M´ ASTER PRESENTADA POR: DIRIGIDA POR: FRANCISCO J. ALVENTOSA RUEDA PEDRO ALONSO JORD´ A Resumen Hoy en d´ ıa, las tabletas y los smartphones est´ an equipados con procesadores de bajo consumo energ´ etico, como las series de arquitecturas ARMv7 y ARMv8. Estos procesadores tienen una unidad de SIMD potente, que permiten explotar el paralelismo intr´ ınseco de los datos que tienen la mayoria de aplicaciones multimedia y de procesamiento de la se˜ nal. En el procesamiento de la se˜ nal ac´ ustica, existen m´ ultiples aplicaciones que requieren el uso de operaciones de filtrado como son entre otros, las ecualizaciones o los sintetizadores de se˜ nal. La mayor´ ıa de estas aplicaciones se pueden ejecutar en dispositivos m´ oviles, sin embargo, una correcta gesti´ on de los diferentes recursos computacionales CPU’s and GPU’s, as´ ı como las unidades SIMD (Simple Instruction - Multiple Data) que incorporan estos, se convierte en una ardua tarea para el programador. En este documento, hablamos sobre la ejecuci´ on eficiente de filtrado multi-canal de se˜ nales de audio. Con este fin, analizamos tres estructuras comunes de filtrado de audio: FIR, IIRI y PIIR. Adem´ as, se ha desarrollado una implementaci´ on en C del algoritmo Beamformer, empleando librer´ ıas de altas prestaciones como el BLAS y LAPACK optimizado (ATLAS), el PLASMA y el CUBLAS. Palabras clave Procesadores de bajo consumo, ARMv7, ARM Cortex-A15, intrinsics NEON, procesamiento de audio, ´ algebra lineal densa, computaci´ on de alto rendimiento, Jetson TK1. Abstract Tablets and smart phones are equipped nowadays with low power processor architectures such as the ARMv7 and the ARMv8 series. These processors integrate powerful SIMD units to exploit the intrinsic data-parallelism of most media and signal processing applications. In audio signal processing, there exist multiple applications that require the use of filtering operations such as equalizations or signal synthesizers, among others. Most of these applications can be executed today on mobile devices via efficient exploitation of computer resources such as CPU’s, GPU’s and CPU feature SIMD (simple-Instruction Multiple-Data) units. In this paper, we target the implementation of multi-channel filtering of audio signals on ARM architectures. To this end, three common audio filter structures are considered: FIR, IIRI and PIIR. In addition, an implementation has been developed in C of the algorithm Beamformer using libraries of high performance such as BLAS and LAPACK optimized (ATLAS), PLASMA and CUBLAS. Key words Low Power Processors, ARMv7, ARM Cortex-A15, NEON Intrinsics, Audio Processing, Dense linear algebra, High-performance computing, Jetson TK1. Agradecimientos Las sisuientes l´ ıneas, las quiero dedicar para dar un sincero agradecimiento a las personas que han puesto todo su empe˜ no, dedicaci´ on y apoyo para que este trabajo haya podido hacerse realidad: Pedro Alonso Jord´ a Antonio Manuel Vidal Maci´ a Alberto Gonz´ alez Salvador Gemma Pi˜ nero Sip´ an Jose Antonio Belloch Rodriguez Sin olvidar a mis personas m´ as allegadas y en especial a mi pareja Tamara Reig, que en estos meses de intenso trabajo me han ayudado y mostrado su apoyo incondicional. Gracias a todos. ´ Indice general 1. Introducci´ on general 11 1.1. Motivaci´ on ............................................ 11 1.2. Motivavi´ on personal ....................................... 12 1.3. Objetivos ............................................. 12 1.4. Hardware ............................................. 12 1.5. Software .............................................. 13 1.6. Librerias .............................................. 14 1.7. Estructura del trabajo ...................................... 14 2. Puesta en marcha y configuraci´ on NVIDIA Jetson TK1 17 2.1. Puesta en marcha ......................................... 17 2.1.1. Conexiones ........................................ 18 2.1.2. Flashing .......................................... 18 2.1.3. Configuraciones b´ asicas recomendadas ......................... 19 2.1.4. Instalaci´ on de CUDA ................................... 19 2.1.5. Instalaci´ on de ATLAS ................................... 20 2.1.6. Instalaci´ on de PLASMA .................................. 21 3. Filtrado de sonido: FIR, IIRI y PIIR 23 3.1. Filtros FIR ............................................. 23 3.2. Filtros IIR ............................................. 24 3.3. Configuraci´ on del sistema e implementaci´ on .......................... 25 3.3.1. Configuraci´ on del sistema ................................ 25 3.3.2. Implementaci´ on ..................................... 26 3.4. An´ alisis de resultados ...................................... 28 4. Beamformer 33 4.1. Modelo de se˜ nal ......................................... 33 4.2. Algoritmos Beamforming ..................................... 34 4.3. Configuraci´ on del sistema e implementaci´ on .......................... 36 4.3.1. Configuraci´ on del sistema ................................ 36 5 12 CAP´ ITULO 1. INTRODUCCI´ ON GENERAL cambio, para optimizar la aplicaci´ on de separaci´ on de se˜ nales, se ha trabajado a un nivel superior. En este caso se han empleado librer´ ıas de altas prestaciones como ATLAS [12], PLASMA [13] y CUBLAS [14] con el fin de mejorar el tiempo de c´ omputo empleando todos los elementos computacionales disponibles en el dispositivo: cores CPU y la GPU. Resaltar que la librer´ ıa ATLAS, hace uso en su implementaci´ on de las instrucciones NEON. 1.2. Motivavi´ on personal A nivel personal, como estudiante de master en computaci´ on paralela y distribuida en la rama de computaci´ on cient´ ıfica y futuro alumno de doctorado en inform´ atica, la elecci´ on de este TFM viene dada por la incre´ ıble oportunidad que se me dio durante mis estudios de m´ aster de poder formar parte en un grupo de investigaci´ on como es el INCO2 del DSIC y poder iniciarme de este modo en el campo de la investigaci´ on. Dicho grupo de investigaci´ on mantiene una estrecha relaci´ on con el instituto iTeam (Instituto de Telecomunicaciones y Aplicaciones Multimedia), especializado en mejorar aplicaciones como las que he tratado en este trabajo. 1.3. Objetivos El objetivo principal de esta investigaci´ on es el de emplear todas las t´ ecnicas y librer´ ıas de computaci´ on de altas prestaciones con el fin de reducir al m´ aximo el tiempo de procesamiento de dos aplicaciones particulares de procesado de se˜ nal. Este objetivo principal se podr´ ıa subdividir y detallar en los siguientes objetivos secundarios: Configurar los dispositivos hardware que se emplean, en este caso el NVIDIA Jetson TK1. Instalar y configurar los paquetes necesarios para emplear las librer´ ıas de altas prestaciones: CUBLAS, ATLAS y PLASMA. Reescribir y reorganizar c´ odigo para mejorar el tiempo de computo. En particular, los c´ odigos facilitados est´ an escritos en Matlab y deben ser reescritos al lenguaje C/C++. Emplear las instrucciones SIMD de la arquitectura ARM (NEON) para reducir el tiempo de c´ omputo. Emplear librer´ ıas de computaci´ on de altas prestaciones como CUBLAS, ATLAS, PLASMA, para reducir el tiempo de c´ omputo. 1.4. Hardware El hardware empleado en la realizaci´ on de este proyecto, han sido las nuevas placas de NVIDIA llamadas Jetson [8]. Estas placas incorporan el chip Tegra K1 de NVIDIA (incluido en una gran variedad de dispositivos m´ oviles como la tableta NVIDIA SHIELD) el cual incluye cuatro n´ ucleos ARM CortexA15, una GPU de 192 n´ ucleos Kepler y un quinto n´ ucleo bastante limitado ARM para cuando la carga del sistema es muy baja reducir al m´ aximo el consumo energ´ etico. Adem´ as, incluye dos GygaBytes de memoria principal compartida por los procesadores ARM y la gr´ afica, adem´ as de una gran variedad de conexiones, que se enumeran con m´ as detalles a continuaci´ on, con el fin de hacer al Jetson de una herramienta lo m´ as ´ util posible. Todas estas caracter´ ısticas, hacen al NVIDIA Jetson una herramienta muy ´ util para una gran variedad de aplicaciones que necesiten de un buen rendimiento computacional con un bajo consumo de energ´ ıa. A continuaci´ on se realiza una enumeraci´ on detallada de las caracter´ ısticas de la placa NVIDIA Jetson TK1 que podemos observar en la Figura 1.1. 1.5. SOFTWARE 13 Figura 1.1: NVIDIA Jetson TK1. Tegra K1. •GPU Kepler 192 n´ ucleos. •CPU 4+1 procesadores ARM Cortex-A15 (ARMv7). Memoria Principal 2GB. Disco Duro 16GB eMMC. 1 mini-PCIE. 1 SSD/MMC. 1 HDMI. 1 USB 2.0 micro AB. 1 USB 3.o. 1 RS232 serie. 1 SATA3. 1 GigaEthernet. 1 ALC5639 Audio in/out . . . 1.5. Software Las aplicaciones desarrolladas se han implementado en lenguaje C y se ha empleado el compilador de GNU gcc, compilador com´ un en los sistemas operativos basados en unix como el que opera en el Jetson. Tambi´ en se ha utilizado el compilador nvcc, compilador propietario de NVIDIA empleado para la compilaci´ on del c´ odigo dise˜ nado para las GPUs de NVIDIA escrito en el lenguaje CUDA [9] y tambi´ en para compilar c´ odigo que incorpora llamadas a rutinas que hacen uso de estas GPUs. 14 CAP´ ITULO 1. INTRODUCCI´ ON GENERAL 1.6. Librerias En el desarrollo de este trabajo, se han empleado cuatro librer´ ıas de altas prestaciones con el fin de mejorar al m´ aximo las prestaciones de las aplicaciones desarrolladas; estas librer´ ıas son: BLAS, LAPACK, PLASMA y CUBLAS. Adem´ as, se ha empleado el paquete ATLAS para la instalaci´ on optimizada de BLAS y LAPACK. BLAS [10] (Basic Linear Algebra Subprograms): Librer´ ıa que provee rutinas que realizan operaciones b´ asicas de ´ algebra lineal num´ erica sobre vectores y matrices. Las rutinas de esta librer´ ıa est´ an agrupadas en tres niveles: BLAS1, que contiene operaciones de coste lineal sobre vectores unidimensionales; BLAS2, que contiene operaciones de coste cuadr´ atico que involucran matrices y vectores; y BLAS3, que contiene operaciones de coste c´ ubico sobre matrices. LAPACK [11] (Linear Algebra PACKage): Librer´ ıa desarrollada para resolver problemas de ´ algebra lineal num´ erica tales como resoluci´ on de sistemas de ecuaciones lineales, resoluci´ on de sistemas de m´ ınimos cuadrados, resoluci´ on de problemas de valores propios y descomposiciones en valores singulares. Las rutinas de LAPACK utilizan llamadas a las rutinas de BLAS simplificando su implementaci´ on y aprovechando la eficiencia del BLAS en caso de que se disponga de una implementaci´ on optimizada. ATLAS [12] (Automatically Tuned Linear Algebra Software): No es una librer´ ıa propiamente dicha, sino que permite configurar de forma emp´ ırica las librer´ ıas BLAS y LAPACK a las caracter´ ısticas del hardware sobre el que se est´ an instalando, maximizando as´ ı el rendimiento de ambas y por lo tanto, reduciendo el tiempo computacional de nuestras aplicaciones. PLASMA [13] (Parallel Linear Algebra for Scalable Multi-core Architectures): Es una implementaci´ on mejorada de la liber´ ıa LAPACK, que permite explotar las caracter´ ısticas de las arquitecturas multi-n´ ucleos en las funciones contenidas en ella. CUBLAS [14] (CUda Basic Linear Algebra Subroutines): Librer´ ıa que implementa las rutinas de BLAS pero empleando la GPU como hardware para los c´ alculos. 1.7. Estructura del trabajo La memoria de la tesis de m´ aster se estructura en los siguientes cap´ ıtulos: Configuraci´ on En el Cap´ ıtulo 2se describen con detalle los pasos necesarios para actualizar y configurar de manera b´ asica el Sistema Operativo que incorpora el NVIDIA Jetson, la distribuci´ on Ubuntu L4T de Linux. Adem´ as, se exponen las instalaciones y configuraciones de las librer´ ıas empleadas en el desarrollo de esta tesis, as´ ı como los problemas que surgieron al instalarlas y configurarlas. Filtrado En el Cap´ ıtulo 3se describen dos estructuras comunes empleadas para el filtrado de se˜ nales, como son el FIR (Finite Impulse Response) y el IIR (Infinite Impulse Response). Adem´ as de han realizado dos implementaciones diferentes de la estructura de filtrado IIR: IRRI (forma directa I) y PIIR (forma paralela). Las implementaciones han sido realizadas empleando instrucciones vectoriales NEON y tres tipos de datos, dos que trabajan con datos enteros y uno que trabaja con n´ umeros reales: INT16,INT32 yFLO32. Al final del cap´ ıtulo se realizan pruebas con las diferentes versiones y se analizan los resultados obtenidos. 1.7. ESTRUCTURA DEL TRABAJO 15 Beamformer En el Cap´ ıtulo 4se describe el algoritmo de Beamformer, empleado en se˜ nales digitales para separar una se˜ nal de otra/s de un conjunto en el que han sido mezcladas varias de ellas por un medio de transmisi´ on con el objeto de extraer la original. Se han desarrollado diferentes versiones empleando librer´ ıas y herramientas de computaci´ on de altas prestaciones. Al final del cap´ ıtulo se muestran los resultados obtenidos por las diferentes versiones y se analizan los mismos. Conclusiones En el Cap´ ıtulo 5se hace una recopilaci´ on de las ideas m´ as relevantes obtenidas durante la elaboraci´ on de la tesis, tratando de enfatizar los aspectos m´ as interesantes, los conceptos aprendidos y la aportaci´ on en la tem´ atica de procesamiento de se˜ nal digital en tiempo real. Adem´ as al final de este cap´ ıtulo, se exponen las posibles mejoras o extensiones ha realizar en ambos trabajos expuestos en esta tesis, as´ ı como las futuras l´ ıneas de investigaci´ on por las que podemos encauzarnos a tenor de los resultados obtenidos. 16 CAP´ ITULO 1. INTRODUCCI´ ON GENERAL Cap´ ıtulo 2 Puesta en marcha y configuraci´ on NVIDIA Jetson TK1 ENeste segundo cap´ ıtulo se abordan los primeros pasos que se han de llevar a cabo una vez tenemos en nuestras manos el Jetson TK1 [15], antes de pasar a implementar y ejecutar los c´ odigos que vamos a desarrollar en los cap´ ıtulos 3y4. A la hora de configurar el NVIDIA Jetson TK1, nos hemos encontrado con un gran n´ umero de problemas que nos han reportado un tiempo extra para poder configurarlo de forma satisfactoria. Uno de los principales problemas que nos hemos encontrado, es que el sistema operativo del Jetson es una versi´ on incompleta del sistema operativo Ubuntu para escritorio (L4T – Linux for Tegra), esto nos ha causado algunos problemas a la hora de configurarlo e instalar las librer´ ıas necesarias, ya que faltan algunos ficheros y/o paquetes de software que s´ ı que se encuentran en una versi´ on completa de Ubuntu. Adem´ as, al ser un producto relativamente nuevo, tiene problemas de madurez en el sistema operativo que los t´ ecnicos de NVIDIA van solucionando en versiones sucesivas seg´ un los van detectando o son informados por parte de la comunidad de usuarios. Todo esto sin nombrar que, al tratarse de un producto nuevo y en fase de adquisici´ on de usuarios y siendo la documentaci´ on bastante escueta en algunos apartados, el feedback a los usuarios a´ un es bastante escaso, lo que resulta en dificultades a la hora de encontrar soluciones a consultas acerca de cualquier problema surgido con el Jetson. Para finalizar, comentar que no ´ unicamente el software del Jetson carece de madurez a la hora de utilizarlo, sino que adem´ as, las librer´ ıas y/o aplicaciones no se encuentran adaptadas a´ un a algunas de sus peculiaridades por lo que surgen problemas a˜ nadidos al instalarlas (si es que se consiguen instalar correctamente en el Jetson), problemas que no se producen en equipos con un Ubuntu o Linux convencional. Por comentar algunos casos, destacar que no se ha podido hacer funcionar correctamente ni el compilador de ARM (armcc) ni la librer´ ıa MAGMA que implementa las rutinas del Lapack haciendo uso de todos los n´ ucleos computacionales disponibles en el hardware, tanto CPU’s como GPU’s. En lo referente a la librer´ ıa CUDA para el Jetson, comentar que la ´ ultima versi´ on disponible en el momento que se redactan estas l´ ıneas es la CUDA-6.5, mientras que para sistemas convencionales est´ a disponible actualmente la CUDA-7.5. Adem´ as, rese˜ nar que el paquete CUDA para el Jetson no dispone de la totalidad de las herramientas disponibles en una versi´ on CUDA convencional. 2.1. Puesta en marcha En esta primera secci´ on, se van a detallar los pasos b´ asicos para poner en marcha el NVIDIA Jetson, as´ ı como para obtener una configuraci´ on b´ asica del mismo a partir de la cual podamos empezar a trabajar. 17 18 CAP´ ITULO 2. PUESTA EN MARCHA Y CONFIGURACI´ ON NVIDIA JETSON TK1 2.1.1. Conexiones Conectar un extremo del cable serie al puerto serie J1A2 UART4 de la placa NVIDIA Jetson TK1, y el otro extremo a un equipo Linux (Preferentemente el equipo Linux debe ser Ubuntu, que es el sistema Linux que incorporan los Jetson, en caso contrario pueden surgir problemas adicionales a los expuestos en este documento). Conectar el cable USB Micro-B (incluido en la caja) a la placa Jetson y por el otro extremo USB al equipo Linux. Conectar teclado y rat´ on al puerto USB de la placa, como la placa ´ unicamente dispone de un ´ unico puerto USB, si se desea conectar ambos dispositivos es necesario disponer de un adaptador que saque m´ as de un puerto USB. Conectar la salida HDMI de la placa Jetson a un monitor HDMI. Conectar un cable de red al puerto Ethernet de la placa por un extremo y por el otro a un switch/- router/etc que le proporcione acceso a internet al Jetson. Conectar el cable de alimentaci´ on del Jetson al puerto AC del mismo. 2.1.2. Flashing La primera tarea que realizamos una vez recibidos los Jetson, fue actualizar el sistema operativo que incorporan a la versi´ on m´ as reciente disponible, tarea a la que llaman flashing. En este punto ´ unicamente hacer un peque˜ no inciso, y es que en la gu´ ıa se nos indica que el flashing se puede realizar desde cualquier equipo Linux, pero aprovechando nuestra experiencia, indicar que los sistemas operativos OSX no sirven para tal cometido, aunque bien sabemos que tambi´ en est´ an basados en Linux (distribuci´ on Darwin). Por lo tanto, decidimos finalmente emplear un equipo donde ten´ ıamos instalado Ubuntu como sistema operativo, que es el mismo sistema operativo en el que se basa NVIDIA para dise˜ nar el de los Jetson. A continuaci´ on vamos a detallar paso por paso el procedimiento a seguir: 1. Descargar la ´ ultima versi´ on del software para el Jetson (L4T) de la url https://developer. nvidia.com/linux-tegra 2. Descomprimir el paquete descargado en una carpeta y ensamblar el sistema de ficheros. sudo tar xpf ${RELEASE NAME} cd Linux_for_Tegra/rootfs/ sudo tar xpf ../../Tegra_Linux_Sample-Root-Filesystem_Rxx.x.x_armhf.tbz2 cd .. sudo ./apply_binaries.sh 3. Copiamos el sistema de ficheros a la memoria interna eMMC del Jetson. a) Poner la placa Jetson en Recovery Mode pulsando los botones reset y recovery de la placa y posteriormente soltando el bot´ on de reset. b) Asegurarse que el equipo linux est´ a conectado a la placa con el cable USB para el flashing. sudo ./flash.sh -S 16GiB ${Jetson-tk1} mmcblk0p1 En unos minutos el sistema del Jetson abr´ a sido reinstalado por completo. 4. Instalamos el sistema gr´ afico, despu´ es de haber configurado la red, y reiniciamos el sistema. sudo apt-get update sudo apt-get install ubuntu-desktop reboot 2.1. PUESTA EN MARCHA 19 2.1.3. Configuraciones b´ asicas recomendadas Una vez tenemos realizado el flashing del Jetson por completo y el sistema ha sido reiniciado, procedemos a realizar unas configuraciones b´ asicas en el sistema antes de proceder a instalar ning´ un otro software. Para ello, debemos abrir un terminal o consola del sistema Ubuntu del Jetson y seguir los pasos siguientes: 1. Es muy importante que se le diga al comando apt que no sobrescriba el archivo libglx.so si vamos a actualizar el sistema. El archivo libglx.so es un archivo de controladores gr´ aficos de NVIDIA que podr´ ıa ser reemplazado por una versi´ on incorrecta al actualizar el sistema Ubuntu del Jetson, y esto causar´ ıa que no pudiera ser arrancado el entorno gr´ afico del Jetson. Por esta raz´ on, es conveniente ejecutar el comando sudo apt-mark hold xserver-xorg-core antes de conectar a Internet y/o actualizar el sistema Ubuntu del Jetson. 2. A˜ nadir a la lista de repositorios del sistema Ubuntu del Jetson el repositorio universe con los comandos sudo apt-add-repository universe ysudo apt-get update, ya que es com´ un que se requieran paquetes contenidos en este repositorio cuando estamos desarrollando c´ odigos. 3. Instalar el paquete bash-completion, permite que se auto-complete los comandos shell y te da sugerencias sobre los paquetes a instalar si se escribe un comando que no se encuentra en el sistema. 4. Instalar los paquetes gfortran ygfortran-multilib ya que las liberias BLAS yLAPACK est´ an escritas en el lenguaje Fortran por lo que es necesario un compilador de fortran para instalarlas con el paquete ATLAS. 2.1.4. Instalaci´ on de CUDA Como siguiente paso una vez hecho el flashing del Jetson a la ´ ultima versi´ on de software disponible, se procede a instalar las herramientas CUDA que nos permitir´ an hacer uso de la GPU incorporada en el chip TK1 del Jetson. A continuaci´ on detallaremos los pasos que se han se seguir para dicha instalaci´ on: 1. Descargar la ´ ultima versi´ on de CUDA disponible para sistemas L4T de la url http://developer. download.nvidia.com/compute/cuda/release_version/rel/installers/cuda-repo-l4t-rXX. X-X-X-prod_X.X-XX_armhf.deb. 2. A˜ nadimos el repositorio descargado al sistema de paquetes del Jetson con el comando sudo dpkg -i cuda-repo-l4t-rXX.X-X-X-prod_X.X-XX_armhf.deb. 3. Actualizamos la lista de paquetes que hay en los repositorios con el comando sudo apt-get update. 4. Instalamos el paquete CUDA-toolkit de la versi´ on descargada con el comando sudo apt-get install cuda-toolkit-X-X. 5. A˜ nadimos los usuarios que deban tener acceso a la GPU al grupo video para que as´ ı puedan acceder con el comando sudo usermod -a -G video usuario. 6. A˜ nadimos los path del CUDA instalado en el fichero bashrc de los usuarios, recargamos el fichero bashrc para que se apliquen las actualzaciones y comprobamos que podemos acceder al compilador de NVIDIA nvcc con los comandos: echo ‘‘# Add CUDA bin & library paths:’’ >> ~/.bashrc echo ‘‘export PATH=/usr/local/cuda/bin:$PATH’’ >> ~/.bashrc echo ‘‘export LD_LIBRARY_PATH=/usr/local/cuda/lib:$LD_LIBRARY_PATH’’ >> ~/.bashrc source ~/.bashrc nvcc -V 20 CAP´ ITULO 2. PUESTA EN MARCHA Y CONFIGURACI´ ON NVIDIA JETSON TK1 2.1.5. Instalaci´ on de ATLAS Una vez tenemos instalado el CUDA-toolkit, procedemos a la instalaci´ on del paquete ATLAS que nos har´ a una instalaci´ on ´ optima para nuestro equipo de las librer´ ıas BLAS yLAPACK. Para ello expondremos a continuaci´ on los pasos necesarios para una correcta configuraci´ on y instalaci´ on de la misma. ´ Unicamente comentar, que si posteriormente se va ha hacer uso de otras librer´ ıas como el PLASMA, es necesario realizar la instalaci´ on del ATLAS con la opci´ on hard, esto a diferencia que la opci´ on soft, indica que la arquitectura dispone de unidades hardware para realizar las operaciones con n´ umeros reales, mientras que la opci´ on soft es para las arquitecturas que carecen de este tipo de unidades implementadas por hardware y lo que hacen es realizar las operaciones con n´ umeros reales por software. En el caso del Jetson, el procesador dispone de estas unidades hardware, por lo que no tenemos mayores problemas para instalar la librer´ ıa con una opci´ on u otra, pero es necesario instalarla con la opci´ on hard si vamos a utilizar otras librer´ ıas como en nuestro caso el PLASMA. A continuaci´ on se detallan con esmero los pasos a seguir para una correcta configuraci´ on e instalaci´ on del paquete ATLAS en el Jetson: 1. Activar el modo performance en el Jetson echo 0 > /sys/devices/system/cpu/cpuquiet/tegra_cpuquiet/enable echo 1 > /sys/devices/system/cpu/cpu0/online echo 1 > /sys/devices/system/cpu/cpu1/online echo 1 > /sys/devices/system/cpu/cpu2/online echo 1 > /sys/devices/system/cpu/cpu3/online echo performance > /sys/devices/system/cpu/cpu0/cpufreq/scaling_governor 2. Crear una carpeta y posicionarnos dentro de ella mkdir ATLAS cd ATLAS 3. Descargar el paquete ATLAS de la url http://sourceforge.net/projects/math-atlas/files/. 4. Descomprimir fichero tar -xjvf atlas3.XX.X.tar.bz2 rm -r atlas3.XX.X.tar.bz2 5. Descargar y configurar opci´ on hard wget http://math-atlas.sourceforge.net/fixes/armhardfp_archdef.tar tar xvf armhardfp_archdef.tar rm -r armhardfp_archdef.tar vi ./CONFIG/src/atlcomp.txt :g/=softfp/s//=hard/g 6. Instalar ATLAS con todas las funciones Lapack wget http://www.netlib.org/lapack/lapack-3.5.0.tgz 7. Configurar, construir, comprobar e instalar ATLAS mkdir srcATLAS cd srcATLAS ../ATLAS/configure -D c -DATL_ARM_HARDFP=1 -Ss ADdir $CURRENT_PATH/ATLAS/ARMHARDFP --with-netlib-lapack-tarfile=$CURRENT_PATH/lapack-3.5.0.tgz sudo make build sudo make check sudo make ptcheck sudo make install 2.1. PUESTA EN MARCHA 21 2.1.6. Instalaci´ on de PLASMA Para finalizar con las librer´ ıas ha instalar, procedimos a la instalaci´ on de la librer´ ıa PLASMA la cual nos proporciona el poder realizar las funciones que incorpora la librer´ ıa LAPACK, haciendo uso de los cuatro n´ ucleos computacionales (CPU’s) del Jetson, y no ´ unicamente de uno. A continuaci´ on mostramos los comandos necesarios para la configuraci´ on e instalaci´ on de la librer´ ıa PLASMA 1. Descargar la librer´ ıa PLASMA desde la url http://icl.cs.utk.edu/projectsfiles/plasma/pubs/ plasma-installer_X.X.X.tar.gz. 2. Descomprimir la librer´ ıa tar -xzvf plasma-installer_X.X.X.tar.gz cd plasma-installer_X.X.X 3. Configuraci´ on e instalaci´ on sudo ./setup.py --prefix=/usr/lib/ --build=/usr/local --cc=gcc --fc=gfortran --cflags=‘‘-O2’’ --ldflags_c=-L/usr/local/atlas/lib/ --ldflags_fc=-L/usr/local/ atlas/lib/ --blaslib=‘‘-lf77blas -lcblas -latlas’’ --lapacklib=‘‘-llapack’’ --lapclib=‘‘-llapacke’’ --nbcores=4 Pulsamos ‘‘d’’ Pulsamos ‘‘d’’ 28 CAP´ ITULO 3. FILTRADO DE SONIDO: FIR, IIRI Y PIIR Tabla 3.1: Instrucciones NEON empleadas con una breve descripci´ on. Instruction Description vmovq_n_s32(value) Initialize int32x4_t vector registers vmovq_n_s64(value) Initialize int64x2_t vector registers vmovq_n_f32(value) Initialize flo32x4_t vector registers vmov_n_s16(value) Initialize int16x4_t vector registers vmov_n_s32(value) Initialize int32x2_t vector registers vld1_s16(*data) Load int16x4_t vector registers vld1_s32(*data) Load int32x4_t vector registers vld1q_f32(*data) Load flo32x4_t vector registers vmls_s16(dest,data1,data2) Vector multiply subtract, int16x4_t dest. vector registers vmls_s32(dest,data1,data2) Vector multiply subtract, int32x2_t dest. vector registers vmlsq_f32(dest,data1,data2) Vector multiply subtract, flo32x4_t dest. vector registers vmlal_s16(dest,data1,data2) Vector multiply accumulate long, int32x4_t dest. vector registers vmlal_s32(dest,data1,data2) Vector multiply accumulate long, int64x2_t dest. vector registers vmlaq_f32(dest,data1,data2) Vector multiply accumulate, flo32x4_t dest. vector registers vst1_s16(dest,source) Store a single vector into memory, *int16_t[4] dest. registers vst1_s32(dest,source) Store a single vector Store a single, *int32_t[4] dest. registers vst1q_f32(dest,source) Store a single vector into memory, *flo32_t[4] dest. registers vget_high_s32(reg) Splitting vector registers, int32x2_t return vector registers vget_high_f32(reg) Splitting vector registers, flo32x2_t return vector registers vget_low_s32(reg) Splitting vector registers, int32x2_t return vector registers vget_low_f32(reg) Splitting vector registers, flo32x2_t return vector registers vpadd_s32(data1,data2) Pairwise add, int32x2_t dest. vector registers vpadd_f32(data1,data2) Pairwise add, flo32x2_t dest. vector registers vget_lane_s32(reg,pos) Extract lanes from a vector, int32_t dest. registers vget_lane_s64(reg,pos) Extract lanes from a vector, int64_t dest. registers vget_lane_f32(reg,pos) Extract lanes from a vector, flo32_t dest. registers 3.4. An´ alisis de resultados Como hemos comentado con anterioridad, en este an´ alisis de resultados estudiamos tres estructuras diferentes de filtros: FIR, IIRI y PIIR, filtros que hemos expuesto al principio del cap´ ıtulo, empleando como tama˜ nos de coeficientes del filtro 52 y 256 valores. Para poder realizar una comparativa equitativa, elegimos los valores de M y N id´ enticos (M=N). Todos los c´ odigos han sido compilados con el compilador GNU GCC empleando las opciones de compilaci´ on -O3,-mfpu=neon y-march=armv7. Para auto-vectorizar un c´ odigo incluyendo rutinas NEON intrinsics, se emplea la opci´ on -ftree-vectorize [22]. El primer resultado, plasmado en la Figura 3.4, muestra el tiempo de ejecuci´ on requerido por los tres tipos de datos utilizados: 16-bit y 32-bit integers y 32-bit floats. Nosotros comparamos la versi´ on optimizada empleando instrucciones intr´ ınsecas NEON con la versi´ on que emplea operaciones comunes (no vectoriales). El resultado mostrado recalca el beneficio en coste computacional que obtenemos si utilizamos el tipo de datos 16-bit iteger comparado con el tipo de datos 32-bit float. Para comparar el speedUp obtenido empleando los diferentes tipos de datos, se ha elegido la comparativa entre los tipos de datos INT16 yFLO32, ya que se consideran los m´ as interesantes de utilizar dependiendo de las necesidades de la aplicaci´ on. En la Figura 3.5, la parte superior del gr´ afico representa el speedUp alcanzado por el tipo de datos INT16, mientras que la parte inferior representa el speedUp obtenido para el tipo de datos FLO32. En el caso del filtro FIR, el speedUp obtenido es muy alto, obteniendose unos valores de 5 y 6,5 para un n´ umero de coeficientes de 52 y 256, respectivamente, cuando el tipo de datos utilizado es el INT16. En cambio, para el tipo de datos FLO32 y el mismo filtro FIR, los va- 3.4. AN´ ALISIS DE RESULTADOS 29 0 10 20 30 40 50 60 10 20 40 80 160 320 Time (ms.) Number of filters FIR of length 52 INT16 INT32 FLO32 0 50 100 150 200 250 300 350 10 20 40 80 160 320 Time (ms.) Number of filters FIR of length 256 INT16 INT32 FLO32 0 5 10 15 20 25 30 10 20 40 80 160 320 Time (ms.) Number of filters IIRI of length 52 INT16 INT32 FLO32 0 20 40 60 80 100 120 10 20 40 80 160 320 Time (ms.) Number of filters IIRI of length 256 INT16 INT32 FLO32 0 10 20 30 40 50 60 10 20 40 80 160 320 Time (ms.) Number of filters Parallel IIR of length 52 INT16 INT32 FLO32 0 50 100 150 200 250 300 350 10 20 40 80 160 320 Time (ms.) Number of filters Parallel IIR of length 256 INT16 INT32 FLO32 Figura 3.4: Rendimiento con respecto al tipo de datos. Figura 3.5: Speedup alcanzado usando intr´ ınsecas NEON. lores de speedUp obtenidos son de 3 y 3,5 aproximadamente para los mismos tama˜ nos de los vectores de coeficientes. Estos resultados nos dicen que obtenemos un gran beneficio en cuanto a tiempo de computo comparado con la versi´ on auto-vectorizada por parte del compilador gcc; obteniendose un mejor beneficio cuando utilizamos el tipo de datos INT16 en vez del FLO32, aunque produci´ endose un p´ erdida de precisi´ on en los resultados del primero comparado con los del segundo (esta p´ erdida es dependiente del rango de los datos de entrada de la aplicaci´ on). Tambi´ en se han implementado versiones paralelas de los c´ odigos (empleando OpenMP) para estos tres tipos de datos y estructuras de filtro. La escalabilidad obtenida en muy razonable en todos los casos. Como se muestra en la Figura 3.6, los speedUp obtenidos para el caso de emplear dos n´ ucleos computacionales (CPU), es de dos para ambas longitudes de los vectores de coeficientes. En cambio para 30 CAP´ ITULO 3. FILTRADO DE SONIDO: FIR, IIRI Y PIIR 0 1 2 3 4 10 20 40 80 160 320 Speedup Number of filters 2 cores. Filter length = 52 4 cores. Filter length = 52 2 cores. Filter length = 256 4 cores. Filter length = 256 Figura 3.6: Speedup del filtro IIRI empleando 2 y 4 cores. cuatro n´ ucleos, el comportamiento es un poco m´ as irregular obteniendose unos valores de speedUp entre tres y cuatro aunque entrando estos dentro de la normalidad. Con cuatro n´ ucleos ya no son tan ideales como con dos n´ ucleos, teniendo en cuenta que en ambos casos no se han contemplado dependencias entre los datos de los diferentes n´ ucleos computacionales. Esta diferencia en el rendimiento deber´ ıa estar relacionada con alg´ un cuello de botella en el hardware del Jetson, presumiblemente en la memoria cach´ e de nivel dos (L2). Se define tbuff como el intervalo de tiempo que se tarda en generar dos buffers consecutivos de muestras de entrada a procesar. Tambi´ en se define tproc como el tiempo que tarda nuestra aplicaci´ on en procesar las muestras del buffer. Por lo tanto, para que nuestra aplicaci´ on funcione en tiempo real, las muestras contenidas en el buffer actual han de ser procesadas antes de que el siguiente buffer est´ e listo para ser procesado, esto es tproc < tbuff . El objetivo de la Figura 3.7 es el de mostrar, para cada una de las estructuras de filtro y para los tipos de datos INT16 y FLO32 empleando en ambos casos 52 y 256 como tama˜ no de los coeficientes del filtro, cu´ al es el n´ umero m´ aximo de filtros que nuestra aplicaci´ on ser´ ıa capaz de procesar en tiempo real, es decir, siendo tproc < tbuff . En la Figura 3.7 se puede apreciar la evoluci´ on de tproc en funci´ on del n´ umero de filtros procesados. Esta figura compara los tipos de datos INT16 y FLO32. Ahora vamos a comentar por encima que l´ ımites se pueden extraer de la gr´ afica para cada una de las estructuras de filtro estudiadas en este trabajo, que permitan a la aplicaci´ on que emplee estas estructuras de filtro trabajar en tiempo real. Para la estructura FIR los l´ ımites para los tipos de datos INT16 y FLO32 con 256 coeficientes son de 170 y 260 filtros, respectivamente. Para la estructura IIRI son de 80 y 125. Por ´ ultimo, para la estructura PIIR son de 50 y 25 filtros aplicados en tiempo real. 3.4. AN´ ALISIS DE RESULTADOS 31 0 5 10 15 20 25 50 100 150 200 250 300 tproc(ms) 52 coef. - INT16 FIR IIRI PIIR tbuff (ms) 0 5 10 15 20 25 50 100 150 200 250 300 tproc(ms) 256 coef. - INT16 FIR IIRI PIIR tbuff (ms) 0 5 10 15 20 25 50 100 150 200 250 300 tproc(ms) 52 coef. - FLO32 FIR IIRI PIIR tbuff (ms) 0 5 10 15 20 25 50 100 150 200 250 300 tproc(ms) 256 coef. - FLO32 FIR IIRI PIIR tbuff (ms) Figura 3.7: Evoluci´ on del tiempo de computo necesario para el c´ alculo en tiempo real. 32 CAP´ ITULO 3. FILTRADO DE SONIDO: FIR, IIRI Y PIIR Cap´ ıtulo 4 Beamformer ESTE cap´ ıtulo aborda el desarrollo del algoritmo del beamforming, algoritmo capaz de, a partir de una se˜ nal recibida que ha sido alterada por ruido, reverberaciones y otras se˜ nales, limpiarla y separarla del ruido a fin de que ´ esta quede como originalmente era antes de haber sido transmitida y alterada por agentes externos. Para ello, la captura de la se˜ nal a procesar se realiza a partir de un array de micr´ ofonos que capturan la se˜ nal desde diferentes puntos del espacio. 4.1. Modelo de se˜ nal En el modelo presentado vamos a asumir un sistema compuesto por Maltavoces que van a emitir M se˜ nales diferentes: s1(k), s2(k), . . . , sM(k), y que van a ser capturadas conjuntamente por Nmicr´ ofonos situados en otro espacio de la habitaci´ on diferente al de los altavoces y que capturar´ an las Mse˜ nales como una se˜ nal ´ unica que, adem´ as, se encontrar´ a alterada por ruido y reverberaciones de la habitaci´ on. El problema que trata de solucionar este algoritmo consiste en c´ omo separar cada una de las se˜ nales s1(k),s2(k),...,sM(k)a partir de las se˜ nales capturadas por los distintos micr´ ofonos del sistema. En la Figura 4.1 se puede apreciar un ejemplo con dos altavoces y tres micr´ ofonos. El enfoque adoptado en este documento hace uso de algoritmos de procesamiento de se˜ nal para el dise˜ no del ancho de banda beamformers (filtros gnen la Figura 4.1), una vez que todos los canales de respuesta de la habitaci´ on (hnm en la Figura 4.1) son conocidos. Este sistema puede ser modelado como un sistema multicanal con dos entradas (altavoces) y tres salidas (micr´ ofonos). La generalizaci´ on a un sistema de Multiples entradas - M´ ultiples salidas (MIMO, Multiple Input - Multiple output) puede ser estudiado en [23]. De acuerdo con la Figura 4.1, la salida del micr´ ofono nth viene dada por la EQ 4.1: xn(k) = M X m=1 Lh X j=1 hnm(j)sm(k−j) + vn(k),(4.1) donde n= 1,2, . . . , N, siendo Nel n´ umero de micr´ ofonos del sistema, y m= 1,2, . . . , M, donde Mes el n´ umero de se˜ nales de entrada (fuentes) o n´ umero de altavoces de la Figura 4.1. El par´ ametro Lhes la longitud del canal de respuesta ac´ ustico hnm m´ as largo de la habitaci´ on, mientras que vn(k)es el ruido de la se˜ nal. En aras de aportar una mayor claridad a la ecuaci´ on del problema, el t´ ermino que aporta ruido, vn(k)en la EQ 4.1, no se considerar´ a en el siguiente modelo de la se˜ nal EQ 4.2: xn(k) = M X m=1 hT nmsm(k),(4.2) siendo sm(k)el vector columna definido como: sm(k) = [sm(k), sm(k−1) . . . sm(k−Lh+ 1)]T, yhnm es el RLh×1vector del canal ac´ ustico del altavoz mal micr´ ofono n. 33 34 CAP´ ITULO 4. BEAMFORMER Figura 4.1: Modelo de la se˜ nal para M=2 altavoces (entradas) y N=3 micr´ ofonos (salidas). Considerando ahora el problema de recuperaci´ on de se˜ nales emitidas sm(k)a partir de las se˜ nales adquiridas por los micr´ ofonos xn(k), el filtro beamforming gnde la Figura 4.1 tiene que estar dise˜ nado de forma que la se˜ nal de salida y(k)constituya una buena estimaci´ on de la se˜ nal sm(k), es decir, que y(k) = ˆsm(k−τ)nos proporcione un error m´ ınimo. Dada una longitud m´ axima de Lgfija para cada uno de los filtros gn, el ancho de banda de la se˜ nal de salida se puede expresar de forma similar a la EQ 4.2: y(k) = N X n=1 gT nxn(k),(4.3) donde gnes el vector RLg×1conteniendo ordenadamente los filtros gnde la Figura 4, y xn(k)es el vector columna definido como xn(k) = [xn(k)xn(k−1) . . . xn(k−Lg+ 1)]T. Con el fin de calcular el vector completo xn(k)utilizado en la EQ 4.3 de forma matricial, la EQ 4.2 ha tenido que ser reescrita en forma compacta redefiniendo hnm como matrices de Sylvester. Para m´ as detalles consultar [23]. 4.2. Algoritmos Beamforming En [23], Benesty et al. presenta un excelente repaso de los principales algoritmos utilizados en procesamiento de se˜ nal. Debido a que ofrece un mejor rendimiento, nosotros nos centramos en este trabajo en un algoritmo basado en correlaci´ on de matrices, como es el LCMV (Linearly Constrained Minimum Variance). El Algoritmo LCMV calcula los filtros beamforming como se puede apreciar en la EQ 4.4, gLCMV =ˆ R−1 xH:m[HT :mˆ R−1 xH:m]−1um,(4.4) donde gLCMV est´ a formado por la concatenaci´ on de los filtros gn, esto es, gLCMV = [gT 1. . . gT N]T, la matriz H:mes una partici´ on de la matriz de impulso del canal que solo incluye las respuestas al impulso de las fuentes mth de los N micr´ ofonos [23] utilizados en la matriz de Sylvester para las dimensiones [NLg·Lg +Lh −1]. La matriz ˆ Rxes la matriz de correlaci´ on de las se˜ nales capturadas y umes un vector de ceros a excepci´ on de una de las componentes del vector con el fin de compensar el retardo de la respuesta al impulso de la habitaci´ on. Buscando la implementaci´ on m´ as eficiente y precisa del LCMV, utilizamos un m´ etodo basado en la 4.2. ALGORITMOS BEAMFORMING 35 descomposici´ on QR de la matriz XT, siendo ´ esta definida como X∈R[NLg·K]: X=1 √k      x1(k)x1(k+ 1) . . . x1(k+K−1) x2(k)x2(k+ 1) . . . x2(k+K−1) . . .. . .. . . . . . xN(k)xN(k+ 1) . . . xN(k+K−1)      ,(4.5) donde K > NLges el n´ umero de muestras utilizadas. Esta forma, XT=Q·R, donde Qes una matriz ortogonal y Res una matriz triangular superior, permite la resoluci´ on r´ apida de sistemas algebraicos lineales empleando las rutinas de la librer´ ıa LAPACK. Nosotros constru´ ımos directamente la matriz XTen la representaci´ on Column Major Order para poder aplicar directamente estas rutinas eficientes. Considerando la descomposici´ on QR de las observaciones de los micr´ ofonos, se puede redefinir ˆ Rx como: ˆ Rx=X·XT=RT·QT·Q·R=RT·R. Ahora nosotros definimos por conveniencia la matriz Wcomo W=ˆ R−1 xH:m, as´ ı que el filtro beamformer LCMV (gLCMV ) definido en la EQ 4.4 queda de la siguiente forma: gLCMV =W[HT :mW]−1um.(4.6) Se define la matriz Zcomo la soluci´ on al sistema lineal mostrado en la EQ 4.7: RTZ=H:m,(4.7) entonces, utilizando la descomposici´ on QR de la matriz X, tenemos: W=ˆ R−1 xH:m= (RTR)−1H:m=R−1R−TH:m=R−1Z, donde se puede apreciar claramente que la matriz Wes la soluci´ on al sistema lineal RW =Z. La soluci´ on para obtener los filtros beamforming se obtiene de resolver el sistema lineal de la EQ 4.8, Abm=um,(4.8) donde A=HT :mW=HT :mR−1Z=ZTZ. Tambi´ en aqu´ ı, la soluci´ on del sistema lineal de la EQ 4.8 es obtenida a partir de la factorizaci´ on QR, en este caso, de la matriz Z. Ahora, Z=Q0R0es la descomposici´ on QR de la matriz Z, entonces, el vector bmpuede ser calculado resolviendo los dos siguientes sistemas lineales triangulares que se muestran en la EQ 4.9 yEQ4.10. R0Ty=um,(4.9) R0bm=y . (4.10) Finalmente, es f´ acil deducir que el c´ alculo del filtro beamformer de la EQ 4.6, puede ser calculado empleando los ´ ultimos c´ alculos realizados, esto es, R,Zybmde la forma que se muestra en la EQ 4.11, gLCMV =R−1Zbm.(4.11) Estos ´ utimos c´ alculos se resuelven realizando un producto matriz-vector y resolviendo un sistema lineal triangular. 36 CAP´ ITULO 4. BEAMFORMER 4.3. Configuraci´ on del sistema e implementaci´ on La placa NVIDIA Jetson incorpora el chip Tegra K1 de NVIDIA, el cual, contiene a su vez una GPU basada en arquitectura Kepler. Por lo tanto, podemos emplear una CPU multi-core (4+1) y una GPU de 192 cores ensamblados en un ´ unico chip y que hacen uso de una memoria RAM com´ un con 2 GB de capacidad. Al estar incorporada la GPU en el mismo chip y emplear la misma memoria RAM, los intercambios de datos entre CPU y GPU son a priori mucho m´ as r´ apidos que en las configuraciones tradicionales donde la GPU se comunica con el host a trav´ es del bus PCIe. Adem´ as, se trata de un dispositivo de bajo consumo con un buen ratio GFlop/wattio, lo que permite desarrollar aplicaciones que despu´ es vayan a ser empleadas en dispositivos m´ oviles. 4.3.1. Configuraci´ on del sistema En este trabajo se han desarrollado diferentes implementaciones del algoritmo para el c´ alculo de los filtros beamformer que emplean la CPU quad-core Cortex-A15 a 2,32 GHz, la GPU Kepler de 192 cores a 0,85 GHz o ambas unidades computacionales (CPU y GPU). La dos unidades computacionales han sido configuradas al m´ aximo de su potencial (modo performance) para obtener el menor tiempo computacional posible en el c´ alculo de los filtros beamformer. El fin que se persigue con las diferentes implementaciones realizadas para el c´ alculo de los filtros beamformer es poder calcularlos en tiempo real, por lo que se han empleado librer´ ıas de altas prestaciones como LAPACK, PLASMA y CUBLAS. En otras palabras, se han explorado todas las posibilidades basadas en librer´ ıas existens de altas prestaciones con el fin de potenciar el rendimiento de la aplicaci´ on y la eficiencia de los n´ ucleos computacionales disponibles en el Jetson. La configuraci´ on del sistema desarrollado en las pruebas es un sistema con dos altavoces situados en un lugar de la habitaci´ on que emiten dos se˜ nales diferentes de la misma duraci´ on (si no fuese as´ ı se alargar´ ıa la m´ as corta rellenando el buffer con ceros), y tres micr´ ofonos en otro lugar de la habitaci´ on que capturan las se˜ nales emitidas por los altavoces. Para finalizar con esta secci´ on, comentar que en este trabajo se analizan implementaciones que emplean para el c´ alculo de los filtros Beamformer: una sola CPU empleando la librer´ ıa LAPACK, varias CPU’s empleando la librer´ ıa PLASMA y, por ´ ultimo, con una CPU y la GPU con LAPACK y CUBLAS conjuntamente. 4.3.2. Implementaci´ on En las implementaciones realizadas tenemos el vector Xmicro, que es el vector que nos proporciona las se˜ nales capturadas por los micr´ ofonos. A partir de los datos contenidos en el vector Xmicro, generamos la matriz X(EQ 4.5), la cual va a ser el punto de partida para calcular los filtros beamformer para cada se˜ nal que se quiera separar. Seguidamente tenemos las matrices H,ZyA, donde Hes la matriz del canal. Existe una por cada altavoz (se˜ nal) que se emita en el sistema. La matriz Zes la matriz resultado de resolver el sistema de ecuaciones con m´ ultiples vectores independientes (EQ 4.7), siendo Rla QR de matrixX. La matriz Aes el producto ZTZ=A. Para finalizar, tenemos los vectores u,byg, siendo uun vector que tiene ceros en todas sus componentes menos en la componente Lg+ 1 que vale uno. El vector bes el vector resultado de resolver dos sistemas de ecuaciones lineales, EQ 4.9 yEQ4.10. El vector g, por su parte, contiene los filtros beamformer y hay tantos vectores gcomo se˜ nales tengamos que separar. Estos vectores se calculan en mediante la EQ 4.11. Para obtener el m´ aximo rendimiento en el c´ alculo de los filtros beamformer, se ha aprovechado el hecho que diferencia el tipo de almacenamiento de los datos que hacen el lenguaje C y el lenguaje Fortran, que es el que se emplea en las librer´ ıas LAPACK y BLAS. Manejando adecuadamente el tipo de almacenamiento hemos evitamos tener que trasponer las matrices XyZa la hora de calcular la QR de las mismas. Esto se consigue construyendo la matriz Xdirectamente traspuesta. Por otro lado, tambi´ en se 4.3. CONFIGURACI´ ON DEL SISTEMA E IMPLEMENTACI´ ON 37 Algoritmo 4.1 Realiza el cociente de la ra´ ız de k con cada uno de los valores de Xmicro a utilizar en la matriz X requiere: ∗Xmicro,Lx,MICROS,Ndatos asegurar: ∗Xmicro 1: float iraizk=1.0\sqrt(k); 2: float *pXmicro = Xmicro; 3: for(i=0; i<MICROS; i++){ 4: for(j=0; j<Ndatos; j++){ 5: pXmicro[j] *= iraizk; 6: } 7: pXmicro += Lx; 8: } Algoritmo 4.2 Construir la traspuesta de la matriz X con los datos de Xmicro requiere: ∗X,∗Xmicro,Lx,MICROS,Ndatos,Lg asegurar: ∗X 1: float *pXmicro = Xmicro+Lg-1; 2: float *pmatrixX = matrixX; 3: for(i=0; i<MICROS; i++){ 4: float *pXmicroaux = pXmicro; 5: for(r=0; r<Lg; r++){ 6: for(j=0; j<K; j++){ 7: *pmatrixX = *(pXmicroaux+j); 8: *pmatrixX++; 9: } 10: pXmicroaux--; 11: } 12: pXmicro += lx; 13: } debe tener almacenada la matriz Hpor columnas, ya que Zse forma a partir de estas dos. En definitiva, hemos comenzado la optimizaci´ on del algoritmo por la base: tener un acceso eficiente a memoria. A continuaci´ on desgranaremos paso a paso las distintas implementaciones realizadas para el c´ alculo de los filtros beamformer en los algoritmos siguientes. En el Algoritmo 4.1 se calcula el cociente entre los valores de los datos capturados por los micr´ ofonos y la ra´ ız de la variable K. Se calcula la inversa de la ra´ ız de Kpara as´ ı emplear multiplicaciones en vez de cocientes. Esto es debido a que las primeras tienen un coste computacional menor. Se calculan ´ unicamente los K+Lgprimeros datos de cada micr´ ofono en caso de no ser iterativo el algoritmo, y se le a˜ nade al valor anterior el producto del n´ umero de repeticiones del algoritmo por el valor de las nuevas muestras (solapamiento) que se procesan por iteraci´ on, es decir, K+Lg +NF rames ∗overLap. A este valor lo llamamos Ndatos. Este c´ alculo se realiza en una funci´ on independiente debido a que, a la hora de formar la matriz X, se repite varias veces el mismo valor, por lo que se gana en eficiencia al calcularlos una ´ unica vez. El valor de la variable Lx es el cociente entre el tama˜ no de Xmicro y el n´ umero de micr´ ofonos MICROS. Una vez ya tenemos los datos de la matriz Xmicro calculados correctamente, se procede a la construcci´ on de la matriz Xcon dichos valores y, como hemos comentado con anterioridad, dicha construcci´ on se realizar´ a con la consideraci´ on de que la matriz debe ser traspuesta antes de emplearla. Esto se muestra en el Algoritmo 4.2. El Algoritmo 4.3 muestra una funci´ on auxiliar que sirve para copiar de forma invertida los elementos de un vector en otro. Estos vectores invertidos son de un tama˜ no dado por un intervalo, por lo que se invierten los datos que componen cada intervalo por separado. Para finalizar con el apartado de implementaci´ on, se muestran los algoritmos encargados del c´ alculo de los filtros beamformer, la descomposici´ on QR de las matrices que lo requieran, y la modificaci´ on para 44 CAP´ ITULO 4. BEAMFORMER Cap´ ıtulo 5 Conclusiones y trabajo futuro ENesta tesis de m´ aster se han tratado dos problemas del procesado de se˜ nales digitales ac´ usticas. Ambos problemas, aunque diferentes, tienen en com´ un su importancia en la b´ usqueda de un mejor ambiente sonoro, pero tambi´ en el dispositivo sobre el que se ha trabajado. Por un lado, se han implementado y estudiado tres estructuras de filtrado (FIR, IIR y PIIR) para se˜ nales digitales. Las operaciones de filtrado son b´ asicas, es decir, operaciones aritm´ eticas esenciales realizadas sobre vectores de mayor o menor dimensi´ on. Para realizar estas operaciones de manera eficiente hemos utilizado las instrucciones intr´ ınsecas NEON de los procesadores ARM. Los resultados comparativos entre la autovectorizaci´ on por parte del compilador gcc y las implementaciones realizadas empleando instrucciones NEON muestran una gran mejora en la implementaci´ on utilizando expl´ ıcitamente el uso de las instrucciones NEON frente a la autovectorizaci´ on que proporciona el compilador. Esta mejora es importante, llegando a ser en varios casos superior a 4 veces el tiempo de la versi´ on autovectorizada. Esto tiene que ver con el hecho de que la autovectorizaci´ on proporcionada por el compilador gcc no ofrece una buena ganancia con respecto a la versi´ on no vectorizada. Nuestra implementaci´ on permite su utilizaci´ on en tiempo real. Por ejemplo, se pueden aplicar aproximadamente 125 filtros IIRI y 260 FIR para un n´ umero de coeficientes de filtro de 256 y con el tipo de datos INT16 antes de que est´ e disponible el siguiente frame de datos de entrada a procesar. Adem´ as, se muestra en los c´ odigos desarrollados que el tipo de datos INT16 es m´ as eficiente que el tipo de datos FLO32, esto es debido a que el coste computacional de una operaci´ on con n´ umeros reales es m´ as costosa que una operaci´ on con n´ umeros enteros. A tenor de estos resultados, se podr´ ıa concluir que es mejor emplear el tipo de datos INT16, siempre y cuando la precisi´ on en los c´ alculos nos lo permita. Esto es as´ ı, no solo porque es m´ as eficiente en cuanto a coste computacional que el tipo de datos FLO32, si no tambi´ en porque es el tipo de datos gen´ erico para formatos de audio digital como los CD’s de m´ usica. Por otra parte, se han desarrollado varias implementaciones del Algoritmo LCMV para el c´ alculo de filtros Beamformer. Este caso es diferente ya que las operaciones que involucra el algoritmo son de m´ as alto nivel, es decir, operaciones de ´ algebra lineal num´ erica sobre matrices (descomposiciones matriciales, resoluci´ on de sistemas lineales, etc.). En este caso, lo apropiado era acudir a la utilizaci´ on de librer´ ıas que resuelven este tipo de problemas de manera eficiente sobre arquitecturas parecidas basadas en multicore y manycore. Esto ha tenido una implicaci´ on natural en el trabajo consistente en la instalaci´ on, adaptaci´ on y comprobaci´ on de las librer´ ıas sobre la m´ aquina objeto de estudio. La “inmadurez” del dispositivo y, especialmente, de su software, ha supuesto un trabajo de investigaci´ on a˜ nadido nada despreciable, pero ha permitido, al mismo tiempo, obtener un buen rendimiento de este novedoso dispositivo en lo que se refiere a la soluci´ on de este tipo de problemas. En particular, hemos comenzado por realizar un estudio del coste computacional de cada una de las partes de las que se compone el algoritmo. Hemos concluido que una gran parte del coste computacional del algoritmo viene dado por el c´ alculo de las factorizaciones QR matriciales previas a la resoluci´ on de los sistemas de ecuaciones lineales subsiguientes. En concreto, la factorizaci´ on QR de la matriz Xes la que mayor coste tiene, estando este coste comprendido entre un 60 % y un 70 % del coste total del algoritmo para un c´ alculo de dos filtros Beamformer (g1yg2). En vista de estos resultados se han implementan diferentes estrategias para calcular eficientemente las factorizaciones QR, utilizando los distintos recursos computacionales del NVIDIA Jetson. Se concluye, a 45 46 CAP´ ITULO 5. CONCLUSIONES Y TRABAJO FUTURO ra´ ız de los tiempos tomados, que lo m´ as eficiente es emplear la librer´ ıa PLASMA que dispondr´ a de los cuatro n´ ucleos computacionales ARM para calcularla. En relaci´ on al Cap´ ıtulo 3y como l´ ınea futura de investigaci´ on, se podr´ ıa estudiar la conveniencia de emplear el tipo de datos INT16 oFLO32 en varias aplicaciones reales con diferentes requisitos de precisi´ on y tiempo computacional, para poder concluir con una aplicaci´ on real que tipo de datos es m´ as conveniente y mostrar resultados concretos para estas aplicaciones. Con respecto al problema tratado en el Cap´ ıtulo 4, cabe decir que una de las tareas pendientes de realizar en este apartado es el realizar un algoritmo que emplee todos los n´ ucleos computacionales disponibles, es decir, las cuatro CPU’s ARM y la GPU Kepler. Esto lo conseguiremos conjugando el calculo de las QR con PLASMA y el c´ alculo de la rutina larfb en GPU; dicha implementaci´ on no ha podido ser incluida en esta memoria por falta de tiempo, pero ya se esta trabajando en ella y en futuros trabajos se mostraran los resultados. Adem´ as, como futura l´ ınea de investigaci´ on, se podr´ ıa proponer un algoritmo colaborativo entre diferentes Jetson que intercambien informaci´ on empleando el est´ andar de intercambio de mensajes MPI, de tal modo que los Jetson colaboren entre si intercambiando parte de la informaci´ on procesada para mejorar la calidad y/o el tiempo de procesado para la separaci´ on de la se˜ nal que se quiere. Bibliograf´ ıa [1] ARM NEON, www.arm.com/, (accessed 2015 February 23). [2] M. FLYNN,Some Computer Organizations and Their Effectiveness, IEEE Trans. Comput C-21 (1972) 948–960. [3] INTEL SSE, https://software.intel.com/sites/landingpage/IntrinsicsGuide/, (accessed 2015 September 17). [4] J. R¨ AMO, V. V¨ ALIM¨ AKI AND BAL´ AZS BANK,High-Precision Parallel Graphic Equalizer, IEEE Transactions on Audio, Speech and Language Processing 22 (2014) 1894–1904. [5] J.C. RISSET,Computer Music Experiments 1964, Computer Music J. 22 (1985) 11–18. [6] Y. HUANG, J. CHEN AND J. BENESTY,Inmerse Audio Schemes, IEEE Signal Processing Magazine 28 (2011) 20–32. [7] J. LORENTE, G. PI˜ NERO, A.M. VIDAL, J.A. BELLOCH AND ALBERTO GONZ´ ALEZ,19th European Signal Procesing Conference, Parallel Implementations of Beamforming Design and Filtering for Microphone Array Applications (29 August - 02 september 2011). [8] NVIDIA JETSON TK1, https://developer.nvidia.com/jetson-tk1/, (accessed 2015 February 10). [9] CUDA, http://www.nvidia.es/object/cuda-parallel-computing-es.html, (accessed 2015 September 17). [10] BLAS Library, http://www.netlib.org/blas/, (accessed 2015 February 15). [11] LAPACK Library, http://www.netlib.org/lapack/, (accessed 2015 February 15). [12] ATLAS Library, http://math-atlas.sourceforge.net/, (accessed 2015 February 15). [13] PLASMA Library, http://icl.cs.utk.edu/plasma/, (accessed 2015 March 15). [14] CUBLAS Library, http://docs.nvidia.com/cuda/cublas/, (accessed 2015 April 25). [15] Nvidia Jetson Quick Start Guide, http://developer.download.nvidia.com/embedded/jetson/ TK1/docs/2_GetStart/Jeston_TK1_User_Guide.pdf, (accessed 2015 September 04). [16] A. V. OPPENHEIM, A. S. WILLSKY,AND S. HAMID Signals and systems, Processing series. Prentice Hall, 2nd edition, 1997. [17] L. R. RABINER AND B. GOLD,Theory and Application of Digital Signal Processing Prentice-Hall, Englewood Cliffs, New Jersey, USA, 1975. [18] ARM NEON intrinsics, gcc.gnu.org/onlinedocs/gcc-4.4.1/gcc/ARM-NEON-Intrinsics.html, (accessed 2015 July 12). 47 48 BIBLIOGRAF´ IA [19] Advanced Linux Sound Architecture (ALSA), www.alsa-project.org/, (accessed 2015 January 08). [20] Steinberg Media Technologies GmbH, www.steinberg.net/, (accessed 2015 Aug. 14). [21] S. B. HOLGERSSON,Optimising IIR filters using ARM NEON, Master Thesis of University of Danemark (2012). [22] ARM NEON auto-vectorization, gcc.gnu.org/onlinedocs/gcc/ARM-Options.html, (accessed 2015 July 22). [23] D. THEODOROPOULOS, G. KUZMANOV AND G. GAYDADJIEV,Multi-core Platforms for Beamforming and Wave Field Synthesis IEEE Transactions on Multimedia 99 (2010). [24] Source DGEQRF.c, http://www.netlib.org/clapack/CLAPACK-3.1.1/SRC/dgeqrf.c, (accessed 2015 September 18). [25] Source DLARFT.c, http://www.netlib.org/clapack/CLAPACK-3.1.1/SRC/dlarft.c, (accessed 2015 September 18). [26] Source DLARFB.c, http://www.netlib.org/clapack/CLAPACK-3.1.1/SRC/dlarfb.c, (accessed 2015 September 18). [27] QUARK Library, http://icl.cs.utk.edu/quark/, (accessed 2015 September 18). [28] QR factoritation PLASMA, http://www.netlib.org/lapack/lawnspdf/lawn222.pdf, (accessed 2015 September 18). [29] QR factoritation MAGMA, http://www.netlib.org/lapack/lawnspdf/lawn233.pdf, (accessed 2015 September 18).