처리 할 요소의 수가 무작위로 늘어날 때 고정 된 global_work_size 및 local_work_size를 유지하는 것은 좋지 않은 생각입니까?

Question

종종 처리해야하는 "요소"의 논리적 인 양과 동일하게 유지하는 것이 좋습니다. 내 응용 프로그램에는 그런 것이 없다. N 요소를 처리해야하는 경우 단일 커널을 통과 한 후 M 요소를 갖게됩니다. 즉, N에 의존하지 않는 완전히 다른 숫자입니다.

이 상황에 대처하기 위해, 나는 같은 루프를 작성할 수

while (elementsToBeProcessed) 
    read "elementsToBeProcessed" variable from device 
    enqueue ND range kernel with global_work_size = elemnetsToBeProcessed 

하지만 그 패스 당 하나 개의 읽기가 필요합니다. GPU 레이아웃과 일치하는 고정 된 global_work_size 및 local_work_size로 enqueueNDRangeKernel을 한 번만 호출 한 다음 모든 스레드를 사용하여 GPU 내부의 모든 작업을 동기화 할 수 있습니다.

제 질문은 간단합니다. 두 번째 옵션이 더 좋다고 내 직감이 맞습니까? 아니면 첫 번째 옵션과 함께 갈 이유가 있습니까?

Answer 1

글로벌 작업 크기는 고정 될 필요가 없습니다. E. g. 128 개의 스트림 프로세서가 있습니다. 따라서 로컬 크기 128의 커널을 만들 수 있습니다. 전역 작업 크기는 256, 4096 등 여러 값을 가질 수 있습니다.

일반적으로 로컬 그룹의 크기는 하드웨어 사양에 따라 결정됩니다. 처리 할 데이터가 더 많은 경우 관련 로컬 그룹의 수를 늘리십시오.

Answer 2

호스트와 장치간에 데이터 전송을 피할 수 있다면 장치에 대해 조금 더 많은 작업이 이루어 지더라도 항상 좋습니다. 많은 애플리케이션에서 데이터 전송은 가장 느린 부분입니다.

시스템 구성에 대한 더 나은 해결책을 찾으려면 두 가지를 모두 테스트해야합니다. 여러 플랫폼을 대상으로하는 경우 두 번째 플랫폼이 더 빠릅니다. 그러나 느리게 만들 수있는 많은 것들이 있습니다. 예를 들어 컴파일러를 최적화하기위한 코드가 더 어려울 수도 있고 데이터 액세스 패턴이 캐시 미스를 더 많이 초래할 수도 있습니다.

OpenCL 2.0을 타겟팅하는 경우 pipes은 이러한 종류의 임의의 요소를 조사하고자하는 내용 일 수 있습니다. (2.0을 지원하지 않는 플랫폼으로 인해 투표를하기 전에 AMD는 올해 2.0 드라이버를 발표했습니다.) 파이프를 사용하면 제작자 커널과 소비자 커널을 만들 수 있습니다. 소비자 커널은 작업 할 항목이있는 즉시 작업을 시작할 수 있습니다. 이렇게하면 모든 리소스를 더 잘 활용할 수 있습니다.

Answer 3

그건 까다로운 문제입니다. 그리고 글로벌 크기 값에 따라 달라지며 시간이 지남에 따라 달라집니다.

A는 패스 당 읽어 (에 더 높은 값 변경)

• 장착 글로벌 크기가 모든 작업 항목의 HW에 대한
• Unfitted 지역의 크기 유용한 작업을 할 것입니다 작업 크기 경우 작은
• 차단
• 쉽게 이해할 수있는 큐의 행동, 나쁜 장치 활용 및 디버그는

고정 커널 발사 크기 : (더 나은 안정하지만 값 변경)

• 취소가-장착 글로벌 크기, 장치에 널 (null) 작업 항목을
• 장착 지방의 크기를 실행 시간을 낭비 할 수
• 비 일부 답변이 이미 말한대로
단지
• 는

