5

私は OpenCL の学習を開始し、現在、単純な骨格アニメーション アルゴリズムのパフォーマンスをどれだけ改善できるかをテストしようとしています。これを行うために、ランダムに生成された頂点と変換行列から骨格アニメーションを 2 回実行するプログラムを作成しました。 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));
    }
}

ここで、ワークアイテムごとに一定数の頂点を処理し、ワークアイテムごとに 1 回だけすべてのボーン マトリックスをローカル メモリにプリフェッチします。複数の頂点のマトリックスをその後、より高速なローカルメモリ。残念ながら、このカーネルは私の最初の試行よりもパフォーマンスが悪く、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;
}

他のグローバル読み取りのいくつかをまとめてバッチ処理するなど、最適化できることが他にもあると思いますが、最初に、この最初の最適化が機能しなかった理由を知りたいと思います。

4

2 に答える 2

-2

計算が開始される前に、ワーク グループ内の各スレッドが同じ 50 個の float をコピーしているようです。これにより、グローバル メモリの帯域幅が飽和します。

これを試して

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

これにより、ワーク グループごとに 1 回だけコピーが実行されます。

于 2012-08-05T13:44:36.583 に答える