OpenCLパフォーマンスの最適化

Question

私はOpenCLの学習を始めました。現在、単純な骨格アニメーションアルゴリズムのパフォーマンスをどれだけ向上させることができるかを試しています。これを行うために、無作為に生成された頂点と変換行列から骨格アニメーションを2回、普通のC++でSSE最適化線形代数ライブラリで1回、GPUで自分のOpenCLカーネルを使って1回実行するプログラムを作成しました。 Nvidia GTX 460）。

私は、各作業項目が正確に1つの頂点を変換し、すべての値がグローバルメモリから読み込まれる単純なカーネルを使い始めました。私はこのカーネルの性能に満足していなかったので、私は少し最適化しようとしました。私の現在のカーネルは、次のようになります。

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

あなたの演習では、2つのことがパフォーマンスに影響します。

1）OpenCLは、インライン関数については何も含まれている、すなわちCLCCコンパイラのどちらかだけinlineキーワードを無視し、通常の呼び出しを行い、またはそれは静かにインライン化をサポートしていませんC99 STDに準拠しています。しかし、その機能をサポートすることは必須ではありません。

したがって、MultiplyMatrixVectorをプリプロセッサマクロとして定義することをお勧めします。これはあなたのケースでは大きな問題ではありませんが。

2）ローカルメモリ（LDM）を間違って脅かしました。

global memoryのレイテンシは、正しくアクセスしたときのレイテンシよりも短くはありますが、local memoryはバンク競合の影響を受けます。

頂点インデックスは、作業項目ごとに100ストライドで計算されます。バンクの数は使用中のGPUに依存するが、通常は16または32である。e。すべての銀行が異なる銀行にある場合、最大16（32）個の4バイトのLDM変数をペナルティなしで1サイクルでアクセスすることができます。それ以外の場合は、bank conflict（2つ以上のスレッドが同じバンクにアクセスするとき）がシリアル化されます。 ワークグループ内の100個のスレッドは、LDMの配列にアクセスしますが、バンクの競合は特にありません。さらに、配列要素はfloat16であり、すなわち、1つの要素がすべての16のバンク（または32のバンクのうちの半分）に及ぶ。したがって、あなたはMultiplyMatrixVector機能の各行に銀行の競合があります。少なくとも16x32（ここでは16はアクセスするベクトル要素の数であり、32は波面または半波の半分のサイズです）と競合する累積値degreeです。ここ

ソリューションはLDMにその配列をコピーするのではなく、（あなたはすでにやっている）CL_MEM_READ_ONLYでホストにそれを割り当て、boneMats引数に__constant指定子を使用してカーネルを宣言することはありません。私はパフォーマンスについて知っているが、何をやっていることは、未定義の結果を持っていると思わないでください

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