UNIVERSIDAD DE BURGOS
Á ea de Tecnología Elec ónica
INTRODUCCIÓN A LA
PROGRAMACIÓN EN CUDA
Césa Rep esa Pé ez
José Ma ía Cáma a Neb eda
Ped o Luis Sánchez O ega
In oducción a la p og amación en CUDA©2016 3.1
Á ea de Tecnología Elec ónica
Depa amen o de Ingenie ía Elec omecánica
Uni e sidad de Bu gos
ÍNDICE
INTRODUCCIÓN: GPU Compu ing.........................................................................5
PRÁCTICA nº 1: Disposi i os CUDA ......................................................................13
PRÁCTICA nº 2: T ans e encia de da os ................................................................21
PRÁCTICA nº 3: Lanzamien o de un Ke nel..........................................................25
PRÁCTICA nº 4: Hilos y Bloques ............................................................................31
PRÁCTICA nº 5: Tempo ización y E o es .............................................................37
PRÁCTICA nº 6: A ays Mul idimensionales.........................................................45
PRÁCTICA nº 7: Memo ia Compa ida...................................................................51
PRÁCTICA nº 8: Memo ia Cons an e .....................................................................57
PRÁCTICA nº 9: G á icos en CUDA........................................................................63
PRÁCTICA nº 10: Eje cicio de aplicación p ác ica .................................................71
En es e manual se han u ilizado las siguien es unciones de CUDA 5.0:
Función Página
cudaGe De iceCoun ()...............................................................................................17
cudaGe De ice()........................................................................................................17
cudaSe De ice()........................................................................................................17
cudaGe De iceP ope ies()......................................................................................17
cudaMalloc().............................................................................................................22
cudaMemcpy().............................................................................................................22
cudaF ee().................................................................................................................23
cudaMemGe In o()......................................................................................................23
cudaE en C ea e()....................................................................................................38
cudaE en Reco d()....................................................................................................39
cudaE en Synch onize() ...........................................................................................39
cudaE en ElapsedTime() ...........................................................................................39
cudaE en Des oy() ..................................................................................................40
cudaGe E o S ing()...............................................................................................41
cudaDe iceSynch onize() .........................................................................................41
cudaGe Las E o () ..................................................................................................41
__sync h eads()........................................................................................................53
cudaMemcpyToSymbol()...............................................................................................58
INTRODUCCIÓN: GPU Compu ing
INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA 5
INTRODUCCIÓN:
GPU Compu ing
1. Compu ación en sis emas he e ogéneos
Podemos de ini la compu ación sob e a je as g á icas (en inglés GPU compu ing) como
el uso de una a je a g á ica (GPU - G aphics P ocessing Uni ) pa a ealiza cálculos
cien í icos de p opósi o gene al. El modelo de compu ación sob e a je as g á icas consis e en
usa conjun amen e una CPU (Cen al P ocessing Uni ) y una GPU de mane a que o men un
modelo de compu ación he e ogéneo (Figu a i.1). Siguiendo es e modelo, la pa e secuencial
de una aplicación se ejecu a ía sob e la CPU (comúnmen e denominada hos ) y la pa e más
cos osa del cálculo se ejecu a ía sob e la GPU (que se denomina de ice). Desde el pun o de
is a del usua io, la aplicación simplemen e se a a ejecu a más ápido po que es á u ilizando
las al as p es aciones de la GPU pa a inc emen a el endimien o.
HOST DEVICE
Figu a i.1. Modelo de compu ación en sis emas he e ogéneos: CPU + GPU.
El p oblema inicial del uso de las a je as g á icas pa a el cálculo cien í ico de p opósi o
gene al (en inglés GPGPU - Gene al-Pu pose Compu ing on G aphics P ocessing Uni s) e a
que se necesi aba usa lenguajes de p og amación especí icos pa a g á icos como el OpenGL
o el Cg pa a p og ama la GPU. Los desa ollado es debían hace que sus aplicaciones
cien í icas pa ecie an aplicaciones g á icas con i iéndolas en p oblemas que dibujaban

INTRODUCCIÓN: GPU Compu ing
6 INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA
iángulos y polígonos. Es o cla amen e limi aba el acceso po pa e del mundo cien í ico al
eno me endimien o de las GPUs.
NVIDIA ue conscien e del po encial que suponía ace ca es e eno me endimien o a la
comunidad cien í ica en gene al y decidió in es iga la o ma de modi ica la a qui ec u a de
sus GPUs pa a que ue an comple amen e p og amables pa a aplicaciones cien í icas además
de añadi sopo e pa a lenguajes de al o ni el como C y C++. De es e modo, en No iemb e de
2009, NVIDIA in odujo pa a sus a je as g á icas la a qui ec u a CUDATM (Compu e Uni ied
De ice A chi ec u e), una nue a a qui ec u a pa a cálculo pa alelo de p opósi o gene al, con
un nue o epe o io de ins ucciones y un nue o modelo de p og amación pa alela, con
sopo e pa a lenguajes de al o ni el (Figu a i.2), y cons i uidas po cien os de núcleos que
pueden p ocesa de mane a concu en e miles de hilos de ejecución. En es a a qui ec u a, cada
núcleo iene cie os ecu sos compa idos, incluyendo egis os y memo ia. La memo ia
compa ida in eg ada en el p opio chip pe mi e que las a eas que se es án ejecu ando en es os
núcleos compa an da os sin ene que en ia los a a és del bus de memo ia del sis ema.
Figu a i.2. CUDA es á diseñado pa a sopo a di e en es lenguajes de p og amación.
2. A qui ec u a CUDA
En la a qui ec u a clásica de una a je a g á ica podemos encon a la p esencia de dos
ipos de p ocesado es, los p ocesado es de é ices y los p ocesado es de agmen os,
dedicados a a eas dis in as e independien es den o del cauce g á ico, y con epe o ios de
ins ucciones di e en es. Es o p esen a dos p oblemas impo an es, po un lado el
desequilib io de ca ga que apa ece en e ambos p ocesado es y po o o la di e encia en e sus
espec i os epe o ios de ins ucciones. De es e modo, la e olución na u al en la a qui ec u a
INTRODUCCIÓN: GPU Compu ing
INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA 7
de una GPU ha sido la búsqueda de una a qui ec u a uni icada donde no se dis inguie a en e
ambos ipos de p ocesado es. Así se llegó a la a qui ec u a CUDATM, donde odos los núcleos
de ejecución necesi an el mismo epe o io de ins ucciones y p ác icamen e los mismos
ecu sos.
Figu a i.3. A qui ec u a de una a je a g á ica CUDA-enabled.
En la Figu a i.3 se mues a la a qui ec u a de una a je a g á ica compa ible con CUDA.
En ella se puede obse a la p esencia de unas unidades de ejecución denominadas S eaming
Mul ip ocesso s (SM), 8 unidades en el ejemplo de la igu a, que es án in e conec adas en e
sí po una zona de memo ia común. Cada SM es á compues o a su ez po unos núcleos de
cómpu o llamados “núcleos CUDA” o S eaming P ocesso s (SP), que son los enca gados de
ejecu a las ins ucciones y que en nues o ejemplo emos que hay 32 núcleos po cada SM,
lo que hace un o al de 256 núcleos de p ocesamien o. Es e diseño ha dwa e pe mi e la
p og amación sencilla de los núcleos de la GPU u ilizando un lenguaje de al o ni el como
puede se el lenguaje C pa a CUDA. De es e modo, el p og amado simplemen e esc ibe un
p og ama secuencial den o del cual se llama a lo que se conoce como ke nel, que puede se
una simple unción o un p og ama comple o. Es e ke nel se ejecu a de o ma pa alela den o
de la GPU como un conjun o de hilos ( h eads) y que el p og amado o ganiza den o de una
je a quía en la que pueden ag upa se en bloques (blocks), y que a su ez se pueden dis ibui
o mando una malla (g id), al como se mues a en la Figu a i.4. Po con eniencia, los
INTRODUCCIÓN: GPU Compu ing
8 INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA
bloques y las mallas pueden ene una, dos o es dimensiones. Exis en mul i ud de
si uaciones en las que los da os con los que se abaja poseen de o ma na u al una es uc u a
de malla, pe o en gene al, descompone los da os en una je a quía de hilos no es una a ea
ácil. Así pues, un bloque de hilos es un conjun o de hilos concu en es que pueden coope a
en e ellos a a és de mecanismos de sinc onización y compa i accesos a un espacio de
memo ia exclusi o de cada bloque. Y una malla es un conjun o de bloques que pueden se
ejecu ados independien emen e y que po lo an o pueden se lanzados en pa alelo en los
S eaming Mul ip ocesso s (SM).
Figu a i.4. Je a quía de hilos en una aplicación CUDA.
Cuando se in oca un ke nel, el p og amado especi ica el núme o de hilos po bloque y el
núme o de bloques que con o man la malla. Una ez en la GPU, a cada hilo se le asigna un
único núme o de iden i icación den o de su bloque, y cada bloque ecibe un iden i icado
den o de la malla. Es o pe mi e que cada hilo decida sob e qué da os iene que abaja , lo
que simpli ica eno memen e el di eccionamien o de memo ia cuando se abaja con da os
mul idimensionales, como es el caso del p ocesado de imágenes o la esolución de ecuaciones
di e enciales en dos y es dimensiones.
O o aspec o a des aca en la a qui ec u a CUDA es la p esencia de una unidad de
dis ibución de abajo que se enca ga de dis ibui los bloques en e los SM disponibles. Los
hilos den o de cada bloque se ejecu an concu en emen e y cuando un bloque e mina, la
unidad de dis ibución lanza nue os bloques sob e los SM lib es. Los SM mapean cada hilo
sob e un núcleo SP, y cada hilo se ejecu a de mane a independien e con su p opio con ado de
p og ama y egis os de es ado. Dado que cada hilo iene asignados sus p opios egis os, no
INTRODUCCIÓN: GPU Compu ing
INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA 9
exis e penalización po los cambio de con ex o, pe o en cambio sí exis e un lími e en el
núme o máximo de hilos ac i os debido a que cada SM iene un nume o de e minado de
egis os.
Una ca ac e ís ica pa icula de la a qui ec u a CUDA es la ag upación de los hilos en
g upos de 32. Un g upo de 32 hilos ecibe el nomb e de wa p, y se puede conside a como la
unidad de ejecución en pa alelo, ya que odos los hilos de un mismo wa p se ejecu an
ísicamen e en pa alelo y po lo an o comienzan en la misma ins ucción (aunque después
son lib es de bi u ca se y ejecu a se independien emen e). Así, cuando se selecciona un
bloque pa a su ejecución den o de un SM, el bloque se di ide en wa ps, se selecciona uno
que es é lis o pa a ejecu a se y se emi e la siguien e ins ucción a odos los hilos que o man
el wa p. Dado que odos ellos ejecu an la misma ins ucción al unísono, la máxima e iciencia
se consigue cuando odos los hilos coinciden en su u a de ejecución (sin bi u caciones).
Aunque el p og amado puede igno a es e compo amien o, con iene ene lo en cuen a si se
p e ende op imiza alguna aplicación.
Figu a i.5. Je a quía de memo ia den o de la a qui ec u a CUDA.
En cuan o a la memo ia, du an e su ejecución los hilos pueden accede a los da os desde
di e en es espacios den o de una je a quía de memo ia (Figu a i.5). Así, cada hilo iene una
zona p i ada de memo ia local y cada bloque iene una zona de memo ia compa ida
PRÁCTICA nº 1: Disposi i os CUDA
16 INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA
1. Copia el con enido de la ca pe a “ CUDA Samples 5.0 0_Simple empla e” en
una nue a ca pe a c eada en la misma u a que el es o de ejemplos y con el nomb e
que deseemos (po ejemplo CUDA Samples 5.0 p ac icas).
2. Elimina los a chi os de p oyec o y solución que no co espondan a nues a e sión de
Visual S udio (po ejemplo, pa a la e sión 2010 nos queda emos con los iche os
“ empla e_ s2010. cxp oj” y “ empla e_ s2010.sln”).
3. Ab i la solución “ empla e_ s2010.sln” con Visual S udio y cons ui la pa a
win32 en modo “debug”.
4. Ejecu a el esul ado “ empla e.exe” desde Visual S udio o bien desde su ubicación
en CUDA Samples 5.0 bin win32 Debug.
Si hemos log ado cons ui y ejecu a el p oyec o sin e o es que á deci que odo el
en o no de NVIDIA pa a CUDA es á ins alado co ec amen e y que podemos comenza a
edi a los di e en es a chi os uen e y modi ica los pa a ealiza nues os p opios cálculos.
Pa a ello comenzamos eliminando de nues o p oyec o los a chi os “ empla e_cpu.cpp”
y “ empla e_ke nel.cu”. De es e modo, nues o p oyec o con end á un único a chi o
uen e llamado “ empla e.cu”, que es el único que debemos u iliza pa a esc ibi nues o
p opio código.
Po o o lado, aunque no es necesa io, podemos es a in e esados en enomb a nues os
a chi os y u iliza o o nomb e, como po ejemplo “p ac ica” . Pa a ello, una ez ce ado
Visual S udio y con los cambios an e io es gua dados, podemos p ocede de la siguien e
o ma:
5. Modi ica los nomb es de los iche os uen e, p oyec o y solución sus i uyendo la
palab a “ empla e” po o o nomb e (po ejemplo, po la palab a “p ac ica”). De es e
modo la ca pe a queda á con 3 a chi os:
“p ac ica.cu”
“p ac ica_ s2010.sln”
“p ac ica_ s2010. cxp oj”
6. Edi a el con enido de los a chi os *.sln y *. cxp oj con un edi o de ex o y
eemplaza odas la ocu encias de la palab a “ empla e” po la palab a escogida
an e io men e (en nues o ejemplo, po la palab a “p ac ica”).

