Saludos a todos los amantes de los algoritmos. Quiero contarles acerca de mi investigación sobre la ordenación en general y profundizar en la consideración de la raíz del género.
Como desarrollador con muchos años de experiencia, comenzó a encontrar cada vez más una extraña tendencia en el desarrollo de software: a
pesar del desarrollo del hardware de las computadoras modernas y las mejoras en los algoritmos, en general, el rendimiento del código no solo no aumentó, sino que en algunos lugares se degradó considerablemente.
Creo que esto se debe a la idea general de dar preferencia a la programación rápida utilizando frameworks cada vez más potentes y lenguajes de scripting de nivel ultra alto. Los lenguajes como Ruby o Python son increíblemente amigables para los desarrolladores. Mucho 'azúcar sintáctico', incluso diría 'Cariño', acelera el desarrollo a veces, si no órdenes de magnitud, pero a qué costo. Como usuario, me molesta la baja eficiencia térmica del código, simplemente guardaré silencio sobre la cantidad de memoria consumida, pero el principal recurso de la humanidad es el tiempo. Desaparece sin dejar rastro en interminables abstracciones, es robado por analizadores de código, limpiado por recolectores de basura inteligentes. No llamo a volver al pasado, abandonando los beneficios del desarrollo moderno, a escribir código "caro",Solo propongo pensar en la posible eliminación de los cuellos de botella de rendimiento cuando sea posible en las tareas típicas. Esto a menudo se puede lograr optimizando secciones de código de alta carga.
La clasificación se puede destacar como una de las tareas básicas de optimización. El tema está tan explorado, arriba y abajo, que parece que es difícil encontrar algo interesante en el camino. Sin embargo, lo intentaremos.
No clasificaremos matrices pequeñas (menos de un millón de elementos). Incluso si es extremadamente ineficiente hacer esto, es bastante difícil sentir las caídas, ya que están niveladas por el rendimiento de los equipos modernos. Grandes cantidades de datos (miles de millones de elementos) son otro asunto; la velocidad de ejecución varía mucho de una selección competente de un algoritmo.
Todos los algoritmos basados en comparaciones generalmente resuelven el problema de clasificación no mejor que O (n * Log n). En n grande, la eficiencia disminuye rápidamente y no es posible cambiar esta situación. Esta tendencia se puede corregir abandonando los métodos basados en la comparación. El más prometedor para mí es el algoritmo de ordenación Radix. Su complejidad computacional es O (k * n), donde k es el número de pasadas a través del arreglo. Si n es lo suficientemente grande, pero k, por el contrario, es muy pequeño, entonces este algoritmo gana sobre O (n * Log n).
En la arquitectura de procesador moderna, k casi siempre se reduce al número de bytes del número ordenado. Por ejemplo, para DWord (int) k = 4, que no es mucho. Esta es una especie de "pozo potencial" de cálculos, que se debe a varios factores:
- Los registros del procesador se agudizan para los operadores de 8 bits a nivel de hardware
- El búfer de conteo utilizado en el algoritmo encaja en una línea de la caché L1: el procesador. (256 * números de 4 bytes)
Puede intentar verificar la verdad de esta afirmación usted mismo. Sin embargo, en el momento actual, dividir por bits es la mejor opción. No excluyo que cuando el caché L1 de procesadores crezca a 256 KB, la opción de dividir a lo largo del límite de 2 bytes será más rentable.
Una implementación eficiente de la clasificación no es solo un algoritmo, sino también un delicado problema de ingeniería para optimizar el código.
En esta solución, el algoritmo consta de varias etapas:
- , ,
- ,
- (LSD),
Aplicamos el algoritmo LSD como más rápido (al menos en mi versión) debido a un procesamiento más suave con varias fluctuaciones de los datos de entrada.
La matriz original completamente ordenada es el peor de los casos para el algoritmo, ya que los datos aún estarán completamente ordenados. Por otro lado, el ordenamiento Radix es extremadamente efectivo en datos aleatorios o mixtos.
Ordenar una matriz simple de números es poco común, normalmente se necesita un diccionario de la forma: clave - valor, donde el valor puede ser un índice o un puntero.
Para la universalización, aplicaremos una estructura de la forma:
typedef struct TNode {
//unsigned long long key;
unsigned int key;
//unsigned short key;
//unsigned char key;
unsigned int value;
//unsigned int value1;
//unsigned int value2;
} TNode;
Naturalmente, cuanto menor sea el bitness de la clave, más rápido funcionará el algoritmo. Dado que el algoritmo no funciona con punteros a una estructura, en realidad la conduce a la memoria. Con un mínimo de campos, obtenemos alta velocidad. Con un aumento en el volumen de campos de datos de la estructura, la eficiencia cae drásticamente.
Anteriormente, ya escribí una nota con la implementación de Radix sort en Pascal, pero la "segregación" de lenguajes de programación está ganando un ritmo sin precedentes entre los habituales de este recurso. Por lo tanto, decidí reescribir parcialmente el código de este artículo en 'si' como más
En contraste con la implementación anterior, utilizando punteros en el direccionamiento de bucle en lugar de operaciones de desplazamiento bit a bit, fue posible obtener un aumento aún mayor del 1-2% en la velocidad del algoritmo.
C
#include <stdio.h>
#include <omp.h>
#include <time.h>
#include <windows.h>
#include <algorithm>
//=============================================================
typedef struct TNode {
//unsigned long long key;
unsigned int key;
//unsigned short key;
//unsigned char key;
unsigned int value;
//unsigned int value1;
//unsigned int value2;
} TNode;
//=============================================================
void RSort_step(TNode *source, TNode *dest, unsigned int n, unsigned int *offset, unsigned char sortable_bit)
{
unsigned char *b = (unsigned char*)&source[n].key + sortable_bit;
TNode *v = &source[n];
while (v >= source)
{
dest[--offset[*b]] = *v--;
b -= sizeof(TNode);
}
}
//=============================================================
void RSort_Node(TNode *m, unsigned int n)
{
//
TNode *m_temp = (TNode*)malloc(sizeof(TNode) * n);
//
unsigned int s[sizeof(m->key) * 256] = {0};
//
unsigned char *b = (unsigned char*)&m[n-1].key;
while (b >= (unsigned char*)&m[0].key)
{
for (unsigned int digit=0; digit< sizeof(m->key); digit++)
{
s[*(b+digit)+256*digit]++;
}
b -= sizeof(TNode);
}
//
for (unsigned int i = 1; i < 256; i++)
{
for (unsigned int digit=0; digit< sizeof(m->key); digit++)
{
s[i+256*digit] += s[i-1+256*digit];
}
}
// (LSD)
for (unsigned int digit=0; digit< sizeof(m->key); digit++)
{
RSort_step(m, m_temp, n-1, &s[256*digit] ,digit);
TNode *temp = m;
m = m_temp;
m_temp = temp;
}
// ,
if (sizeof(m->key)==1)
{
TNode *temp = m;
m = m_temp;
m_temp = temp;
memcpy(m, m_temp, n * sizeof(TNode));
}
free(m_temp);
}
//=============================================================
int main()
{
unsigned int n=10000000;
LARGE_INTEGER frequency;
LARGE_INTEGER t1, t2, t3, t4;
double elapsedTime;
TNode *m1 = (TNode*)malloc(sizeof(TNode) * n);
TNode *m2 = (TNode*)malloc(sizeof(TNode) * n);
srand(time(NULL));
for (unsigned int i=0; i<n; i++)
{
m1[i].key = rand()*RAND_MAX+rand();
m2[i].key = m1[i].key;
}
QueryPerformanceFrequency(&frequency);
QueryPerformanceCounter(&t1);
RSort_Node(m1, n);
QueryPerformanceCounter(&t2);
elapsedTime=(float)(t2.QuadPart-t1.QuadPart)/frequency.QuadPart;
printf("The RSort: %.5f seconds\n", elapsedTime);
QueryPerformanceFrequency(&frequency);
QueryPerformanceCounter(&t3);
std::sort(m2, m2+n,[](const TNode &a, const TNode &b){return a.key < b.key;});
QueryPerformanceCounter(&t4);
elapsedTime=(float)(t4.QuadPart-t3.QuadPart)/frequency.QuadPart;
printf("The std::sort: %.5f seconds\n", elapsedTime);
for (unsigned int i=0; i<n; i++)
{
if (m1[i].key!=m2[i].key)
{
printf("\n\n!!!!!\n");
break;
}
}
free(m1);
free(m2);
return 0;
}
Pascal
program SORT;
uses
SysUtils, Windows;
//=============================================================
type TNode = record
key : Longword;
//value : Longword;
end;
type ATNode = array of TNode;
//=============================================================
procedure RSort_Node(var m: array of TNode);
//------------------------------------------------------------------------------
procedure Sort_step(var source, dest: array of TNode; len : Longword; offset: PLongword; const num: Byte);
var b : ^Byte;
v : ^TNode;
begin
b:=@source[len];
v:=@source[len];
inc(b,num);
while v >= @source do
begin
dec(offset[b^]);
dest[offset[b^]] := v^;
dec(b,SizeOf(TNode));
dec(v);
end;
end;
//------------------------------------------------------------------------------
var // ,
s: array[0..1023] of Longword =(
0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0);
i : Longword;
b : ^Byte;
p : Pointer;
begin
GetMem(p, SizeOf(TNode)*Length(m));
//
b:=@m[High(m)];
while (b >= @m[0]) do
begin
Inc(s[(b+3)^+256*3]);
Inc(s[(b+2)^+256*2]);
Inc(s[(b+1)^+256*1]);
Inc(s[(b+0)^+256*0]);
dec(b,SizeOf(TNode));
end;
//
for i := 1 to 255 do
begin
Inc(s[i+256*0], s[i-1+256*0]);
Inc(s[i+256*1], s[i-1+256*1]);
Inc(s[i+256*2], s[i-1+256*2]);
Inc(s[i+256*3], s[i-1+256*3]);
end;
//
Sort_step(m, ATNode(p), High(m), @s[256*0], 0);
Sort_step(ATNode(p), m, High(m), @s[256*1], 1);
Sort_step(m, ATNode(p), High(m), @s[256*2], 2);
Sort_step(ATNode(p), m, High(m), @s[256*3], 3);
FreeMem(p);
end;
//=============================================================
procedure test();
const
n = 10000000;
var
m1: array of TNode;
i,j,k,j1,j2: Longword;
iCounterPerSec: TLargeInteger;
T1, T2: TLargeInteger; //
begin
SetLength(m1,n);
for i := 0 to n - 1 do
begin
m1[i].key := Random(65536 * 65536);
end;
QueryPerformanceFrequency(iCounterPerSec);//
QueryPerformanceCounter(T11); //
RSort_Node(m1);
QueryPerformanceCounter(T22);//
WRITELN('1='+FormatFloat('0.0000', (T22 - T11)/iCounterPerSec) + ' sec.');//
SetLength(m, 0);
end;
//------------------------------------------------------------------------------
begin
test();
Readln();
exit;
end.
Medité en este código durante mucho tiempo, pero no pude escribir mejor. Quizás pueda decirme cómo hacer que este código sea más rápido.
¿Y si quieres un poco más rápido?
El siguiente paso lógico, como lo vi, es usar una tarjeta de video. En Internet, vi muchos razonamientos de que Radix sort es perfectamente paralelo en muchos núcleos de tarjetas de video (casi el mejor método de clasificación no es una tarjeta de video).
Habiendo sucumbido a la tentación de obtener un rendimiento adicional, implementé varias opciones de clasificación usando OpenCL que conozco. Desafortunadamente, no tengo la oportunidad de verificar la implementación en tarjetas de video de gama alta, pero en mi GEFORCE GTX 750 TI, el algoritmo se perdió en la implementación de un solo subproceso en la CPU, debido al hecho de que los datos tenían que enviarse a la tarjeta de video y luego recuperarse. Si no corrió los datos de un lado a otro en el autobús, la velocidad sería aceptable, pero aún así no una fuente. Hay una observación más. En OpenCL, los hilos de ejecución no son síncronos (los grupos de trabajo se ejecutan en un orden arbitrario, hasta donde yo sé, este no es el caso en CUDA, correcto, quién sabe), lo que impide escribir código más eficiente en este caso.
Código de la serie: ¿es posible crear un trolebús ... procesando en OpenCL en Delphi ... pero por qué?
Fui demasiado vago para reescribirlo en 'C', lo publico como está.
program project1;
uses Cl, SysUtils, Windows, Math;
//------------------------------------------------------------------------------
function OCL_Get_Prog(context: cl_context; pdevice_id: Pcl_device_id; name: PChar; S:AnsiString):cl_kernel;
var Tex:PCHAR;
Len:QWORD;
PLen:PQWORD;
Prog:cl_program;
kernel:cl_kernel;
Ret:cl_int;
begin
Tex:=@S[1];
Len:=Length(S);
PLen:=@LEN;
prog:=nil;
kernel:=nil;
//
prog:= clCreateProgramWithSource(context, 1, @Tex, @Len, ret);
if CL_SUCCESS<>ret then writeln('clCreateProgramWithSource Error ',ret);
//
ret:= clBuildProgram(prog, 1, pdevice_id, nil, nil, nil);
if CL_SUCCESS<>ret then writeln('clBuildProgram Error ',ret);
//
kernel:= clCreateKernel(prog, name, ret);
if ret<>CL_SUCCESS then writeln('clCreateKernel Error ',ret);
//
clReleaseProgram(prog);
OCL_Get_Prog:=kernel;
end;
//------------------------------------------------------------------------------
var
context:cl_context;
kernel1, kernel2, kernel3, kernel4 :cl_kernel;
Ret:cl_int;
valueSize : QWord =0;
s0 : PChar;
platform_id:cl_platform_id;
ret_num_platforms:cl_uint;
ret_num_devices:cl_uint;
device_id:cl_device_id;
command_queue:cl_command_queue;
S1,S2,S3,S4 : AnsiString;
memobj1, memobj2, memobj3 :cl_mem;
mem:Array of LongWord;
size:LongWord;
g_works, l_works :LongInt;
iCounterPerSec: TLargeInteger;
T1, T2, T3, T4: TLargeInteger; //
i,j,step:LongWord;
procedure exchange_memobj(var a,b:cl_mem);
var c:cl_mem;
begin
c:=a;
a:=b;
b:=c;
end;
begin
//---------------------------------------------------------------------------
S1 :=
// 1 (O )
'__kernel void sort1(__global uint* m1) {'+
' uint g_id = get_global_id(0);'+
' m1[g_id] = 0;'+
'}';
//---------------------------------------------------------------------------
S2 :=
// 2 ( )
'__kernel void sort2(__global uint* m1, __global uint* m2, const uint len, const uint digit) {'+
' uint g_id = get_global_id(0);'+
' uint size = get_global_size(0);'+
' uint a = g_id / len;'+
' uchar key = (m1[g_id] >> digit);'+
' atomic_inc(&m2[key]);'+
' atomic_inc(&m2[256 * a + key + 256]);'+
'}';
//---------------------------------------------------------------------------
S3 :=
// 3 ( )
'__kernel void sort3(__global uint* m1, const uint len) {'+
' uint l_id = get_global_id(0);'+
' for (uint i = 0; i < 8; i++) {'+
' uint offset = 1 << i;'+
' uint add = (l_id>=offset) ? m1[l_id - offset] : 0;'+
' barrier(CLK_GLOBAL_MEM_FENCE);'+
' m1[l_id] += add;'+
' barrier(CLK_GLOBAL_MEM_FENCE);'+
' }'+
' for (uint i = 1; i < 1024; i++) {'+
' m1[i*256+l_id + 256] += m1[(i-1)*256+l_id + 256];'+
' }'+
' barrier(CLK_GLOBAL_MEM_FENCE);'+
' if (l_id>0) {'+
' for (uint i = 0; i < 1024; i++) {'+
' m1[i*256+l_id + 256] += m1[l_id-1];'+
' }'+
' }'+
'}';
//---------------------------------------------------------------------------
S4 :=
// 4
'__kernel void sort4(__global uint* m1, __global uint* m2, __global uint* m3, const uint digit, const uint len) {'+
' uint g_id = get_global_id(0);'+
' for (int i = len-1; i >= 0; i--) {'+ // !
' uchar key = (m1[g_id*len+i] >> digit);'+
' m2[--m3[g_id*256 + key + 256]] = m1[g_id*len+i];'+
' }'+
'}';
//---------------------------------------------------------------------------
//
ret := clGetPlatformIDs(1,@platform_id,@ret_num_platforms);
if CL_SUCCESS<>ret then writeln('clGetPlatformIDs Error ',ret);
//
ret := clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_GPU, 1, @device_id, @ret_num_devices);
if CL_SUCCESS<>ret then writeln('clGetDeviceIDs Error ',ret);
clGetDeviceInfo(device_id, CL_DEVICE_NAME, 0, nil, valueSize);
GetMem(s0, valueSize);
clGetDeviceInfo(device_id, CL_DEVICE_NAME, valueSize, s0, valueSize);
Writeln('DEVICE_NAME: '+s0);
FreeMem(s0);
//
context:= clCreateContext(nil, 1, @device_id, nil, nil, ret);
if CL_SUCCESS<>ret then writeln('clCreateContext Error ',ret);
//
command_queue := clCreateCommandQueue(context, device_id, 0, ret);
if CL_SUCCESS<>ret then writeln('clCreateContext Error ',ret);
//-------------------------------------------------------------
kernel1 := OCL_Get_Prog(context, @device_id, 'sort1', S1);
kernel2 := OCL_Get_Prog(context, @device_id, 'sort2', S2);
kernel3 := OCL_Get_Prog(context, @device_id, 'sort3', S3);
kernel4 := OCL_Get_Prog(context, @device_id, 'sort4', S4);
//-------------------------------------------------------------
size:=256*256*16*10;
g_works := size;
l_works := 256;
Randomize;
SetLength(mem, size);
for i:=0 to size-1 do mem[i]:=random(256*256*256*256);
//
memobj1 := clCreateBuffer(context, CL_MEM_READ_WRITE, size * sizeof(cl_uint), nil, ret);
if ret<>CL_SUCCESS then writeln('clCreateBuffer1 Error ',ret);
memobj2 := clCreateBuffer(context, CL_MEM_READ_WRITE, size * sizeof(cl_uint), nil, ret);
if ret<>CL_SUCCESS then writeln('clCreateBuffer2 Error ',ret);
memobj3 := clCreateBuffer(context, CL_MEM_READ_WRITE, (256+256*1024) * sizeof(cl_uint), nil, ret);
if ret<>CL_SUCCESS then writeln('clCreateBuffer3 Error ',ret);
QueryPerformanceFrequency(iCounterPerSec); //
QueryPerformanceCounter(T1); //
//
ret := clEnqueueWriteBuffer(command_queue, memobj1, CL_TRUE, 0, size * sizeof(cl_int), @mem[0], 0, nil, nil);
if ret<>CL_SUCCESS then writeln('clEnqueueWriteBuffer Error ',ret);
QueryPerformanceCounter(T2);//
Writeln('write '+FormatFloat('0.0000', (T2 - T1)/iCounterPerSec) + ' second.');//
QueryPerformanceFrequency(iCounterPerSec); //
QueryPerformanceCounter(T3); //
for step:=0 to 3 do
begin
//-------------------------------------------------------------
QueryPerformanceFrequency(iCounterPerSec); //
QueryPerformanceCounter(T1); //
//-------------------------------------------------------------
// 1 ( )
ret:= clSetKernelArg(kernel1, 0, sizeof(cl_mem), @memobj3 );
if ret<>CL_SUCCESS then writeln('clSetKernelArg_1_1 Error ',ret);
i := 256+256*1024;
ret:= clEnqueueNDRangeKernel(command_queue, kernel1, 1, nil, @i, nil, 0, nil, nil);
if ret<>CL_SUCCESS then writeln('clEnqueueNDRangeKernel_1 Error ',ret);
//-------------------------------------------------------------
clFinish(command_queue); //
QueryPerformanceCounter(T2); //
Writeln('step 1 '+FormatFloat('0.0000', (T2 - T1)/iCounterPerSec) + ' second.');//
QueryPerformanceFrequency(iCounterPerSec); //
QueryPerformanceCounter(T1); //
//-------------------------------------------------------------
// 2 ( )
ret:= clSetKernelArg(kernel2, 0, sizeof(cl_mem), @memobj1 );
if ret<>CL_SUCCESS then writeln('clSetKernelArg_2_1 Error ',ret);
ret:= clSetKernelArg(kernel2, 1, sizeof(cl_mem), @memobj3 );
if ret<>CL_SUCCESS then writeln('clSetKernelArg_2_2 Error ',ret);
j := size div (1024);
ret:= clSetKernelArg(kernel2, 2, sizeof(j), @j );
if ret<>CL_SUCCESS then writeln('clSetKernelArg_2_3 Error ',ret);
j:=step*8;
ret:= clSetKernelArg(kernel2, 3, sizeof(j), @j );
if ret<>CL_SUCCESS then writeln('clSetKernelArg_2_4 Error ',ret);
i := size;
ret:= clEnqueueNDRangeKernel(command_queue, kernel2, 1, nil, @i, nil, 0, nil, nil);
if ret<>CL_SUCCESS then writeln('clEnqueueNDRangeKernel_2 Error ',ret);
//-------------------------------------------------------------
clFinish(command_queue); //
QueryPerformanceCounter(T2);//
Writeln('step 2 '+FormatFloat('0.0000', (T2 - T1)/iCounterPerSec) + ' second.');//
QueryPerformanceFrequency(iCounterPerSec);//
QueryPerformanceCounter(T1); //
//-------------------------------------------------------------
// 3 ( )
ret:= clSetKernelArg(kernel3, 0, sizeof(cl_mem), @memobj3 );
if ret<>CL_SUCCESS then writeln('clSetKernelArg_3_1 Error ',ret);
j := size;
ret:= clSetKernelArg(kernel3, 1, sizeof(j), @j );
if ret<>CL_SUCCESS then writeln('clSetKernelArg_3_3 Error ',ret);
j := 256;
ret:= clEnqueueNDRangeKernel(command_queue, kernel3, 1, nil, @j, @j, 0, nil, nil);
if ret<>CL_SUCCESS then writeln('clEnqueueNDRangeKernel_3 Error ',ret);
//-------------------------------------------------------------
clFinish(command_queue); //
QueryPerformanceCounter(T2);//
Writeln('step 3 '+FormatFloat('0.0000', (T2 - T1)/iCounterPerSec) + ' second.');//
QueryPerformanceFrequency(iCounterPerSec);//
QueryPerformanceCounter(T1); //
//-------------------------------------------------------------
// 4 ()
ret:= clSetKernelArg(kernel4, 0, sizeof(cl_mem), @memobj1 );
if ret<>CL_SUCCESS then writeln('clSetKernelArg_4_1 Error ',ret);
ret:= clSetKernelArg(kernel4, 1, sizeof(cl_mem), @memobj2 );
if ret<>CL_SUCCESS then writeln('clSetKernelArg_4_2 Error ',ret);
ret:= clSetKernelArg(kernel4, 2, sizeof(cl_mem), @memobj3 );
if ret<>CL_SUCCESS then writeln('clSetKernelArg_4_3 Error ',ret);
j:=step*8;
ret:= clSetKernelArg(kernel4, 3, sizeof(j), @j );
if ret<>CL_SUCCESS then writeln('clSetKernelArg_4_4 Error ',ret);
j := size div (1024);
ret:= clSetKernelArg(kernel4, 4, sizeof(j), @j );
if ret<>CL_SUCCESS then writeln('clSetKernelArg_4_5 Error ',ret);
i := (1024); //
ret:= clEnqueueNDRangeKernel(command_queue, kernel4, 1, nil, @i, nil, 0, nil, nil);
if ret<>CL_SUCCESS then writeln('clEnqueueNDRangeKernel_4 Error ',ret);
clFinish(command_queue);
//
// memobj1
exchange_memobj(memobj2, memobj1);
//-------------------------------------------------------------
clFinish(command_queue); //
QueryPerformanceCounter(T2);//
Writeln('step 4 '+FormatFloat('0.0000', (T2 - T1)/iCounterPerSec) + ' second.');//
//-------------------------------------------------------------
end;
QueryPerformanceCounter(T4);//
Writeln('all not R/W '+FormatFloat('0.0000', (T4 - T3)/iCounterPerSec) + ' second.');//
QueryPerformanceFrequency(iCounterPerSec);//
QueryPerformanceCounter(T1); //
//
ret:= clEnqueueReadBuffer(command_queue, memobj1, CL_TRUE, 0, size * sizeof(cl_int), @mem[0], 0, nil, nil);
if ret<>CL_SUCCESS then writeln('clEnqueueReadBuffer Error ',ret);
QueryPerformanceCounter(T2);//
Writeln('Read '+FormatFloat('0.0000', (T2 - T1)/iCounterPerSec) + ' second.');//
//
clReleaseMemObject(memobj1); //
clReleaseMemObject(memobj2);
clReleaseMemObject(memobj3);
clReleaseKernel(kernel1); //
clReleaseKernel(kernel2);
clReleaseKernel(kernel3);
clReleaseKernel(kernel4);
clReleaseCommandQueue(command_queue); //
clReleaseContext(context); //
//-------------------------------------------------------------
SetLength(mem, 0);
readln;
end.
Queda una opción más sin probar: el subproceso múltiple. Ya usando solo C desnudo y la biblioteca OpenMP, decidí averiguar qué efecto tendría el uso de múltiples núcleos de CPU.
Inicialmente, la idea era dividir la matriz original en partes iguales, transferirlas a secuencias separadas y luego fusionarlas (fusionar ordenación). La clasificación no salió mal, pero la fusión ralentizó mucho toda la estructura, cada encolado equivale a una pasada adicional a través de la matriz. El efecto fue peor que trabajar en un hilo. Rechazó la implementación por no ser práctica.
Como resultado, apliqué una clasificación paralela muy similar a la que se usa en la GPU. Con ella todo resultó mucho mejor.
Características del procesamiento paralelo:
En la implementación actual, evitó los problemas de sincronización de subprocesos tanto como fue posible (las funciones atómicas tampoco se utilizan debido al hecho de que diferentes subprocesos leen bloques de memoria diferentes, aunque a veces vecinos). Las transmisiones no son algo gratuito; el procesador gasta microsegundos preciosos en su creación y sincronización. Si mantiene las barreras al mínimo, puede ahorrar un poco. Desafortunadamente, dado que todos los subprocesos usan la misma memoria caché L3 y RAM, la ganancia general del algoritmo no es tan significativa debido a la ley de Amdahl , con un aumento en el número de subprocesos.
C
#include <stdio.h>
#include <omp.h>
//=============================================================
typedef struct TNode {
//unsigned long long key;
unsigned int key;
//unsigned short key;
//unsigned char key;
unsigned int value;
//unsigned int value1;
//unsigned int value2;
} TNode;
//=============================================================
void RSort_Parallel(TNode *m, unsigned int n)
{
//
unsigned char threads = omp_get_num_procs();
//
TNode *m_temp = (TNode*)malloc(sizeof(TNode) * n);
unsigned int *s = (unsigned int*)malloc(sizeof(unsigned int) * 256 * threads);
#pragma omp parallel num_threads(threads)
{
TNode *source = m;
TNode *dest = m_temp;
unsigned int l = omp_get_thread_num();
unsigned int div = n / omp_get_num_threads();
unsigned int mod = n % omp_get_num_threads();
unsigned int left_index = l < mod ? (div + (mod == 0 ? 0 : 1)) * l : n - (omp_get_num_threads() - l) * div;
unsigned int right_index = left_index + div - (mod > l ? 0 : 1);
for (unsigned int digit=0; digit< sizeof(m->key); digit++)
{
unsigned int s_sum[256] = {0};
unsigned int s0[256] = {0};
unsigned char *b1 = (unsigned char*)&source[right_index].key;
unsigned char *b2 = (unsigned char*)&source[left_index].key;
while (b1 >= b2)
{
s0[*(b1+digit)]++;
b1 -= sizeof(TNode);
}
for (unsigned int i=0; i<256; i++)
{
s[i+256*l] = s0[i];
}
#pragma omp barrier
for (unsigned int j=0; j<threads; j++)
{
for (unsigned int i=0; i<256; i++)
{
s_sum[i] += s[i+256*j];
if (j<l)
{
s0[i] += s[i+256*j];
}
}
}
for (unsigned int i=1; i<256; i++)
{
s_sum[i] += s_sum[i-1];
s0[i] += s_sum[i-1];
}
unsigned char *b = (unsigned char*)&source[right_index].key + digit;
TNode *v1 = &source[right_index];
TNode *v2 = &source[left_index];
while (v1 >= v2)
{
dest[--s0[*b]] = *v1--;
b -= sizeof(TNode);
}
#pragma omp barrier
TNode *temp = source;
source = dest;
dest = temp;
}
}
// ,
if (sizeof(m->key)==1)
{
memcpy(m, m_temp, n * sizeof(TNode));
}
free(s);
free(m_temp);
}
//=============================================================
int main()
{
unsigned int n=10000000;
TNode *m1 = (TNode*)malloc(sizeof(TNode) * n);
srand(time(NULL));
for (unsigned int i=0; i<n; i++)
{
m1[i].key = rand()*RAND_MAX+rand();
}
RSort_Parallel(m1, n);
free(m1);
return 0;
}
Espero que haya sido interesante.
Saludos cordiales, Atentamente, Rebuilder.
PS
Comparación de algoritmos en términos de velocidad deliberadamente no lo es. Los resultados de la comparación, las sugerencias y las críticas son bienvenidos.
PPS
