Nuestro kernel básicamente consiste en realizar
la suma de una serie de elementos y guardarla en el primero de ellos. Cada
ejecución se encargará de una posición de memoria diferente, con lo que no es necesario
preocuparnos de implementar ningún mecanismo de sincronización. El número de
elementos a sumar viene dado por el factor de desenrollado (UF).
Los parámetros
de entrada del Algoritmo son: Datos: puntero a la estructura que
contiene los datos a sumar, Elementos: número de elementos que se suman,
Impar: indica si el número de elementos es impar (para evitar la
realización del cálculo del módulo, dentro del kernel, ya que es
computacionalmente costoso), y finalmente UF, que se utilizará para indicar el número de elementos a
sumar dentro de cada hilo.
Algoritmo 3: kernel
de reducción de vectores
|
||
1
|
tid ← threadIdx.x + blockIdx.x * blockDim.x
|
|
2
|
t ← (ene/UF)
|
|
3
|
si tid < t hacer
|
|
4
|
para i←1 hasta UF hacer
|
|
5
|
a[tid] ← a[tid] + a[(tid+(t*i))]
|
|
6
|
Si tid = 0 y impar > 0 hacer
|
|
7
|
para i←1 hasta impar+1
hacer
|
|
8
|
a[tid] ← a[tid]
+ a[(ene-i)]
|
|
Algoritmo 1 Kernel de reducción de vectores
1.-La primera línea del algoritmo,
“tid ← threadIdx.x + blockIdx.x *
blockDim.x”, es la típica asignación para el cálculo de la posición global
de ejecución del kernel. Así se identifica exactamente al hilo que se está
referenciando en cada uno de los kernels y es la forma que tenemos para
organizar el proceso en paralelo. Como en nuestro caso estamos trabajando únicamente
con una dimensión, sólo se referencia al valor del campo “x” de cada variable (threadIdx, blockIdx y blockDim), pero podrían ser hasta tridimensionales, para ello se utilizan los
campos “y” y
“z” para referirnos a la segunda y
tercera dimensión respectivamente.
2.- Dado
que el número real de cálculos que tenemos que ejecutar no siempre coincide con
el total de hilos ejecutados en los Kernels, debemos controlar esta situación para que el
dispositivo no haga trabajo inútil (o redundante). Por este motivo la segunda
línea del algoritmo calcula el número total de threads a emplear y en la
tercera sólo se permite la ejecución del proceso de reducción a los hilos que
están dentro del rango necesario.
4.-El siguiente paso define al bucle que empleamos
para la suma de elementos del vector, que viene definido por el factor de
desenrollado que empleamos en la llamada.
5.-La siguiente línea realiza la suma de elementos
y la acumula en el primer elemento.
6.-Finalmente, tenemos una condición especial: sólo
si se trata del primer hilo y el número de elementos a procesar es impar,
adicionamos esos elementos impares también al primero.
Las operaciones realizadas en el siguiente
ejemplo son: (3+6=9), (8+5=13), (4+2=6), y en la siguiente iteración
(9+13+6=28).
En la siguiente Figura se puede observar
gráficamente el proceso, para un factor UF de 2:
Las propiedades fundamentales de este método de reducción que acabamos de
proponer son:
·
No hay relación entre bloques ni hilos del mismo stream, con lo que no es
necesario establecer ningún tipo de mecanismo para sincronizar el algoritmo en
ningún punto. Es un algoritmo orientado al hilo y no al bloque.
·
Adecúa la petición de recursos del sistema en función al estado de su
ejecución.
·
Sólo se realiza un desplazamiento de datos de memoria central a memoria de
dispositivo por el total de elementos a sumar, y además solo se necesita desplazar
el primer elemento a memoria principal para obtener el resultado de la
reducción.
·
Con el control de UF se reduce el
efecto de la latencia de memoria global del dispositivo, no siendo por tanto
necesario realizar las lecturas alineadas con el orden de los hilos.
Ya hemos presentado un algoritmo para la reducción de vectores, que como sabréis, es una de las operaciones más típicas en multitud de escenarios, tanto en análisis de datos, como en optimización. El próximo post presentará los datos obtenidos en las pruebas.
