왜이 간단한 OpenCL 커널이 너무 느리게 실행됩니까?

Question

OpenCL을 조사해 보니이 커널이 너무 천천히 실행되는 이유에 대해 조금 혼란 스럽습니다. 나는 임의의 값으로 포인터의 초기화 한 것을 제외하고,

char* out = new char[2048*2048]; 
cl::Buffer(
    context, 
    CL_MEM_USE_HOST_PTR | CL_MEM_WRITE_ONLY, 
    length, 
    out); 

저두 입력 버퍼 : 나는 다음과 같은 방법으로 버퍼를 만든

__kernel void copy(
    const __global char* pSrc, 
    __global __write_only char* pDst, 
    int length) 
{ 
    const int tid = get_global_id(0); 

    if(tid < length) { 
    pDst[tid] = pSrc[tid]; 
    } 
} 

: 여기에 커널입니다. 마지막으로, 커널이 방법을 실행 : 인텔 i5-3450과

cl_ulong startTime = event.getProfilingInfo<CL_PROFILING_COMMAND_START>(); 
cl_ulong endTime = event.getProfilingInfo<CL_PROFILING_COMMAND_END>(); 
std::cout << (endTime - startTime) * SECONDS_PER_NANO/SECONDS_PER_MILLI << "\n"; 

나는, 윈도우 7을 실행 해요 : 평균적으로

cl::Event event; 
queue.enqueueNDRangeKernel(
    kernel, 
    cl::NullRange, 
    cl::NDRange(length), 
    cl::NDRange(1), 
    NULL, 
    &event); 

event.wait(); 

을에 의해 계산 된 시간은 약 75 밀리 초 칩 (샌디 브릿지 아키텍처). 비교를 위해 복사를 수행하는 "직접적인"방법은 5 밀리 초 미만입니다. 나는 event.getProfilingInfo가 호스트와 장치 사이의 통신 시간을 포함한다고 생각하지 않는다. 생각?

는 편집 : ananthonline의 제안에서

, 내가 대신 문자의 float4s를 사용하도록 커널을 변경하고 약 50 밀리의 평균 실행 시간을 떨어졌다. 아직도 내가 기대했던 것만 큼 빠르지는 않지만 개선점. 감사합니다 ananthonline!

Answer 1

주요 문제는 사용중인 2048 * 2048 개의 작업 그룹이라고 생각합니다. 이 단일 항목 작업 그룹이 많은 경우 시스템의 opencl 드라이버가 훨씬 많은 오버 헤드를 관리해야합니다. 이것은 gpu를 사용하여이 프로그램을 실행하는 경우 특히 위험합니다. 왜냐하면 하드웨어의 채도가 매우 낮기 때문입니다.

최적화 : 더 큰 작업 그룹으로 커널을 호출하십시오. 기존 커널을 ​​변경할 필요조차 없습니다. see question: What should this size be? 예를 들어 아래 64 개를 사용했습니다. 64는 대부분의 하드웨어에서 적절한 숫자입니다.

cl::size_t myOptimalGroupSize = 64; 
cl::Event event; 
queue.enqueueNDRangeKernel(
    kernel, 
    cl::NullRange, 
    cl::NDRange(length), 
    cl::NDRange(myOptimalGroupSize), 
    NULL, 
    &event); 

event.wait(); 

또한 커널이 단일 값을 복사하는 것 이상을 수행해야합니다. 전역 메모리에 관한 비슷한 질문에 대한 대답을했습니다. over here.

Answer 2

CPU는 GPU와 매우 다릅니다. x86 CPU에서이를 실행하는 경우, 괜찮은 성능을 얻는 가장 좋은 방법은 char 또는 float4 대신에 double16 (가장 큰 데이터 형식)을 사용하는 것입니다 (다른 사람이 제안한 것처럼).

OpenCL을 사용한 CPU 사용 경험이 적 으면서도 OpenMP 병렬 처리로 얻을 수있는 성능 수준에 이르지 못했습니다. CPU와 병렬로 복사본을 만드는 가장 좋은 방법은 작은 수의 큰 하위 블록으로 복사하도록 블록을 나누고 각 스레드가 하위 블록을 복사하도록하는 것입니다. GPU 접근법은 직각입니다. 각 스레드는 동일한 블록의 복사본에 참여합니다. GPU에서는 서로 다른 스레드가 인접한 메모리 영역에 효율적으로 액세스 할 수 있기 때문에 (병합) 이러한 현상이 발생합니다.

OpenCL을 사용하여 CPU에서 효율적인 복사를 수행하려면 커널 내부 루프를 사용하여 인접한 데이터를 복사하십시오. 사용 가능한 코어 수보다 크지 않은 작업 그룹 크기를 사용하십시오.

Answer 3

나는 그것이 단일 항목 작업 그룹을 사용하도록 런타임에 알려주는 cl :: NDRange (1)라고 생각한다. 이것은 효율적이지 않습니다.C API에서이 값을 NULL로 설정하면 작업 그룹 크기를 런타임까지 남겨 둘 수 있습니다. C++ API에서도 그렇게 할 수있는 방법이 있어야합니다 (아마도 NULL 일 수도 있습니다). 이것은 CPU에서 더 빨라야합니다. 확실히 GPU에있을 것입니다.