Uso de tarjetas GPU para acelerar el procesado de señales

TRABAJO FINAL DE GRADO
TÍTULO DEL TFG: Uso de a je as GPU pa a acele a el p ocesado de
señales
TITULACIÓN: G ado en Ingenie ía de Sis emas de Telecomunicación
AUTOR: Da id Ama Sanz
DIRECTOR: Gab iel Mon o o López
CODIRECTOR: Pe e Lluis Gilabe Pinal
FECHA: 05 de Julio de 2017
Tí ulo: Uso de a je as GPU pa a acele a el p ocesado de señales
Au o : Da id Ama Sanz
Di ec o : Gab iel Mon o o López
Codi ec o : Pe e Lluis Gilabe Pinal
Fecha: 05 de Julio de 2017
Resumen
Es e T abajo de Final de G ado p e ende analiza y pone en p ác ica o a
o ma de p ocesa señales digi ales, mejo ando su endimien o y elocidad de
ejecución.
Los DSP y FPGA son los elemen os más usados en la ac ualidad pa a
cualquie ipo de p ocesado de señales. Es e p oyec o se cen a en el uso de
las a je as g á icas (GPU) pa a explo a al máximo el pa alelismo del que se
dispone hoy en día.
Los p ocesado es ac uales (CPU) cuen an con unos pocos núcleos y abajan
secuencialmen e lo que puede compo a ele ado consumo de iempo si se
es á p ocesando can idades muy g andes de da os.
La en aja de las GPU es que disponen de miles de núcleos, que, pese a se
menos po en es indi idualmen e, si abajan de o ma simul ánea pueden
o ece una no able acele ación espec o a la CPU.
El algo i mo Leas Mean Squa e, usado en il os adap a i os, es el escogido
pa a lle a a cabo la pa alelización.
Su complejidad compu acional es ideal pa a es a op imización, pues, como
hay que encon a los coe icien es de adap ación median e un g an núme o de
i e aciones, amos a pode ealiza odas es as mul iplicaciones a la ez con
CUDA C, que es el lenguaje que se u iliza en compu ación pa alela pa a
a je as g á icas de la ma ca N idia.
Ti le: Use o GPU ca ds o speed up signal p ocessing
Au ho : Da id Ama Sanz
Di ec o : Gab iel Mon o o López
Codi ec o : Pe e Lluis Gilabe Pinal
Da e: July 05 h 2017
O e iew
This Bachelo 's Deg ee Final P ojec aims o analyze and implemen ano he
way o p ocess digi al signals, imp o ing hei pe o mance and speed o
execu ion.
DSP and FPGA a e he mos commonly used elemen s o any kind o signal
p ocessing. This p ojec ocuses on he use o g aphics ca ds (GPU) o exploi
o he maximum he pa allelism ha is a ailable oday.
Cu en p ocesso s (CPUs) ha e a ew co es and wo k sequen ially which can
be e y ime consuming i la ge amoun s o da a a e being p ocessed.
The ad an age o GPUs is ha hey ha e housands o co es, which, al hough
hey a e less powe ul indi idually, i hey wo k simul aneously hey can o e a
ema kable accele a ion compa ed o he CPU.
The Leas Mean Squa e algo i hm, used in adap i e il e s, is he one chosen
o pe o m he pa alleliza ion.
I s compu a ional complexi y is ideal o his op imiza ion, because, as we ha e
o ind he adap a ion coe icien s h ough a la ge numbe o i e a ions, we will
be able o pe o m all hese mul iplica ions a he same ime wi h CUDA C,
which is he language used in pa allel compu ing o N idia g aphics ca ds.
Quie o ag adece a oda mi amilia, amigos y a
mi pa eja po apoya me desde el inicio.
Su mo i ación me ha ayudado a log a los
obje i os que pe seguía. También ag adece
an o a mi u o Gab iel Mon o o como a mi
co u o Pe e Lluis Gilabe po la ayuda o ecida
du an e odo el desa ollo del p oyec o.

ÍNDICE
INTRODUCCIÓN ............................................................................................... 1
CAPÍTULO 1. TIPOS DE FILTROS Y ELECCIÓN DEL ADAPTATIVO ............ 2
1.1. Fil o elec ónico................................................................................................................ 2
1.1.1. Fil os analógicos .................................................................................................... 2
1.1.2. Fil os digi ales ........................................................................................................ 2
1.1.3. Respues a de Impulso In ini a (IIR) ........................................................................ 3
1.1.4. Respues a de Impulso Fini a (FIR) ......................................................................... 3
1.2. Fil o Adap a i o ................................................................................................................ 4
1.3. In oducción al algo i mo LMS ......................................................................................... 5
1.4. Desc ipción del LMS ......................................................................................................... 6
1.5. Ecuación del algo i mo LMS ............................................................................................ 9
CAPÍTULO 2. USO DE LA GPGPU PARA LA PARALELIZACIÓN ............... 11
2.1. Uso de la GPU an e el DSP ............................................................................................. 11
2.2. Ven ajas del uso de la GPU ............................................................................................ 12
2.3. GPU – CUDA s OpenCL ................................................................................................ 13
2.4. Aplicaciones de cálculo en la GPU (CUDA) .................................................................. 14
2.5. Pa alelismo e icien e ....................................................................................................... 17
2.6. Compa ación GPU-CPU .................................................................................................. 18
2.7. Ca ac e ís icas del Ha dwa e ......................................................................................... 20
CAPÍTULO 3. CUDA Y PARALELIZACIÓN DEL ALGORITMO LMS ............ 22
3.1. His o ia de CUDA ............................................................................................................. 22
3.2. Ven ajas de CUDA ........................................................................................................... 23
3.3. Modelo de p og amación ................................................................................................ 24
3.4. Acele ación de MATLAB usando MEX .......................................................................... 28
3.5. Funcionamien o de mexcuda ......................................................................................... 30
3.6. Pa alelización del algo i mo LMS .................................................................................. 33
3.7. Implemen ación del LMS en CUDA ................................................................................ 35
3.8. Resul ados ob enidos ..................................................................................................... 40
CONCLUSIONES ............................................................................................ 46
BIBLIOGRAFÍA ............................................................................................... 47
ANEXOS .......................................................................................................... 49
In oducción 1
INTRODUCCIÓN
En la ac ualidad los il os adap a i os cons i uyen una impo an e pa e del
p ocesamien o digi al de señales.
El uso de es os il os o ece una a ac i a solución al p oblema que usualmen e
no pueden sol en a los il os de coe icien es ijos con encionales.
Los il os de coe icien es cons an es no suelen se su icien es a la ho a de
a a señales que son a ian es en el iempo o que son pa e de en o nos
con aminados, debido a que abajan en ecuencias ijas y po al mo i o se
en limi ados en el p ocesamien o adecuado de la señal.
Combina un il o con un algo i mo adap a i o, gene a posibilidades de que el
sis ema de il ado se ajus e a cualquie señal de en ada, es o pe mi e
ecalcula los coe icien es del il o de una o ma dinámica y e ec i a,
con i iéndolo de es a mane a en un il o adap a i o.
En es e abajo se p opone una e sión pa alela del algo i mo LMS (Leas
Mean Squa e Algo i hm), que se u iliza pa a p ocesamien o de señales digi ales
como la eliminación de eco y educción de uido pa a la ob ención de una señal
deseada.
El pa alelismo en la GPU (Unidad de p ocesamien o g á ico) pe mi e la
descomposición de un p oblema en una se ie de p oblemas más pequeños y
que pueden se calculados más ápidamen e.
Los esul ados ob enidos, especialmen e el aumen o de la elocidad y la
e iciencia, mues an que el mé odo pa alelo implemen ado en la GPU es mucho
más ápido que o os p ocedimien os exis en es que se lle an a cabo
solamen e en la CPU, pues es a solo abaja con unos cuan os núcleos
op imizados pa a el p ocesamien o en se ie secuencial.
8 Uso de a je as GPU pa a acele a el p ocesado de señales
La siguien e igu a p esen a los de alles del mecanismo de con ol adap a i o
pa a los coe icien es, especí icamen e, una e sión escala del p oduc o in e no
de la es imación del e o “e(n)” y la en ada “x(n-k)” es calculada pa a k=0,1, 2,
..., M-2, M-1.
Fig. 1.8 Mecanismo de con ol adap a i o pa a los coe icien es
El esul ado así ob enido de ine la co elación “ ” aplicada a los alo es
de los coe icien es “ ” en la i e ación n+1.
El ac o de escala u ilizado en es e cálculo se designa po “µ” y es conocido
como pa áme o del amaño de paso de adap ación o elocidad de adap ación

