OpenCL Performance Optimization

Question

Ich habe angefangen, OpenCL zu lernen und ich versuche derzeit zu testen, wie sehr ich die Leistung für einen einfachen Skelettanimationsalgorithmus verbessern kann. Um dies zu tun, habe ich ein Programm geschrieben, das Skelettanimation aus zufällig generierten Scheitelpunkten und Transformationsmatrizen zweimal durchführt, einmal mit einer SSE-optimierten linearen Algebra-Bibliothek in reinem C++ und einmal mit meinem eigenen OpenCL-Kernel auf GPU (ich teste auf einem Nvidia GTX 460).

Ich begann mit einem einfachen Kernel, bei dem jedes Arbeitselement genau einen Knoten transformiert, wobei alle Werte aus dem globalen Speicher gelesen werden. Da ich mit der Leistung dieses Kernels nicht zufrieden war, versuchte ich ein wenig zu optimieren. Mein aktueller Kernel sieht wie folgt aus:

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

Jetzt habe ich eine konstante Anzahl von Eckpunkten pro Arbeitsposten verarbeiten, und ich Prefetch alle Knochen Matrizen in den lokalen Speicher nur einmal für jede Arbeitsposition, die ich glaubte, würde dazu führen, zu einer besseren Leistung, da die Matrizen für mehrere Vertices anschließend aus dem schnelleren lokalen Speicher gelesen werden können. Leider ist dieser Kernel schlechter als mein erster Versuch und sogar schlechter als die reine CPU-Implementierung.

Warum ist die Leistung so schlecht mit dieser sollte Optimierung sein?

Wenn es hilft, ist hier, wie ich den Kernel ausführen:

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

Ich denke, es gibt mehr Dinge, die optimiert werden könnte, vielleicht einige Dosieren der anderen globalen liest zusammen, aber zuerst würde ich wirklich wie zu wissen, warum diese erste Optimierung nicht funktioniert hat.

Answer 1

Es sieht so aus, als würde jeder Thread in einer Arbeitsgruppe die gleichen 50 Floats kopieren, bevor die Berechnung beginnt. Dies wird die Global Memory-Bandbreite sättigen.

versuchen diese

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

Diese die Kopie tut nur einmal pro Arbeitsgruppe.

Answer 2

Haben Sie den Grund für die Verlangsamung Ihres Kernels herausgefunden?

Vielleicht liege ich falsch, aber ich denke, dass alle Arbeitsaufgaben innerhalb einer Arbeitsgruppe, die auf denselben lokalen Speicher zugreifen, zu einem Engpass führen können.

Answer 3

Zwei Dinge beeinflussen die Leistung in Ihrer Übung.

1) OpenCL entspricht C99 std, die nichts über Inline-Funktionen, das heißt den CLCC Compiler enthalten entweder ignoriert einfach das inline Schlüsselwort und macht einen regulären Anruf, oder es unterstützt den inlining geräuschlos. Es ist jedoch nicht vorgeschrieben, diese Funktion zu unterstützen.

Also, besser definieren Sie Ihre MultiplyMatrixVector als Pre-Prozessor-Makro. Obwohl dies in Ihrem Fall kein großes Problem ist.

2) Sie gefährden fälschlicherweise den lokalen Speicher (LDM).

Obwohl seine Latenzzeiten weniger als die Latenz der global memory, wenn es richtig zugegriffen hat, unterliegt die local memory Bankkonflikte.

Ihr Vertex-Index wird mit Schritt 100 pro Arbeitselement berechnet. Die Anzahl der Bänke hängt von der verwendeten GPU ab, aber normalerweise ist sie 16 oder 32, d.h.e. Sie können bis zu 16 (32) vier Byte LDM Variablen in einem Zyklus ohne Strafe zugreifen, wenn alle in verschiedenen Banken sind. Andernfalls erhalten Sie eine bank conflict (wenn zwei oder mehr Threads auf dieselbe Bank zugreifen), die serialisiert ist. Ihre 100 Threads in einer Arbeitsgruppe greifen auf das Array in LDM ohne besondere Vereinbarung über Bankkonflikte zu. Darüber hinaus sind die Array-Elemente float16, d. H. Ein einzelnes Element überspannt alle 16 Bänke (oder die Hälfte von 32 Bänken). Somit haben Sie in jeder Zeile der MultiplyMatrixVector Funktion einen Bankkonflikt. Die kumulative degree, die mindestens 16x32 Konflikt (hier 16 ist die Anzahl der Vektorelemente, die Sie zugreifen und 32 ist eine halbe Wellenfront oder Halfwarp).

Die Lösung ist hier nicht das Array zu LDM kopieren, aber es in den Host mit CL_MEM_READ_ONLY (die Sie bereits getan haben) und erklären Sie Ihren Kernel mit __constant Spezifizierer für boneMats Argument zuzuweisen. Dann wird die OpenCL Bibliothek würde den Speicher in dem konstanten Bereich innerhalb GPU und der Zugang zu diesem Array wäre schnell zuweisen:

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