PRÁCTICA nº 1: Disposi i os CUDA
INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA 17
3. P opiedades de un disposi i o CUDA
Además de las ca ac e ís icas gene ales de una a je a g á ica con a qui ec u a CUDA
exis en o as ca ac e ís icas que dependen de cada disposi i o en pa icula . Po ese mo i o,
es muy impo an e conoce las ca ac e ís icas ac uales del disposi i o o disposi i os con los
que amos a abaja , como po ejemplo cuán a memo ia y qué capacidad de cómpu o posee.
Po o o lado, si disponemos de más de un disposi i o, ambién es con enien e pode
selecciona el que mejo se adap e a nues as necesidades.
Así pues, lo p ime o que enemos que hace es sabe cuán os disposi i os CUDA enemos
ins alados en nues o sis ema. Pa a ello llamamos a la unción cudaGe De iceCoun ():
cudaGe De iceCoun (in *coun );
Es a unción de uel e en la a iable coun el núme o de disposi i os con capacidad de
cómpu o igual o supe io a 1.0 (que es la especi icación mínima de un disposi i o con
a qui ec u a CUDA). Los dis in os disposi i os encon ados se iden i ican median e un
núme o en e o en el ango que a desde 0 has a coun -1. Una ez que conocemos el núme o
o al de disposi i os, podemos i e a a a és de ellos y ob ene su in o mación u ilizando la
unción cudaGe De iceP ope ies():
cudaGe De iceP ope ies(cudaDe iceP op *p opiedades, in de iceID);
Con cada llamada a es a unción, odas las p opiedades del disposi i o iden i icado con el
núme o de iceID quedan almacenadas en p opiedades, que es una es uc u a del ipo
cudaDe iceP op y que con iene la in o mación o ganizada en di e en es campos al como se
indica en la TABLA 3.
Si es amos abajando en un en o no con a ios disposi i os ins alados, o as unciones
que esul an de u ilidad son la unción cudaGe De ice() y la unción cudaSe De ice():
cudaGe De ice(in *ge ID);
cudaSe De ice(in se ID);
La p ime a de uel e en la a iable ge ID el núme o de iden i icación del disposi i o
seleccionado mien as que la segunda nos pe mi e selecciona un disposi i o cuyo núme o de
iden i icación sea se ID.
PRÁCTICA nº 1: Disposi i os CUDA
18 INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA
TABLA 3. Campos con enidos en la es uc u a cudaDe iceP op.
PRÁCTICA nº 1: Disposi i os CUDA
INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA 19
REALIZACIÓN PRÁCTICA
1. Comp oba que el en o no de p og amación en CUDA es á en o den y unciona
co ec amen e esc ibiendo el ejemplo inicial po excelencia, aquel en el que se mues a po
pan alla el mensaje de bien enida: “¡Hola, mundo!”:
///////////////////////////////////////////////////////////////////////////
// includes
///////////////////////////////////////////////////////////////////////////
#include <s dio.h>
#include <s dlib.h>
#include <cuda_ un ime.h>
///////////////////////////////////////////////////////////////////////////
// de ines
///////////////////////////////////////////////////////////////////////////
///////////////////////////////////////////////////////////////////////////
// decla acion de unciones
///////////////////////////////////////////////////////////////////////////
///////////////////////////////////////////////////////////////////////////
// u ina p incipal ejecu ada en el hos
in main(in a gc, cha ** a g )
{
// cue po del p og ama
p in ("Hola, mundo!! n");
// salida del p og ama
p in (" npulsa INTRO pa a inaliza ...");
lush(s din);
cha ecla = ge cha ();
e u n 0;
}
///////////////////////////////////////////////////////////////////////////
2. Busca odos los disposi i os ins alados en el sis ema que sean compa ibles con CUDA y
mos a las siguien es p opiedades de cada uno de ellos:
» Nomb e del disposi i o.
» Capacidad de cómpu o.
» Núme o de mul ip ocesado es, S eaming Mul ip ocesso s (SM).
» Núme o de núcleos de cómpu o, S eaming P ocesso s (SP)1.
» Memo ia global (en MiB).
1 Consul a TABLA 2.
PRÁCTICA nº 2: T ans e encia de da os
INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA 21
PRÁCTICA nº 2:
T ans e encia de da os
1. Manejo de memo ia en CUDA
Como ya hemos mencionado an e io men e, el modelo de p og amación CUDA asume
un sis ema compues o po un hos y un de ice, y cada uno de ellos con su espacio de
memo ia. Un ke nel sólo pueden ope a sob e la memo ia del disposi i o, po lo que
necesi a emos unciones especí icas pa a ese a y libe a la memo ia del disposi i o, así
como unciones pa a la ans e encia da os en e la memo ia del hos y del de ice. En la
Figu a 2.1 se han ep esen ado los di e en es ni eles den o de la je a quía de memo ia.
Vemos que las únicas zonas de memo ia accesibles desde el hos son la memo ia global
(Global Memo y) y la memo ia cons an e (Cons an Memo y) y desde es as zonas de
memo ia el ke nel puede ans e i da os al es o de ni eles. Se puede obse a cómo odos
los hilos pueden accede a la zona de memo ia global/cons an e (lo que esul a en un canal de
comunicación en e odos ellos) mien as que únicamen e sólo los hilos pe enecien es a un
mismo bloque pueden accede a una zona de memo ia denominada memo ia compa ida
Figu a 2.1. Je a quía de memo ia en CUDA.