Tipos de il os y elección del adap a i o 9
1.5. Ecuación del algo i mo LMS
El algo i mo LMS, pa a un il o de o den “M”, puede esumi se de la siguien e
mane a:
Pa áme os:
= O den del il o = Tamaño del paso
Inicialización: (Po de ec o “0” si no hay da os del ec o de coe icien es del
il o)
(1.1)
Da os:
Señal de en ada en el ins an e “n”:
(1.2)
: Señal deseada a la salida del il o.
A calcula :
Es imación del ec o de coe icien es (pesos) del il o en el ins an e “n+1”
(1.3)
Cómpu o:
Pa a n=0, 1, 2, ..., calcula :
Señal de e o :
(1.4)
Adap ación de los coe icien es del il o:
(1.5)
Salida del il o:
(1.6)
10 Uso de a je as GPU pa a acele a el p ocesado de señales
deno a asposición, deno a aspues a conjugada, el as e isco
deno a conjugación.
El ac o “µ” que de e mina la elocidad de con e gencia, es di ec amen e
p opo cional a la elocidad de con e gencia e in e samen e p opo cional al
e o cuad á ico medio mínimo.
Po lo an o, exis e un comp omiso en e elocidad de con e gencia y e o
cuad á ico p omedio mínimo, es en es e pun o donde adica es e algo i mo
p esen a su mayo incon enien e.
Uso de la GPGPU pa a la pa alelización 11
CAPÍTULO 2. Uso de la GPGPU pa a la pa alelización
2.1. Uso de la GPU an e el DSP
El p ocesado de señales mul idimensionales es un p oblema común en la
in es igación cien í ica. No malmen e, la complejidad de cálculo de un
p oblema de DSP c ece exponencialmen e con el núme o de dimensiones.
Sin emba go, con un al o g ado de complejidad de almacenamien o y iempo,
es ex emadamen e di ícil p ocesa señales mul idimensionales en iempo eal.
Aunque se han p opues o muchos algo i mos ápidos (po ejemplo, FFT) pa a
p oblemas de DSP 1-D ( ec o es de 1 dimensión), oda ía no son lo
su icien emen e e icaces pa a adap a se a p oblemas con mayo núme o de
dimensiones.
Toda ía es di ícil ob ene los esul ados de cálculo deseados con p ocesado es
de señales digi ales (DSPs). Es po eso que se necesi an mejo es algo i mos y
a qui ec u a de ha dwa e pa a acele a los cálculos DSP mul idimensionales.
Las unidades de p ocesamien o g á ico de p opósi o gene al (GPGPU)
mode nas se conside a que ienen un endimien o excelen e en ope aciones
ec o iales y manipulaciones numé icas a a és de un g an núme o de
cálculos pa alelos.
El é mino GPGPU es concep o ecien e den o de in o má ica que a a de
es udia y ap o echa las capacidades de cómpu o de una GPU.
Mien as que el p ocesamien o de señales digi ales, en pa icula las señales
mul idimensionales, a menudo implica una se ie de ope aciones de ec o en
una g an can idad de mues as de da os independien es, las GPGPUs aho a
son comúnmen e u ilizadas pa a acele a los DSP mul idimensionales, ales
como el p ocesamien o de imágenes, señales, códecs de ídeo y ul asonido.
Concep ualmen e, el uso de disposi i os GPGPU, pa a ealiza las a eas que
ha ía el DSP mul idimensional, es capaz de educi d ás icamen e la
complejidad compu acional en compa ación con las unidades cen ales de
p ocesamien o (CPUs), p ocesado es de señales digi ales (DSP) u o os
acele ado es FPGA.
Fig. 2.1 A qui ec u a in e na de una FPGA
12 Uso de a je as GPU pa a acele a el p ocesado de señales
2.2. Ven ajas del uso de la GPU
La demanda en el me cado de consumo, de ha dwa e g á ico que acele a el
ende izado de imágenes en 3D, ha dado luga a a je as g á icas que son
capaces de o ece so p enden es ni eles de endimien o.
Es os esul ados se log a on adap ando especí icamen e el ha dwa e pa a esa
inalidad. Los acele ado es g á icos es án ol iéndose cada ez más
p og amables po lo que es e endimien o los ha con e ido en un obje i o
a ac i o pa a o os dominios.
Las Unidades de P ocesamien o G á ico (GPU) p opo cionan una a qui ec u a
de compu ación pa alela de bajo cos o.
Es posible log a un pa alelismo masi o con la écnica del SIMT (Single
Ins uc ion Mul iple Th ead) en la Unidad de P ocesamien o de G á icos de Uso
Gene al (GPGPU) in eg ada con la Unidad Cen al de P ocesamien o (CPU).
Las en ajas de u iliza SIMT (Ins ucción Única pa a Múl iples Hilos) pueden
explica se median e un ejemplo: cambiando el b illo de una imagen.
Cada píxel de una imagen cons a de es alo es pa a el b illo siendo los
colo es: ojo (R), e de (G) y azul (B).
No malmen e se cambia ía el alo de cada pixel secuencialmen e, lo que
du a ía un cie o iempo.
Fig. 2.2 Ins ucciones de o ma clásica e Ins ucciones u ilizando SIMT
Con un p ocesado que u ilice SIMT hay dos mejo as en es e p oceso.
P ime amen e, los da os se o ganizan en bloques, y una se ie de alo es se
pueden se ca gados de una ez.
En luga de usa una se ie de ins ucciones que digan: " ecupe a es e píxel,
aho a ecupe a el píxel siguien e, e c", un p ocesado SIMT end á una sola
ins ucción que e ec i amen e di á: " ecupe a n píxeles" (donde n es un
núme o que a ía de diseño a diseño).
Po un g an núme o de azones, es o puede a da mucho menos iempo que
ecupe a cada píxel indi idualmen e, como esul a ía con el diseño adicional
de la CPU.
Uso de la GPGPU pa a la pa alelización 13
2.3. GPU – CUDA s OpenCL
Los a chi os con g an núme o de da os consumen una eno me can idad de
iempo en se p ocesados po la CPU.
Se equie e mucho iempo pa a ealiza muchas ope aciones en ella, lo que
esul a en la disminución del endimien o.
CUDA es un modelo de p og amación pa alela po el cual se mejo a el
endimien o en é minos de iempo pa a cualquie ipo de compu ación.
Su al e na i a, OpenCL (Open Compu ing Language), se di e encia pues es un
es ánda de p og amación abie o lo que le pe mi e se usado en cualquie
GPU.
AMD es su pa ocinado y lo que p e ende es que no sea necesa io u iliza
exclusi amen e disposi i o NVIDIA po lo que sopo a a ias clases de
disposi i os como pueden se : CPUs, GPUs, DSPs, elé onos mó iles, e c.
Pa a abaja con CUDA, es an sencillo como desca ga se desde su misma
web, un paque e que incluye el ki de he amien as CUDA, con olado es de
desa ollado y el SDK.
En cambio, pa a OpenCL, las he amien as necesa ias y el SDK son
p opo cionados po cada p o eedo de sus disposi i os compa ibles.
OpenCL cons a de una se ie de he amien as y una API (In e az de
p og amación de aplicaciones) a di e encia de CUDA que ep esen a oda una
a qui ec u a como e emos más adelan e.
A día de hoy CUDA sigue o eciendo mayo endimien o po a ios mo i os:
• Puede usa se más de una GPU pa a GPGPU.
• CUDA es compa ible con más aplicaciones (Adobe CC, P emie e P o,
e c).
• Mayo op imización pues es á ejecu ándose en su p opio ha dwa e.
La o ma más e icien e de abajo es el uso conjun o de CPU y GPU median e
la a qui ec u a de cálculo pa alelo CUDA.
Po ello, más adelan e, e emos la di e encia en e el p ocesado en CPU y en
GPU, que eside p incipalmen e en el ha dwa e u ilizado y en la o ma que
cada núcleo p ocesa los da os.

