OpenCL on Xeon Phi : 2D 컨볼 루션 경험 - OpenCL 대 OpenMP

Question

opnecl에서 2D 컨볼 루션으로 벤치 마크 된 Xeon Phi의 성능은 컴파일러 사용 벡터화에서도 openmp 구현보다 훨씬 우수합니다. Openmp 버전은 파이 네이티브 모드로 실행되었으며, 타이밍은 계산 부분 : For-loop 만 측정되었습니다. opencl 구현에서 타이밍은 커널 계산을 위해서만 사용되었으며 데이터 전송은 포함되지 않았습니다. OpenMp-enbaled 버전은 2,4,60,120,240 개의 스레드로 테스트되었습니다. - 240 개의 스레드가 균형 잡힌 스레드 선호도 설정을위한 최상의 성능을 제공했습니다. 그러나 Opencl은 240 스레드의 오픈 엠프베이스 라인에서도 pragma-enbled vectorization이 소스 코드로 17 배 정도 향상되었습니다. 입력 이미지 크기는 1024x1024에서 최대 16384x16384이고 필터 크기는 3x3에서 최대 17x17입니다. 호출이 실행될 때 opencl은 openmp보다 낫습니다. 이것은 opencl의 예상되는 속도인가 ?? 진실한 너무 좋은 것 같다.

EDIT :

편집 (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

인텔의 OpenCL 구현은 벡터 부동 소수점 단위를 활용하기 위해 "암시 적 벡터화"라고하는 것을 사용합니다. 여기에는 작업 항목을 SIMD 레인에 매핑하는 작업이 포함됩니다. 예에서 각 작업 항목은 단일 픽셀을 처리하고 있습니다. 즉, 각 하드웨어 스레드는 Xeon Phi의 512 비트 벡터 단위를 사용하여 한 번에 16 픽셀을 처리합니다.

반면 OpenMP 코드는 여러 픽셀에서 병렬 처리를 수행 한 다음 픽셀 내에서 계산을 벡터화합니다. 성능 차이가 발생하는 곳은 거의 확실합니다.

암시 벡터화 오픈 CL 코드와 유사한 방식으로 OpenMP의 코드를 벡터화 ICC을 취득하기 위하여는, 당신은 가장 안쪽의 루프에서 #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로 이것을 컴파일하면 원하는 루프가 성공적으로 벡터 라이 제이션되고 있음을 알립니다.

Answer 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를 초과 할 수 있습니다 : 여기 내 결과와 비교입니까?

Answer 3

OpenMP 프로그램은 이미지 행에 하나의 스레드를 사용합니다. 동일한 행의 픽셀이 벡터화됩니다. OpenCL에서 하나의 차원 작업 그룹과 같습니다. 각 작업 그룹은 한 행의 이미지를 처리합니다. 하지만 OpenCL 코드에서는 2 차원 작업 그룹이있는 것처럼 보입니다. 각 작업 그룹 (phi의 한 스레드로 매핑 됨)은 이미지의 행을 처리하는 것이 아니라 이미지의 BLOCK을 처리하는 것입니다. 캐시 적중은 다릅니다.