首页 文章

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

提问于
浏览 次
3

在opnecl中使用2D卷积进行基准测试的Xeon Phi的性能似乎比使用编译器启用的矢量化的openmp实现要好得多 . Openmp版本以phi native模式运行，并且定时仅测量计算部分：For-loop . 对于opencl实现，时序也仅用于内核计算：不包括数据传输 . OpenMp-enbaled版本使用2,4,60,120,240个线程进行了测试 . - 240个线程为 balancer 线程亲和性设置提供了最佳性能 . 但Opencl大约是17倍，即使对于240线程的openmp基线，pragma-enbled矢量化也是源代码 . 输入图像尺寸为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总线上的数据传输时间 .

opencl openmp convolution xeon-phi

3 回答

  • 1

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

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

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

    #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编译它时，它报告它成功地向所需的循环进行矢量化 .

    回复于
  • 2

    以前:(对于内部最内部循环使用 #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更快吗？

    回复于
  • 1

    您的OpenMP程序使用一个线程作为一行图像 . 同一行中的像素是矢量化的 . 它等于你在OpenCL中有一个维度的工作组 . 每个工作组处理一行图像 . 但在您的OpenCL代码中，您似乎有一个二维工作组 . 每个工作组（映射到phi上的一个线程）正在处理图像的BLOCK，而不是图像的ROW . 缓存命中将是不同的 .

    回复于

相关问题