Xeon Phi 上的 OpenCL : 2D Convolution Experience - OpenCL vs OpenMP

Question

Xeon Phi 在 opnecl 中以 2D 卷积为基准的性能似乎比 openmp 实现要好得多，即使使用启用编译器的矢量化也是如此。 Openmp 版本以 phi native 模式运行，计时只测量计算部分:For 循环。对于 opencl 实现，计时也仅用于内核计算:不包括数据传输。 OpenMp-enbaled 版本使用 2,4,60,120,240 个线程进行了测试。 - 240 个线程为平衡的线程关联设置提供了最佳性能。但 Opencl 甚至对于 240 线程 openmp 基线和 pragma-enbled 矢量化是源代码也好 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 总线的数据传输时间。

Answer 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 快这么多？