OpenCL Performance Optimization

Question

Ho iniziato ad imparare OpenCL e attualmente sto provando a testare quanto posso migliorare le prestazioni per un semplice algoritmo di animazione scheletrica. Per fare questo ho scritto un programma che esegue l'animazione scheletrica da vertici generati casualmente e matrici di trasformazione due volte, una volta con una libreria algebra lineare ottimizzata SSE in C++ semplice, e una volta usando il mio kernel OpenCL su GPU (sto testando su un Nvidia GTX 460).

Ho iniziato con un semplice kernel in cui ogni oggetto di lavoro trasforma esattamente un vertice, con tutti i valori letti dalla memoria globale. Poiché non ero soddisfatto delle prestazioni di questo kernel, ho cercato di ottimizzare un po '. Il mio kernel attuale si presenta così:

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)); 
    } 
} 

Ora mi elaborare un numero costante di vertici per il lavoro-voce, e io prefetch tutte le matrici ossee nella memoria locale solo una volta per ogni opera-voce, che ho creduto porterebbe per migliorare le prestazioni in quanto le matrici per più vertici potrebbero essere letti dalla memoria locale più veloce in seguito. Sfortunatamente, questo kernel ha prestazioni peggiori del mio primo tentativo e persino peggiore rispetto all'implementazione della sola CPU.

Perché le prestazioni sono così negative con questa ottimizzazione dovrebbe essere?

Se aiuta, ecco come ho eseguire il 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; 
} 

Credo che ci sono altre cose che potrebbero essere ottimizzati, forse il dosaggio alcune delle altre globale legge insieme, ma prima mi piacerebbe davvero per sapere perché questa prima ottimizzazione non ha funzionato.

Answer 1

Sembra che OGNI thread in un gruppo di lavoro stia copiando gli stessi 50 float prima dell'avvio del calcolo. Questo saturerà la larghezza di banda della memoria globale.

provare questo

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

Questo fa la copia solo una volta per ogni gruppo di lavoro.

Answer 2

Hai scoperto il motivo del rallentamento del kernel?

Forse ho sbagliato, ma penso che avere tutti gli elementi di lavoro all'interno di un gruppo di lavoro che accede alla stessa memoria locale possa portare a un collo di bottiglia.

Answer 3

Due cose stanno influenzando la prestazione nel vostro esercizio.

1) OpenCL conforme alle C99 std che non contiene nulla di funzioni inline, vale a dire il compilatore CLCC o solo ignora la parola inline e fa una chiamata normale, o sostiene la messa in linea in silenzio. Ma non è obbligatorio per supportare questa funzionalità.

Quindi, è meglio definire il MultiplyMatrixVector come una macro pre-processore. Anche se questo non è un grosso problema nel tuo caso.

2) In modo errato minaccia la memoria locale (LDM).

Sebbene la latenza sia inferiore alla latenza dello global memory quando accede correttamente, lo local memory è soggetto a conflitti bancari.

L'indice dei vertici viene calcolato con passo 100 per articolo di lavoro. Il numero di banchi dipende dalla GPU in uso ma solitamente è 16 o 32, i.e. è possibile accedere a un massimo di 16 (32) quattro byte LDM variabili in un ciclo senza penalità se tutte sono in banche diverse. Altrimenti, si ottiene bank conflict (quando due o più thread accedono allo stesso banco) che è serializzato. I 100 thread in un gruppo di lavoro accedono all'array in LDM senza alcun accordo speciale sui conflitti bancari. Inoltre, gli elementi dell'array sono float16, cioè un singolo elemento si estende su tutti i 16 banchi (o metà su 32 banchi). Pertanto, si ha un conflitto bancario in ogni riga della funzione MultiplyMatrixVector. Il cummativo degree che confligge almeno 16x32 (qui 16 è il numero degli elementi vettoriali a cui si accede e 32 è una dimensione di mezzo fronte d'onda o halfwarp).

La soluzione non è quello di copiare tale matrice a LDM, ma di allocare nel host con CL_MEM_READ_ONLY (che avete già fatto) e dichiarare il kernel usando __constant identificatore per boneMats argomento. Poi la libreria OpenCL sarebbe allocare la memoria nella zona costante all'interno GPU e l'accesso a tale matrice sarebbe veloce:

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