14 Uso de a je as GPU pa a acele a el p ocesado de señales
2.4. Aplicaciones de cálculo en la GPU (CUDA)
Desde que debu ó, di e sidad de emp esas y aplicaciones han dis u ado
mucho de un g an éxi o po habe escogido CUDA C como pla a o ma pa a
c ea sus aplicaciones.
Es os bene icios a menudo incluyen aumen os signi ica i os de endimien o en
los úl imos pasos de la implemen ación de la écnica. Además, las aplicaciones
que se ejecu an sob e p ocesado es g á icos NVidia, gozan de un endimien o
supe io po dóla y po a ios, que las implemen aciones cons uidas
exclusi amen e pa a ecnologías de CPU.
Los sis emas de ada suelen eque i econs ui una can idad nume osa de
mues as de da os en 3-D o 4-D en iempo eal. T adicionalmen e y
especialmen e en el campo mili a , es o necesi a el apoyo de los
supe o denado es.
Fig. 2.3 P ocesado GPU en sis emas ada
Hoy en día, las GPGPUs se emplean pa a eemplaza supe o denado es pa a
p ocesa señales de ada . Po ejemplo, pa a señales de ada de ape u a
sin é ica (SAR), que no malmen e implica cálculos mul idimensionales de FFT.
Los GPGPUs se pueden u iliza pa a ealiza ápidamen e FFT y / o iFFT en
es e ipo de aplicaciones.
Uso de la GPGPU pa a la pa alelización 15
Muchos au omó iles de conducción au omá ica aplican écnicas de
econocimien o de imagen 3D pa a con ola au omá icamen e los ehículos.
Es e iden e que, pa a adap a se al en o no ex e io que cambia ápidamen e,
los p ocesos de econocimien o y decisión deben ealiza se en iempo eal.
Fig. 2.4 Ejemplo de coche de conducción au omá ica
O os usos de las GPU (CUDA) pueden se :
• Bioin o má ica
Pa a la secuencia y el acoplamien o de p o eínas se equie en p ocesos
in o má icos muy in ensos.
Fig. 2.5 Compa ación GPU-CPU en el alineamien o de secuencias de ADN
• Ciencia de los da os, analí ica y bases de da os
Se usa pa a analiza can idades masi as de da os (Big Da a) y mejo a así la
oma de decisiones de negocio en iempo eal.
16 Uso de a je as GPU pa a acele a el p ocesado de señales
• De ensa e In eligencia
Hay más de 100 millones de huellas dac ila es y o og a ías almacenadas en la
base de da os del FBI. Las GPUs acele an odas las unciones que in e ienen
en el p ocesamien o de esas imágenes, lo que incluye la geo ec i icación, los
algo i mos de il ado, la de ección de cambios y la econs ucción en 3D.
• Cálculo inancie o
La eno me acele ación que p opo cionan las GPUs de NVIDIA a aplicaciones
de cálculo an u ilizadas como las simulaciones Mon eca lo ep esen a una
g an en aja compe i i a pa a las emp esas de se icios inancie os.
Pode calcula el p ecio y el iesgo de opciones complejas y de i ados OTC
(me cados donde se negocian acciones, bonos, e c.) en segundos en luga de
ho as pe mi e ejecu a más simulaciones y, po an o, mejo a la calidad y
iabilidad de los esul ados
• Ap endizaje au omá ico
Es una ama cada ez más impo an e de la in eligencia a i icial (IA) y es la
ciencia que pe mi e a los o denado es ac ua sin habe sido exp esamen e
p og amados pa a ello. (Machine Lea ning).
Las edes neu onales basadas en sis emas in o má icos pueden ap ende a
imi a el compo amien o del ce eb o, lo que incluye las capacidades que és e
posee pa a econoce obje os, pe sonas, oces y sonidos.
Fig. 2.6 Compa ación en e se ido es CPU y GPU
NVIDIA ha colabo ado con un equipo de in es igado es de la Uni e sidad de
S an o d pa a c ea la ed neu onal a i icial más g ande del mundo con el
obje i o de ep oduci la o ma en que ap ende la men e humana.
La ed es 6,5 eces más g ande que la an e io poseedo a del éco d,
desa ollada po Google en 2012
En e o as muchas aplicaciones, es as son las que más u ilidad es án
apo ando:
• Dinámica de luidos compu acional (CFD), Química cuán ica, Explo ación
sísmica, Modelos me eo ológicos y climá icos, p ocesado de señales.
Uso de la GPGPU pa a la pa alelización 17
2.5. Pa alelismo e icien e
Una de las en ajas de los FPGAs (Field P og ammable Ga e A ay) y GPUs es
el pa alelismo a qui ec ónico que con ienen es os disposi i os.
El p oblema, es que no malmen e ienen ecuencias de eloj más bajas que las
CPUs con encionales. Es o signi ica que, pa a log a una acele ación, el
pa alelismo del disposi i o debe se lo más e icien e posible, iden i icando y
sol en ando el p oblema a esol e .
Teó icamen e, si uno dobla el núme o de p ocesado es, el iempo de ejecución
debe ía educi se a la mi ad, pe o no es así y e emos el po qué:
La ley de Amdahl es un modelo ma emá ico que desc ibe la elación en e la
acele ación espe ada de la implemen ación pa alela de un algo i mo y la
implemen ación se ial del mismo algo i mo.
Supongamos que una po ción P de un cálculo puede se pa alelizada y que
exis en N u as de da os pa alelas disponibles (núme o de p ocesado es).
La ley de Amdahl es ablece que la máxima elocidad alcanzable a a és de la
pa alelización es:
Cada ez que se dobla el núme o de p ocesado es la acele ación disminuye,
de es a mane a se iende al lími e siguien e:
Si el 95% de un p og ama es pa alelizable, la máxima acele ación ob enida
se á de 20x como emos en el siguien e g á ico:
Fig. 2.7 G á ico de la ley de Amdahl sob e la pa alelización
Es impo an e señala que la ley de Amdahl p opo ciona una acele ación
máxima alcanzable, no la acele ación que se ob end á.
La acele ación eal puede es a limi ada po el ha dwa e u ilizado y el iempo de
comunicación en e la CPU y la GPU.
24 Uso de a je as GPU pa a acele a el p ocesado de señales
3.3. Modelo de p og amación
El modelo de p og amación CUDA asume que los hilos ( h eads) CUDA se
ejecu an en una unidad ísica dis in a que ac úa como cop ocesado (de ice o
GPU) al p ocesado (hos o CPU) donde se ejecu a el p og ama.
CUDA C es una ex ensión del lenguaje de p og amación C, que pe mi e al
p og amado de ini unciones C, llamadas ke nels, que, al se llamadas, se
ejecu an en pa alelo po N hilos di e en es. Los ke nels se ejecu an de o ma
concu en e en el de ice.
Como ejemplo, el siguien e código mues a cómo se de ine un ke nel y cómo
se llama desde un p og ama:
Exis e una je a quía pe ec amen e de inida sob e los hilos de CUDA. Los hilos
se ag upan en ec o es a los que se les llama bloques, es os ec o es pueden
se de una, dos o es dimensiones, de o ma que de inen bloques de hilos de
una, dos o es dimensiones.
Hilos del mismo bloque pueden coope a en e sí, compa iendo da os y
sinc onizando sus ejecuciones
Sin emba go, hilos de dis in os bloques no pueden coope a . Los bloques a su
ez, se o ganizan en un g id de bloques. Es e g id, de nue o puede se de una
o dos dimensiones.
Los alo es en e <<< ... >>> que apa ecen en código an e io se conocen
como la con igu ación del ke nel, y de inen la dimensión (1) del g id y el núme o
de hilos de cada bloque (N).
// Ke nel de ini ion
__global_- oid VecAdd( loa * A, loa * B, loa * C)
{
in i = h eadIdx . x;
C[ i ] = A[ i ] + B[ i ];
}
in main ( )
{ // Ke nel in oca ion wi h N h eads
VecAdd<<<1, N>>>(A, B, C);
}

