3

opnecl の 2D 畳み込みでベンチマークされた Xeon Phi のパフォーマンスは、コンパイラー対応のベクトル化を使用した場合でも、openmp 実装よりもはるかに優れているようです。Openmp 版は phi ネイティブ モードで実行し、タイミングは計算部分 (For ループ) のみを測定しました。opencl の実装では、タイミングはカーネル計算のみであり、データ転送は含まれていませんでした。OpenMp 対応バージョンは、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

1

OpenMP プログラムは、画像の行に対して 1 つのスレッドを使用します。同じ行のピクセルはベクトル化されます。OpenCL に 1 つのディメンション ワークグループがあることと同じです。各ワークグループは、1 行の画像を処理します。しかし、あなたの OpenCL コードでは、2 次元のワークグループがあるようです。各ワークグループ (phi の 1 つのスレッドにマップされている) は、画像の ROW ではなく、画像のブロックを処理しています。キャッシュヒットは異なります。

于 2014-08-22T18:40:26.910 に答える
1

Intel の OpenCL 実装は、ベクトル浮動小数点ユニットを利用するために、「暗黙的なベクトル化」と呼ばれるものを使用します。これには、作業項目を SIMD レーンにマッピングすることが含まれます。あなたの例では、各作業項目は単一のピクセルを処理しています。つまり、各ハードウェア スレッドは、Xeon Phi の 512 ビット ベクトル ユニットを使用して一度に 16 ピクセルを処理します。

対照的に、OpenMP コードはピクセル間で並列化し、ピクセル内で計算をベクトル化しています。これはほぼ確実に、パフォーマンスの違いがどこから来ているかです。

暗黙的にベクトル化された OpenCL コードと同様の方法で ICC に OpenMP コードをベクトル化させるには、最も内側のループから#pragma ivdepand#pragma vector alignedステートメントを削除し、代わりに#pragma simd水平ピクセル ループの前にa を配置する必要があります。

#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 に答える