5

我已经开始学习 OpenCL，目前我正在尝试测试我可以在多大程度上提高简单骨骼动画算法的性能。为此，我编写了一个程序，从随机生成的顶点和变换矩阵执行骨骼动画两次，一次使用纯 C++ 中的 SSE 优化线性代数库，一次使用我自己在 GPU 上的 OpenCL 内核（我正在测试英伟达 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;
}

我想还有更多可以优化的东西，也许将其他一些全局读取批处理在一起，但首先我真的很想知道为什么第一次优化不起作用。

4

2 回答 2

1

有两件事会影响你的锻炼表现。

1)OpenCL符合C99不包含任何有关内联函数的 std，即 clcc 编译器要么只是忽略inline关键字并进行常规调用，要么它默默地支持内联。但没有强制要求支持该功能。

因此，最好将您的定义MultiplyMatrixVector为预处理器宏。尽管在您的情况下这不是主要问题。

2）您错误地威胁本地内存（LDM）。

虽然它的延迟时间小于global memory它正确访问时的延迟时间，但它local memory会受到银行冲突的影响。

您的顶点索引以每个工作项的步幅 100 计算。bank 的数量取决于所使用的 GPU，但通常为 16 或 32，即LDM如果所有变量都位于不同的 bank 中，则您可以在一个周期内访问多达 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)
于 2016-10-31T06:12:27.333 回答
-2

看起来工作组中的每个线程在计算开始之前都在复制相同的 50 个浮点数。这将使全局内存带宽饱和。

尝试这个

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

每个工作组只复制一次。

于 2012-08-05T13:44:36.583 回答