PRÁCTICA nº 2: T ans e encia de da os
22 INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA
(Sha ed Memo y).
Pa a pode ese a espacio en la zona de memo ia global del de ice y pode accede a
ella desde el hos se u iliza la unción cudaMalloc(), que iene un compo amien o simila a
la co espondien e unción es ánda de C:
cudaMalloc( oid **de P , size_ size);
El p ime a gumen o de es a unción es un doble pun e o, que co esponde a la di ección
del pun e o (de P ) en el que amos a almacena la di ección de la memo ia ese ada en el
disposi i o. El segundo a gumen o es la can idad de memo ia exp esada en by es que
que emos ese a . De es a o ma podemos ese a size by es de memo ia lineal den o de
la memo ia global de la a je a g á ica.
Una ez que enemos el espacio de memo ia ese ado en la memo ia global de nues o
disposi i o, lo siguien e que podemos que hace es ans e i da os en e es a memo ia y la
memo ia de nues a CPU. Pa a ello u ilizamos o a unción pa ecida a la que disponemos en C
es ánda pa a al e ec o, sólo que en es e caso end emos algún pa áme o adicional que nos
pe mi a especi ica el o igen y el des ino de los da os. Es a unción es cudaMemcpy():
cudaMemcpy( oid *ds , oid *s c, size_ coun , cudaMemcpyKind kind);
El p ime pa áme o (ds ) co esponde al pun e o con la di ección de des ino de los
da os, el segundo (s c) es el pun e o con la di ección de o igen, es deci , donde se encuen an
los da os que que emos copia , coun es el núme o de by es que amos a ans e i y kind es
el ipo de ans e encia que amos a ealiza . En la TABLA 4 enemos las cua o
posibilidades que exis en pa a kind y co esponden a los dis in os sen idos de ans e encia
de da os en e el hos y el de ice.
Po úl imo, la memo ia ese ada po cudaMalloc() ambién se puede libe a , y pa a ello
se u iliza la unción cudaF ee():
TABLA 4. Tipos de ans e encias de da os en CUDA.
Tipo de ans e encia Sen ido de la ans e encia
cudaMemcpyHos ToHos hos  hos
cudaMemcpyHos ToDe ice hos  de ice
cudaMemcpyDe iceToHos de ice  hos
cudaMemcpyDe iceToDe ice de ice  de ice
PRÁCTICA nº 2: T ans e encia de da os
INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA 23
cudaF ee( oid *de P );
Con es a llamada libe amos el espacio de memo ia apun ado po de P , el cual debe se
un pun e o de uel o po la unción cudaMalloc() en una llamada p e ia.
O a unción que puede se de u ilidad a lo la go de la ejecución de un p og ama es la
unción cudaMemGe In o():
cudaMemGe In o(size_ *memLib e, size_ *memTo al);
que de uel e en las a iables memLib e y memTo al el alo en by es de la can idad de
memo ia lib e en el disposi i o así como el alo de memo ia o al, espec i amen e.
2. Ejemplo
En el siguien e ejemplo se mues an las di e encias y las simili udes que exis en a la ho a
de ese a memo ia an o en el hos como en el de ice. En es e ejemplo se ese a espacio
pa a una ma iz cuad ada de NN elemen os, se inicializa en el hos con alo es alea o ios2
(en e 0 y 9) de ipo loa y después se ans ie en los da os desde el hos has a el de ice:
// includes
#include <s dio.h>
#include <s dlib.h>
#include < ime.h>
#include <cuda_ un ime.h>
#de ine N 8
// MAIN: u ina p incipal ejecu ada en el hos
in main(in a gc, cha ** a g )
{
// decla acion
loa *hs _ma iz;
loa *de _ma iz;
// ese a en el hos
hs _ma iz = ( loa *)malloc( N*N*sizeo ( loa ) );
// ese a en el de ice
cudaMalloc( ( oid**)&de _ma iz, N*N*sizeo ( loa ) );
// inicializacion de da os
s and ( (in ) ime(NULL) );
o (in i=0; i<N*N; i++)
{
hs _ma iz[i] = ( loa )( and() % 10 );
}
// copia de da os
cudaMemcpy(de _ma iz, hs _ma iz, N*N*sizeo ( loa ), cudaMemcpyHos ToDe ice);
// salida
cudaF ee( de _ma iz )
p in (" npulsa INTRO pa a inaliza ...");
lush(s din);
cha ecla = ge cha ();
e u n 0;
}
2 La unción es ánda in and() gene a núme os alea o ios comp endidos en e 0 y un alo máximo de inido
en una cons an e llamada RAND_MAX. Pa a gene a núme os dis in os cada ez que se ejecu a el p og ama es
necesa io u iliza la unción s and(in semilla) donde semilla debe se un alo dis in o en cada ejecución.
PRÁCTICA nº 2: T ans e encia de da os
24 INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA
REALIZACIÓN PRÁCTICA
1. Esc ibi un p og ama que enga como obje i o la ans e encia de da os en e el hos y el
de ice, al como se mues a en la Figu a 2.2.
2. El a ay hs _A debe se decla ado de ipo loa e inicializado con N elemen os alea o ios
comp endidos en e 0 y 1.
3. Imp imi po pan alla el con enido de los a ays hs _A y hs _B u ilizando sólo dos ci as
decimales y comp oba que son iguales.
4. U iliza la unción cudaMemGe In o() pa a calcula la can idad de memo ia ese ada en
el de ice exp esada en MiB.
Figu a 2.2. T ans e encia de da os en e hos y de ice.
HOST DEVICE
hs _A de _A
de _B hs _B
PRÁCTICA nº 3: Lanzamien o de un Ke nel
INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA 25
PRÁCTICA nº 3:
Lanzamien o de un Ke nel
1. Modelo de p og amación he e ogéneo
Pa a un p og amado de CUDA, el sis ema de cómpu o consis e en una CPU adicional
(hos ), y una o más a je as g á icas (de ice) que son p ocesado es masi amen e pa alelos,
equipados con un g an núme o de unidades a i mé ico-lógicas. En la mayo ía de las
aplicaciones ac uales, a menudo hay secciones del p og ama donde los da os p esen an un
cla o pa alelismo, es deci , que sob e esos da os se pueden ealiza muchas ope aciones
a i mé icas de mane a simul ánea. En ese escena io, los disposi i os CUDA pueden acele a
Figu a 3.1. Modelo de p og amación he e ogéneo.
PRÁCTICA nº 4: Hilos y Bloques
32 INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA
Cuando an es lanzamos el ke nel median e la sin axis an e io , especi icamos que
que íamos una malla o mada po un único bloque y N hilos pa alelos. Es o le dice al sis ema
que que emos una malla unidimensional (los alo es escala es se in e p e an como
unidimensionales) con los hilos epa idos a lo la go del eje x (Figu a 4.1), ya que po de ec o
se conside a es e eje como la dimensión de abajo (de ahí que se añada el su ijo “.x” pa a
indica dicha di ección). De es e modo, cada uno de los hilos end á un alo dis in o de
h eadIdx.x que i á desde 0 has a N1. Es deci , cada hilo end á su p opio alo de myID
que pe mi i á al p og amado decidi sob e qué da os debe abaja cada uno de ellos.
Figu a 4.1. Malla (g id) o mada po un único bloque (block) de N hilos pa alelos
( h eads) epa idos a lo la go del eje x. Cada hilo se puede iden i ica median e la
a iable h eadIdx.x.
El ha dwa e de la GPU limi a el núme o de hilos po bloque con que podemos lanza un
ke nel. Es e núme o puede se mayo o meno dependiendo de la capacidad de cómpu o de
nues a GPU y en pa icula no puede supe a el alo dado po maxTh eadsPe Block, que es
uno de los campos que o man pa e de la es uc u a de p opiedades cudaDe iceP op. Pa a
las a qui ec u as con capacidad de cómpu o 1.0 es e lími e es de 512 hilos po bloque.
O a al e na i a pa a lanza un ke nel de N hilos consis i ía en lanza N bloques pe o de
un hilo cada uno. En es e caso la sin axis pa a el lanzamien o se ía:
myKe nel<<<N,1>>>(a g_1,a g_2,...,a g_n);
Median e es a llamada end íamos una malla o mada po N bloques epa idos a lo la go
del eje x y con un hilo cada uno (Figu a 4.2).
Figu a 4.2. Malla o mada po N bloques pa alelos de 1 hilo epa idos a lo la go
del eje x. Cada bloque se puede iden i ica median e la a iable blockIdx.x.