CUDA y pa alelización del algo i mo LMS 25
La siguien e igu a mues a cómo se o ganizan los hilos en un g id de 2x3
bloques y 3x4 hilos cada uno.
Fig. 3.4 Je a quía de los hilos en la GPU
Como puede e se, cada hilo queda pe ec amen e iden i icado po un ID de
bloque y el ID del p opio hilo den o del bloque. Es os IDs suelen usa se como
índices pa a de ini qué po ciones de los da os p ocesa cada hilo.
Fig. 3.5 Bloques e hilos en GPU
Los hilos de un bloque compa en la “sha ed memo y” (memo ia compa ida) y
pueden sinc oniza se median e llamadas a "synch h eads()".
Al asigna el amaño de los hilos y bloques hay que ene en cuen a las
es icciones del ha dwa e que es amos u ilizando pues cada GPU,
dependiendo de su gene ación, posee un núme o máximo de bloques
ecu en es.
26 Uso de a je as GPU pa a acele a el p ocesado de señales
La a qui ec u a Pascal, que ac ualmen e es la úl ima, es á limi ada a 1024 hilos
po bloque y 32 bloques concu en es.
Es o puede obse a se si ejecu amos la unción “gpuDe ice” en Ma lab lo que
nos de ol e á la in o mación de la a je a g á ica que engamos seleccionada.
En mi caso, uso una N idia GTX 970 que es á en ocada a los ideojuegos,
pe o que como es bas an e nue a o ece á un buen endimien o g acias al
núme o de “MaxTh eadsPe Block”.
El g an a ance se ha p oducido con la úl ima a qui ec u a desa ollada po
N idia, Pascal, que pe mi e una concu encia de 32 bloques y o eciendo
mayo po encia de p ocesado pa alelo, a di e encia de la an e io que solo
pe mi ía 16 bloques concu en es.
Es muy impo an e ema ca que a medida que apa ecen nue as a je as
g á icas, el núme o de núcleos CUDA que con ienen aumen a
conside ablemen e.
El p oblema es que en cada gene ación de GPUs la a qui ec u a cambia, po lo
que, aunque una nue a GTX Ti an Xp (Pascal) posea 3840 CUDA Co es y la
an igua GTX Ti an Z (Keple ) posea 5760, el endimien o se e á ealmen e
a ec ado po la capacidad de p ocesado (FLOPS) en la que la nue a a je a es
mucho supe io a la Ti an Z. (12 TFOPS con a 8 TFLOPS).
CUDA y pa alelización del algo i mo LMS 27
Las limi aciones al escoge los amaños de hilo ( h ead) y bloque pa a un
mul ip ocesado Pascal son:
1- 32 bloques concu en es.
2- 1024 hilos/bloque.
3- 2048 hilos en o al.
Ejemplos de limi ación:
 1 bloque de 2048 hilos
o No lo pe mi e [2].
 2 bloques de 1024 hilos.
o Posible en el mismo mul ip ocesado . (2048 hilos en o al)
 4 bloques de 512 hilos.
o Posible en el mismo mul ip ocesado . (2048 hilos en o al)
 4 bloques de 1024 hilos.
o No lo pe mi e [3] en el mismo mul ip ocesado , pe o es posible
usando dos mul ip ocesado es. (4096 hilos en o al)
 8 bloques de 256 hilos.
o Posible en el mismo mul ip ocesado . (2048 hilos en o al)
 256 bloques de 8 hilos.
o No lo pe mi e [1] en el mismo mul ip ocesado , pe o es posible
usando 8 mul ip ocesado es (8 * 32 bloques).
Fig. 3.6 Pe spec i a del so wa e pa a el bloque de hilos
28 Uso de a je as GPU pa a acele a el p ocesado de señales
3.4. Acele ación de MATLAB usando MEX
MATLAB p opo ciona un sc ip , llamado mex, pa a compila un a chi o MEX a
un obje o compa ido o dll que se puede ca ga y ejecu a den o de un
p og ama de MATLAB.
Es e sc ip es capaz de analiza el código C, C++ y FORTRAN y puede llama
al código CUDA pa a c ea una mayo op imización en las GPUs.
A pesa de que MATLAB llama a biblio ecas op imizadas in e namen e, oda ía
hay espacio pa a una mayo op imización.
Los a chi os MEX se han u ilizado en el pasado como una o ma de ca ga
biblio ecas mul ip oceso o ec o izadas y se usan en es e caso pa a ejecu a
código en la GPU y pa a maneja la ans e encia de da os en e el hos (CPU)
y el de ice (GPU).
Todos los a chi os MEX deben inclui 4 elemen os:
1. “#include mex.h” (pa a C y C++)
2. La u ina “ga eway” a cada a chi o MEX se llama “mexFunc ion”.
Es e es el pun o de en ada que MATLAB u iliza pa a accede a la DLL.
En C/C++, siemp e se usa de la siguien e mane a:
“MexFunc ion (in nlhs, mxA ay * plhs [], in n hs, cons mxA ay * p hs [])”
Dónde:
Nlhs = núme o de mxA ays espe ados (Le Hand Side)
Plhs = ec o de pun e os a la salida espe ada
N hs = núme o de en adas (Righ Hand Side)
P hs = ec o de pun e os a los da os de en ada.
(Los da os de en ada son de sólo lec u a)
3. “mxA ay”:
Es una es uc u a especial que con iene los da os de MATLAB. Es la
ep esen ación C de un a ay MATLAB. Todos los ipos de ma ices
MATLAB (escala es, ec o es, a ays, s ings, e c.) son mxA ays.
4. Funciones API (como asignación de memo ia y libe ación de es a).
CUDA y pa alelización del algo i mo LMS 29
Pa a pode ejecu a es e ipo de p og amas hay que esc ibi el siguien e
comando en Ma lab: “mex ejemplo.c”.
Es e comando gene a un a chi o compilado MEX. (El su ijo p oducido depende
del sis ema ope a i o):
ejemplo.mexw32 (Windows 32 bi )
ejemplo.mexw64 (Windows 64 bi )
ejemplo.mexglx (Linux 32 bi )
Fig. 3.7 Ca pe a as compilación del código
Pa a que MATLAB pueda ejecu a sus unciones en C, hay que pone los
a chi os compilados MEX que con ienen esas unciones en un di ec o io en la
u a MATLAB o ejecu a MATLAB en el di ec o io en el que se hallan.
Pa a la mayo ía de los casos, el a chi o MEX con iene código CUDA
( unciones del ke nel, con igu ación) y pa a se p ocesado po el compilado
CUDA, el a chi o necesi a ene un su ijo ".cu".
Es muy pa ecido al an e io ejemplo de “mex ejemplo.c”, pe o en es e caso se
u iliza “mexcuda ejemplo.cu” pa a pode ejecu a con MATLAB odo el código
que con enga CUDA.
Los a chi os CUDA MEX se pa ecen a los MEX (C, C++ y FORTRAN) pe o
di ie en a la ho a de asigna los da os en la memo ia pues hay que hace lo en
la GPU y ambién hay que ealiza los p ocesos que se encuen an en e CPU-
GPU, como es la llamada a los ke nels.

