OpenCL Performance Optimization

Question

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).

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ó.

Answer 1

Parece que CADA hilo en un grupo de trabajo está copiando los mismos 50 flotantes antes de que comience el cálculo. Esto saturará el ancho de banda de la memoria global.

probar esto

if (lid == 0) 
{ 
    async_work_group_copy(lBoneMats, boneMats, NUM_BONES, 0); 
} 

Esto hace la copia de una sola vez por cada grupo de trabajo.

Answer 2

¿Has descubierto el motivo de la ralentización de tu kernel?

Tal vez estoy equivocado, pero creo que tener todos los elementos de trabajo dentro de un grupo de trabajo accediendo a la misma memoria local puede conducir a un cuello de botella.

Answer 3

Dos cosas están afectando el rendimiento en su ejercicio.

1) OpenCL se ajusta a C99 std que no contiene nada acerca de las funciones en línea, es decir, el compilador CLCC ya sea simplemente ignora la palabra clave inline y hace una llamada normal, o que apoya la expansión en línea en silencio. Pero no está obligado a admitir esa característica.

Por lo tanto, mejor defina su MultiplyMatrixVector como una macro de pre-procesador. Aunque este no es un problema importante en su caso.

2) Usted amenaza incorrectamente la memoria local (LDM).

Aunque su latencia es inferior a la latencia del global memory cuando se accedió correctamente, el local memory está sujeto a conflictos bancarios.

Su índice de vértice se calcula con zancada 100 por elemento de trabajo. El número de bancos depende de la GPU en uso, pero generalmente es 16 o 32, i.mi. puede acceder a hasta 16 (32) variables de cuatro bytes LDM en un ciclo sin penalización si todas están en bancos diferentes. De lo contrario, obtienes un bank conflict (cuando dos o más subprocesos tienen acceso al mismo banco) que se serializa. Sus 100 hilos en un grupo de trabajo accede a la matriz en LDM sin ningún acuerdo especial sobre conflictos bancarios. Además, los elementos del conjunto son float16, es decir, un único elemento abarca los 16 bancos (o la mitad de 32 bancos). Por lo tanto, tiene un conflicto bancario en cada fila de la función MultiplyMatrixVector. El cumulativo degree que entra en conflicto al menos 16x32 (aquí 16 es la cantidad de elementos vectoriales a los que accede y 32 tiene un tamaño de mitad de frente de onda o medio).

La solución a este problema no es copiar esa matriz a LDM, pero asignarlo en el huésped con CL_MEM_READ_ONLY (que ya se hizo) y declarar el núcleo usando __constant especificador de boneMats argumento. A continuación, la biblioteca OpenCL asignaría la memoria en el área constante en el interior GPU y el acceso a esa matriz sería rápida:

kernel void skelanim(__constant const float16* boneMats, 
        global const float4* vertices, 
        global const float4* weights, 
        global const uint4* indices, 
        global float4* resVertices)