Comencé a aprender OpenCL y actualmente trato de probar cuánto puedo mejorar el rendimiento de un algoritmo de animación esquelética simple. Para hacer esto, he escrito un programa que realiza animaciones esqueléticas desde vértices generados aleatoriamente y matrices de transformación dos veces, una con una biblioteca de álgebra lineal optimizada para SSE en C++ simple y otra usando mi propio kernel OpenCL en GPU (estoy probando un Nvidia GTX 460).OpenCL Performance Optimization
Empecé con un kernel simple donde cada elemento de trabajo transforma exactamente un vértice, con todos los valores leídos de la memoria global. Como no estaba satisfecho con el rendimiento de este kernel, traté de optimizar un poco. Mi núcleo actual se parece a esto:
inline float4 MultiplyMatrixVector(float16 m, float4 v)
{
return (float4) (
dot(m.s048C, v),
dot(m.s159D, v),
dot(m.s26AE, v),
dot(m.s37BF, v)
);
}
kernel void skelanim(global const float16* boneMats, global const float4* vertices, global const float4* weights, global const uint4* indices, global float4* resVertices)
{
int gid = get_global_id(0);
int lid = get_local_id(0);
local float16 lBoneMats[NUM_BONES];
async_work_group_copy(lBoneMats, boneMats, NUM_BONES, 0);
barrier(CLK_LOCAL_MEM_FENCE);
for (int i = 0 ; i < NUM_VERTICES_PER_WORK_ITEM ; i++) {
int vidx = gid*NUM_VERTICES_PER_WORK_ITEM + i;
float4 vertex = vertices[vidx];
float4 w = weights[vidx];
uint4 idx = indices[vidx];
resVertices[vidx] = (MultiplyMatrixVector(lBoneMats[idx.x], vertex * w.x)
+ MultiplyMatrixVector(lBoneMats[idx.y], vertex * w.y)
+ MultiplyMatrixVector(lBoneMats[idx.z], vertex * w.z)
+ MultiplyMatrixVector(lBoneMats[idx.w], vertex * w.w));
}
}
Ahora procesar un número constante de vértices por el trabajo a punto, y yo captación previa todas las matrices óseas en la memoria local solamente una vez para cada-elemento de trabajo, lo que yo creía que llevaría a un rendimiento mucho mejor porque las matrices para múltiples vértices podrían leerse luego en la memoria local más rápida. Desafortunadamente, este núcleo funciona peor que mi primer intento, e incluso peor que la implementación de solo CPU.
¿Por qué el rendimiento es tan malo con esta optimización que debería ser?
Si ayuda, aquí es cómo ejecutar el kernel:
#define NUM_BONES 50
#define NUM_VERTICES 30000
#define NUM_VERTICES_PER_WORK_ITEM 100
#define NUM_ANIM_REPEAT 1000
uint64_t PerformOpenCLSkeletalAnimation(Matrix4* boneMats, Vector4* vertices, float* weights, uint32_t* indices, Vector4* resVertices)
{
File kernelFile("/home/alemariusnexus/test/skelanim.cl");
char opts[256];
sprintf(opts, "-D NUM_VERTICES=%u -D NUM_REPEAT=%u -D NUM_BONES=%u -D NUM_VERTICES_PER_WORK_ITEM=%u", NUM_VERTICES, NUM_ANIM_REPEAT, NUM_BONES, NUM_VERTICES_PER_WORK_ITEM);
cl_program prog = BuildOpenCLProgram(kernelFile, opts);
cl_kernel kernel = clCreateKernel(prog, "skelanim", NULL);
cl_mem boneMatBuf = clCreateBuffer(ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, NUM_BONES*sizeof(Matrix4), boneMats, NULL);
cl_mem vertexBuf = clCreateBuffer(ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, NUM_VERTICES*sizeof(Vector4), vertices, NULL);
cl_mem weightBuf = clCreateBuffer(ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, NUM_VERTICES*4*sizeof(float), weights, NULL);
cl_mem indexBuf = clCreateBuffer(ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, NUM_VERTICES*4*sizeof(uint32_t), indices, NULL);
cl_mem resVertexBuf = clCreateBuffer(ctx, CL_MEM_WRITE_ONLY | CL_MEM_ALLOC_HOST_PTR, NUM_VERTICES*sizeof(Vector4), NULL, NULL);
uint64_t s, e;
s = GetTickcount();
clSetKernelArg(kernel, 0, sizeof(cl_mem), &boneMatBuf);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &vertexBuf);
clSetKernelArg(kernel, 2, sizeof(cl_mem), &weightBuf);
clSetKernelArg(kernel, 3, sizeof(cl_mem), &indexBuf);
clSetKernelArg(kernel, 4, sizeof(cl_mem), &resVertexBuf);
size_t globalWorkSize[] = { NUM_VERTICES/NUM_VERTICES_PER_WORK_ITEM };
size_t localWorkSize[] = { NUM_BONES };
for (size_t i = 0 ; i < NUM_ANIM_REPEAT ; i++) {
clEnqueueNDRangeKernel(cq, kernel, 1, NULL, globalWorkSize, localWorkSize, 0, NULL, NULL);
}
clEnqueueReadBuffer(cq, resVertexBuf, CL_TRUE, 0, NUM_VERTICES*sizeof(Vector4), resVertices, 0, NULL, NULL);
e = GetTickcount();
return e-s;
}
supongo que hay más cosas que podrían ser optimizados, tal vez dosificadora de algunas de las otras mundial lee juntos, pero primero me gusta mucho para saber por qué esta primera optimización no funcionó.
no sé sobre el rendimiento, pero lo que está haciendo parece tener resultados indefinidos . Utiliza una operación async_copy seguida de una barrera. La barrera no esperará a que finalice la copia asíncrona, sino que continuará tan pronto como todos los elementos de trabajo hayan llegado a ese punto. De acuerdo con la especificación, debe usar la función wait_group_events en su núcleo después de una async_copy, o los resultados no están definidos. Esto tiene sentido, porque la async_copy está ocurriendo mientras el resto del kernel se está ejecutando, por lo que wait_group_events forzará al kernel a asegurarse de que la copia de la memoria esté lista. –