References
- https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html
- Programming Massively Parallel Processors
Contents
- Dynamic Parallelism Overview
- Memory Data Visibility
- Execution Environment
- Synchronization, Streams, and Events
CUDA Dynamic Parallelism(동적 병렬)은 CUDA 프로그래밍 모델의 확장이며, CUDA 커널이 새로운 커널을 launch함으로써 새로운 스레드 그리드를 만들 수 있게 해줍니다. 동적 병렬은 Kepler 아키텍처에서 도입되었고, GK110 칩에서 처음 선보였습니다. 과거 CUDA 시스템, 커널은 오직 Host Code에서만 커널을 실행할 수 있었습니다. 재귀나, 불규칙한 루프 구조, 시공간 변화 등 단일 레벨의 병렬화에 적합하지 않은 이러한 알고리즘은 여러 커널이 실행되어야 했고, 이는 호스트의 작업, 호스트-디바이스 간 통신, 총 실행 시간을 증가시켰습니다.
동적 병렬 지원은 동적으로 새로운 작업들을 발견하는 알고리즘들이 host에 부담을 주지 않으면서 커널을 준비하고 실행할 수 있도록 해줍니다. 이번 포스팅에서는 CUDA 프로그래밍 모델에서 동적 병렬을 가능하게 하는 확장된 CUDA 아키텍처 기능에 대해 알아보겠습니다.
(조금 이해하기 어려운 부분들이 있어 어색한 해석과 개인 의견이 포함되어 있으니 혹시 잘못되거나 다른 의견이 있으시면 언제든지 댓글로 지적해주세요. 지적은 언제나 환영입니다.. !)
Dynamic Parallelism은 compute capability 3.5 이상의 디바이스에서만 지원합니다.
Background
실제 세상의 많은 어플리케이션들은 공간 전체에서 다양한 작업량을 가지거나 시간에 따라 동적으로 변화하는 작업량을 갖는 알고리즘을 사용합니다.
Graph Search (Breadth-First Search)
Graph Search (Breadth-First Search)
References Programming Massively Parallel Processors Contents Breadth-First Search (BFS) A Sequential BFS Function A Parallel BFS Function Optimization 이번에 살펴 볼 병렬 패턴은 Graph Search입니다..
junstar92.tistory.com
이전 포스팅에서 살펴본 그래프 탐색에서 봤듯이, 각 정점들을 처리할 때 수행되는 작업의 양은 소셜 네트워크와 같은 그래프에서처럼 다양합니다.
또 다른 예시로, 아래 그림은 요구되는 모델링 세부사항의 수준이 시공간에 따라 변화하는 난류(turbulence) 시뮬레이션을 보여줍니다.
combustion flow가 왼쪽에서 오른쪽으로 움직일 수록, 활동량과 강도의 수준은 증가합니다. 모델의 오른쪽 부분에서 요구되는 세부사항의 수준은 모델의 왼쪽보다 훨씬 높습니다.
고정된 미세한 그리드를 사용하는 것은 모델의 왼쪽 부분에서는 도움이 되지 않는 쓸데없는 작업을 너무 많이 수행합니다. 반면 고정된 더 큰 그리드를 사용하면 모델 오른쪽 부분에 대한 정확도가 낮아질 수 있습니다.
모델의 세부사항이 필요한 부분에는 미세한 그리드를 사용하고 그렇지 않은 부분은 큰 그리드를 사용하는 것이 이상적입니다.
이전 CUDA 시스템에서는 모든 커널을 호스트 코드에서 실행해야 합니다. 스레드 그리드에 의한 작업의 양은 커널이 실행되는 동안 미리 결정됩니다. 커널 코드를 SPMD 프로그래밍 스타일로 다른 그리드 간격을 사용하는 스레드 블록을 갖는 것이 매우 어렵지는 않지만 단조롭습니다. 이러한 제한은 고정 그리드 시스템에서 유리합니다. 원하는 정확도를 달성하기 위해서 Fig.1의 오른쪽 위의 모델에서 설명하는 이러한 고정 그리드 접근 방법은 불필요한 추가 작업을 수행해야합니다.
보다 바람직한 접근 방법은 Fig.1의 오른쪽 아래의 dynamic grid를 사용하는 것입니다. 시뮬레이션 알고리즘은 모델의 일부 영역에서 빠르게 변화하는 시뮬레이션 양을 감지하기 때문에 해당 영역의 그리드를 조정하여 원하는 수준의 정확도를 달성합니다. 활동량이 적은 부분에 대해서는 이러한 개선이 이루어질 필요가 없습니다. 이러한 방법으로 알고리즘은 동적으로 연산이 더 많이 필요한 부분에 그리드를 더 할당할 수 있습니다.
아래 그림은 Fig.1의 시뮬레이션 모델에 대해 동적 병렬이 없는 시스템과 동적 병렬이 있는 시스템 사이의 동작을 차이를 보여줍니다.
동적 병렬이 없다면 호스트 코드에서 모든 커널을 시작해야 합니다. 커널을 실행하는 동안 모델 영역의 그리드를 개선할 필요가 있다면, 커널은 스스로를 종료하고 호스트 코드로 피드백을 주어 새로운 커널을 시작하도록 합니다. 이는 위 그림의 (A)에 해당하고, 호스트는 커널을 실행하고, 커널로 부터 정보를 받습니다. 이 커널이 종료된 후에, 완료된 커널로부터 발견된 모든 새로운 작업에 대해 다음 커널을 시작합니다.
Fig.2의 (B)는 동적 병렬 처리를 통해 새로운 작업을 발견한 스레드가 스스로 그 작업을 수행하기 위해 커널을 실행할 수 있다는 것을 보여줍니다. 위 예시에서 스레드가 모델의 영역을 세분화할 필요가 있다는 것을 발견하면, 커널을 종료하고 호스트에 전달하고 호스트에서 새로운 커널을 시작하는 오버헤드없이, 스스로 새로운 커널을 실행할 수 있습니다.
Dynamic Parallelism Overview
프로그래머 관점에서 동적 병렬은 커널 내에서 새로운 커널 실행문을 작성하는 것을 의미합니다.
위 그림에서 main 함수(host code)는 A, B, C라는 3개의 커널을 실행합니다. 이는 기존 CUDA 모델에서 커널 실행입니다. 한 가지 다른 점은 커널 B에서 X, Y, Z라는 3개의 커널을 실행합니다. 이는 이전 CUDA 시스템에서 불가능한 것이었습니다.
커널을 실행하기 위한 문법은 host code에서 커널을 실행하는 것과 동일합니다.
kernel_name<<< Dg, Db, Ns, S>>>([kernel arguments])- Dg는 dim3 타입이며, 그리드의 차원과 크기를 지정합니다.
- Db 또한 dim3 타입이며, 각 스레드 블록의 차원과 크기를 지정합니다.
- Ns는 size_t 타입이며, 각 스레드 블록에 동적으로 할당되는 공유 메모리 크기(byte 단위)입니다. 이는 정적으로 할당된 메모리에 추가되며, Ns의 기본값은 0입니다.
- S는 cudaStream_t 타입이며 커널 호출과 연관된 스트림을 지정합니다. 스트림은 호출이 생성되는 곳에 위치하는 동일한 스레드 블록에서 할당되어야 합니다. S는 optional 인자이며, 기본값은 0(Null Stream)입니다.
A Simple Example
두 가지 스타일의 예제 코드를 살펴보겠습니다. 하나는 기존의 CUDA 스타일이고, 다른 하나는 동적 병렬 스타일입니다.
아래 코드는 동적 병렬이 없는 간단한 예제 커널을 보여줍니다.
__global__ void kernel(unsigned int* start, unsigned int* end, float* someData, float* moreData)
{
unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
doSomeWork(someData[i]);
for (unsigned int j = start[i]; j < end[i]; ++j)
doMoreWork(moreData[j]);
}커널의 각 스레드는 someData에 대해 어떠한 연산(line 4)를 수행하고, for문을 반복하면서 moreData의 데이터 요소의 연산을 수행(line 6-7)합니다.
여기에는 두 가지 문제점이 있습니다.
첫 번째는 커널 함수의 for 루프는 병렬로 수행될 수 있지만, 어플리케이션에서 병렬화를 더 이끌어낼 수 있는 기회를 놓치고 있습니다.
두 번째 문제점은 만약 동일한 워프(warp) 내에서 각 스레드 간의 루프 범위가 상당히 다르다면, 이는 control divergence를 유발하고 결과적으로 프로그램의 성능을 감소시킨다는 것입니다.
아래 코드는 동적 병렬을 사용하여 위의 커널과 동일한 역할을 수행하는 커널 함수입니다. 위에서 살펴본 커널 함수가 parent 커널과 child 커널로 분리되었습니다.
__global__ void kernel_parent(unsigned int* start, unsigned int* end, float* someData, float* moreData)
{
unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
doSomeWork(someData[i]);
kernel_child<<<ceil((end[i]-start[i])/256.0), 256>>>(start[i], end[i], moreData);
}
__global__ void kernel_child(unsigned int start, unsigned int end, float* moreData)
{
unsigned int j = start + blockIdx.x*blockDim.x + threadIdx.x;
if (j < end)
doMoreWork(moreData[j]);
}parent kernel의 시작은 원래 커널과 동일합니다. 대신 루프를 반복하지 않고 child 커널을 실행합니다. child 커널은 child 그리드에서 호출된 다른 스레드들에 의해서 수행되며, child 커널에서 수행하는 작업은 원래 커널의 for 루프의 body와 동일합니다.
동적 병렬을 사용한 프로그램을 작성하면 원래 코드에 언급된 두 가지 문제를 모두 해결할 수 있습니다.
첫째, 루프 반복은 원래 커널 스레드에 의해 순차적으로 수행되는 대신 child 커널 스레드에 의해서 병렬로 수행됩니다. 따라서, 프로그램에서 더 많은 병렬 처리를 이끌어냈습니다.
둘째, 각 스레드는 이제 단일 루프 반복을 실행하고 로드 밸런스(load balance)를 개선하고 control divergence를 제거합니다.
이러한 두 가지 목표는 커널을 조금 다르게 작성하여 달성할 수 있지만, 이러한 변환은 어색하고 복잡하며 오류가 발생하기 쉽습니다. 하지만 동적 병렬은 이러한 연산 패턴을 쉽게 표현할 수 있도록 해줍니다.
아래 이어지는 내용은 동적 병렬 처리를 사용하는 프로그램의 동작을 좌우하는 몇 가지 중요한 세부 사항을 간략하게 설명합니다.
Memory Data Visibility
먼저 메모리 데이터의 visibility(가시성)에 대한 규칙을 살펴보겠습니다. 이러한 규칙은 parent 그리드의 데이터 객체가 child 그리드의 스레드에서 액세스할 수 있는 방법을 알려줍니다. 이 규칙들은 non-dynamic parallelism 프로그램에서 동일한 그리드의 스레드 vs 다른 그리드의 스레드 간 데이터 일관성(data consistency) 규칙의 확장입니다.
예를 들어, 한 그리드 내의 스레드들에 의해 쓰여진 전역 메모리 데이터는 명시적인 메모리 동기화나 커널이 종료될 때까지 다른 스레드가 볼 수 있다는 것을 보장하지 않습니다.
Global Memory
Parent 스레드와 이 스레드의 child 그리드는 서로가 전역 메모리 데이터를 볼 수 있는데, child와 parent 간의 약간 일관성을 보장합니다. parent 스레드와 child 그리드에서의 메모리 view는 만약 각자에게 서로의 메모리 연산의 효과가 보이는 경우에 일치한다고 말할 수 있습니다.
여기에 parent 스레드와 메모리 view가 일치할 때, child 그리드 실행에서 두 가지 포인트가 있습니다.
- parent 스레드에 의해서 child 그리드가 생성될 때 입니다. 이는 child 그리드 생성 이전에 발생하는 parent 스레드에서의 모든 전역 메모리 연산은 child 그리드가 볼 수 있다는 의미입니다.
- parent 스레드에서 동기화 API 호출이 완료됨으로써 child 그리드가 완료될 때입니다. 이는 child 그리드에서의 모든 메모리 연산이 parent에서 child 그리드의 완료에 대해 동기화된 이후에 parent가 볼 수 있다는 의미입니다.
아래 예제 코드를 통해 자세히 살펴보겠습니다.
__global__ void child_launch(int *data) {
data[threadIdx.x] = data[threadIdx.x]+1;
}
__global__ void parent_launch(int *data) {
data[threadIdx.x] = threadIdx.x;
__syncthreads();
if (threadIdx.x == 0) {
child_launch<<< 1, 256 >>>(data);
cudaDeviceSynchronize();
}
__syncthreads();
}
void host_launch(int *data) {
parent_launch<<< 1, 256 >>>(data);
}child_launch를 수행하는 child 그리드는 오직 child 그리드가 실행되기 전의 'data' 수정을 확인할 수 있도록 보장됩니다. parent의 스레드 0이 하위 커널을 수행하기 때문에 child는 parent의 스레드 0에서 보는 메모리와 일치합니다. 첫 번째 __syncthreads() 호출로 인해 child는 data[0]=0, data[1]=1, ..., data[255]=255로 수정된 것을 볼 수 있습니다. 만약 첫 번째 __syncthreads()가 없다면 오직 data[0]만이 child에서 확인할 수 있도록 보장됩니다. child 그리드가 반환될 때, 스레드 0은 child 그리드의 스레드들에 의한 data 변경을 확인할 수 있도록 보장됩니다. 다른 스레드들이 이러한 변경을 확인할 수 있다고는 보장할 수 없으며, 두 번째 __syncthreads() 호출 이후에서야 parent의 다른 스레드들도 변경된 data를 확인할 수 있습니다.
Zero-Copy Memory
Zero-copy system memory는 전역 메모리와 동일한 일관성을 보장하며, 위에서 설명한 것과 같은 규칙을 따릅니다. 그러나 커널은 zero-copy 메모리를 할당하거나 해제할 수 없지만 host code로부터 전달받은 포인터를 사용할 수 있습니다.
Constant Memory
상수는 커널에 의해서 쓰여질 수 없으며, 동적 병렬 커널 실행 간에도 불가능합니다. 즉, 모든 __constant__ 변수의 값은 처음 커널이 실행되기 전에 host에서 설정되어야 합니다. 상수 메모리 변수는 모든 커널에서 전역으로 접근할 수 있습니다. 스레드 내에서 상수 메모리 객체의 주소를 사용하는 것은 non-dynamic-parallelism과 동일하고, parent에서 child로 포인터를 전달하는 것과 반대의 경우도 지원합니다.
Local Memory
지역 메모리(local memory)는 스레드에 private한 저장 공간이며, 그 스레드 외에는 볼 수 없습니다. child 커널을 실행할 때 커널 argument로 지역 메모리의 포인터를 전달하는 것은 illegal 입니다. child로부터 지역 메모리 포인터를 역참조하는 동작의 결과는 정의되어 있지 않습니다. 예를 들어, 아래의 child_launch 커널을 실행하는 스레드가 x_array에 액세스하는 것은 정의되지 않은 동작을 발생시키며, illegal 합니다.
int x_array[10]; // Creates x_array in parent's local memory
child_launch<<<1, 1>>>(x_array);
컴파일러에 의해 변수가 지역 메모리에 배치되는 시기를 프로그래머가 인지하는 것은 어렵습니다. 일반적으로 child 커널에 전달되는 모든 storage는 global-memory Heap에서 cudaMalloc() 또는 new()를 사용하거나 전역 scope에서 __device__ storage를 선언하여 명시적으로 할당해야 합니다.
예시로, 아래의 (A)는 전역 메모리 변수에 대한 포인터를 child 커널의 argument로 전달할 때, 유효한 커널 실행을 보여줍니다. 반면, (B)는 지역 메모리(auto) 변수에 대한 포인터를 child 커널로 전달하는 유효하지 않은 코드입니다.
NVIDIA CUDA C 컴파일러는 만약 지역 메모리의 포인터가 커널의 argument로 전달되는 것을 감지하면 경고를 출력합니다. 그러나 이러한 감지가 보장되지는 않습니다.
Shared Memory
공유 메모리(Shared Memory)는 실행 중인 스레드 블록에 private한 storage이며, 데이터는 해당 스레드 블록 외부에서 볼 수 없습니다. 공유 메모리 변수의 포인터를 메모리를 통해서나 argument로 child 커널에 전달하는 것은 정의되지 않은 동작입니다.
NVIDIA 컴파일러는 커널 실행에서 공유 메모리를 가리키는 포인터를 argument로 전달하는 것을 감지하면, 경고를 출력합니다. 런타임에서 프로그래머는 __isGlobla() 을 사용하여 해당 포인터가 전역 메모리를 참조하는지 확인할 수 있고, 이를 사용하면 안전하게 child launch로 포인터를 전달할 수 있습니다.
cudaMemcpy*Async() 또는 cudaMemset*Async()를 호출하는 것은 stream semantics를 보존하기 위해 디바이스에서 새로운 child 커널을 실행할 수 있습니다. 따라서, 이 API에 공유 메모리나 지역 메모리 포인터를 전달하는 것은 illegal이며 에러를 반환할 것 입니다.
Texture Memory
텍스처가 매핑된 전역 메모리 영역에 대한 쓰기는 텍스처 액세스 측면에서 일관성이 없습니다. 텍스처 메모리의 일관성은 child 그리드를 호출할 때 그리고 child 그리드가 완료될 때 이루어집니다. 이는 child 커널이 실행되기 전에 발생한 메모리 쓰기는 child 커널의 텍스처 메모리 액세스에 반영이 됩니다. 마찬가지로 child에 의한 메모리 쓰기는 parent에 의한 텍스처 메모리 액세스에 반영이 되지만, 오직 child 커널 완료에 대해 parent가 동기화된 이후에만 반영됩니다. parent와 child의 동시 액세스는 데이터의 불일치를 야기할 수 있습니다.
(텍스처 메모리에 자세한 내용은 추후 다른 포스팅에서 알아보겠습니다.)
Execution Environment
CUDA 실행 모델은 threads, thread blocks, grid의 기본 요소를 기반으로 하며 스레드 블록과 그리드 내의 개별 스레드에 의해 실행되는 프로그램을 정의하는 커널 함수를 가지고 있습니다. 커널 함수가 호출될 때 그리드의 속성은 CUDA의 특별한 문법을 가진 execution configuration(<<<...>>>)에 의해 설명됩니다. CUDA의 동적 병렬 지원은 디바이스에서 실행 중인 스레드에서 새로운 그리드를 configure, launch, synchronize 하는 기능을 확장합니다.
Parent and Child Grids
parent 그리드에 속하여 새로운 그리드를 구성하고 실행하는 디바이스 스레드와 생성되는 그리드를 child 그리드라고 합니다. child 그리드의 생성과 완료는 적절하게 중첩되어 있는데, 이는 parent 그리드는 parent의 스레드에 의해서 생성된 모든 child 그리드가 완료되기 전까지 완료되었다고 간주되지 않는다는 의미입니다. 비록 실행 중인 스레드가 실행된 child 그리드에서 명시적으로 동기화되지 않더라도 런타임은 parent와 child 간의 암시적 동기화를 보장합니다.
Scope of CUDA Primitives
Host와 Device에서 CUDA 런타임은 커널을 실행하고, 실행된 작업이 완료되기를 기다리고, streams과 events를 통해 실행(launches) 간의 의존성을 추적하는 API를 제공합니다. Host 시스템에서 실행 상태와 stream과 event를 참조하는 CUDA 기본 요소는 프로세스 내의 모든 스레드 사이에서 공유됩니다. 그러나 프로세스는 독립적으로 실행되고 CUDA 객체를 공유하지 않을 수 있습니다.
디바이스에도 비슷한 계층이 존재합니다. 실행된 커널과 CUDA 객체는 스레드 블록 내의 모든 스레드에서 볼 수 있습니다만, 스레드 블록 간에는 독립적입니다. 이는 하나의 스레드에서 스트림이 생성되고, 동일 스레드 블록의 다른 스레드들이 이 스트림을 사용할 수 있지만, 다른 스레드 블록의 스레드와는 공유할 수 없습니다.
Synchronization
커널 실행을 포함하여 어떤 스레드에서의 CUDA 런타임 명령은 스레드 블록 간에서 visible합니다. 이는 parent 그리드에서 실행중인 스레드는 그 스레드에 의해서, 스레드 블록 내의 다른 스레드에 의해서, 또는 같은 스레드 블록 내에서 생성된 스트림에서 실행된 그리드에 대한 동기화를 수행한다는 것을 의미합니다. 스레드 블록의 수행은 블록의 모든 스레드에 의한 모든 launches가 완료될 때까지 완료된 것으로 간주되지 않습니다. 만약 블록의 모든 스레드가 모든 child launches가 완료되기 전에 종료된다면, 동기화 명령은 자동으로 트리거됩니다.
(NVIDIA 프로그래밍 가이드에서의 설명을 번역했지만, 의미 전달이 조금 부족한 것 같습니다 ㅠ.ㅠ 동기화는 아래에서 조금 더 자세하게 살펴보겠습니다.)
Streams and Events
CUDA Streams과 Events는 그리드 launches 간의 의존성에 대한 제어를 허용합니다. 동일한 스트림에서 실행되는 그리드는 순서대로 수행되며, 이벤트는 스트림간의 의존성을 생성하기 위해 사용될 수 있습니다. 디바이스에서 생성된 스트림과 이벤트는 동일하게 사용됩니다.
그리드 내에서 생성된 스트림과 이벤트는 스레드 블록 범위 내에서 존재하지만 생성된 스레드 블록 외부에서 사용될 경우 정의되지 않은 동작이 발생합니다. 위에서 설명한 바와 같이 스레드 블록에 의해서 시작된 모든 작업은 블록이 종료될 때 암시적으로 동기화됩니다. 따라서, 스레드 블록 범위 외부에서 수정된 스트림에 대한 동작은 정의되지 않습니다.
Perent 그리드에서 생성된 스트림과 이벤트가 child 그리드에서 사용될 때 정의되지 않은 동작이 발생하는 것과 같이 Host에서 생성된 스트림과 이벤트가 어느 커널 내부에서 사용될 때 정의되지 않은 동작을 발생시킵니다.
Ordering and Concurrency
디바이스 런타임으로부터 커널 실행의 순서는 CUDA Stream 순서 semantics를 따릅니다. 스레드 블록 내에서 동일한 스트림으로 실행된 모든 커널은 순서대로 수행됩니다. 동일한 스레드 블록의 여러 스레드이 동일한 스트림으로 실행되면, 스트림에서의 순서는 블록 내의 스레드 스케줄링에 의존하며, 이는 __syncthreads()와 같은 동기화 기본 요소에 의해 제어될 수 있습니다.
스레드 블록 내의 모든 스레드들에 의해서 스트림이 공유되기 때문에 암시적인 널(NULL) 스트림 또한 공유됩니다. 만약 한 스레드 블록의 여러 스레드가 널 스트림에서 실행된다면, 그 실행은 순서대로 수행됩니다. 만약 동시에 수행되기를 원한다면, 명시적인 스트림이 사용되어야 합니다.
아래 코드에서 A,B,C 커널이 널 스트림에서 수행되는데, A,B,C가 순차적으로 수행된다는 것을 의미합니다.
Device Management
디바이스 런타임으로부터 다중 GPU는 지원하지 않으며, 디바이스 런타임은 현재 실행 중인 디바이스에서만 동작할 수 있습니다. 그러나 시스템에서 CUDA를 사용할 수 있는 디바이스를 쿼리하는 것은 허용됩니다.
Synchronization, Streams, and Events
Synchronization
호스트로부터 커널을 실행하는 것처럼, 디바이스에서의 커널 실행은 non-blocking 입니다. 만약 parent 스레드에서 child 커널이 완료되기를 기다렸다가 진행되는 것을 원한다면 명시적으로 동기화를 수행해주어야 합니다.
parent 스레드에서 child 커널에 동기화를 수행하는 한 가지 방법은 cudaDeviceSynchronize() API를 호출하는 것입니다. 이를 호출한 스레드는 스레드 블록(thread-block)의 어떠한 스레드로부터 실행된 모든 커널이 완료될 때까지 대기합니다. 그러나, 이는 블록의 모든 스레드들이 대기한다는 것을 의미하지 않습니다. 따라서, 만약 block-wide(블록 전체)의 동기화를 원한다면, 한 블록의 모든 스레드에서 __syncthreads() 가 뒤이어 호출되어야 합니다.
만약 parent 커널이 다른 child 커널을 실행했지만, 명시적으로 동기화를 수행하지 않았다면 런타임은 parent 커널 종료 전에 암시적으로 동기화를 수행합니다. 이는 parent와 child 커널이 적절히 중첩되어 있다는 것을 확실시해주고, 어떠한 커널도 child 커널이 완료되기 전까지 완료할 수 없습니다.
위에서 봤던 그림인데, 아래 그림이 parent 커널과 child 커널의 실행이 어떤 순서로 어떻게 완료되는지 잘 보여주고 있습니다.
간단하게 실제로 어떻게 동기화가 이루어지는지 살펴보겠습니다. 아래 코드는 위에서 살펴본 parent_launch, child_launch 커널 함수와 host_launch 함수에 각 함수의 진입과 종료를 출력하도록 printf를 추가하였습니다.
__global__ void child_launch(int *data)
{
if (threadIdx.x == 0)
printf("-- child kernel launched\n");
data[threadIdx.x] = data[threadIdx.x] + 1;
if (threadIdx.x == 0)
printf("-- child kernel end\n");
}
__global__ void parent_launch(int *data)
{
if (threadIdx.x == 0)
printf("- parent kernel launched\n");
data[threadIdx.x] += threadIdx.x;
__syncthreads();
if (threadIdx.x == 0)
{
child_launch<<<1, 256>>>(data);
cudaDeviceSynchronize();
}
__syncthreads();
if (threadIdx.x == 0)
printf("- parent kernel end\n");
}
void host_launch(int *data)
{
printf("host_launch called\n");
parent_launch<<<1, 256>>>(data);
printf("host_launch is end\n");
}위 코드를 간단하게 아래 main 함수로 실행하면,
int main()
{
int *h_data = (int *)malloc(256 * sizeof(int));
int *d_data;
CUDA_CHECK(cudaMalloc((void **)&d_data, 256 * sizeof(int)));
host_launch(d_data);
printf("copy from device to host\n");
CUDA_CHECK(cudaMemcpy(h_data, d_data, 256 * sizeof(int), cudaMemcpyDeviceToHost));
printf("Data: \n");
for (int i = 0; i < 256; i++)
printf("%d ", h_data[i]);
printf("\n");
return 0;
}다음의 결과를 확인하실 수 있습니다.
호스트 코드에서 host_launch 실행은 non-blocking이기 때문에 'host_launch called'가 출력된 이후에 바로 'host_launch is end'와 'copy from device to host'가 호출됩니다. 하지만 Data 출력은 커널 실행이 완료된 이후에 수행되는데, 이는 cudaMemcpy 함수를 시행하면 내부적으로 cudaDeviceSynchronize()의 역할을 수행하기 때문입니다. 따라서, cudaMemcpy 함수를 호출했을 때, 실행 중인 디바이스 커널이 완료된 후에 메모리 카피를 호스트에서 수행하게 됩니다.
parent_launch와 child_launch에 동작은 위에서 본 것과 같습니다.
처음에 말씀드렸듯이 동적 병렬은 compute capability 3.5 이상의 디바이스에서만 지원하기 때문에 nvcc로 컴파일 시 -arch 옵션을 지정해주어야 합니다. 그리고 저의 경우에는 -rdc=true 옵션도 지정해주어야 했는데, 이 옵션은 relocatable device code 생성을 활성화합니다.
Synchronization Depth
parent 커널의 child 커널에 대한 명시적인 동기화를 수행한다면, child 커널이 완료되는 것을 기다리기 위해 실행 중이던 작업이 swap됩니다. 이러한 이유로 parent 커널 상태를 저장하기 위한 저장 공간이 메모리에 할당되어 있어야 합니다. parent 커널의 parent 커널 또한 swap 되어 대기 중일 수 있습니다. 그러므로 이러한 저장 공간은 동기화가 수행되는 가장 깊은 nesting level까지의 모든 커널을 저장할만큼 충분히 커야합니다.
이렇게 가장 깊은 nesting level을 synchronization depth라고 합니다.
각 레벨의 동기화 depth를 위해 할당되어야하는 저장 공간 메모리의 양은 디바이스에서 지원하는 실행 중인 최대 개수의 스레드 상태를 저장할 수 있을만큼 충분히 커야합니다. 현재 디바이스에서 이 메모리의 양은 level당 ~150MB 정도입니다. 최대 synchronization depth는 소프트웨어에서 이 저장 공간을 위해 할당되는 메모리 양에 의해 제한되고, 이는 하드웨어에서 의해 정해진 최대 nesting depth 이상으로 제한될 수 있습니다.
이 저장 공간을 위해 예약된 기본 메모리 양은 2단계의 동기화 depth에 충분한 양입니다. 그러나 이 저장 공간 크기는 호스트 함수에서 cudaDeviceSetLimit() API 호출(with cudaLimitDevRuntimeSyncDepth 파라미터)를 통해 더 큰 값으로 설정할 수 있습니다.
Streams
named 와 unnamed(NULL) stream은 디바이스 런타임에서 모두 사용가능합니다. named 스트림은 스레드 블록 내에서 어떤 스레드에 의해서든 사용할 수 있지만, 스트림 핸들(stream handles)은 다른 블록이나 child/parent 커널에 전달되어서는 안됩니다. 다시 말하자면, 스트림은 스트림이 생성된 블록에서 private로 취급되어야 합니다. 스트림 핸들은 블록 간에 유일한 것이라고 보장되지 않으며, 블록 내에서 할당되지 않은 스트림 핸들을 사용하는 것은 정의되지 않은 결과를 유발합니다. (호스트에서 생성된 스트림을 커널 내에서 사용하거나, parent 그리드에서 생성된 스트림을 child 그리드에서 사용하면 정의되지 않은 동작이 발생합니다.)
커널을 실행할 때 스트림을 지정하지 않는다면, 기본으로 NULL 스트림이 모든 스레드에 사용됩니다.
호스트 쪽의 launch와 유사하게, 분리된 스트림에서 실행되는 작업은 동시에(concurrently) 수행될 수 있지만, 실제로 동시성(concurrency)은 보장되지 않습니다. child 커널 간의 동시성에 요구하는 프로그램은 CUDA 프로그래밍 모델에서 지원되지 않으며, 정의되지 않은 동작을 일으킵니다. 블록에서 named 스트림의 개수는 제한되어 있지 않으나, 플랫폼에서 maximum concurrency는 제한됩니다. 만약 지원하는 동시 실행보다 더 많은 스트림이 생성되면, 이들 중 일부는 순차적으로 실행됩니다.
호스트 쪽의 널 스트림에서 전역 동기화는 동적 병렬에서 지원되지 않습니다. 호스트와 동적 병렬이 있는 디바이스 간에 스트림 동작을 다르게 하려면, 커널에서 생성되는 모든 스트림을 반드시 cudaStreamNonBlocking 플래그를 사용한 cudaStreamCreateWithFlags() API를 호출하여 생성해야 합니다. 커널로부터 cudaStreamCreate() 호출은 컴파일 할 때, "unrecognized function call" 에러가 발생합니다.
cudaStreamSynchronize() API 와 cudaStreamQuery() API는 커널 내에서 사용할 수 없으며, stream에서 수행되는 child kernel과 동기화를 수행하기 위해서는 오직 cudaDeviceSynchronize()만 사용할 수 있습니다.
아직 스트림에 관해 자세히는 알지 못하기 때문에, 개인적인 호기심으로 몇 가지 테스트를 해봤습니다. 방금 위에서 살펴본 parent_launch와 child_launch 커널 함수에서 main 함수를 다음과 같이 작성하고 실행해봤습니다. host_launch를 연속해서 두 번 실행합니다.
int main()
{
int *h_data = (int *)malloc(256 * sizeof(int));
int *d_data;
CUDA_CHECK(cudaMalloc((void **)&d_data, 256 * sizeof(int)));
host_launch(d_data);
host_launch(d_data);
printf("copy from device to host\n");
CUDA_CHECK(cudaMemcpy(h_data, d_data, 256 * sizeof(int), cudaMemcpyDeviceToHost));
printf("Data: \n");
for (int i = 0; i < 256; i++)
printf("%d ", h_data[i]);
printf("\n");
return 0;
}
그 결과, 첫 번째 host_launch에서 실행되는 커널들이 모두 수행된 후에 두 번째 host_launch에서 실행되는 커널이 순차적으로 실행되는 것을 확인할 수 있습니다.
이제 host_launch 함수와 main 함수를 스트림을 사용하도록 아래와 같이 변경하고 실행해봤습니다.
void host_launch(int *data, cudaStream_t* s)
{
printf("host_launch called\n");
parent_launch<<<1, 256, 0, *s>>>(data);
printf("host_launch is end\n");
}
int main()
{
int *h_data = (int *)malloc(256 * sizeof(int));
int *d_data;
CUDA_CHECK(cudaMalloc((void **)&d_data, 256 * sizeof(int)));
cudaStream_t s1, s2;
CUDA_CHECK(cudaStreamCreate(&s1));
CUDA_CHECK(cudaStreamCreate(&s2));
host_launch(d_data, &s1);
host_launch(d_data, &s2);
printf("copy from device to host\n");
CUDA_CHECK(cudaMemcpy(h_data, d_data, 256 * sizeof(int), cudaMemcpyDeviceToHost));
printf("Data: \n");
for (int i = 0; i < 256; i++)
printf("%d ", h_data[i]);
printf("\n");
CUDA_CHECK(cudaStreamDestroy(s1));
CUDA_CHECK(cudaStreamDestroy(s2));
return 0;
}그 결과,
각 host_launch에서 실행되는 커널이 각각 동시에 실행되는 것을 볼 수 있습니다. 따라서, 같은 data 배열을 사용하기 때문에 의도한 결과가 아닌 이상한 결과를 출력합니다.
아래의 코드도 위와 같은 결과를 출력합니다.
void host_launch(int *data)
{
cudaStream_t s;
CUDA_CHECK(cudaStreamCreate(&s));
printf("host_launch called\n");
parent_launch<<<1, 256, 0, s>>>(data);
printf("host_launch is end\n");
}
int main()
{
int *h_data = (int *)malloc(256 * sizeof(int));
int *d_data;
CUDA_CHECK(cudaMalloc((void **)&d_data, 256 * sizeof(int)));
host_launch(d_data);
host_launch(d_data);
printf("copy from device to host\n");
CUDA_CHECK(cudaMemcpy(h_data, d_data, 256 * sizeof(int), cudaMemcpyDeviceToHost));
printf("Data: \n");
for (int i = 0; i < 256; i++)
printf("%d ", h_data[i]);
printf("\n");
return 0;
}
이번에는 parent_launch에서 child_launch를 스트림없이 두 번 실행하도록 했습니다.
__global__ void parent_launch(int *data)
{
if (threadIdx.x == 0)
printf("- parent kernel launched\n");
data[threadIdx.x] += threadIdx.x;
__syncthreads();
if (threadIdx.x == 0)
{
child_launch<<<1, 256>>>(data);
child_launch<<<1, 256>>>(data);
cudaDeviceSynchronize();
}
__syncthreads();
if (threadIdx.x == 0)
printf("- parent kernel end\n");
}
int main()
{
int *h_data = (int *)malloc(256 * sizeof(int));
int *d_data;
CUDA_CHECK(cudaMalloc((void **)&d_data, 256 * sizeof(int)));
host_launch(d_data);
printf("copy from device to host\n");
CUDA_CHECK(cudaMemcpy(h_data, d_data, 256 * sizeof(int), cudaMemcpyDeviceToHost));
printf("Data: \n");
for (int i = 0; i < 256; i++)
printf("%d ", h_data[i]);
printf("\n");
return 0;
}
child 커널이 널 스트림을 사용하여, 각각 순차적으로 실행되는 것을 볼 수 있습니다.
parent 커널 내에서 각각의 스트림(cudaStreamDefault)으로 child 커널을 실행하면,
(아까 언급했듯이 커널 함수 내에서는 cudaCreateStream() API는 지원되지 않습니다.)
__global__ void parent_launch(int *data)
{
if (threadIdx.x == 0)
printf("- parent kernel launched\n");
data[threadIdx.x] += threadIdx.x;
__syncthreads();
if (threadIdx.x == 0)
{
cudaStream_t s1, s2;
cudaStreamCreateWithFlags(&s1, cudaStreamDefault);
cudaStreamCreateWithFlags(&s2, cudaStreamDefault);
child_launch<<<1, 256, 0, s1>>>(data);
child_launch<<<1, 256, 0, s2>>>(data);
cudaDeviceSynchronize();
cudaStreamDestroy(s1);
cudaStreamDestroy(s2);
}
__syncthreads();
if (threadIdx.x == 0)
printf("- parent kernel end\n");
}child 커널이 순차적으로 실행됩니다.
만약 cudaStreamNonBlocking을 플래그로 전달하여 실행해도 동일한 결과를 보여주고 있습니다.
__global__ void parent_launch(int *data)
{
if (threadIdx.x == 0)
printf("- parent kernel launched\n");
data[threadIdx.x] += threadIdx.x;
__syncthreads();
if (threadIdx.x == 0)
{
cudaStream_t s1, s2;
cudaStreamCreateWithFlags(&s1, cudaStreamNonBlocking);
cudaStreamCreateWithFlags(&s2, cudaStreamNonBlocking);
child_launch<<<1, 256, 0, s1>>>(data);
child_launch<<<1, 256, 0, s2>>>(data);
cudaDeviceSynchronize();
cudaStreamDestroy(s1);
cudaStreamDestroy(s2);
}
__syncthreads();
if (threadIdx.x == 0)
printf("- parent kernel end\n");
}
사실 child 커널이 동시에 실행될 것이라고 생각했지만, 순차적으로 실행되고 있습니다...
이것이 CUDA 동적 병렬이 child launch에서의 동시 실행이 수행될 수는 있지만, 보장하지는 않는다는 의미인 것 같습니다.
Events
커널 함수에서는 오직 CUDA Events의 스트림간 동기화 기능만 지원됩니다. 즉, cudaStreamWithEvent()는 지원되지만, cudaEventSynchronize(), cudaEventElapsedTime(), cudaEventQuery()는 지원되지 않습니다. cudaEventElapsedTime()이 지원되지 않기 때문에 cudaEvent는 반드시 cudaEventCreateWithFlags()를 cudaEventDisableTiming 플래그로 호출하여 생성해야 합니다.
CUDA Event를 사용하는 예제는 따로 첨부하지 않습니다. Event를 사용한 예제는 아래 코드를 참조하시면 조금 도움이 되실 것 같습니다 !
https://github.com/junstar92/parallel_programming_study/blob/master/CUDA/matrixMul/matrixMul.cu
GitHub - junstar92/parallel_programming_study: Study parallel programming - CUDA, OpenMP, MPI, Pthread
Study parallel programming - CUDA, OpenMP, MPI, Pthread - GitHub - junstar92/parallel_programming_study: Study parallel programming - CUDA, OpenMP, MPI, Pthread
github.com
모든 디바이스 런타임 객체와 마찬가지로 Event 객체는 스레드 블록 내의 모든 스레드 간에 공유될 수 있지만, 해당 스레드 블록에 local(지역 변수)이므로, child/parent 커널에 전달되어서는 안됩니다. 해당 블록에 할당되지 않은 이벤트를 사용한 이벤트 핸들은 정의되지 않은 동작을 야기합니다.
블록당 이벤트의 수를 제한되지 않지만, 이벤트는 디바이스 메모리를 소비합니다. 따라서 너무 많은 이벤트가 생성되어 리소스 한계에 도달하면, 디바이스에서 실행된 그리드에서 동시에 실행되는 스레드가 감소할 수 있습니다만, 결과에 문제가 있는 것은 아닙니다.
지금까지 동적 병렬과 이에 관련된 내용들을 조금 부족하지만 살펴봤습니다.
다음 포스팅에서는 동적 병렬을 사용하는 알고리즘이나 다른 어플리케이션을 직접 구현해보는 내용으로 찾아오겠습니다.
'NVIDIA > CUDA' 카테고리의 다른 글
| Nsight Compute로 Warp 성능 측정하기 (0) | 2022.01.07 |
|---|---|
| WARP Execution (3) | 2022.01.05 |
| Graph Search (Breadth-First Search) (0) | 2021.12.30 |
| Parallel Merge Sort (merge operation) (0) | 2021.12.24 |
| Sparse Matrix Computation (0) | 2021.12.21 |
댓글