Xeon Phi 上的 OpenCL：2D 卷积体验 - OpenCL 与 OpenMP答案

【问题标题】：OpenCL on Xeon Phi: 2D Convolution Experience - OpenCL vs OpenMPXeon Phi 上的 OpenCL：2D 卷积体验 - OpenCL 与 OpenMP
【发布时间】：2014-04-24 04:33:04
【问题描述】：

在 opnecl 中使用 2D 卷积进行基准测试的 Xeon Phi 的性能似乎比 openmp 实现要好得多，即使使用了启用编译器的矢量化也是如此。 Openmp 版本在 phi 原生模式下运行，时间仅测量计算部分：For-loop。对于 opencl 实现，计时也仅用于内核计算：不包括数据传输。 OpenMp-enbaled 版本使用 2,4,60,120,240 个线程进行了测试。 - 240 个线程为平衡线程关联设置提供了最佳性能。但是，即使对于 240 线程的 openmp 基线，Opencl 的性能也提高了大约 17 倍，其中启用了编译指示的矢量化是源代码。输入图像尺寸为 1024x1024 到 16384x16384，过滤器尺寸为 3x3 到 17x17。在调用运行中，opencl 优于 openmp。这是opencl的预期加速吗？好得令人难以置信。

编辑：

编译（openmp）

icc Convolve.cpp -fopenmp -mmic -O3 -vec-report1 -o conv.mic
Convolve.cpp(71): (col. 17) remark: LOOP WAS VECTORIZED

来源（Convole.cpp）：

void Convolution_Threaded(float * pInput, float * pFilter, float * pOutput,
          const int nInWidth, const int nWidth, const int nHeight,
          const int nFilterWidth, const int nNumThreads)
{
    #pragma omp parallel for num_threads(nNumThreads)
    for (int yOut = 0; yOut < nHeight; yOut++)
    {
        const int yInTopLeft = yOut;

        for (int xOut = 0; xOut < nWidth; xOut++)
        {
            const int xInTopLeft = xOut;

            float sum = 0;
            for (int r = 0; r < nFilterWidth; r++)
            {
                const int idxFtmp = r * nFilterWidth;

                const int yIn = yInTopLeft + r;
                const int idxIntmp = yIn * nInWidth + xInTopLeft;

                #pragma ivdep           //discards any data dependencies assumed by compiler                                        
                #pragma vector aligned      //all data accessed in the loop is properly aligned
                for (int c = 0; c < nFilterWidth; c++)
                {
                    const int idxF  = idxFtmp  + c;
                    const int idxIn = idxIntmp + c;    
                    sum += pFilter[idxF]*pInput[idxIn];
                }
            } 

            const int idxOut = yOut * nWidth + xOut;
            pOutput[idxOut] = sum;
        } 
    } 
}

源 2 (convolve.cl)

    __kernel void Convolve(const __global  float * pInput,
                        __constant float * pFilter,
                        __global  float * pOutput,
                        const int nInWidth,
                        const int nFilterWidth)
{
    const int nWidth = get_global_size(0);

    const int xOut = get_global_id(0);
    const int yOut = get_global_id(1);

    const int xInTopLeft = xOut;
    const int yInTopLeft = yOut;

    float sum = 0;
    for (int r = 0; r < nFilterWidth; r++)
    {
        const int idxFtmp = r * nFilterWidth;

        const int yIn = yInTopLeft + r;
        const int idxIntmp = yIn * nInWidth + xInTopLeft;

        for (int c = 0; c < nFilterWidth; c++)
        {
            const int idxF  = idxFtmp  + c;
            const int idxIn = idxIntmp + c;
            sum += pFilter[idxF]*pInput[idxIn];
        }
    }
    const int idxOut = yOut * nWidth + xOut;
    pOutput[idxOut] = sum;
}

OpenMP 的结果（与 OpenCL 相比）：

            image filter  exec Time (ms)
OpenMP  2048x2048   3x3   23.4
OpenCL  2048x2048   3x3   1.04*

*原始内核执行时间。不包括通过 PCI 总线的数据传输时间。

【问题讨论】：

