OpenCL - 벡터화 vs 스레드 내 루프

Question

알려진 수의 스레드를 병렬 처리 (위대한)해야하지만 각 스레드의 내부 반복 수가 크게 다를 수있는 문제가 있습니다. . 내 마음 속에서, 이것은 더 나은이 같은 커널 방식을 할 수 있습니다 : 반드시 FULL 가능한 위해 실행됩니다

__kernel something(whatever) 
{ 
    unsigned int glIDx = get_global_id(0); 
    unsigned int glIDy = get_global_id(1); // max "unroll dimension" 

    if(glIDy_meets_condition) 
     do_something(); 
    else 
     dont_do_anything(); 

} 

:보다는 ID가 (0) 사전에 알려져있다

__kernel something(whatever) 
{ 
    unsigned int glIDx = get_global_id(0); 

    for(condition_from_whatever) 
    { 

    }//alternatively, do while 

} 

,

Killing OpenCL Kernels

나에 대한 특정 정보를 찾을 수 없습니다 : 등이 논의에 따라, 사전에 종료 할 수있는 방법으로 glIDy의 범위, 커널 내에서 동적 크기의 forloops/do-while 문에 대한 비용은 Nvidia 및 AMD SDK의 커널에서 어디에서나 볼 수 있습니다. 커널 내부 조건 분기가 비 주기적 일수록 성능이 어떻게 떨어지는지를 읽었던 것을 기억합니다.

실제 질문 : 제가 제안 제 1 방식보다 GPU 아키텍처에 대처하기 위해보다 효율적인 방법이

있습니까?

이 주제에 대한 일반 정보는 공개되어 있습니다.

감사합니다.

Answer 1

나는 그 질문에 부여 할 수있는 일반적인 대답이 있다고 생각하지 않습니다. 그것은 정말로 당신의 문제에 달려 있습니다. 루프

/else 문이 나 커널의 성능에 영향을주지 않을 수 있습니다 경우 :

그러나 여기에서이 주제에 대한 몇 가지 고려 사항입니다. 사실 성능 비용은 커널 수준이 아니라 작업 그룹 수준입니다. 작업 그룹은 하나 이상의 워프 (NVIDIA)/파면 (AMD)으로 구성됩니다. 이 warps (NVIDIA 용어는 그대로 유지할 것이지만 정확히 AMD와 동일합니다)는 잠금 단계에서 실행됩니다.

if else (또는 반복 횟수가 다른 for 루프)로 인해 워프 내에서 분기가 발생하면 실행이 일련 화됩니다.즉, 첫 번째 경로를 따르는이 워프 내의 스레드는 작업을 수행하고 다른 스레드는 유휴 상태가됩니다. 작업이 완료되면이 스레드는 유휴 상태가되고 다른 스레드는 작업을 시작합니다.

스레드를 장벽과 동기화해야하는 경우 이러한 문에 또 다른 문제가 발생합니다. 모든 스레드가 장벽에 부딪치지 않으면 정의되지 않은 동작을 보입니다.

이제 특정 문제에 따라 작업 그룹 내에서 분기가 발생하지 않도록 스레드를 그룹화 할 수 있습니다. 단, 작업 그룹간에 차이가있을 수 있습니다 (아니요. 거기에 영향을 미침).

워프가 32 개의 스레드와 64 개의 웨이브 프런트로 구성되어 있다는 것을 알고 있다면 (잘 모르겠지만 오래된 AMD GPU에서는 아닐 수도 있음) 잘 조직 된 작업 그룹의 크기를이 숫자의 배수로 만들 수 있습니다. 다른 문제도 고려해야하기 때문에 매우 간단합니다. 예를 들어 this question을보고 Chanakya.sun이 제공 한 답변을 봅니다 (어쩌면 더 많은 정보를 얻을 수 있습니다).

방금 ​​설명한대로 문제를 구성 할 수없는 경우에는 분기를 처리하는 CPU에 OpenCL을 사용하는 것이 좋습니다. 내가 잘 회상한다면 일반적으로 작업 그룹당 하나의 작업 항목을 갖게 될 것입니다. 이 경우 CPU에 대한 Intel 및 AMD의 설명서를 확인하는 것이 좋습니다. 또한 프로그래밍 할 때 OCL과 GPU를 사용하는 것과 CPU를 사용할 때의 차이점을 설명하는 Heterogeneous Computing with OpenCL 장을 매우 좋아합니다.

나는 this article도 좋아합니다. 주로 GPU (단순한 문제) 감소에 대한 성능 향상에 대한 토론이지만이 기사의 마지막 부분에서는 CPU 성능에 대해서도 살펴 봅니다.

마지막으로, @Oak이 제공하는 "내부 장치 스레드 대기열 지원"에 대한 귀하의 의견과 관련하여 실제로 동적 병렬 처리라고합니다. 이 기능을 사용하면 분명히 문제가 해결되지만 CUDA를 사용하더라도 3.5 이상의 기능을 갖춘 장치가 필요합니다. 따라서 Kepler GK104 아키텍처를 사용하는 NVIDIA GPU도 지원하지 않습니다 (기능 3.0). OCL의 경우, 동적 병렬 처리는 표준 버전 2.0의 일부입니다. (내가 아직까지 구현이 없다는 것을 알고있는 한).

Answer 2

for은 반복 사이에 거짓 의존성을 삽입하므로 더 많은 두 번째 버전을 좋아합니다. 내부 반복이 독립적 인 경우, 각각을 다른 작업 항목으로 보내고 OpenCL 구현이이를 실행하는 최선의 방법을 분류하도록하십시오.

두 가지주의 사항 : 반복의 평균은 반복의 최대 수보다 현저히 낮은 경우

• ,이 여분의 더미 작업 항목 가치가 없을 수도 있습니다.
• 작업 항목이 더 많으며 각 조건을 계산해야합니다. 조건을 계산하는 것이 복잡한 경우이 방법은 좋은 생각이 아닙니다.
  • 또는 인덱스를 x 차원으로 병합하고 모든 반복을 동일한 작업 그룹으로 그룹화 한 다음 작업 그룹 당 한 번만 조건을 계산하고 로컬 메모리 + 장벽을 사용하여 조건을 동기화 할 수 있습니다.