hoy quiero ilustraros un método de ordenación que implementé para la ordenación de vectores en GPU, la verdad es que primero lo ideé y como suele pasar en estas cosas luego busqué en internet y efectivamente el método ya existía y para mi sorpresa desde hace muchos años, en este caso aplicado a lás máquinas matriciales.
La solución de la que hablo está basada en el método
de ordenación paralelo Parallel Neighbor-Sort propuesto por Nico Habermann(1), tambien conocido como odd-even sort, se diseño inicialmente para computación en arrays de procesadores a principios de los 70. Se trata realizar
comparaciones de los elementos del vector dos a dos, siendo estos vecinos,
realizando paralelamente Elementos/2 comparaciones en cada ejecución del
kernel, las llamadas se intercalarán variando la paridad del elemento
comparado, repitiendo este proceso hasta que no se produzca ningún cambio, esto
determinará el final de la ordenación.
A continuación os ilustro el kernel para la ordenación en CUDA, está ordenando 3 vectoresunas reglas especificas demás complejo del que ya os hablaré, sirva ahora de ejemplo:
__global__ void ordenar_especies(int par,int
*cambios,int *contar, int M, int D,
float *Puntos, int *Nivel, int *Orden, float *Valores) {
__shared__ int cuenta;
int tid =
threadIdx.x + blockIdx.x * blockDim.x;
int elem =
((3*M*3)+M)/2;
float
puntoA, puntoB, aux = 0.0f;
int ordenA, ordenB = 0;
int
nivelA, nivelB = 0;
int
cambiar_puntos = 0;
if (tid
== 0){
cambios = 0;
if
(contar[0] == 0){
cuenta = 0;
}
}
tid = (tid*2)+par;
if (tid <
elem) {
puntoA
= tex1Dfetch(valoresTex,tid);
puntoB
= tex1Dfetch(valoresTex,tid+1);
ordenA
= tex1Dfetch(ordenTex,tid);
ordenB
= tex1Dfetch(ordenTex,tid+1);
nivelA
= tex1Dfetch(nivelTex,tid);
nivelB
= tex1Dfetch(nivelTex,tid+1);
if (ordenB > 0) {
if (ordenA == 0) {
cambiar_puntos = 1;
} else {
if (nivelB > nivelA){
cambiar_puntos = 1;
} else if ((nivelB == nivelA) && (puntoB > puntoA)){
cambiar_puntos = 1;
}
}
}
if
(cambiar_puntos == 1) {
for (int x=0;x<D;++x){
aux = Puntos[tid+(x*elem)];
Puntos[tid+(x*elem)] = Puntos[tid+1+(x*elem)];
Puntos[tid+1+(x*elem)] = aux;
}
Valores[tid] = puntoB;
Valores[tid+1] = puntoA;
Orden[tid] = ordenB;
Orden[tid+1] = ordenA;
Nivel[tid] = nivelB;
Nivel[tid+1] = nivelA;
cambios[0] = 1;
}
if
(contar[0] == 0){
if
((ordenA == 0) && (ordenB == 0)) {
__syncthreads();
cuenta = cuenta + 2;
}
else if ((ordenA
== 0) || (ordenB == 0)) {
__syncthreads();
cuenta = cuenta + 1;
}
}
}
if (tid
== (elem-1)){ //el ultimo que ponga la
cuenta de su bloque
if (contar[0] == 0){
__syncthreads();
contar[blockIdx.x] = elem-cuenta;
}
}
}
Hasta la próxima!!
____________________
(1)Parallel Neighbor-Sort (or the Glory of the Induction Principle). 5285 Port Royal Rd Sprigfield VA : National Technical Information Service, US Department of Commerce, 1972. AD0759248.
No hay comentarios:
Publicar un comentario