PRÁCTICA nº 4: Hilos y Bloques
INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA 33
Si aho a que emos iden i ica los hilos pa a pode epa i a eas, no podemos u iliza la
a iable an e io ( h eadIdx.x), ya que es a a iable sólo pe mi e iden i ica a los hilos que
pe enecen al mismo bloque. Es o quie e deci que den o de cada bloque odos los hilos
comienzan a nume a se a pa i del núme o 0, y po lo an o, en nues o ejemplo, al habe un
único hilo en cada bloque, odos los hilos end ían el mismo alo de h eadIdx.x y po
an o en odos ellos end íamos myID = 0. Aho a la iden i icación de los hilos pasa po
iden i ica el bloque en el que se encuen a cada hilo. Es o lo podemos hace median e o a
a iable simila a la an e io denominada blockIdx:
in myID = blockIdx.x;
Al igual que an es, cada uno de los hilos end á un alo dis in o de blockIdx.x que i á
desde 0 has a N1. También exis e un lími e impues o po el ha dwa e en el núme o máximo
de bloques que podemos lanza , que pa a la mayo ía de las a qui ec u as es de 65.535
bloques. El alo pa icula pa a nues a GPU lo podemos a e igua a pa i del campo
maxG idSize[i] de nues a es uc u a de p opiedades, y que nos da el máximo núme o de
bloques pe mi idos en cada una de las di ecciones del espacio (i = 0, 1, 2  ejes x, y, y z
espec i amen e).
Po úl imo, el caso más gene al se ía el lanzamien o de un ke nel con M bloques y N
hilos po bloque (Figu a 4.3), en cuyo caso end íamos un o al de MN hilos . La sin axis en
es e caso se ía:
myKe nel<<<M,N>>>(a g_1,a g_2,...,a g_n);
Figu a 4.3. Malla o mada po M bloques pa alelos (g idDim.x) de N hilos cada
uno (blockDim.x) epa idos a lo la go del eje x.
Aho a hab á un o al de MN hilos ejecu ándose en pa alelo que pa a iden i ica los se á
necesa io hace uso conjun o de las dos a iables an e io es y de dos nue as cons an es que
pe mi an a la GPU conoce en iempo de ejecución las dimensiones del ke nel que hemos
lanzado. Es as dos cons an es son g idDim.x y blockDim.x. La p ime a nos da el núme o de
bloques (en nues o caso se ía M) y la segunda el núme o de hilos que iene cada bloque (N en
PRÁCTICA nº 4: Hilos y Bloques
34 INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA
nues o ejemplo). De es e modo, den o del ke nel cada hilo se puede iden i ica de o ma
uní oca median e la siguien e exp esión ( igu a 4.4):
in myID = h eadIdx.x + blockDim.x * blockIdx.x;
Figu a 4.4. Iden i icación de los hilos den o de una malla o mada po cua o bloques de cua o hilos cada uno
y epa idos a lo la go del eje x.
2. Ejemplo
En el siguien e ejemplo se mues a o ma de ap o echa el pa alelismo de da os
p og amando un ke nel que ealice la suma de dos ec o es de longi ud N inicializados con
alo es alea o ios comp endidos en e 0 y 1:
// includes
#include <s dio.h>
#include <s dlib.h>
#include <cuda_ un ime.h>
// de ines
#de ine N 16 // amano de los ec o es
#de ine BLOCK 5 // amano del bloque
// decla acion de unciones
// GLOBAL: uncion llamada desde el hos y ejecu ada en el de ice (ke nel)
__global__ oid suma( loa *a, loa *b, loa *c )
{
in myID = h eadIdx.x + blockDim.x * blockIdx.x;
// Solo abajan N hilos
i (myID < N)
{
c[myID] = a[myID] + b[myID];
}
}
// MAIN: u ina p incipal ejecu ada en el hos
in main(in a gc, cha ** a g )
{
// decla aciones
loa * ec o 1, * ec o 2, * esul ado;
loa *de _ ec o 1, *de _ ec o 2, *de _ esul ado;
// ese a en el hos
ec o 1 = ( loa *)malloc(N*sizeo ( loa ));
ec o 2 = ( loa *)malloc(N*sizeo ( loa ));
esul ado = ( loa *)malloc(N*sizeo ( loa ));
// ese a en el de ice
cudaMalloc( ( oid**)&de _ ec o 1, N*sizeo ( loa ));
cudaMalloc( ( oid**)&de _ ec o 2, N*sizeo ( loa ));
cudaMalloc( ( oid**)&de _ esul ado, N*sizeo ( loa ));
// inicializacion de ec o es
o (in i = 0; i < N; i++)
{
ec o 1[i] = ( loa ) and() / RAND_MAX;
ec o 2[i] = ( loa ) and() / RAND_MAX;
}
PRÁCTICA nº 4: Hilos y Bloques
INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA 35
// copia de da os hacia el de ice
cudaMemcpy(de _ ec o 1, ec o 1, N*sizeo ( loa ), cudaMemcpyHos ToDe ice);
cudaMemcpy(de _ ec o 2, ec o 2, N*sizeo ( loa ), cudaMemcpyHos ToDe ice);
// lanzamien o del ke nel
// calculamos el nume o de bloques necesa io pa a un amaño de bloque ijo
in nBloques = N/BLOCK;
i (N%BLOCK != 0)
{
nBloques = nBloques + 1;
}
in hilosB = BLOCK;
p in ("Vec o de %d elemen os n", N);
p in ("Lanzamien o con %d bloques (%d hilos) n", nBloques, nBloques*hilosB);
suma<<< nBloques, hilosB >>>( de _ ec o 1, de _ ec o 2, de _ esul ado );
// ecogida de da os desde el de ice
cudaMemcpy( esul ado, de _ esul ado, N*sizeo ( loa ), cudaMemcpyDe iceToHos );
// imp esion de esul ados
p in ( "> ec o 1: n");
o (in i = 0; i < N; i++)
{
p in ("%.2 ", ec o 1[i]);
}
p in (" n");
p in ( "> ec o 2: n");
o (in i = 0; i < N; i++)
{
p in ("%.2 ", ec o 2[i]);
}
p in (" n");
p in ( "> SUMA: n");
o (in i = 0; i < N; i++)
{
p in ("%.2 ", esul ado[i]);
}
p in (" n");
// libe amos memo ia en el de ice
cudaF ee( de _ ec o 1 );
cudaF ee( de _ ec o 2 );
cudaF ee( de _ esul ado );
// salida
p in (" npulsa INTRO pa a inaliza ...");
lush(s din);
cha ecla = ge cha ();
e u n 0;
}
PRÁCTICA nº 4: Hilos y Bloques
36 INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA
REALIZACIÓN PRÁCTICA
1. Ejecu a un ke nel compues o po 24 hilos que se enca gue de ellena es a ays con los
índices de iden i icación de cada hilo (un a ay po cada ipo de índice).
2. Cada hilo esc ibi á en el co espondien e a ay su índice de hilo ( h eadIdx.x), su índice
de bloque (blockIdx.x) y su índice global ( h eadIdx.x + blockDim.x * blockIdx.x).
3. El ke nel se debe ejecu a es eces, cada una de ellas con di e en es opciones de
o ganización:
» 1 bloque de 24 hilos.
» 24 bloques de 1 hilo.
» 4 bloques de 6 hilos.
4. En cada ejecución se debe mos a po pan alla el con enido de los es a ays:
>> Opcion 1: 1 bloque 24 hilos
indice de hilo:
0 1 2 3 4 5 6 7 ...
indice de bloque:
0 0 0 0 0 0 0 0 ...
indice global:
0 1 2 3 4 5 6 7 ...
>> Opcion 2: 24 bloques 1 hilo
...
...
...
>> Opcion 3: 4 bloques 6 hilos
...
...
...
PRÁCTICA nº 5: Tempo ización y E o es
INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA 37
PRÁCTICA nº 5:
Tempo ización y E o es
1. Co ien es (S eams)
Has a aho a hemos is o cómo el p ocesamien o de da os en pa alelo en una GPU puede
p opo ciona g andes mejo as en el endimien o en compa ación con el mismo código
ejecu ado en una CPU. Sin emba go, hay o a clase de pa alelismo pa a se explo ado en los
p ocesado es g á icos NVIDIA. Es e pa alelismo es simila al pa alelismo de a eas ( ask
pa allelism) que se encuen a en las aplicaciones de una CPU mul ip oceso. En luga de
calcula simul áneamen e la misma unción sob e una g an can idad de da os como se hace en
el pa alelismo de da os, el pa alelismo de a eas implica hace dos o más a eas
comple amen e di e en es en pa alelo.
En es e con ex o de pa alelismo, una a ea puede se cualquie ac i idad. Po ejemplo,
una aplicación pod ía es a ealizando dos a eas: ol e a dibuja la pan alla de un moni o
mien as se desca ga una ac ualización en la ed. Es as a eas se ejecu an en pa alelo, a pesa
Figu a 5.1. C onog ama de ejecución de una aplicación con dos co ien es independien es.

