OpenCL性能優化

Question

我已經開始學習OpenCL，並且我目前正在嘗試測試我可以提高多少程度來提高簡單的骨骼動畫算法的性能。爲此，我編寫了一個程序，用於從隨機生成的頂點和變換矩陣執行兩次骨骼動畫，一次使用純C++中的SSE優化線性代數庫，一次在GPU上使用我自己的OpenCL內核（我正在測試Nvidia GTX 460）。

我從一個簡單的內核開始，每個工作項目只轉換一個頂點，所有值都從全局內存中讀取。因爲我對這個內核的性能不滿意，我試着優化一下。我現在的內核是這樣的：

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

現在我處理每個工作項目頂點的常數，我預取所有的骨基質到本地內存只有一次，每個工作項目，我相信會導致以提高性能，因爲之後可以從更快的本地存儲器讀取用於多個頂點的矩陣。不幸的是，這個內核比我第一次嘗試的性能差，甚至比僅有CPU的實現更差。

爲什麼性能如此糟糕以至於應該進行優化？

如果有幫助，這是我如何執行內核：

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

我想有更多的事情可以優化，也許配料的其他一些全球讀取在一起，但首先我真的想要知道爲什麼這第一次優化不起作用。

Answer 1

看起來像工作組中的每個線程都在計算開始之前複製相同的50個浮點數。這會使全局內存帶寬飽和。

試試這個

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

這每個工作組做副本只有一次。

Answer 2

您是否發現內核放緩的原因？

也許我錯了，但我認爲讓一個工作組中的所有工作項訪問相同的本地內存可能會導致瓶頸。

Answer 3

兩件事影響你的鍛鍊表現。

1）OpenCL符合C99 STD不包含有關內聯函數什麼，即CLCC編譯器要麼只是忽略inline關鍵字和不正常通話，或者它只是默默地支持內聯。但它沒有強制支持該功能。

因此，最好將您的MultiplyMatrixVector定義爲預處理器宏。雖然這不是你的情況的主要問題。

2）您錯誤地威脅本地內存（LDM）。

雖然它的延遲時間小於global memory正確訪問時的延遲時間，但local memory會受到銀行衝突的影響。

您的頂點索引是用每個工作項目的步幅100來計算的。銀行的數量取決於正在使用的GPU，但通常是16或32，即。即您可以在一個週期內訪問多達16（32）個四字節的變量，如果它們全部位於不同的銀行，則不會受到懲罰。否則，你會得到一個bank conflict（當兩個或更多的線程訪問同一個存儲區時）被序列化。 工作組中的100個線程訪問LDM中的數組，而沒有關於銀行衝突的特殊安排。而且，陣列元素是float16，即單個元素跨越全部16個組（或32個組的一半）。因此，您在每行MultiplyMatrixVector函數中存在銀行衝突。累積degree衝突至少16x32（這裏16是您訪問的向量元素的數量，32是半波前或半扭曲的大小）。

這裏的解決方案是不是該數組複製到LDM，而是將其與CL_MEM_READ_ONLY分配的主機（你已經做了），並使用__constant符爲boneMats參數聲明你的內核。 然後OpenCL庫將分配在固定區域內部的存儲器GPU和訪問該陣列是快：

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