c - 为什么面向 OpenCL 的英特尔内核构建器告诉我我的内核未矢量化？

Question

我要写一个内核来在有限区域内添加两个 3 维矩阵。我有我的代码

#define PREC float

typedef struct _clParameter clParameter;
struct _clParameter {
    size_t width;
    size_t minWidth;
    size_t maxWidth;
    size_t height;
    size_t minHeight;
    size_t maxHeight;
    size_t depth;
    size_t minDepth;
    size_t maxDepth;
};

__kernel void clMatrixBasicOperate1Add(
    __global const PREC * restrict in1,
    __global const PREC * restrict in2,
    __global PREC * restrict out,
    __private const clParameter par) {

    size_t sizeOfXY = par.width * par.height;

    // 3-Dimension matrix

    size_t X = get_global_size(0);
    size_t x = get_global_id(0);

    size_t Y = get_global_size(1);
    size_t y = get_global_id(1);

    size_t Z = get_global_size(2);
    size_t z = get_global_id(2);

    size_t endX = (par.maxWidth - par.minWidth + 1)     / X;
    size_t endY = (par.maxHeight - par.minHeight + 1)   / Y;
    size_t endZ = (par.maxDepth - par.minDepth + 1)     / Z;

    if(x<( (par.maxWidth    - par.minWidth  + 1) % X) )     endX += 1;
    if(y<( (par.maxHeight   - par.minHeight + 1) % Y) )     endY += 1;
    if(z<( (par.maxDepth    - par.minDepth  + 1) % Z) )     endZ += 1;

    for(size_t k=0;k<endZ;k++)
    for(size_t j=0;j<endY;j++)
    for(size_t i=0;i<endX;i++) {
        size_t index = (par.minDepth + k*Z+z) * sizeOfXY + (par.minHeight + j*Y+y) * par.width + (par.  minWidth + i*X +x);
        out[index] = in1[index] + in2[index];
    }

    // return
}

当我使用 Intel Kernel Builder For OpenCL API 构建它时，它告诉我

Setting target instruction set architecture to: Default (Advanced Vector Extension (AVX))
OpenCL Intel CPU device was found!
Device name: Intel(R) Core(TM) i7-2630QM CPU @ 2.00GHz
Device version: OpenCL 1.2 (Build 83073)
Device vendor: Intel(R) Corporation
Device profile: FULL_PROFILE
Compilation started
Compilation done
Linking started
Linking done
Device build started
Device build done
Kernel <clMatrixBasicOperate1Add> was not vectorized
Done.
Build succeeded!

我想知道为什么 clMatrixBasicOperate1Add 没有向量化。

Answer 1

您的内核无法矢量化，部分原因是 for 循环中的终止条件。这些条件都依赖于根据内核输入计算的变量。因此，在内核编译时，英特尔 OpenCL C 编译器不知道这些循环将执行多少次迭代，因此根本无法优化它们。如果您将内部循环从 for(size_t i=0;i<endX;i++) 更改为至 for(size_t i=0;i<4;i++)然后内核被矢量化。当然，此更改不会执行您想要的操作，但至少您的内核已矢量化 :)。

我认为您想要尝试的策略是沿线程网格的 X 维度进行矢量化。这意味着您将沿着 X 启动 1/2 数量的线程，而是使用 vload2 和 vstore2 函数来读取和写入全局内存。您也可以使用 4、8 或 16 个元素 vector ，在这种情况下，您将分别沿 X 维度启动当前线程数的 1/4、1/8 或 1/16。

由于您使用的是第二代 Core i7 和 float 据，您可能希望使用 float8、vload8 和 vstore8，因为您的 CPU 支持同时对 8 个浮点值进行操作的 AVX 指令。请注意，这不是性能可移植的，例如一些 GPU 在 float2 下运行良好，但在使用 float4/8/16 时性能下降。使用 AMD CPU 运行时的旧 CPU 无法访问 AVX 指令，只能访问使用 4 元素浮点 vector 的 SSE。因此，您应该使用诸如“-D vectype=float4”之类的字符串，通过在 clBuildProgram 的选项中传递的宏，使 vector 大小成为可调参数。