PRÁCTICA nº 5: Tempo ización y E o es
38 INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA
de no ene nada en común. Aunque el pa alelismo de a eas en una GPU no es ac ualmen e
an lexible como en un p ocesado de p opósi o gene al, puede se una opo unidad pa a
ex ae mayo endimien o a cie as aplicaciones. La idea que subyace es ene siemp e
ocupada la GPU ealizando a eas.
En CUDA, una co ien e o s eam ep esen a una lis a de ope aciones que se ejecu an en
la GPU en un o den especí ico. Es as ope aciones pueden se lanzamien os de un ke nel o
copias de memo ia desde la CPU a la GPU o ice e sa (Figu a 5.1). El o den en que las
ope aciones se añaden a la co ien e especi ica el o den en el que se án ejecu adas. Podemos
pensa en una co ien e o s eam como una lis a de a eas en la GPU, y hay que ap o echa
cualquie opo unidad pa a que es as a eas se puedan ejecu a en pa alelo. La capacidad de
una GPU pa a pode ejecu a es as a eas de o ma concu en e depende de su capacidad de
cómpu o, y po an o es algo que debemos a e igua a pa i de la es uc u a de p opiedades
cudaDe iceP op de nues o disposi i o. En pa icula , hay que consul a el alo del campo
de iceO e lap, que es una a iable bina ia que nos in o ma si nues o disposi i o sopo a la
ejecución de un ke nel simul áneamen e con la ans e encia de da os ( ue) o no lo sopo a
( alse). Pa a nues os p opósi os de momen o nos bas a á con u iliza la GPU empleando una
única co ien e de a eas.
2. E en os (E en s)
Un e en o o e en en CUDA es esencialmen e una ma ca de iempo en la GPU que
noso os podemos g aba en un de e minado ins an e de iempo. El hecho de que es a ma ca
de iempo se g abe en la p opia GPU elimina una g an can idad de p oblemas que pod íamos
encon a cuando in en á amos medi el iempo de ejecución de un ke nel u ilizando los
empo izado es de la CPU. La o ma de medi un iempo de ejecución es ela i amen e ácil,
ya que básicamen e consis e en c ea e en os pa a después g aba en ellos ma cas empo ales
con el in de calcula la di e encia de iempos en e las ma cas. Los e en os se almacenan en
un nue o ipo de da o denominado cudaE en _ , y la unción pa a c ea un e en o es
cudaE en C ea e():
cudaE en C ea e (cudaE en _ *ma ca);
Debemos c ea an os e en os como ma cas empo ales ayamos a necesi a , que en
gene al se án dos, una de inicio y o a de in. Una ez c eados los e en os donde amos a
PRÁCTICA nº 5: Tempo ización y E o es
INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA 39
gua da las ma cas inicial y inal, pod emos g aba dichas ma cas empo ales en el momen o
que deseemos. Es a g abación se ealiza u ilizando la unción cudaE en Reco d():
cudaE en Reco d(cudaE en _ ma ca, cudaS eam_ s eam);
donde ma ca es alguna de las a iables decla adas an e io men e pa a almacena las ma cas
empo ales de inicio o in y s eam es la co ien e en la que g aba el e en o. Habi ualmen e
se pone un 0 pa a ene en cuen a odas las co ien es (si es que hay más de una). A
con inuación lanzamos el abajo que que emos empo iza sob e la GPU y después g abamos
sob e la ma ca inal el iempo anscu ido.
Lo siguien e se ía calcula el iempo anscu ido en e dos e en os. Sin emba go, es
necesa io añadi llamada adicional después del úl imo cudaE en Reco d() y que iene como
obje i o sinc oniza el hos con el de ice. La unción eque ida pa a ello es:
cudaE en Synch onize(cudaE en _ ma ca);
Es a unción de iene la CPU has a la inalización de odo el abajo pendien e en la GPU
y que p ecede a la llamada más ecien e a cudaE en Reco d(). El po qué de la necesidad de
inclui es a unción iene que e con el hecho de que algunas de las llamadas que hacemos en
CUDA ealmen e son asínc onas. Po ejemplo, cuando lanzamos un ke nel, la GPU comienza
a ejecu a su co espondien e código, pe o la CPU con inua ejecu ando la siguien e línea de
nues o p og ama p incipal (main) an es de que la GPU haya e minado. Es o es muy
in e esan e desde el pun o de is a del endimien o po que signi ica que podemos ene
abajando simul áneamen e a la CPU y a la GPU, pe o po o o lado complica la medida del
iempo ya que una llamada a la unción cudaE en Reco d() no implica que se g abe el
iempo sino que queda pendien e en las a eas de la GPU g aba ese iempo. Po eso, has a
que no u ilizamos la unción cudaE en Synch onize() no podemos es a segu os de ob ene
la ma ca de iempo co ec a.
T as la llamada a es a unción sabemos que odo el abajo an e io se hab á comple ado y
po lo an o ambién se á segu o lee su co espondien e ma ca de iempo. Po an o, una ez
que enemos las dos ma cas empo ales de inicio y in, podemos calcula el iempo
anscu ido en e ambos e en os u ilizando la unción cudaE en ElapsedTime():
cudaE en ElapsedTime( loa * iempo, cudaE en _ ma ca1, cudaE en _ ma ca2);
PRÁCTICA nº 5: Tempo ización y E o es
40 INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA
que nos de uel e en la a iable iempo, de ipo loa , el iempo en milisegundos
anscu ido en e los e en os ma ca1 y ma ca2. Cabe menciona que dado que los e en os se
implemen an di ec amen e sob e la GPU, sólo si en pa a empo iza código pa a la GPU y
nunca pa a la CPU.
Pa a e mina podemos deci que los e en os c eados ambién pueden des ui se
u ilizando una unción denominada cudaE en Des oy():
cudaE en Des oy(cudaE en _ ma ca);
y de es e modo libe amos los ecu sos asociados con el e en o ma ca. En esumen, si
que emos medi el iempo de ejecución en la GPU es necesa io inclui en nues o código las
unciones elacionadas con la ges ión de e en os de mane a simila a como se mues a en el
siguien e ejemplo:
// MAIN: u ina p incipal ejecu ada en el hos
in main(in a gc, cha ** a g )
{
// ...
// decla acion de e en os
cudaE en _ s a ;
cudaE en _ s op;
// c eacion de e en os
cudaE en C ea e(&s a );
cudaE en C ea e(&s op);
// ma ca de inicio
cudaE en Reco d(s a ,0);
// codigo a empo iza en el de ice
//...<<< , >> ...
// ma ca de inal
cudaE en Reco d(s op,0);
// sinc onizacion GPU-CPU
cudaE en Synch onize(s op);
// calculo del iempo en milisegundos
loa elapsedTime;
cudaE en ElapsedTime(&elapsedTime,s a ,s op);
// imp esion de esul ados
p in ("> Tiempo de ejecucion: % ms n",elapsedTime);
// libe acion de ecu sos
cudaE en Des oy(s a );
cudaE en Des oy(s op);
// salida
p in (" npulsa INTRO pa a inaliza ...");
lush(s din);
cha ecla = ge cha ();
e u n 0;
}
3. Manejo de e o es
A la ho a de depu a un p og ama esul a muy ú il dispone de algún mecanismo pa a
localiza los di e en es e o es que se hayan ido p oduciendo du an e su ejecución. Con el in
de acili a dicho manejo de e o es, en CUDA odas las unciones (excep o los lanzamien os
de un ke nel) de uel en un código de e o del ipo cudaE o _ . Es e código es
PRÁCTICA nº 5: Tempo ización y E o es
INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA 41
simplemen e un alo en e o y oma dis in os alo es dependiendo del ipo de e o
encon ado. Pa a una llamada inalizada con éxi o, el código de e o oma un alo de inido
como cudaSuccess. De es a o ma, comp obando el código de e o de uel o as la llamada
a una unción podemos sabe si és a ha e minado con éxi o o no. El p incipal incon enien e
es que es e p ocedimien o puede esul a edioso al ene que comp oba con inuamen e el
código de e o de uel o po cada unción e imp imi el mensaje de e o co espondien e,
como po ejemplo:
// ...
cudaE o _ e o ;
// ...
e o = cudaMalloc( . . . );
i (e o != cudaSuccess)
{
p in (" nERROR en cudaMalloc: %s n", cudaGe E o S ing(e o ) );
p in (" npulsa INTRO pa a inaliza ...");
lush(s din);
cha ecla = ge cha ();
exi (-1);
}
// ...
donde la unción cudaGe E o S ing() se enca ga de de ol e un mensaje de ex o
explicando el ipo de e o que ha sucedido:
cha * cudaGe E o S ing(cudeE o _ codigo);
Sin emba go, esul a mucho más ú il de ini una unción dedicada al chequeo de e o es,
con un único pa áme o de en ada que sea el mensaje de e o que deseamos que apa ezca
po pan alla, y coloca la jus o después de la llamada a la unción cuyo código de e o
que amos sabe . Un ejemplo de es e ipo de unción puede se el siguien e:
__hos __ oid check_CUDA_E o (cons cha *mensaje)
{
cudaE o _ e o ;
cudaDe iceSynch onize();
e o = cudaGe Las E o ();
i (e o != cudaSuccess)
{
p in ("ERROR %d: %s (%s) n", e o , cudaGe E o S ing(e o ), mensaje);
p in (" npulsa INTRO pa a inaliza ...");
lush(s din);
cha ecla = ge cha ();
exi (-1);
}
}
En la de inición de es a unción se han u ilizado dos nue as unciones de CUDA que son
cudaDe iceSynch onize() y cudaGe Las E o (). Como ya hemos dicho an e io men e,
algunas de las llamadas que hacemos en CUDA ealmen e son asínc onas, po eso la p ime a
unción es necesa ia pa a que la CPU no con inúe con la ejecución del p og ama has a que la
GPU no haya e minado (en ealidad sólo se ía necesa ia pa a comp oba los e o es después
del lanzamien o de un ke nel). En cuan o a la segunda unción, és a se enca ga de cap u a el
PRÁCTICA nº 6: A ays Mul idimensionales
48 INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA
2. Ejemplo
El siguien e ejemplo demues a el g an po encial de una GPU pa a ap o echa el
pa alelismo de da os. El ejemplo consis e en un ke nel bidimensional que ealiza la suma de
dos ma ices cuad adas de amaño NN. El lanzamien o se ha ealizado u ilizando un único
bloque con NN hilos epa idos a lo la go de los ejes x e y de mane a análoga al p oblema que
se p e ende esol e , es deci , cada hilo se a a enca ga de calcula un elemen o de la ma iz
inal. Hay que ene p esen e que el espacio de memo ia ese ada en el de ice con
cudaMalloc() es un espacio lineal, po lo que el acceso a los da os hay que ealiza lo
median e un índice lineal ob enido a pa i de los índices co espondien es a los ejes x e y:
// includes
#include <s dio.h>
#include <s dlib.h>
#include <cuda_ un ime.h>
// de ines
#de ine N 20
// decla acion de unciones
// GLOBAL: uncion llamada desde el hos y ejecu ada en el de ice (ke nel)
__global__ oid suma_gpu( loa *A, loa *B, loa *C )
{
// indice de columna
in columna = h eadIdx.x;
// indice de ila
in ila = h eadIdx.y;
// indice lineal
in myID = columna + ila * blockDim.x;
// sumamos cada elemen o
C[myID] = A[myID] + B[myID];
}
// MAIN: u ina p incipal ejecu ada en el hos
in main(in a gc, cha ** a g )
{
// decla aciones
loa *hs _A, *hs _B, *hs _C;
loa *de _A, *de _B, *de _C;
// ese a en el hos
hs _A = ( loa *)malloc(N*N*sizeo ( loa ));
. . .
// ese a en el de ice
cudaMalloc( ( oid**)&de _A, N*N*sizeo ( loa ));
. . .
// incializacion
o (in i=0;i<N*N;i++)
{
hs _A[i] = ( loa )( and() % 10 );
. . .
}
// copia de da os
cudaMemcpy( de _A, hs _A, N*N*sizeo ( loa ), cudaMemcpyHos ToDe ice );
. . .
// dimensiones del ke nel
dim3 Nbloques(1);
dim3 hilosB(N,N);
// llamada al ke nel bidimensional de NxN hilos
suma_gpu<<<Nbloques,hilosB>>>(de _A, de _B, de _C);

PRÁCTICA nº 6: A ays Mul idimensionales
INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA 49
// ecogida de da os
cudaMemcpy( hs _C, de _C, N*N*sizeo ( loa ), cudaMemcpyDe iceToHos );
// imp esion de esul ados
p in ("A: n");
. . .
p in ("B: n");
. . .
p in ("C: n");
o (in i=0;i<N;i++)
{
o (in j=0;j<N;j++)
{
p in ("%2.0 ",hs _C[j+i*N]);
}
p in (" n");
}
// salida
p in (" npulsa INTRO pa a inaliza ...");
lush(s din);
cha ecla = ge cha ();
e u n 0;
}
PRÁCTICA nº 6: A ays Mul idimensionales
50 INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA
REALIZACIÓN PRÁCTICA
1. Ejecu a un ke nel que ope e sob e una ma iz de 1616 elemen os.
2. La ope ación consis e en sus i ui cada elemen o de la ma iz o iginal po la suma de los
elemen os adyacen es ( igu a 6.2).
3. Los elemen os si uados en los bo des no deben modi ica se.
4. La ma iz o iginal se debe inicializa en el hos con alo es alea o ios que sean 0 ó 1.
5. U iliza una unción auxilia del ipo __de ice__ pa a calcula el índice lineal a pa i de
los índices de ila y columna:
__de ice__ in index (in ila, in columna);
6. El p og ama debe mos a po pan alla:
» Ma iz inicial y ma iz inal.
» Tiempo de ejecución.
7. Inclui en el código del p og ama llamadas a la unción check_CUDA_E o () pa a
de ec a posibles e o es.
Figu a 6.2. Cada elemen o de la ma iz B se ob iene como la suma de los
elemen os de la ma iz A ubicados en posiciones adyacen es.
PRÁCTICA nº 7: Memo ia Compa ida
INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA 51
PRÁCTICA nº 7:
Memo ia Compa ida
1. Je a quía de memo ia
Como ya hemos is o en p ác icas an e io es, los hilos de CUDA pueden accede a los
da os desde di e en es espacios de memo ia du an e su ejecución. (Figu a 7.1). Sabemos que
cada hilo iene sus p opios egis os y su p opia memo ia local. También que cada hilo iene
una memo ia compa ida que es isible pa a odos los hilos de su mismo bloque y que odos
los hilos ienen acceso a la misma memo ia global.
Figu a 7.1. Di e en es espacios de memo ia en CUDA.
La ca ac e ís ica más impo an e de la memo ia compa ida (sha ed memo y) es que
és a se encuen a ce ca de cada uno de los núcleos de p ocesamien o (semejan e a una
memo ia cache de ni el 1), lo que hace que sea una memo ia con muy baja la encia. Debido a
que es á en el mismo chip que el p ocesado , la memo ia compa ida es mucho más ápida
que los espacios de memo ia local y global. Po ello, cualquie opo unidad de eemplaza los
accesos a memo ia global po accesos a memo ia compa ida debe se ap o echada al
máximo. En es e sen ido, lo más habi ual se á u iliza es a zona de memo ia pa a ealiza
PRÁCTICA nº 7: Memo ia Compa ida
52 INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA
cálculos in e medios, pa a coloca da os que son ecuen emen e accedidos, o bien pa a lee o
esc ibi esul ados que o os hilos del mismo bloque necesi an pa a su ejecución.
Desde el pun o de is a del ha dwa e, pa a alcanza es e ele ado ancho de banda, la
memo ia compa ida se di ide en módulos de memo ia de igual amaño denominados
“bancos”, los cuales pueden se accedidos de o ma simul ánea. Es o quie e deci que
cualquie acceso a n di ecciones de memo ia que co espondan a n bancos di e en es pueden
se a endidos de mane a simul ánea, esul ando en un ancho de banda que es n eces el
co espondien e a un solo módulo. El p oblema es á cuando dos accesos coinciden con el
mismo banco, en cuyo caso los accesos de ealizan de o ma secuencial y el ancho de banda
se e se iamen e educido.
El hecho de que sólo los hilos pe enecien es al mismo bloque puedan compa i es e
espacio de memo ia es o o pa áme o de decisión a la ho a de lanza un ke nel y o ganiza lo
en di e en es bloques e hilos, además de las conocidas limi aciones del ha dwa e.
La sin axis pa a ese a memo ia en es e espacio se ealiza a a és del iden i icado
__sha ed__ den o del código de la unción __global__ (el ke nel). Po ejemplo, podemos
ese a espacio pa a un a ay de 32 elemen os de ipo loa de la siguien e o ma:
__sha ed__ loa ec o [32];
Es o c ea un espacio de almacenamien o pa a cada uno de los bloques que hemos lanzado
sob e la GPU, pe o con la ca ac e ís ica de que los hilos de un bloque no pueden e ni
modi ica los da os de o os bloques, lo que cons i uye un excelen e medio po el cual los
hilos de un mismo bloque pueden comunica se y colabo a en la ealización de un cálculo.
2. Sinc onización
El g an po encial de cálculo que nos p opo ciona CUDA al di idi una a ea en cien os o
miles de hilos se basa en la posibilidad de que odos los hilos se puedan ejecu a de mane a
independien e y simul ánea (o casi) ac uando sob e sus p opios da os. Sin emba go, dado que
los hilos son independien es, si espe amos que los hilos de un mismo bloque puedan coope a
compa iendo da os a a és de alguna zona de memo ia, ambién necesi amos algún
mecanismo de sinc onización en e ellos, es deci , necesi amos sinc oniza su ejecución pa a
coo dina los accesos a memo ia. Po ejemplo, si un hilo A esc ibe un alo en una zona de
memo ia compa ida y que emos que o o hilo B haga algo con ese alo , no podemos hace
que el hilo B comience su abajo has a que no sepamos que la esc i u a del hilo A ha
PRÁCTICA nº 7: Memo ia Compa ida
INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA 53
e minado. Sin una sinc onización, enemos un iesgo de da os del ipo RAW (Read A e
W i e) donde la co ección del esul ado depende de aspec os no de e minis as del ha dwa e.
Pa a e i a es e p oblema, en CUDA podemos especi ica pun os de sinc onización
den o del código del ke nel median e llamadas a la unción:
__sync h eads();
Es a unción ac úa como ba e a en la que odos los hilos de un mismo bloque deben
espe a an es de pode con inua con su ejecución. Con es a llamada ga an izamos que odos
los hilos del bloque han comple ado las ins ucciones p eceden es a __sync h eads() an es
de que el ha dwa e lance la siguien e ins ucción pa a cualquie o o hilo. De es e modo
sabemos que cuando un hilo ejecu a la p ime a ins ucción pos e io a __sync h eads(),
odos los demás hilos ambién han e minado de ejecu a sus ins ucciones has a ese pun o.
3. Ejemplo: educción pa alela
Una aplicación donde se pone de mani ies o la necesidad de sinc oniza la ejecución de
los hilos es la conocida como “ educción pa alela”. En gene al, un algo i mo de educción es
aquel donde el amaño del ec o de salida es más pequeño que el ec o de en ada (se ha
educido). En nues o caso, la aplicación que amos a es udia consis e en la suma de los
componen es de un ec o donde odos los hilos colabo an pa a ealiza la suma. El pun o
cla e es a á en la sinc onización de los hilos pa a que los esul ados in e medios sean
co ec os.
La o ma secuencial de ealiza es a suma se ía i e a a lo la go del ec o e i sumando
de o ma acumula i a cada elemen o. Sin emba go, en nues o ejemplo amos a ap o echa la
p esencia de a ios hilos de ejecución pa a epa i el abajo. Pa a implemen a nues o
algo i mo de educción pa alela necesi amos que la longi ud n del ec o sea una po encia de
2. Si no ue a así, bas a ía con comple a el ec o con ce os has a alcanza la longi ud
deseada. El núme o de pasos que necesi amos pa a implemen a es e algo i mo es de log2 (n),
y en cada uno de ellos enemos que hace que cada hilo sume los alo es co espondien es a
su posición y su posición más la mi ad (de ahí la necesidad de que n sea una po encia de 2),
gua dando los esul ados pa ciales en sus espec i as posiciones de memo ia al como se
mues a en la Figu a 7.2. El ec o de da os es a á almacenado en una zona de memo ia
compa ida po odos los hilos y en cada paso los hilos que abajan son la mi ad que en el
paso an e io . Vemos que en cada paso es muy impo an e sinc oniza la ejecución de odos