30 Uso de a je as GPU pa a acele a el p ocesado de señales
3.5. Funcionamien o de mexcuda
La pa e secuencial de una aplicación se ejecu a sob e la CPU (comúnmen e
denominada hos ) y la pa e más cos osa del cálculo se ejecu a sob e la GPU
(que se denomina de ice).
Dada la na u aleza he e ogénea del modelo de p og amación CUDA, una
secuencia ípica de ope aciones pa a un p og ama CUDA C es:
• Decla a y asigna la memo ia del hos y del de ice.
• Inicializa los da os del hos .
• T ans e i da os desde el hos al de ice.
• Ejecu a uno o más ke nels (llamadas a la GPU usando CUDA).
• T ans e i los esul ados del de ice al hos .
Pa a comp ende un poco mejo su uncionamien o el mejo ejemplo es el
p og ama “TimesTwo”, que mul iplica po dos cada elemen o, del a ay que le
en a, en un hilo di e en e de la GPU.
Lo p ime o es llama a la unción mxIni GPU en la en ada de su a chi o MEX.
Es o ga an iza que el disposi i o GPU se inicialice co ec amen e y sea
conocido po MATLAB.
oid mexFunc ion(in nlhs, mxA ay *plhs[],
in n hs, mxA ay cons *p hs[])
{
/* Decla amos a iables. */
in cons h eadsPe Block = 256;
mxIni GPU();
/* Ve i icamos que A ealmen e es una a ay de double an es de
ex ae el pun e o. */
i (mxGPUGe ClassID(A) != mxDOUBLE_CLASS) {
mexE MsgIdAndTx (e Id, e Msg);
}
/* Ex aemos el pun e o al da o de en ada en el de ice. */
d_A = (double cons *)(mxGPUGe Da aReadOnly(A));
N = (in )(mxGPUGe Numbe O Elemen s(A));
/* Llamamos al ke nel pa a que ejecu e la uncion "TimesTwo" */
TimesTwo<<<1, h eadsPe Block>>>(d_A, d_B, N);
/* Con ie e el esul ado ob enido en gpuA ay como la que usa
Ma lab pa a su en ío. */
plhs[0] = mxGPUC ea eMxA ayOnGPU(B);
/* Los pun e os de mxGPUA ay son es uc u as que se elacionan con
los da os que se ob iene del de ice. Han de se des uidos an es de
sali de la unción MEX. */
mxGPUDes oyGPUA ay(A);
mxGPUDes oyGPUA ay(B);
}
CUDA y pa alelización del algo i mo LMS 31
Como emos as inicializa la GPU, comp obamos que los da os que en an
engan el o ma o co ec o y p ocedemos a ex ae los pun e os pa a después
ob ene los da os calculados en la GPU.
T as escoge el núme o de hilos po bloque (256 en ese ejemplo), se p ocede a
la llamada del ke nel con la unción “TimesTwo”, con la que en ia emos “d_A” y
ecibi emos “d_B” que se á el a ay mul iplicado po dos.
En el siguien e ozo de código puede e se como unciona la ejecución de un
ke nel y como se usan los hilos.
“TimesTwo” es el p og ama que a a ejecu a se en la GPU po lo que hay que
pone asigna le la unción “__global__” que indica que ha sido in ocada desde
el Hos (CPU).
En es e caso, las a iables "i" y "N" se án almacenadas po cada hilo en un
egis o, y los pun e os "A" e "B" se án pun e os asignados a la memo ia del
de ice.
Sólo hay dos líneas en nues o ke nel “TimesTwo”. Como se ha mencionado
an e io men e, el ke nel es ejecu ado po múl iples hilos en pa alelo. Si
que emos que cada hilo p ocese un elemen o de la ma iz esul an e, en onces
necesi amos un medio pa a dis ingui e iden i ica cada hilo.
CUDA de ine las a iables "blockDim", "blockIdx" y " h eadIdx".
La a iable p ede inida "blockDim" con iene las dimensiones de cada bloque de
hilos como se especi ica en el segundo pa áme o de la llamada al ke nel
(Th eadsPe Block).
Las a iables p ede inidas " h eadIdx" y "blockIdx" con ienen el índice del hilo
den o de su bloque de hilos y el bloque de hilos den o del g id,
espec i amen e.
32 Uso de a je as GPU pa a acele a el p ocesado de señales
An es de u iliza es e índice pa a accede a los elemen os del a ay, su alo se
comp ueba con el núme o de elemen os “n”, pa a asegu a se que no hay
accesos a la memo ia ue a de los lími es pe mi idos.
“i (i < N)”
Es a comp obación es necesa ia pa a los casos en los que el núme o de
elemen os en un a ay no sea di isible po el amaño del bloque de hilos. Es o
pod ía esul a en que el núme o de hilos ejecu ados po el ke nel ue a
supe io al amaño del a ay.
“B[i] = 2.0 * A[i];”
La segunda línea del ke nel ealiza el abajo de "TimesTwo", mul iplicando
cada elemen o po dos.
El esul ado se ía el siguien e:
CUDA y pa alelización del algo i mo LMS 33
3.6. Pa alelización del algo i mo LMS
Los il os LMS se basan en la minimización del e o cuad á ico medio.
Es os il os son es ables y áciles de implemen a .
Desa o unadamen e, la pa alelización de es e algo i mo, especialmen e en los
sis emas in o má icos pa alelos de memo ia dis ibuida (Clus e s), no es an
ob ia.
Una des en aja p incipal del algo i mo LMS es la con e gencia len a de es e
plan eamien o.
Hay a ias a ian es de LMS incluyendo PNLMS (P opo ional No malized
Leas Mean Squa e) que se cen an en mejo a la débil con e gencia del
mé odo LMS o iginal, disminuyendo a su ez el e o ob enido, el p oblema es
que son más complejos.
El p ocedimien o de adap ación del il o equie e un cálculo signi ica i o y un
cos e de iempo, que debe se minimizado.
La con e gencia más ápida del algo i mo necesi a un mayo amaño de los
ec o es u ilizados den o del il o (miles de elemen os).
El elemen o más complejo del p oceso compu acional es el p ocedimien o de la
mul iplicación de ma ices.
T as ealiza la pa alelización ob enemos un algo i mo concu en e que
unciona como el secuencial, pe o mucho más ápido.
El algo i mo pa alelo se ha diseñado en C con el uso de la a qui ec u a CUDA
pa a se ejecu ado en GPUs nVidia.
Es e mé odo se denomina pa alelización de un solo paso y consis e en u iliza
la o denación del cálculo pa alelo pa a acele a el elemen o (ope aciones) que
más ecu sos consuma.
En el il o LMS, es e elemen o ep esen a el p oduc o de ma ices.
Fig. 3.8 Mul iplicación ma icial
Hay que calcula cada elemen o en “C”, y cada uno de los o os es
independien e, po lo que pod emos pa aleliza lo de mane a e icaz.
40 Uso de a je as GPU pa a acele a el p ocesado de señales
3.8. Resul ados ob enidos
Ejemplo de Algo i mo LMS en MATLAB (CPU):
Es e ejemplo de Leas Mean Squa e Algo i hm es p opiedad del u o del
p oyec o. La unción “gml_LS_Mgene a ion” de uel e una ma iz compleja de
12288x210 elemen os que si en pa a calcula los coe icien es necesa ios pa a
es e il o adap a i o. Hemos escogido 210 coe icien es pa a es e p oyec o.
clea all; clc; close all;
oo pa h=cd;
addpa h( oo pa h);
addpa h([ oo pa h ' oolbox']);
load('PAc eeLTE20M.ma ');
Ncoe _poly=10;
delays_poly=[-10:1:10];
X=gml_LS_Mgene a ion('MEMPOLY2',xBB,Ncoe _poly,delays_poly,0,0);
[N ows,Ncols]=size(X)
%%%%% LS sol ed wi h LMS in .m
mu=0.015;
disp(' '); disp('LS: Doing LMS');
w=complex(ze os(Ncols,1));
yBB_es =ze os(N ows,1);
e o =ze os(N ows,1);
Ts a = ic;
o n ow=1:N ows
%%% call-1 CUDA
uH=X(n ow,:);
yBB_es (n ow)=uH*w;
e o (n ow)=yBB(n ow)-yBB_es (n ow);
del aw=uH'*e o (n ow);
w=w+del aw*mu;
%%%% e u n om CUDA
end
T_LMS= oc(Ts a );
p in ('-> ime1= % n',T_LMS);
Ts a = ic;
%%% call-2 CUDA
yBB_es _pos =X*w; % es ima ion using he las coe icien s
%%% e u n om CUDA
T_pos = oc(Ts a );
p in ('-> ime2= % n',T_pos );
p in ('->NMSE es ima ion pos = % n',dpd_Qmeasu emen s(yBB,yBB_es _pos ,'NMSE'));
plo (abs(xBB),abs(yBB),'.b')
hold on
plo (abs(xBB),abs(yBB_es _pos ),'. ')

CUDA y pa alelización del algo i mo LMS 41
Resul ado del LMS en MATLAB (CPU):
--- plo (abs(xBB),abs(yBB),'.b')
--- plo (abs(xBB),abs(yBB_es _pos ),'. ')
El alo de NSME (no malized mean squa e e o ) encon ado es una
es imación de las des iaciones o ales en e los alo es que espe ábamos y los
medidos.
Cuan o meno sea el alo de NSME, mayo se á el endimien o de es e
modelo que es amos u ilizando.
Haciendo p uebas en MATLAB la es imación encon ada es de -28.62438 dB.
El iempo de ejecución del algo i mo es de 0.290911 segundos.
He ealizado el mismo algo i mo, pe o u ilizando la GPU y p og amándola con
CUDA C pa a e la mejo ía de endimien o.
Mi compañe o Rahul ha p og amado es e mismo algo i mo pa a su TFG, pe o
en lenguaje .C y u ilizando la CPU, po ello le es oy ag adecido pues ha
pe mi ido que lo compa e con mis esul ados.
Las es adap aciones del algo i mo LMS ob ienen el mismo esul ado de
NSME po lo que puede cons a a se que uncionan co ec amen e y que los
alo es ob enidos son idén icos.
El siguien e paso es p oba los y de e mina qué e sión es la más óp ima pa a
g andes can idades de da os y coe icien es.
El algo i mo ealiza dos llamadas a CUDA, en la p ime a se ob ienen los
alo es de los coe icien es (pesos) que an a adap a la señal y en la segunda
se mul iplican es os alo es po la señal de en ada y se ob iene la es imación
inal de la señal que deseamos. En las siguien es mediciones obse a emos la
du ación de es os dos p ocesos.
42 Uso de a je as GPU pa a acele a el p ocesado de señales
T as a ias p uebas u ilizando 210 coe icien es, se ha ob enido los siguien es
esul ados:
Tabla 3.1. Tiempos de ejecución del algo i mo en dis in os ipos de p ocesado
Tiempos de
ejecución (segundos)
Ve sión Ma lab
(CPU)
Ve sión CUDA
(GPU)
Ve sión .C
(CPU)
PC Labo a o io 225
(GTX 660)
0.443244 (1ª)
0.007568 (2ª)
0.165281 (1ª)
0.003169 (2ª)
0.091177
( o al)
O denado Po á il
(GeFo ce 930M)
0.290911 (1ª)
0.006496 (2ª)
0.146715 (1ª)
0.002908 (2ª)
0.064149
( o al)
PC sob emesa
pe sonal (GTX 970)
0.225325 (1ª)
0.004054 (2ª)
0.124273 (1ª)
0.002026 (2ª)
0.055408
( o al)
La p ime a y segunda llamada a CUDA co esponden a: (1ª) y (2ª).
En .C odo el código se ejecu a en una sola llamada: ( o al), po lo que el
iempo de ejecución se á equi alen e a la suma de (1ª) y (2ª).
Podemos obse a que los esul ados ob enidos son g a amen e posi i os pues
el algo i mo se es á ejecu ando, como mínimo, el doble de ápido.
Hay que ma iza a ios de alles espec o a los esul ados:
 Donde mayo di e encia de iempo se obse a es en el caso del PC
