为什么 Intel Kernel Builder for OpenCL 告诉我我的内核没有矢量化？答案

【问题标题】：Why Intel Kernel Builder for OpenCL tell me that my kernel was not vectorized?为什么 Intel Kernel Builder for OpenCL 告诉我我的内核没有矢量化？
【发布时间】：2014-05-13 07:27:37
【问题描述】：

我打算编写一个内核来在有限的区域内添加两个 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 没有矢量化。

【问题讨论】：

好吧，因为它根本没有矢量化！根本不是错误，只是编译器信息。您输入的 OpenCL 内核是 scalalr 内核，即使编译器有能力将其转换为矢量化内核，它也没有。为什么？因为它决定不这样做，或者因为它不知道怎么做。
我看过一些资料说，如果几个工作项访问一个连续的内存空间，可以合并访问。但是，编译器并没有在这里合并我的操作。为什么？
“如果多个工作项访问一个连续的内存空间，则可以组合访问” - 但在您的情况下，工作项之间的访问不是连续的，因为每个工作项都有自己的循环。跨度>
是否意味着循环访问不能合并？
结合内存访问和内核向量化是两个不同的东西，仅在向量化内核需要更少的内存组合的意义上相关，因为它将进行更广泛的读取。查看英特尔文档，但很可能您的内核中的某些内容过于复杂而无法矢量化。

标签： c linux opencl intel

【解决方案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 个元素向量，在这种情况下，您将分别沿 X 维度启动当前线程数的 1/4、1/8 或 1/16。

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

【讨论】：