PRÁCTICA nº 7: Memo ia Compa ida
54 INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA
los hilos ya que es necesa io que cada hilo haya ealizado su suma an es de con inua con la
suma siguien e. Al inal de es e p oceso, sólo uno de los hilos ealiza á la úl ima suma, el
p ime o, dejando el esul ado inal en la p ime a posición de memo ia del ec o o iginal.
Figu a 7.2. P ime paso en el algo i mo de educción pa alela.
El código del ke nel que implemen a es a educción pa alela es el siguien e:
// . . .
// de ines
#de ine N 16
// ke nel
__global__ oid educcion( loa * ec o , loa *suma )
{
// Rese a de espacio en la zona de memo ia compa ida
__sha ed__ loa empo al[N];
// Indice local de cada hilo -> ke nel con un solo bloque
in myID = h eadIdx.x;
// Copiamos en ' empo al' el ec o y sinc onizamos
empo al[myID] = ec o [myID];
__sync h eads();
//Reduccion pa alela
in sal o = N/2;
// Realizamos log2(N) i e aciones
while(sal o)
{
// Solo abajan la mi ad de los hilos
i (myID < sal o)
{
empo al[myID] = empo al[myID] + empo al[myID+sal o];
}
__sync h eads();
sal o = sal o/2;
}
// El hilo no.'0' esc ibe el esul ado inal en la memo ia global
i (myID==0)
{
*suma = empo al[myID];
}
}
// . . .
PRÁCTICA nº 7: Memo ia Compa ida
INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA 55
REALIZACIÓN PRÁCTICA
1. Calcula de o ma ap oximada el alo del núme o π u ilizando la siguien e elación en e
el núme o π y la se ie o mada po los in e sos de los cuad ados de los núme os na u ales:
  22222