을 디버깅, 100 % 장치 사용 행위를 차단, 오픈 CL 2.0 솔루션입니다 , 파이프를 사용하여. 그러나 또 다른 OpenCL 2.0 기능인 커널 내부에서 커널 호출이 가능합니다. 그래서 여러분의 커널은 CPU 개입없이 커널의 다음 배치를 시작할 수 있습니다.

Answer 4

단점 : 리드 백을 수행 할 때 성능이 저하되는 이유는 GPU가 작업 대기 상태 일 때 유휴 상태가되는 반면, 여러 커널을 대기열에 추가하면 바쁠뿐입니다.

단순 : 답변은 elementsToBeProcessed가 얼마나 달라지는 지에 달려 있다고 생각합니다. 실행 순서가 (예를 들어) 20000, 19760, 15789, 19345 일 경우에는 항상 20000을 실행하고 유휴 작업 항목이 몇 개 있습니다. 반면에, 전형적인 패턴이 20000, 4236, 1234, 9000이면 나는 elementsToBeProcessed를 다시 읽고 필요한 것만 커널을 대기열에 넣습니다.

고급 : 패턴이 단조롭게 감소하면 커널 엔큐와 함께 리드 백을 인터리빙 할 수 있으므로 항상 GPU를 사용중인 상태로 유지하면서 이동 중에도 크기를 작게 만듭니다. 모든 커널 enqueue 사이에서 elementsToBeProcessed 사본의 비동기 이중 버퍼 읽기 백을 시작하고 다음에 커널을 대기열로 사용하려면 에 사용하십시오.이와 같이

:

1. elementsToBeProcessedA = 시작 값
2. elementsToBeProcessedB = 시작 값
3. eventA = NULL
4. eventB = NULL
elementsToBeProcessedA의 NDRange와
5. 엔큐 커널
6. 비에게 elementsToBeProcessedA에 대한 clEnqueueReadBuffer 차단, eventA 가져 오기
7. null이 아닌 경우,
8. 비 차단, eventA에서 대기, eventB
9. null이 아닌 경우를 고려 elementsToBeProcessedB에 대한 clEnqueueReadBuffer을 출시 이벤트
elementsToBeProcessedB의 NDRange로 출시 이벤트, eventB에
10. 인큐 커널을 기다립니다
11. goto 5

이렇게하면 GPU가 완전히 포화 상태가되지만 더 작은 요소를 사용하게됩니다. 처리 된대로 처리됩니다. 이 아닌은 elementsToBeProcessed 이으로 증가하는 경우 처리합니다. 따라서이 경우에는 이런 식으로 처리하지 마십시오.

Answer 5

다른 해결책 : 항상 GPU를 채울만큼 충분히 고정 된 수의 글로벌 작업 항목을 실행하십시오. 각 작업 항목은이 패스 (elementsToBeProcessed)에 대해 수행 할 총 항목 수를 확인한 다음 합계 부분을 수행해야합니다.

uint elementsToBeProcessed = <read from global memory> 
uint step = get_global_size(0); 

for (uint i = get_global_id(0); i < elementsToBeProcessed; i += step) 
{ 
    <process item "i"> 
} 

단순화 예 : 5 일 글로벌 크기 (예를 들면 작은 인위적) elementsToBeProcessed = 19 : 1 루프 요소 0-4를 통과하는 처리, 제 5-9 패스, 세 번째 패스 10-14 앞으로 패스 15-18.

고정 된 글로벌 작업 크기를 하드웨어와 정확히 일치하도록 조정할 수 있습니다 (계산 단위 * 최대 작업 그룹 크기 또는 그 일부분).

작업 그룹 크기에 관계없이 작업 항목이 협력하여 공유 로컬 메모리에 데이터를 복사하는 알고리즘과 다르지 않습니다.