Dynamic Parallelism 이란?

케플러 이전의 GPU에서는 오직 CPU만이 디바이스에 작업을 할당할 수 있었다. 하지만 dynamic parallelism이 도입되면서 GPU의 작업 중인 스레드가 새로운 그리드를 생성하고 작업을 할당할 수 있게 되었다.

실행 중인 CUDA 커널은 새롭게 중첩된(nested) 작업을 생성하고 동기화 할 수 있다. 이 때 원래의 커널을 부모 커널(Parent CUDA Kernel)이라고 하고, 새롭게 생성된 커널을 자식 커널(Child CUDA Kernel)이라고 한다. 부모 커널은 자식 커널에 의해 생성된 결과를 사용할 수 있으며, 이 때 CPU가 개입될 필요가 없다. 또한 재귀 호출도 가능하다. 

 

dynamic parallelism을 이용한다면, 한 커널의 결과를 입력으로 하는 새로운 커널을 시작할 때 데이터를 제어하고 다시 전송하는 번거로움을 감소시킬 수 있다. 또한 재귀 표현식과 불규칙한 루프 구조를 제거하려고 수정해야만 했던 알고리즘을 표현할 수 있게 되었다. 

 

2번째 참고자료에는 dynamic parallelism을 설명하기 좋은 예제가 나와있다. 

각 행마다 같은 수의 원소를 갖는 데이터는 기존처럼 쉽게 병렬화할 수 있다. 단순하게 생각하면 M개의 스레드를 갖는 N개의 스레드 블록을 통해 수행할 수 있을 것이다.

 

반면, 위의 그림과 같이 각 행마다 연산할 데이터의 수가 다른 경우에는 N개의 스레드블록 마다 max(x[i])만큼의 스레드를 초과할당(oversubsription)하거나 순차작업을 통해 처리할 수밖에 없다. 하지만 이 작업을 dynamic parallelism을 통해 처리한다면, 연산 자원의 낭비나 병렬성의 포기를 방지할 수 있다.

부모 커널을 호출할 때, 각 행의 데이터 수에 맞춰 블록 수와 스레드 수를 결정할 오직 하나의 스레드만 호출한다. 이 스레드가 각각 처리할 데이터 수에 맞도록 자식 커널을 호출하는 방식으로 동작하게 하는 것이다.

dynamic parallelism의 전체적인 개념도는 다음 그림과 같다. 

dynamic parallelism을 이용하면 CPU 또한 더이상 개별적인 커널 시작에 사용되지 않고, 다른 용도를 위해 release될 수 있다. 따라서 좀 더 쉽게 동시성을 향상시킬 수 있다. 

 

CUDA 샘플 코드에서는 CUDA Dynamic Parallelism의 예로서 cdp-가 붙은 1) cdpSimplePrint 와 2) cdpSimpleQuicksort 를 제공한다. 여기서는 cdpSimplePrint만 살펴보고 다음 포스트에서 cdpSimpleQuicksort를 살펴보기로 한다.

 

&

주어진 코드는 자기 자신을 dynamic parallelism을 이용하여 재귀호출하는 커널로서, 각 블록의 첫번째 스레드(threadIdx.x==0)가 블록에 대한 정보를 출력한다. 현재 블록이 몇번째 생성된 블록인지는 atomicAdd()에 의해 계산되고 있다.

블록의 모든 스레드는 현재의 깊이가 최대 깊이가 될 때까지 지정된 수의 블록과 스레드만큼 (여기서는 각 스레드 2개씩 가지는 2개의 스레드블록) 재귀호출한다. 최대 깊이가 2로 설정되었기 때문에, 호스트가 호출한 2개의 블록의 모든 스레드( 2*2 = 4 )는 각각 2개의 블록씩 재귀호출한다. 따라서 전체적으로 10개이 블록이, dynamic parallelism에 의해 8개의 블록이 생성되었음을 아래 출력에서 확인할 수 있다.

제시된 코드에서 알 수 있지만, /usr/local/cuda/samples/common/inc/helper_string.h 에 정의된 getCmdLineArgumentInt()를 이용해서 max_depth를 새로 설정할 수 있다. max_depth를 3으로 설정했을 때의 결과는 다음과 같다.

# ./cdpSimplePrint depth=3

 

참고)

http://developer.download.nvidia.com/assets/cuda/files/CUDADownloads/TechBrief_Dynamic_Parallelism_in_CUDA.pdf

http://on-demand.gputechconf.com/gtc/2012/presentations/S0338-GTC2012-CUDA-Programming-Model.pdf