21
4
1
3
1
2
1
1
1
6n

2. Cada hilo de ejecución del ke nel debe gene a un é mino de la exp esión an e io .
3. El p og ama debe mos a po pan alla:
» Can idad de memo ia compa ida disponible en KiB.
» Máximo núme o de hilos disponibles.
» Núme o de hilos lanzados.
» Valo ap oximado del núme o π.
» Tiempo de ejecución.
4. Inclui en el código del p og ama llamadas a la unción check_CUDA_E o () pa a
de ec a posibles e o es.
PRÁCTICA nº 8: Memo ia Cons an e
INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA 57
PRÁCTICA nº 8:
Memo ia Cons an e
1. Je a quía de memo ia
La ca ac e ís ica p imo dial de una GPU ac ual es su eno me po encia de cálculo
(de e minada en úl ima ins ancia po su capacidad de cómpu o). De hecho, es a supe io idad
desde el pun o de is a compu acional de las GPUs sob e las CPUs ha ayudado a p ecipi a el
in e és en usa los p ocesado es g á icos pa a la compu ación de p opósi o gene al. Sin
emba go, la p esencia de cien os de núcleos de p ocesamien o en una GPU hace que el cuello
de bo ella no sea el cálculo en sí, sino el ancho de banda de la memo ia. Es deci , hay al
can idad de núcleos de p ocesamien o que es di ícil consegui que la en ada de da os sea lo
su icien emen e ápida como pa a man ene el ele ado i mo de cálculo. Po es e mo i o, se
hace necesa io dispone de mecanismos que pe mi an educi el á ico de memo ia eque ido
po un de e minado p oblema.
Con el in de ali ia el p oblema del ancho de banda de la memo ia, ya hemos is o que
los hilos de ejecución de una aplicación CUDA pueden accede a los da os desde di e en es
espacios de memo ia. (Figu a 8.1).
Figu a 8.1. Je a quía de memo ia en CUDA.
PRÁCTICA nº 9: G á icos en CUDA
64 INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA
luminosa y depende del núme o de bi s empleados. Po ejemplo, si u ilizamos k bi s
po cada mues a pod emos ep esen a cada una de ellas con 2k ni eles de g is
dis in os. Po an o, u ilizando 8 bi s el núme o de ni eles disponibles es 256.
Vemos que as el p oceso de mues eo y cuan i icación de una imagen lo que ob enemos
es una ma iz de núme os, es deci , que una imagen digi al no es más que la ep esen ación de
una imagen eal a pa i de una ma iz numé ica (Figu a 9.2).
Figu a 9.2. Rep esen ación de una imagen digi al como una ma iz de núme os.
Es o nos hace e que la o ma más adecuada pa a abaja con imágenes digi ales es
u ilizando un ha dwa e que nos pe mi a ope a simul áneamen e con g andes ma ices de
núme os. En es e con ex o, emos que nues o en o no de p og amación CUDA encaja
pe ec amen e con dicho p opósi o ya que nos pe mi e lanza un ke nel bidimimensional con
hilos epa idos en ambas di ecciones del espacio y haciendo co esponde a cada hilo con un
píxel de nues a imagen digi al.
2. Imágenes en colo
Desde el pun o de is a de la in o mación almacenada en una imagen digi al, podemos
clasi ica las imágenes digi ales como:
 Imágenes de in ensidad (escala de g ises), donde cada elemen o de la ma iz
ep esen a un ni el de g is den o del ango de e minado po el núme o de bi s
empleado en la cuan i icación (Figu a 9.2).
 Imágenes en colo , donde pa a es e ipo de imágenes se necesi a ene in o mación de
es colo es p ima ios, de inidos po el espacio de colo que se es é u ilizando.

