3

在 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 总线的数据传输时间。

4

3 回答 3

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 快吗？

于 2014-04-24T18:54:21.170 回答
1

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

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

为了让 ICC 以类似于隐式矢量化 OpenCL 代码的方式矢量化您的 OpenMP 代码，您应该从最里面的循环中删除您的#pragma ivdepand#pragma vector aligned语句，而只需将 a#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 编译它时，它报告它正在成功地矢量化所需的循环。

于 2014-04-24T17:36:49.450 回答
1

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

于 2014-08-22T18:40:26.910 回答