Labo a o io 225, en el que la e sión de CUDA se ejecu a casi es eces
más ápido que Ma lab.
Es o puede da se po la con igu ación ha dwa e que posee es e o denado
pues la a je a g á ica “GTX 660” pese a ene unos cuan os años (2012) es
una a je a de gama media-al a po lo que su endimien o es muy ele ado.
En compa ación con el p ocesado “In el i7 920” (2008) que es de gama
media, la capacidad de compu ación es bas an e in e io ya que en cua o
años la ecnología ha e olucionado mucho.
 U ilizando componen es in o má icos de simila ca ego ía y años emos
que el PC de sob emesa pe sonal ob iene los mejo es esul ados en cuan o
a elocidad de ejecución.
Es a a je a g á ica “GTX 970” posee 1664 núcleos CUDA, óp imos pa a el
pa alelismo, que aumen an po dos el endimien o que ob end ía el
p ocesado “i5-6500”.
 El o denado po á il o ece ambién buenos esul ados pues o que posee
componen es ac uales y de simila ca ego ía.
CUDA y pa alelización del algo i mo LMS 43
 El iempo de ejecución del p og ama en .C es muy supe io a las p uebas
ealizadas en CUDA. Es e so p enden e esul ado se explica po que
es amos usando pocos coe icien es y el “Co e Speed” o ecuencia básica
del CPU (3.20 GHz), es mucho más po en e que la de la GPU (1.05 GHz).
 Una posible solución es mon a o denado es dedicados exclusi amen e a
la compu ación GPU pues las a je as g á icas es a ían lib es de ca ga (no
hab ía ningún da o p ecompilado ni se ha ían ca go de hace unciona el
moni o ).
Es e algo i mo es á abajando con ma ices de 12288x210 alo es complejos.
La pa alelización ha enido éxi o, pe o si que emos obse a mejo es
esul ados y, po lo an o, mayo aumen o de elocidad, enemos que abaja
con más coe icien es ya que en la CPU es os cálculos se an a ealiza
secuencialmen e y a mayo núme o de elemen os, más iempo de p ocesado.
En la GPU en cambio al abaja con an os hilos en pa alelo, es os cálculos se
ealizan al mismo iempo en g andes bloques.
T as p oba es e algo i mo en CUDA u ilizando can idades dis in as de
coe icien es obse amos lo siguien e:
Tabla 3.2. Tiempos de ejecución del algo i mo con dis in os coe icien es
Nº Coe icien es
Du ación con MATLAB (s)
Du ación con CUDA (s)
Speedup
21
0,139905
0,139693
1,002
64
0,161046
0,129507
1,244
128
0,178201
0,135671
1,313
256
0,226118
0,149717
1,510
512
0,364946
0,168932
2,160
1024
0,715634
0,223281
3,205
Cuando hablamos de que se es á u ilizando 21 o 1024 coe icien es, lo que
CUDA es á p ocesando son ma ices de 12288xcoe icien es po lo que, en el
úl imo ejemplo de 1024 coe icien es, la ma iz p ocesada se ía es a:
Fig. 3.10 Ma iz p ocesada po la GPU
44 Uso de a je as GPU pa a acele a el p ocesado de señales
El siguien e g á ico nos mues a la acele ación de CUDA con espec o al uso
de la CPU con MATLAB.
Fig. 3.11 Mejo a de CUDA espec o a la CPU
Comp obamos que a mayo núme o de coe icien es, ease mayo amaño de
ma ices, el iempo de ejecución en MATLAB aumen a conside ablemen e
has a casi seis eces en e los 21 coe icien es y los 1024.
En cambio g acias a CUDA el iempo de p ocesado apenas a ía. El único
p oblema de cuda es que con ma ices pequeñas y poca can idad de da os a
p ocesa , es e apenas o ece mejo ía pues apenas se usa la pa alelización.
Vemos que con 21 coe icien es (ma ices de 12288x210), la CPU es
su icien emen e po en e como pa a iguala el endimien o de la GPU.
De o o lado, con 1024 coe icien es, la CPU iplica el iempo de p ocesado
espec o a la GPU con lo que es ealmen e ú il el uso de es a úl ima.
Aho a ol iendo al código en .C que me ha p es ado Rahul, si hacemos
p uebas pe o en es e caso con 1024 coe icien es, el algo i mo en CUDA se
uel e más e icaz que en .C.
CUDA y pa alelización del algo i mo LMS 45
Tabla 3.3. Tiempos de ejecución del algo i mo con 1024 coe icien es
Tiempos de
ejecución (segundos)
Ve sión Ma lab
(CPU)
Ve sión CUDA
(GPU)
Ve sión .C
(CPU)
PC Labo a o io 225
(GTX 660)
1.248812
0.235045
0.318960
O denado Po á il
(GeFo ce 930M)
0.826053
0.230159
0.284348
PC sob emesa
pe sonal (GTX 970)
0.715634
0.223281
0.260961
Puede obse a se que la acele ación de CUDA en el PC del Labo a o io es de
5 eces espec o a Ma lab y es á po encima de .C.
Como se ha comen ado an es, las mejo ías en el PC de sob emesa y el po a il
son algo meno es pues an o la CPU como la GPU son muy po en es y no se
ap ecia an o la di e encia.
Algo a ene en cuen a con el algo i mo LMS es que pa a calcula la es imación
de la señal de salida y los coe icien es, es necesa io calcula los alo es
u ilizando los an e io es y así an as eces como da os de en ada engamos.
Es o puede penaliza le emen e al pa alelismo de CUDA ya que la mane a
óp ima de u ilización se ob iene ealizando muchas i e aciones en un bucles en
los que los da os no dependen de su an e io i e ación.
También hay que ene en cuen a que se ha abajado con un iche o de 12288
da os complejos po lo que u ilizando mayo can idad de da os y algún mé odo
de op imización del código en CUDA, los esul ados y la acele ación se ían
mucho supe io es a los encon ados.