PRÁCTICA nº 9: G á icos en CUDA
INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA 65
El espacio de colo que se u iliza pa a aplicaciones ales como la adquisición o
gene ación de imágenes en colo es el RGB, donde la combinación adi i a de los colo es
p ima ios ojo (Red), e de (G een) y azul (Blue) p oduce odo el ango de colo es
ep esen ables en dicho espacio (Figu a 9.3).
Figu a 9.3. Rep esen ación del espacio de colo RGB u ilizando alo es de 8 bi s.
Una imagen en escala de g ises ambién puede conside a se como una imagen en colo en
la que sus es componen es son iguales. Así pues, pa a ep esen a una imagen en colo se
necesi an como mínimo es ma ices co espondien es a las componen es oja, e de y azul
de cada píxel indi idual (Figu a 9.4).
Figu a 9.4. De alle de una imagen en colo RGB.
3. Imágenes en mapa de bi s (bi map)
Un bi map (o mapa de bi s) es la o ma na u al de o ganiza en memo ia la in o mación
de una imagen digi al. Así, dado que una imagen digi al se ep esen a como una ma iz
ec angula de píxeles o pun os de colo , en un mapa de bi s la in o mación de cada píxel se
almacena en posiciones consecu i as de la memo ia o mando un a ay cuyo amaño depende
PRÁCTICA nº 9: G á icos en CUDA
66 INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA
de la al u a y la anchu a de la imagen (núme o de píxeles) y de la in o mación de colo
con enida en cada píxel (bi s po píxel).
La in o mación de colo de cada píxel se codi ica u ilizando canales sepa ados, cada uno
de los cuales co esponde a uno de los colo es p ima ios del espacio de colo u ilizado
(gene almen e RGB). A eces, se puede añadi o o canal que ep esen a la anspa encia del
colo espec o del ondo de la imagen y se denomina canal α (modelo RGBA). Po lo an o, el
amaño necesa io pa a almacena en memo ia un mapa de bi s de una imagen en colo de
MN píxeles, con 8 bi s po canal y cua o canales (R, G, B y α) es de 4MN by es. Y es a
in o mación se almacena en memo ia p incipal ocupando posiciones consecu i as de memo ia
(Figu a 9.5).
Figu a 9.5. Ubicación en memo ia p incipal de un mapa de bi s de 4 canales y 8 bi s po canal.
4. OpenGL
Con el in de pode isualiza una imagen digi al en nues o sis ema necesi amos
unciones que hagan la labo de “ab i una en ana” y de “dibuja ” en ella. La o ma más
sencilla es u iliza una API (Applica ion P og amming In e ace - In e az de p og amación
de aplicaciones) des inada a al e ec o como es OpenGL.
OpenGL (Open G aphics Lib a y) es una especi icación es ánda , es deci , un
documen o que desc ibe un conjun o de unciones y el compo amien o exac o que deben
ene (pa iendo de ella, los ab ican es de ha dwa e pueden c ea implemen aciones). Cons a
de más de 250 unciones di e en es que pueden usa se pa a dibuja escenas idimensionales
complejas a pa i de p imi i as geomé icas simples, ales como pun os, líneas y iángulos.
El uncionamien o básico de OpenGL consis e en acep a p imi i as ales como pun os,
líneas y polígonos, y con e i las en píxeles. Dado que es á basado en p ocedimien os de bajo
ni el, equie e que el p og amado dic e los pasos exac os necesa ios pa a “ ende iza ” una
PRÁCTICA nº 9: G á icos en CUDA
INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA 67
escena. Es o con as a con o as in e aces desc ip i as, donde un p og amado sólo debe
desc ibi la escena y puede deja que la biblio eca con ole los de alles pa a ep esen a la. El
diseño de bajo ni el de OpenGL equie e que los p og amado es conozcan en p o undidad la
pipeline g á ica, a cambio de da les libe ad pa a implemen a algo i mos g á icos no edosos.
5. Ejemplos
El siguien e ejemplo mues a una o ma muy sencilla de gene a un bi map y de
mos a lo po pan alla. Las llamadas a las unciones de OpenGL necesa ias pa a su
isualización es án encapsuladas den o del iche o de cabece a cpu_bi map.h:
// Includes
#include <s dio.h>
#include <s dlib.h>
#include <cuda_ un ime.h>
#include "cpu_bi map.h"
// De ines
#de ine DIM 1024 // Dimensiones del Bi map
// GLOBAL: uncion llamada desde el hos y ejecu ada en el de ice (ke nel)
__global__ oid ke nel( unsigned cha *imagen )
{
// coo denada ho izon al
in x = h eadIdx.x + blockIdx.x * blockDim.x;
// coo denada e ical
in y = h eadIdx.y + blockIdx.y * blockDim.y;
// coo denada global de cada pixel
in pixel = x + y * blockDim.x * g idDim.x;
// cada hilo pin a un pixel con un colo a bi a io
imagen[pixel *4 + 0] = 255*x/( blockDim.x * g idDim.x); // canal R
imagen[pixel *4 + 1] = 255*y/( blockDim.y * g idDim.y); // canal G
imagen[pixel *4 + 2] = 2*blockIdx.x + 2*blockIdx.y; // canal B
imagen[pixel *4 + 3] = 255; // canal al a
}
// MAIN: u ina p incipal ejecu ada en el hos
in main(in a gc, cha ** a g )
{
// decla acion del bi map
CPUBi map bi map( DIM, DIM );
// amaño en by es
size_ size = bi map.image_size();
// ese a en el hos
unsigned cha *hos _bi map = bi map.ge _p ();
// ese a en el de ice
unsigned cha *de _bi map;
cudaMalloc( ( oid**)&de _bi map, size );
// gene amos el bi map
dim3 Nbloques(DIM/16,DIM/16);
dim3 hilosB(16,16);
ke nel<<<Nbloques,hilosB>>>( de _bi map );
// ecogemos el bi map desde la GPU pa a isualiza lo
cudaMemcpy( hos _bi map, de _bi map, size, cudaMemcpyDe iceToHos );
// libe acion de ecu sos
cudaF ee( de _bi map );
// isualizacion y salida
p in (" n...pulsa ESC pa a inaliza ...");
bi map.display_and_exi ();
e u n 0;
}
PRÁCTICA nº 9: G á icos en CUDA
68 INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA
El código an e io dibuja un bi map donde cada hilo se enca ga de da colo a un píxel,
asignado a cada canal un alo elacionado con su posición espacial. El bi map se dibuja en la
pan alla de izquie da a de echa y de abajo a a iba, es deci , el píxel de coo denadas (0, 0) se
si úa en la esquina in e io izquie da.
PRÁCTICA nº 9: G á icos en CUDA
INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA 69
REALIZACIÓN PRÁCTICA
1. Lanza un ke nel que gene e el bi map co espondien e al able o de ajed ez mos ado en la
igu a 9.6.
2. U iliza bloques de 1616 hilos.
3. Inclui en el código del p og ama llamadas a la unción check_CUDA_E o () pa a
de ec a posibles e o es.
Figu a 9.6. Table o de ajed ez.

ANEXO: Solución a los eje cicios
INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA 71
PRÁCTICA nº 10:
Eje cicio de aplicación p ác ica
1. Memo ia Compa ida
Den o de la je a quía de memo ia de una GPU con a qui ec u a CUDA hemos is o que
la ca ac e ís ica más impo an e de la memo ia compa ida (sha ed memo y) es que se
encuen a den o de cada uno de los S eaming Mul ip ocesso s (SM) semejan e a una
memo ia cache de ni el 1, y al es a en el mismo chip que el p ocesado , es mucho más ápida
que los espacios de memo ia local y global ( igu a 10.1). Po ello, cualquie opo unidad de
eemplaza los accesos a memo ia global po accesos a memo ia compa ida debe se
ap o echada al máximo. En es e sen ido, lo más habi ual es u iliza es a zona de memo ia
pa a ealiza cálculos in e medios.
Figu a 10.1. Di e en es espacios de memo ia en CUDA.
La sin axis pa a ese a memo ia en es e espacio se ealiza den o del cue po de una
unción de ipo __global__ a a és del iden i icado __sha ed__. Po ejemplo, la siguien e
decla ación ese a espacio pa a un a ay de 32 elemen os de ipo loa :
__sha ed__ loa ec o [32];
ANEXO: Solución a los eje cicios
72 INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA
Es a decla ación c ea de o ma es á ica un espacio de almacenamien o en la memo ia
compa ida de cada uno de los bloques lanzados sob e la GPU, pe o obliga a conoce en
iempo de compilación el amaño de las a iables. Sin emba go, exis e la posibilidad de
ese a dinámicamen e espacio en la memo ia compa ida. Es o se consigue añadiendo el
iden i icado ex e n en la decla ación. En el ejemplo an e io , si hacemos lo siguien e:
ex e n __sha ed__ loa ec o [];
aho a la a iable ec o sigue esidiendo en la zona de memo ia compa ida pe o con el
amaño aún sin especi ica . El amaño de ini i o de la a iable se esuel e en iempo de
ejecución especi icándolo den o de la llamada al ke nel, añadiendo un e ce pa áme o jun o
a las dimensiones del ke nel del siguien e modo:
myKe nel<<<blocks, h eads,Sha edMem>>>(a g_1,a g_2,...,a g_n);
siendo el alo de Sha edMem el amaño en by es asignado a la a iable ubicada en memo ia
compa ida.
2. El núme o 
Una aplicación muy in e esan e pa a pone en p ác ica la po encialidad de la
p og amación pa alela median e una GPU es la ob ención del núme o π a pa i de la siguien e
exp esión:


1
0
2
1
4dx
x

La exp esión an e io ep esen a la in eg al de inida de una unción en el in e alo [0, 1],
y su alo es igual al á ea limi ada po la g á ica de la unción, el eje de abscisas, y las ec as
e icales x = 0 y x = 1 (Figu a 10.2).
Los mé odos pa a e alua de o ma ap oximada dicho á ea eciben el nomb e de mé odos
de in eg ación numé ica, y se basan en ap oxima la unción a in eg a po o a unción de la
cual se conoce la in eg al exac a (gene almen e unciones polinómicas). No obs an e, los
mé odos más sencillos di iden la egión cuyo á ea deseamos calcula en cuad ilá e os
( ec ángulos o apecios) y después suman el á ea de odos ellos pa a es ima el á ea
ence ada bajo la cu a (Figu a 10.3). En onces:
ANEXO: Solución a los eje cicios
INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA 73
n
AAA





21

donde Ai es el á ea del cuad ilá e o i. De es e modo, cuan o mayo sea el núme o de
subin e alos en los que di idimos el in e alo de in eg ación, mayo p ecisión end emos en
el esul ado inal.
Figu a 10.2. G á ica de la unción y = 4/(1 + x2). El á ea de la egión somb eada es igual a π.
Figu a 10.3. Ap oximación del á ea ence ada po la cu a median e ec ángulos.