2013-11-24 13 views
1

저는 mpeg4dst 참조 오디오 인코더를 최적화하려고 시도하면서 OpenCL을 가르칩니다. CPU에서 벡터 명령을 사용하여 3 배의 속도 향상을 달성했지만 GPU가 더 잘 수행 할 수 있다고 생각했습니다.OpenCL slow - 확실하지 않은 이유

첫 번째 개선 영역으로 OpenCL에서 자동 상관 벡터를 계산하는 데 집중하고 있습니다. CPU 코드는 다음과 같습니다.

for (int i = 0; i < NrOfChannels; i++) { 
    for (int shift = 0; shift <= PredOrder[ChannelFilter[i]]; shift++) 
     vDSP_dotpr(Signal[i] + shift, 1, Signal[i], 1, &out, NrOfChannelBits - shift); 
} 
NrOfChannels = 6 
PredOrder = 129 
NrOfChannelBits = 150528. 

내 테스트 파일에서이 함수는 약 188ms 걸립니다. 약 54ms 메모리 복사 오버 헤드로,

gcl_memcpy(gpu_signal_in, Signal, sizeof(float) * NrOfChannels * MAXCHBITS); 

for (int i = 0; i < NrOfChannels; i++) { 
    size_t sz = PredOrder[ChannelFilter[i]] + 1; 
    cl_ndrange range = { 1, { 0, 0, 0 }, { sz, 0, 0}, { 0, 0, 0 } }; 

    calculateAutocorrelation_kernel(&range, i * MAXCHBITS, (cl_float *)gpu_signal_in, (cl_float *)gpu_out, NrOfChannelBits); 
    gcl_memcpy(out, gpu_out, sizeof(float) * sz); 
} 

악기에 따르면, 내 오픈 CL 구현이 약을 13ms을 보인다 :

kernel void calculateAutocorrelation(size_t offset, 
           global const float *input, 
           global float *output, 
           size_t size) { 
size_t index = get_global_id(0); 
size_t end = size - index; 
float sum = 0.0; 

for (size_t i = 0; i < end; i++) 
    sum += input[i + offset] * input[i + offset + index]; 

output[index] = sum; 
} 

이 그것이라고하는 방법입니다

여기 내 오픈 CL 방법입니다 (gcl_memcpy).

OpenCL 코드의 측정 된 성능이 같지만 2 채널 음악 대 1 분, 6 채널 1 초의 훨씬 더 큰 테스트 파일을 사용할 때 CPU 사용량은 약 50 전체 프로그램은 실행하는 데 약 2 배 더 오래 걸립니다.

인스트루먼트에서이 문제의 원인을 찾을 수 없으며 OpenCL에서 매우 많은 오버 헤드 스위칭이 필요하다고 제안하는 내용을 아직 읽지 않았습니다.

+0

시스템 사양은 무엇입니까? GPU? OS? 일반적으로 커널 실행 오버 헤드는 개별 GPU에서 상당히 큽니다. 그리고 당신이 얻을 수있는 최악의 실행 오버 헤드는 윈도우 비스타의 새로운 GPU (드라이버 모델로 인해)입니다. 또한 OpenCL 커널이 상대적으로 적은 스레드 (128 개의 스레드가 GPU에 대해 실제로 부족함)로 실행되는 것 같습니다. 이상적으로 문제는 수천 개의 별개의 스레드 (ndrange)를 가져야합니다. 그 메모리 복사의 꼭대기에는 PciE 버스보다 꽤 비싸다. 루프 반복마다 memcpy 대신 모든 커널 호출 후에 전체 결과를 복사하도록 작성하십시오. – sharpneli

+0

NVidia GT 650M 그래픽 카드가 장착 된 MacBook Pro의 OSX 10.9를 사용 중입니다. 두 번째 gcl_memcpy와 루프를 제거했습니다. cl_ndrange에 대해 배우고 GPU에서 더 많은 스레드를 얻으려고합니다. 그러나 내 벤치 마크에서 gcl_memcpy (양방향에서 PCIe 전송이어야 함)에서 커널과 50ms를 실행하는 데 10ms가 걸립니다. 이렇게하면 좋은 속도 향상, ~ 60ms ~ 190ms가 제공됩니다. 대신 CPU 사용률이 40 %에 불과하고 전체 프로세스가 이전에 수행했던 시간의 거의 두 배가 걸린 시간이 사라지는 것을 보았습니다. – Tim

답변

2

커널 코드를 올바르게 읽는다면, 각 작업 항목은 그 위치에서 끝까지의 모든 데이터를 반복합니다. 이것은 효율적이지 않을 것입니다. 하나 (그리고 주요 성능 문제)의 경우 메모리 액세스가 병합되지 않으므로 전체 메모리 대역폭이되지 않습니다. 둘째, 각 작업 항목의 작업량이 다르기 때문에 작업 그룹 내에서 분기 차이가 발생하여 일부 스레드가 다른 스레드를 기다리고있을 수 있습니다.

이것은 감소 문제와 공통점이 많은 것처럼 보이며 병렬 처리와 관련하여 몇 가지 힌트를 얻으려는 "병렬 감소"에 대해 읽어 볼 것을 제안합니다.

메모리를 읽는 방법을 보려면 16 개의 작업 항목 (예 : global_id 0 ~ 15)이 각 단계의 데이터를 읽는 방법을 알아야합니다.

작업 그룹의 모든 작업 항목이 동일한 메모리에 액세스하는 경우 하드웨어에서 수행 할 수있는 "브로드 캐스트"최적화가 있습니다. 따라서 루프의 순서를 바꾸면 상황이 개선 될 수 있습니다.