dynamic parallelism을 활용한 quick sort

Dynamic Parallelism의 간단한 개념 설명)     https://hayunjong83.tistory.com/22

 

CUDA 샘플 코드에서는 dynamic parallelism을 설명하기 위해, quick sort 코드를 제공하고 있다.

depth가 너무 커지거나 일정 임계값 이하의 원소를 정렬할 때는 선택정렬(selection sort)를 사용한다.

 

 

퀵 정렬 알고리즘 자체는 기존의 흔한 구현과 다르지 않다. 피벗을 기준으로 왼쪽과 오른쪽을 다시 정렬할 스레드를 선언하여 재귀적으로 다시 커널을 호출한다. 이 때 dynamic parallelism을 이용하므로, 중간 결과를 다시 호스트로 memcpy할 필요가 없으며,  stream을 이용하므로 각 부분의 정렬이 다른 부분의 정렬과의 무의미한 동기화를 위해 대기하는 일을 피할 수 있다. 이러한 이유에서 아래 첫번째 링크의 블로그에서는 CPU가 제어하여 GPU에서 수행하는 퀵 정렬에 비해, dynamic parallelism이 2배 이상의 성능향상을 보인다고 기술하고 있다.

퀵 정렬 기본 예제는 부모 그리드가 자식 그리드를 만든 후에 추가적인 쓰기를 수행하지 않기 때문에, memory consistency에 관해 고려할 부분이 딱히 많지 않다. 그래서 이 부분은 cudaDeviceSynchronize()를 이용한 명시적인 동기화 선언과 함께 추후에 알아보기로 하고, 이 글에서는 스트림의 이용에 관하여 살펴보려 한다.

 

  블록 안의 스레드들이 여러 개의 자식 그리드를 생성하더라도, 동일한 스트림을 이용하면 자식 그리드 간의 실행은 순차적으로 일어날 수밖에 없다. 하지만 스트림 간의 실행은 non-blocking이기 때문에 여러 스트림을 이용한다면 동시성을 향상시킬 수 있다. 

( 물론 다른 스트림을 사용한다고해서 GPU에서 동시 실행(concurrent execution)이 보장되는 것은 아니다. 단지 한 스트림의 데이터 전송이나 다른 스트림의 실행이 중첩되게 하는 것처럼 GPU 자원의 활용률을 높이는 것이 스트림 사용의 원래 목적이다.)

물론 서로 다른 스레드블록이 생성한 스트림은 기본 Null 스트림이더라도 서로 다른 스트림으로 고려된다. 하지만 위의 퀵 정렬은 하나의 스레드 블록의 한 스레드에서 커널을 시작하고, 상황에 맞게 두 번의 커널  호출이 일어나므로 서로 다른 스트림을 생성하는 것이 좋다. 커널을 호출할 때는 이 스트림을 지정해서 호출한다.

            cdp_simple_quicksort<<<1, 1, 0, s>>>

            (세 번째 인자는 shared_memory 크기로서 보통 0을 이용한다.)

참고로 CUDA 런타임에서는 스트림 안의 동기화를 제공하지는 않는다. 위의 퀵 정렬에서는 자식 그리드의 실행이 종료되고 나서 스트림을 없애기때문에 조심할 필요가 크게 없어보인다. 하지만 그렇지 않은 경우에는 cudaDeviceSynchronize()를 사용해서 먼저 스트림에서 실행되는 커널을 동기화해야지만, 나중에 cudaStreamDestroy()를 통해서 사용 중인 자원이 강제 해제되는 일을 막을 수 있다.

 

원래의 샘플 코드를 조금 수정해서, verbose 값이 true이면 check_results()에서 결과 배열을 출력하도록 만들었다.

 

참고)

https://parallelcomputingweb.wordpress.com/2017/04/12/now-quicksort-can-be-done-using-cuda/

https://devblogs.nvidia.com/cuda-dynamic-parallelism-api-principles/