您正在使用两组不同的代码，它们可能已进行了不同的优化。你是自己写的 OpenMP 吗？贴出代码。它可能没有优化。 OpenCL 是否使用矢量类型（例如 float4）？这些将使用 SSE/AVX。 OpenMP 只处理线程，如果你想使用 SSE/AVX，你必须自己做。
此外，Xeon Phi 有自己的 SIMD (AVX512)，它是 512 位宽，OpenCL 可以利用它。 OpenMP 不会为您执行此操作。
英特尔很可能对某些内核执行水平向量化，这意味着单个 Xeon Phi 内核实际上可以在单个内核上同时运行 16 个线程（假设一个线程在 32 位 vals 上运行），每个矢量元素。这自然取决于内核，并非所有都可以轻松修改为这种处理。
opencl 和 openmp 基准测试是来自同一来源的相同代码：AMD。 opencl 代码不是手动分解的——纯标量代码。该框架进行矢量化。对于 openmp 编译，编译器报告所有循环都已矢量化。 (--vec-report2)。即使二进制文件包含编译器报告的向量指令，代码是否仍有可能不使用 SSE/AVX？
是的，Opencl 确实很容易利用 SIMD 单元。你如何确保 openmp 做到这一点（使用编译指示在我的代码中指定的编译器提示启用）？我正在将代码添加到 EDIT 帖子中。

标签： opencl openmp convolution xeon-phi

【解决方案1】：

以前：（#pragma ivdep 和 #pragma vector aligned 用于最内层循环）：

Compiler output: 
Convolve.cpp(24): (col. 17) remark: LOOP WAS VECTORIZED

Program output:
120 Cores: 0.0087 ms

根据@jprice 的建议（在水平数据上使用#pragma simd）：

Compiler output:
Convolve.cpp(24): (col. 9) remark: **SIMD** LOOP WAS VECTORIZED

Program output:
120 Cores: 0.00305

OpenMP 现在2.8X 比之前的执行速度更快。现在可以与 OpenCL 进行公平的比较！感谢 jprice 和所有做出贡献的人。从你们那里学到了很多教训。

编辑：这是我的结果和比较：

            image   filter  exec Time (ms)
OpenMP  2048x2048   3x3     4.3
OpenCL  2048x2048   3x3     1.04

Speedup: 4.1X

OpenCL 真的可以比 OpenMP 快吗？

【讨论】：

您也可以尝试在 x 和 y 外部循环上并行化。对于非常大的图像，这无关紧要，但对于 240 个线程上的 1024x1024，它可能会。您必须对编译器进行一项微不足道的更改才能接受这一点：#pragma omp parallel for collapse(2) num_threads(nNumThreads) for (int yOut = 0; yOut

【解决方案2】：

英特尔的 OpenCL 实施将使用他们所谓的“隐式矢量化”来利用矢量浮点单元。这涉及将工作项映射到 SIMD 通道。在您的示例中，每个工作项正在处理一个像素，这意味着每个硬件线程将使用 Xeon Phi 的 512 位矢量单元一次处理 16 个像素。

相比之下，您的 OpenMP 代码是跨像素并行化，然后在一个像素内对计算进行矢量化处理。几乎可以肯定，这就是性能差异的来源。

为了让 ICC 以类似于隐式矢量化 OpenCL 代码的方式矢量化您的 OpenMP 代码，您应该从最内层循环中删除您的 #pragma ivdep 和 #pragma vector aligned 语句，而只需放置一个 @987654323 @在水平像素循环前面：

#pragma omp parallel for num_threads(nNumThreads)
for (int yOut = 0; yOut < nHeight; yOut++)
{
    const int yInTopLeft = yOut;

    #pragma simd
    for (int xOut = 0; xOut < nWidth; xOut++)
    {

当我用 ICC 编译它时，它报告它正在成功地矢量化所需的循环。

【讨论】：

【解决方案3】：

您的 OpenMP 程序对一行图像使用一个线程。同一行中的像素被矢量化。它等于您在 OpenCL 中拥有一维工作组。每个工作组处理一行图像。但是在您的 OpenCL 代码中，您似乎有一个二维工作组。每个工作组（映射到 phi 上的一个线程）正在处理图像的块，而不是图像的行。缓存命中会有所不同。

【讨论】：