46 Uso de a je as GPU pa a acele a el p ocesado de señales
Conclusiones
Es e p oyec o ha se ido pa a e el po encial de la a qui ec u a CUDA y
en ende su uncionamien o aplicándola al algo i mo de un il o adap a i o.
El lenguaje CUDA C puede aplica se en cualquie sec o como se ha explicado
en capí ulos an e io es, cualquie código que enga elemen os pa alelizables es
pe ec o pa a su con e sión.
El e dade o desa ío con la op imización de los algo i mos es de e mina qué
pa es se pueden ealiza en la GPU, y qué pa es se deben ealiza en la CPU.
El algo i mo LMS iene el p opósi o de educi el uido p esen e en la señal de
en ada de al modo que la señal de salida del il o se ap oxime lo más posible
a una señal deseada. Es o lo con ie e en un buen ejemplo pa a aplica la
pa alelización pues necesi a mayo núme o de i e aciones pa a llega a la
con e gencia y a su ez mayo núme o de coe icien es.
Es e algo i mo ha sido o almen e adap ado a CUDA y se ha demos ado que, a
mayo exigencia compu acional, mejo inde la GPU en e a la CPU.
Pa a ob ene mejo es esul ados pod ía u iliza se un se ido GPU dedicado,
pues se dispond ía de la mejo ecnología ac ual. En nues os ejemplos la
a je a g á ica es á abajando ambién con el moni o , po lo que no es á
o almen e lib e de abajo.
Me ha esul ado muy in e esan e ap ende desde ce o el uncionamien o de
es a a qui ec u a, pues has a el momen o desconocía o almen e sus en ajas y
eo que puedo aplica la a lo que quie a, desde el p ocesado de señales que
hemos es udiado a lo la go de la ca e a, como al sec o inancie o o de
de ensa.
La compu ación pa alela es á cada ez más en auge, las p incipales emp esas
la u ilizan y como se ha mencionado an es, N idia ha c eado la ed neu onal
a i icial más g ande del mundo median e el uso de múl iples GPUs.
Tengo cla o que an a apa ece ecnologías que desbanquen a es a, pe o po
el momen o CUDA es el u u o y oy a segui adqui iendo sus conocimien os.
Bibliog a ía 47
Bibliog a ía
[1] N idia, CUDA C p og amming guide. PG-02829-001_ 8.0, Janua y 2017.
[2] Jason Laska, W i ing C Func ions in MATLAB (MEX-Files).
[3] KV, NSmi h, Accele a ing MATLAB wi h CUDA™ Using MEX Files, WP-
03495-001_ 01, Sep embe 2007.
[4] Joss Knigh , Calling CUDA-accele a ed Lib a ies om MATLAb: A
Compu e Vision Example, July 29, 2014.
[5] N idia, Accele a e you Applica ion wi h CUDA C.
[6] Ki k, D., Hwu, P og amming Massi ely Pa allel P ocesso s: A Hands-on
App oach. 1 edn. Mo gan Kau mann (Feb ua y 2010)
[7] Wojciech Bożejko, And zej Dob ucki, Maciej Walczyński, Pa allelizing o
digi al signal p ocessing wi h using GPU, 25 Sep . 2010.
[8] Nikolaos Ploskas Nikolaos Sama as, GPU P og amming in MATLAB,
chap e 8, “MATLAB MEX unc ions con aining CUDA code”, pp. 219-239,
edn. Mo gan Kau mann, 28 h July 2016.
[9] Pascal Ge eue , W i ing MATLAB C/MEX Code, Ap il 2010.
[10] h p://www.n idia.es/objec /gpu-compu ing-applica ions-es.h ml,
Aplicaciones acele adas en la GPU.
[11] h p://www.n idia.es/objec /cuda-pa allel-compu ing-es.h ml,
P ocesamien o pa alelo CUDA.
[12] N idia, CUBLAS LIBRARY, DU-06702-001_ 8.0 | Janua y 2017.
[13] Jagdamb Beha i S i as a a, R.K. pandey, Ji end a Jain, Implemen a ion
o Digi al Signal P ocessing Algo i hm in Gene al Pu pose G aphics
P ocessing Uni (GPGPU), Vol. 1, Issue 4, June 2013.
[14] Maciej WALCZYŃSKI, Wojciech BOŻEJKO, Noise educ ion wi h using
pa allel algo i hms, Zamek Książ - Wałb zych, Poland, 6-9 June 2010.
[15] Egil Fykse, Pe o mance Compa ison o GPU, DSP and FPGA
implemen a ions o image p ocessing and compu e ision algo i hms in
embedded sys ems, June 2013.
48 Uso de a je as GPU pa a acele a el p ocesado de señales
[16] Manuel Alejand o Ayala, Nadezhda Rocío Có do a, Diseño e
implemen ación de un il o con algo i mo adap a i o en pga pa a la
cancelación de uido, año 2016.
[17] Jo ge Lo en e, Miguel Fe e , Real- ime adap i e algo i hms using a
G aphics P ocessing Uni
[18] B. Wid ow, S. D. S ea ns, Adap i e Signal P ocessing, P en ice-Hall
Signal P ocessing Se ies, 1985.
[19] A. Gonzalez, C. Gonzalez, F. J. Ma inezZaldi a , C. Rami o, S. Roge ,
A. M. Vidal, The impac o GPU/Mul ico e in Signal P ocessing: a
quan i a i e app oach, in Wa es, ol. 3, pp. 96-106, 2011.
Anexos 49
Anexos
Ejemplo TimesTwo en CUDA C:
/* Example o how o use he mxGPUA ay API in a MEX ile. This example shows how o w i e a
MEX unc ion ha akes a gpuA ay inpu and e u ns a gpuA ay ou pu , e.g. B=mexFunc ion(A).
* Copy igh 2012 The Ma hWo ks, Inc. */
#include "mex.h"
#include "gpu/mxGPUA ay.h"
/* De ice code */
oid __global__ TimesTwo(double cons * cons A,
double * cons B,
in cons N)
{
/* Calcula e he global linea index, assuming a 1-d g id. */
in cons i = blockDim.x * blockIdx.x + h eadIdx.x;
i (i < N) {
B[i] = 2.0 * A[i];
}
}
/* Hos code */
oid mexFunc ion(in nlhs, mxA ay *plhs[],
in n hs, mxA ay cons *p hs[])
{
/* Decla e all a iables.*/
mxGPUA ay cons *A;
mxGPUA ay *B;
double cons *d_A;
double *d_B;
in N;
cha cons * cons e Id = "pa allel:gpu:mexGPUExample:In alidInpu ";
cha cons * cons e Msg = "In alid inpu o MEX ile.";
/* Choose a easonably sized numbe o h eads o he block. */
in cons h eadsPe Block = 256;
in blocksPe G id;
/* Ini ialize he Ma hWo ks GPU API. */
mxIni GPU();
/* Th ow an e o i he inpu is no a GPU a ay. */
i ((n hs!=1) || !(mxIsGPUA ay(p hs[0]))) {
mexE MsgIdAndTx (e Id, e Msg);
}
A = mxGPUC ea eF omMxA ay(p hs[0]);
/* Ve i y ha A eally is a double a ay be o e ex ac ing he poin e . */
i (mxGPUGe ClassID(A) != mxDOUBLE_CLASS) {
mexE MsgIdAndTx (e Id, e Msg);
}
/*
* Now ha we ha e e i ied he da a ype, ex ac a poin e o he inpu
* da a on he de ice. */
d_A = (double cons *)(mxGPUGe Da aReadOnly(A));
56 Uso de a je as GPU pa a acele a el p ocesado de señales
P ime a llamada a CUDA, “cuda_lms_complex.cu”:
#include "mex.h"
#include "gpu/mxGPUA ay.h"
#include <cuComplex.h>
/* De ice code */
__global__ oid yBB_Es ima ion(cuDoubleComplex *yes , cuDoubleComplex *w, cuDoubleComplex cons *X, in
ncoe s)
{
__sha ed__ cuDoubleComplex lyes [256];
unsigned in id = h eadIdx.x;
unsigned in i = h eadIdx.x;
lyes [ id].x = 0; lyes [ id].y = 0;
i (i<ncoe s)
lyes [ id] = cuCmul( w[ id] , X[ id]);
__sync h eads();
o (unsigned in s=blockDim.x/2; s>0; s>>=1) // Reduc ion
{
i ( id < s) {
lyes [ id] = cuCadd(lyes [ id], lyes [ id+s]);
}
__sync h eads();
}
i ( id == 0)
yes [0] = lyes [0]; // yBB_es (n ow); in MATLAB
}
__global__ oid e o _Es ima ion(cuDoubleComplex *del aw, cuDoubleComplex cons *yBB, cuDoubleComplex
*yes , cuDoubleComplex cons *X, in ncoe s, cuDoubleComplex *w, double mu, cuDoubleComplex *d_e o )
{
in idx = h eadIdx.x;
cuDoubleComplex e o = cuCsub(*yBB, *yes ); // e o (n ow)=yBB(n ow)-yBB_es (n ow); in MATLAB
i (idx<ncoe s)
{
del aw[idx] = cuCmul(cuConj(X[idx]) , e o ); // del aw=X'*e o (n ow); in MATLAB
del aw[idx].x = del aw[idx].x*mu;
del aw[idx].y = del aw[idx].y*mu;
w[idx] = cuCadd( w[idx] , del aw[idx]); // w=w+del aw*mu; in MATLAB
}
i (idx==0)
{
d_e o [0] = e o ;
}
}
/* Hos code */
oid mexFunc ion(in nlhs, mxA ay *plhs[],
in n hs, mxA ay cons *p hs[])
{
/* Decla e all a iables.*/
mxGPUA ay cons *X, *yBB;
cuDoubleComplex cons *d_X, *d_yBB;
cuDoubleComplex *d_e o , *d_yBB_es , *del aw, *d_w;
cuDoubleComplex *w_e o , *w;
in N ows, ncoe s;
cha cons * cons e Id = "pa allel:gpu:mexGPUExample:In alidInpu ";
cha cons * cons e Msg = "In alid inpu o MEX ile.";
/* Ini ialize he Ma hWo ks GPU API. */
mxIni GPU();
/* Th ow an e o i he inpu is no a GPU a ay. */
i ((n hs!=2) || !(mxIsGPUA ay(p hs[0])))
{
mexE MsgIdAndTx (e Id, e Msg);
}

Anexos 57
X = mxGPUC ea eF omMxA ay(p hs[0]);
yBB = mxGPUC ea eF omMxA ay(p hs[1]);
/* Ve i y ha A eally is a double a ay be o e ex ac ing he poin e . */
i (mxGPUGe ClassID(X) != mxDOUBLE_CLASS)
{
mexE MsgIdAndTx (e Id, e Msg);
}
/* Now ha we ha e e i ied he da a ype, ex ac a poin e o he inpu da a on he de ice. */
d_X = (cuDoubleComplex cons *)(mxGPUGe Da aReadOnly(X));
d_yBB = (cuDoubleComplex cons *)(mxGPUGe Da aReadOnly(yBB));
N ows = (in )(mxGPUGe Numbe O Elemen s(yBB));
ncoe s = (in )(mxGPUGe Numbe O Elemen s(X))/N ows;
/* C ea e a GPUA ay o hold he esul and ge i s unde lying poin e . */
w_e o = (cuDoubleComplex*) malloc(sizeo (cuDoubleComplex) * N ows);
w = (cuDoubleComplex*) malloc(sizeo (cuDoubleComplex) * ncoe s);
cudaMalloc(( oid**)&d_w, sizeo (cuDoubleComplex) * ncoe s);
cudaMalloc(( oid**)&del aw, sizeo (cuDoubleComplex) * ncoe s);
cudaMalloc(( oid**)&d_e o , sizeo (cuDoubleComplex) * N ows);
cudaMalloc(( oid**)&d_yBB_es , sizeo (cuDoubleComplex));
double *w_ , *w_i;
double *e o _ , *e o _i;
plhs[0]=mxC ea eDoubleMa ix(ncoe s, 1, mxCOMPLEX);
w_ = mxGe P (plhs[0]);
w_i = mxGe Pi(plhs[0]);
plhs[1]=mxC ea eDoubleMa ix(N ows, 1, mxCOMPLEX);
e o _ = mxGe P (plhs[1]);
e o _i = mxGe Pi(plhs[1]);
double mu = 0.015;
o (in i=0; i<ncoe s; i++)
{
w[i].x = 0;
w[i].y = 0;
}
cudaMemcpy(d_w, w, sizeo (cuDoubleComplex) * ncoe s, cudaMemcpyHos ToDe ice);
cudaE en _ s a , s op;
cudaE en C ea e(&s a );
cudaE en C ea e(&s op);
cudaE en Reco d(s a );
o (in i=0; i<N ows; i++)
{
yBB_Es ima ion<<<1,256>>>(d_yBB_es , d_w, &d_X[i*ncoe s], ncoe s);
e o _Es ima ion<<<1,256>>>(del aw, &d_yBB[i] , d_yBB_es , &d_X[i*ncoe s], ncoe s, d_w,
mu, &d_e o [i] );
}
cudaMemcpy(w_e o , d_e o , sizeo (cuDoubleComplex) * N ows, cudaMemcpyDe iceToHos );
cudaMemcpy(w, d_w, sizeo (cuDoubleComplex) * ncoe s, cudaMemcpyDe iceToHos );
cudaE en Reco d(s op);
cudaE en Synch onize(s op);
loa milliseconds = 0;
cudaE en ElapsedTime(&milliseconds, s a , s op);
p in ("Execu ion ime elapsed: % ms n", milliseconds);
o (in i=0; i<ncoe s; i++)
{
w_ [i] = w[i].x;
w_i[i] = w[i].y;
}
o (in i=0; i<N ows; i++)
{
e o _ [i] = w_e o [i].x;
e o _i[i] = w_e o [i].y;
}
mxGPUDes oyGPUA ay(yBB);
}
58 Uso de a je as GPU pa a acele a el p ocesado de señales
Segunda llamada a CUDA, “cuda_mul iply.cu”:
#include "mex.h"
#include "gpu/mxGPUA ay.h"
#include <cuComplex.h>
__global__ oid yBB_Es ima ion_Pos (cuDoubleComplex *yes _pos , cuDoubleComplex cons *X,
cuDoubleComplex cons *W, in ncoe s)
{
__sha ed__ cuDoubleComplex lyes _pos [256];
unsigned in id = h eadIdx.x;
unsigned in s id = blockIdx.x*ncoe s + id;
lyes _pos [ id].x = 0; lyes _pos [ id].y = 0;
i ( id<ncoe s)
lyes _pos [ id] = cuCmul( X[s id] , W[ id]); // yBB_es _pos =X*w; in MATLAB
__sync h eads();
o (unsigned in s=blockDim.x/2; s>0; s>>=1) // Reduc ion
{
i ( id < s) {
lyes _pos [ id] = cuCadd(lyes _pos [ id], lyes _pos [ id+s]);
}
__sync h eads();
}
i ( id == 0)
yes _pos [blockIdx.x] = lyes _pos [0];
}
/* Hos code */
oid mexFunc ion(in nlhs, mxA ay *plhs[],
in n hs, mxA ay cons *p hs[])
{
/* Decla e all a iables.*/
mxGPUA ay cons *g_X, *g_w;
cuDoubleComplex cons *d_g_X, *d_g_w;
cuDoubleComplex *d_yBB_es _pos ;
cuDoubleComplex *yBB_es _pos ;
in N ows, ncoe s;
cha cons * cons e Id = "pa allel:gpu:mexGPUExample:In alidInpu ";
cha cons * cons e Msg = "In alid inpu o MEX ile.";
mxIni GPU(); /* Ini ialize he Ma hWo ks GPU API. */
/* Th ow an e o i he inpu is no a GPU a ay. */
i ((n hs!=2) || !(mxIsGPUA ay(p hs[0])))
{
mexE MsgIdAndTx (e Id, e Msg);
}
g_X = mxGPUC ea eF omMxA ay(p hs[0]);
g_w = mxGPUC ea eF omMxA ay(p hs[1]);
/* Ve i y ha g_X eally is a double a ay be o e ex ac ing he poin e . */
i (mxGPUGe ClassID(g_X) != mxDOUBLE_CLASS || mxGPUGe ClassID(g_w) != mxDOUBLE_CLASS)
{
mexE MsgIdAndTx (e Id, e Msg);
}
/* Now ha we ha e e i ied he da a ype, ex ac a poin e o he inpu on he de ice. */
d_g_X = (cuDoubleComplex cons *)(mxGPUGe Da aReadOnly(g_X));
d_g_w = (cuDoubleComplex cons *)(mxGPUGe Da aReadOnly(g_w));
ncoe s = (in )(mxGPUGe Numbe O Elemen s(g_w));
N ows = (in )(mxGPUGe Numbe O Elemen s(g_X))/ncoe s;
/* C ea e a GPUA ay o hold he esul and ge i s unde lying poin e . */
yBB_es _pos = (cuDoubleComplex*) malloc(sizeo (cuDoubleComplex) * N ows);
cudaMalloc(( oid**)&d_yBB_es _pos , sizeo (cuDoubleComplex) * N ows);
Anexos 59
double *yBB_es _pos _ , *yBB_es _pos _i;
plhs[0]=mxC ea eDoubleMa ix(N ows, 1, mxCOMPLEX);
yBB_es _pos _ = mxGe P (plhs[0]);
yBB_es _pos _i = mxGe Pi(plhs[0]);
cudaE en _ s a , s op;
cudaE en C ea e(&s a );
cudaE en C ea e(&s op);
cudaE en Reco d(s a );
yBB_Es ima ion_Pos <<<N ows,256>>>(d_yBB_es _pos , d_g_X , d_g_w, ncoe s);
cudaMemcpy(yBB_es _pos , d_yBB_es _pos , sizeo (cuDoubleComplex) * N ows,
cudaMemcpyDe iceToHos );
cudaE en Reco d(s op);
cudaE en Synch onize(s op);
loa milliseconds = 0;
cudaE en ElapsedTime(&milliseconds, s a , s op);
p in ("Execu ion ime elapsed: % ms n", milliseconds);
o (in i=0; i<N ows; i++)
{
yBB_es _pos _ [i] = yBB_es _pos [i].x;
yBB_es _pos _i[i] = yBB_es _pos [i].y;
}
mxGPUDes oyGPUA ay(g_X);
mxGPUDes oyGPUA ay(g_w);
}