Streams and Events (1)

References

• Professional CUDA C Programming

Contents

• CUDA Streams
• CUDA Events
• Stream Synchronization

---

CUDA Dynamic Parallelism (동적 병렬)

CUDA Dynamic Parallelism (동적 병렬)

동적 병렬에 대한 이전 포스팅에서 Execution Envrionment에 대해서 살펴보면서 CUDA의 stream과 event에 대해서 간단하게 알아봤었는데, 이번 포스팅을 통해서 조금 더 자세하게 알아보도록 하겠습니다 !

 

시작하기 전에 먼저 CUDA의 동시성(concurrency)에 대해서 짧게 언급하고 시작하도록 하겠습니다.

CUDA C 프로그래밍에서는 2가지 수준의 동시성이 있습니다.

• Kernel level concurrency
• Grid level concurrency

일반적으로는 하나의 task 또는 kernel에만 초점을 맞추어서, GPU의 많은 스레드에 의해서 병렬로 수행되는 kernel level의 concurrency에 익숙합니다. 이번 포스팅에서는 grid level의 concurrency에 대해 살펴보려고 합니다. 이는 여러 커널의 실행이 하나의 GPU 장치에서 동시에 실행되어 디바이스 활용도를 높일 수 있습니다.

 

---

Introduce Streams and Events

CUDA stream(스트림)은 host code에 의해서 issue된 순서에 따라 디바이스에서 실행되는 비동기 CUDA operation의 순서를 가리킵니다. Stream은 이러한 operation들을 캡슐화하고 순서를 유지하며, 모든 이전 operation들이 실행된 후에 실행되는 operation을 대기열에 들어가도록 하고, 대기열에 있는 operation의 상태를 쿼리할 수 있도록 합니다. 이러한 operation에는 host-device data transfer, kernel launches, 그리고 host에서 호출되지만 device에서 처리되는 다른 커맨드들을 포함합니다. Stream에서 operation의 실행은 host와 항상 비동기입니다. CUDA 런타임은 그 operation이 디바이스에서 실행될 수 있는 시점을 결정합니다. 어떤 비동기 operation의 결과를 사용하기 전에 이 operation이 완료되었는지 CUDA API를 사용하는 것은 개발자의 책임입니다. 동일한 CUDA Stream에서의 operation들은 엄격한 순서를 가지고 있지만, 다른 Stream에서의 operation들은 실행 순서에 제한이 없습니다. 따라서, 여러 Stream을 사용하여 여러 개의 커널을 실행하면 grid level의 concurrency를 구현할 수 있습니다.

 

CUDA Stream에서 쿼리되는 모든 operation들은 비동기이기 때문에, host-device system에서 다른 operation들을 오버랩하는 것이 가능합니다. 이렇게 하면 다른 작업들을 동시에 수행하여 operation의 cost를 숨길 수 있습니다.

 

일반적으로 CUDA 프로그래밍의 전형적인 패턴은 다음과 같습니다.

1. Move input data from the host to the device
2. Execute a kernel on the device
3. Move the result from the device back to the host

많은 경우에 data를 전송하는 것보다 커널을 실행하는 데 더 많은 시간을 사용합니다. 이러한 상황에서 완전히 CPU-GPU communication latency를 숨길 수도 있습니다. 커널 실행과 데이터 전달을 다른 스트림에 dispatch함으로써, 이 operation들은 오버랩될 수 있고, 프로그램의 총 수행 시간은 짧아집니다. 

 

CUDA API  함수들은 일반적으로 동기(synchronous) or 비동기(asynchronous)로 분류될 수 있습니다. 동기로 동작하는 함수는 완료될 때까지 host 스레드를 블락합니다. 비동기로 동작하는 함수는 호출된 후 즉시 host로 제어를 반환합니다. 비동기 함수와 스트림은 CUDA에서 grid-level concurrency를 구현하는 기본적인 요소입니다.

 

하지만 software의 관점에서 다른 스트림에서의 CUDA operation이 동시에 실행되는 것으로 보이지만, 물리적인 하드웨어에서도 항상 그런 것은 아닙니다. PCIe BUS에 대한 경쟁이나 SM별 리소스의 가용성에 따라 서로 다른 CUDA 스트림이 완료되기를 기다려야할 수도 있습니다.

 

---

CUDA Streams

모든 CUDA operations (kernels and data transfers)는 명시적 또는 암시적으로 하나의 스트림에서 실행됩니다.

여기에는 두 종류의 스트림이 있습니다.

• Implicitly declared stream (NULL stream)
• Explicitly declared strea (non-NULL stream)

NULL Stream은 디폴트 스트림이며, 만약 명시적으로 스트림을 지정하지 않으면 kernel launch나 data transfers는 디폴트 스트림을 사용합니다. 제가 작성한 포스팅에서 살펴본 대부분의 예제 코드는 디폴트 스트림(NULL Stream)을 사용했습니다.

 

반면, non-null stream은 명시적으로 생성되고 관리됩니다. 만약 다른 CUDA operation을 오버랩하고 싶다면, 반드시 non-null 스트림을 사용해야합니다. 비동기이며 stream-based kernel launch나 data transfers는 다음의 coarse-grain concurrency를 가능하게 합니다.

• Overlapped host computation and device computation
• Overlapped host computation and host-device data transfer
• Overlapped host-device data transfer and device computation
• Concurrent device computation

 

아래의 코드 일부는 디폴트 스트림을 사용하는 코드입니다.

cudaMemcpy(..., cudaMemcpyHostToDevice);
kernel<<<grid, block>>>(...);
cudaMemcpy(..., cudaMemcpyDeviceToHost);

CUDA 프로그램의 동작을 이해하기 위해서, 항상 device와 host의 관점에서 살펴봐야합니다.

device 관점에서 이 3개의 operation은 default stream에서 issue되고, issue된 순서대로 실행됩니다. device는 수행 중인 다른 host operation을 인식하지 못합니다.

host 관점에서는 각 data transfer가 동기화되며 이를 완료할 때까지 host를 idle 상태로 강제합니다. 커널의 실행은 비동기이므로 host application은 커널의 완료 여부에 상관없이 기다리지 않고 다음 operation을 재개합니다. 이러한 kernal launch의 비동기 동작은 device와 host computation을 쉽게 오버랩할 수 있도록 해줍니다.

 

Data transfer 또한 비동기로 실행할 수 있습니다. 그러나, 반드시 명시적으로 CUDA Stream을 지정해주어야 합니다. CUDA 런타임은 비동기로 데이터를 전달할 수 있는, 다음과 같은 cudaMemcpy의 비동기 버전을 제공합니다.

cudaError_t cudaMemcpyAsync(void* dst, const void* src, size_t count, cudaMemcpyKind kind, cudaStream_t stream = 0);

5번째 파라미터에 스트림 식별자가 추가되었습니다. 기본적으로 스트림 식별자는 디폴트 스트림(NULL)로 설정됩니다. 이 함수는 host에 비동기이므로, 이 함수가 호출된 즉시 제어는 다시 host로 리턴됩니다. 

쉽게 비동기로 동작하도록 할 수 있지만, 먼저 사용할 non-null 스트림을 생성해주어야 합니다.

cudaError_t cudaStreamCreate(cudaStream_t* pStream);

cudaStreamCreate는 명시적으로 관리할 non-null 스트림을 생성합니다. pStream으로 리턴되는 스트림은 cudaMemcpyAsync이나 다른 비동기 CUDA API의 스트림 인자로 사용될 수 있습니다. 비동기 CUDA 함수를 사용할 때 한 가지 혼동할 수 있는 것은 이전에 수행된 operation에서 에러 코드를 리턴할 수 있다는 것입니다.

(에러를 리턴하는 API 호출이 반드시 에러를 발생시킨 것은 아닙니다.)

 

비동기 데이터 전달을 수행할 때, 반드시 pinned(or non-pageable) host memory를 사용해야 합니다. Pinned Memory는 cudaMallocHost나 cudaHostAlloc을 사용하여 할당할 수 있습니다.

cudaError_t cudaMallocHost(void** ptr, size_t size);
cudaError_t cudaHostAlloc(void** pHost, size_t size, unsigned int flags);

host의 가상 메모리에 할당을 고정하면 CPU 메모리의 물리적인 위치를 application이 실행되는 동안 일정하게 유지시킬 수 있습니다. 그렇지 않다면, OS는 언제든지 host 가상 메모리의 물리적 위치를 자유롭게 변경할 수 있습니다. 만약 비동기 CUDA transfer가 pinned host memory없이 수행된다면 OS는 CUDA 런타임에서 device로 데이터가 전송되는 동안 물리적으로 배열을 이동하는 것이 가능하고, 이는 정의되지 않은 동작을 발생시킵니다.

 

non-default 스트림으로 커널을 실행시키기 위해서는 반드시 스트림 식별자를 execution configuration의 4번째 파라미터로 지정해주어야 합니다.

kernel_name<<<grid, block, sharedMemSize, stream>>>(...);

non-default 스트림은 다음과 같이 선언됩니다.

cudaStream_t stream;

그리고, 아래와 같이 스트림을 생성하고, 

cudaStreamCreate(&stream);

다음의 API를 통해 사용을 마친 스트림의 리소스를 해제할 수 있습니다.

cudaError_t cudaStreamDestroy(cudaStream_t stream);

만약 cudaStreamDestroy로 해제되는 스트림에 아직 pending 중인 작업이 있다면, cudaStreamDestroy는 즉시 반환되지만 스트림의 모든 작업이 완료되면 스트림의 리소스가 자동으로 해제됩니다.

 

모든 CUDA 스트림 operation은 비동기이므로, CUDA API는 스트림에서의 모든 operation이 완료되었는지 체크하기 위한 두 개의 API 함수를 제공합니다.

cudaError_t cudaStreamSynchronize(cudaStream_t stream);
cudaError_t cudaStreamQuery(cudaStream_t stream);

cudaStreamSynchronize는 제공된 스트림의 모든 operation이 완료될 때까지 host를 강제로 블락합니다. cudaStreamQuery는 스트림의 모든 operation이 완료되었는지 체크합니다. 만약 모든 operation이 완료되었다면 cudaSuccess를 리턴하고, 하나 이상의 실행 중이거나 pending된 것이 있다면 cudaErrorNotReady를 리턴합니다.

 

실제로 CUDA 스트림이 어떻게 사용되는지 살펴보기 위해서 여러 스트림에서 CUDA operation을 dispatch하는 일반적인 패턴이 아래에 있습니다.

for (int i = 0; i < nStreams; i++) {
	int offset = i * bytesPerStream;
    cudaMemcpyAsync(&d_a[offset], &a[offset], bytesPerStream, streams[i]);
    kernel<<<grid, block, 0, streams[i]>>>(&d_a[offset]);
    cudaMemcpyAsync(&a[offset], &d_a[offset], bytesPerStream, streams[i]);
}

for (int i = 0; i < nStreams; i++)
	cudaStreamSynchronize(streams[i]);

아래 그림은 위 코드에서 3개의 스트림을 사용하는 경우의 timeline을 보여줍니다.

위 그림을 살펴보면, 비록 다른 스트림에서 데이터 전송 operation이 issue되었지만 동시에 수행되지 않고 있는 것을 확인할 수 있습니다. 이는 shared resource(PCIe BUS)에 대한 경쟁 때문에 발생됩니다. 프로그래밍 모델의 관점에서는 이 operation들이 독립적이지만, 이들은 동일한 하드웨어 리소스를 공유하기 때문에 데이터 전송은 시리얼로 실행됩니다. 만약 이중 PCIe BUS가 있는 디바이스라면 두 개의 데이터 전송을 오버랩할 수 있지만, 다른 스트림을 사용해야 하고 데이터 전달의 방향 또한 달라야 합니다. 위 그림에서 D2H 전송과 H2D 전송이 서로 오버랩되어 있는 것을 확인할 수 있습니다.

 

동시에 실행되는 커널의 최대 개수는 디바이스에 따라 다릅니다. Fermi의 경우에는 16-way concurrency를 지원하고, Kepler는 32-way concurrency를 지원합니다. concurrent kernel의 수는 shared memory나 레지스터와 같은 디바이스의 사용가능한 compute 리소스에 의해서도 제한됩니다.

 

NVIDIA의 공식 문서에서 Compute Capability에 따른 concurrent kernel의 최대 개수는 다음과 같습니다.

(출처 : link)

---

Stream Scheduling

개념상, 모든 스트림은 동시에 실행될 수 있습니다. 그러나 스트림을 물리적 하드웨어에 매핑할 때 항상 그런 것은 아닙니다. 이번에는 다중 CUDA 스트림에서의 concurrent kernel operation이 어떻게 하드웨어에 의해 스케쥴링되는지 살펴보겠습니다.

 

False Dependecies

비록 Fermi GPU는 16-way concurrency(한번에 16개의 그리드를 실행)를 지원하지만, 모든 스트림은 궁극적으로 하나의 하드웨어 work queue로 다중화됩니다. 실행할 그리드를 선택할 때, queue의 front에 있는 task는 CUDA 런타임에 의해서 스케쥴링됩니다. 런타임은 task dependencies를 체크하고, 이 task가 의존하는 다른 task들이 아직 실행 중이라면 대기시킵니다. 이후에 모든 dependencies가 만족되면 새로운 task가 사용가능한 SM들에게 dispatch 됩니다.

 

이 single pipeline은 false dependecy를 유발합니다.

아래에 그림에서 보여주는 것처럼, 다른 모든 그리드를 실행하기 전에 런타임이 블락하기 때문에 결국에는 빨간색으로 동그라미 친 작업들만 동시에 실행됩니다.

queue에서 블락된 operation은 서로 다른 스트림에 속해있더라도 queue의 모든 후속 작업들을 블락합니다.

 

Hyper-Q

Kepler에서는 Hyper-Q라고 불리는 기술인 다중 하드웨어 work queue를 사용하여 False Dependencies가 감소되었습니다. Hyper-Q는 host와 device 간의 다중 hardware-managed connections를 두어 다중 CPU 스레드 또는 프로세스가 동시에 단일 GPU에서 작업을 실행할 수 있도록 해줍니다. Fermi의 False Dependencies로 인해 제한되었던 기존 어플리케이션은 코드를 변경하지 않고도 큰 성능 향상을 얻었습니다. Kepler GPU는 32개의 hardware work queue를 사용하며 스트림당 하나의 work queue를 할당합니다. 만약 32개 이상의 스트림이 생성되면, 여러 스트림이 단일 hardware work queue를 공유합니다.

아래 그림은 3개의 work queue에 3개의 스트림이 있는 예제 케이스를 보여줍니다.

 

Stream Priorities

Compute capability 3.5 이상의 디바이스에서 스트림에 우선순위를 할당할 수 있습니다. 다음 함수를 통해 지정된 우선순위를 가진 스트림을 생성할 수 있습니다.

cudaError_t cudaStreamCreateWithPriority(cudaStream_t* pStream, unsigned int flags, int priority);

이 함수는 지정된 정수의 우선순위를 가진 스트림을 생성하고 pStream에 핸들을 리턴합니다. 이 우선순위는 pStream에서 스케줄링되는 작업과 연관됩니다. 높은 우선순위 스트림에서 대기 중인 그리드는 낮은 우선순위 스트림에서 실행 중인 작업보다 먼저 수행될 수 있습니다. 스트림의 우선순위는 데이터 전송에는 영향을 미치지 않으며 오직 compute kernel에만 영향을 미칩니다. 만약 지정된 우선순위가 디바이스에서 관리되는 범위를 벗어나면 해당 범위에서 가장 낮거나 가장 높은 수로 자동으로 설정됩니다.

주어진 디바이스에서 허용되는 우선순위 범위는 다음의 함수를 통해 쿼리할 수 있습니다.

cudaError_t cudaDeviceGetStreamPriorityRange(int *leastPriority, int *greatestPriority);

이 함수는 leastPriority와 greatestPriority에 현재 디바이스에서 각각 가장 낮은 우선순위와 높은 우선순위의 값을 반환합니다. 숫자가 낮을수록 더 높은 우선순위를 갖습니다. 만약 현재 디바이스에서 스트림 우선순위를 지원하지 않으면, 모두 0을 반환합니다.

 

---

CUDA Events

CUDA의 이벤트는 본질적으로 CUDA 스트림에서 operation flow의 특정 지점과 관련된 스트림의 마커입니다. 이벤트를 사용하면 다음의 두 가지 기본 작업을 수행할 수 있습니다.

• Synchronize stream execution
• Monitor device progress

CUDA API는 스트림의 특정 포인트에 이벤트를 삽입하고 이벤트가 완료되었는지 쿼리할 수 있는 함수를 제공합니다. 주어진 스트림에서 기록된 이벤트는 오직 동일한 스트림에서의 모든 preceding operation들이 완료되었을 때만 만족합니다. 디폴트 스트림으로 지정된 이벤트는 모든 CUDA 스트림의 모든 preceding operation들에 적용됩니다.

 

Creation and Destruction

이벤트는 다음과 같이 선언합니다.

cudaEvent_t event;

선언되면, 다음의 함수를 통해 이벤트를 생성할 수 있습니다.

cudaError_t cudaEventCreate(cudaEvent_t* event);

그리고 다음의 함수를 통해 생성된 이벤트를 없앨 수 있습니다.

cudaError_t cudaEventDestroy(cudaEvent_t event);

만약 cudaEventDestroy가 호출되었을 때 이벤트가 아직 제거될 준비가 되지 않았다면, 이 호출은 즉시 반환되고 이 이벤트가 완료될 때, 관련된 리소스들이 자동으로 릴리즈됩니다.

 

Recoding Events and Measuring Elapsed Time

이벤트는 스트림의 실행의 한 포인트를 마킹합니다. 이것으로 실행 중인 스트림 operation이 해당 지점에 도달했는지 체크할 수 있습니다. 이러한 작업은 work queue의 front에서 pop될 때 호스트 측에 flag를 set하여 완료되었는지 알려주는 작업이라고 생각할 수 있습니다.

이벤트는 다음의 함수를 통해 CUDA Stream에 queue 합니다.

cudaError_t cudaEventRecord(cudaEvent_t event, cudaStream_t stream = 0);

파라미터로 전달된 이벤트는 지정된 스트림의 모든 preceding operation이 완료되기를 기다리거나 테스트하는 데 사용될 수 있습니다. 이벤트를 기다리는 것은 host thread를 블락하며, 다음의 함수를 사용하여 수행됩니다.

cudaError_t cudaEventSynchronize(cudaEvent_t event);

cudaEventSynchronize는 스트림에서 cudaStreamSynchronize와 유사하지만 host가 스트림 실행의 중간 지점에서 기다릴 수 있다는 차이점이 있습니다. 또한 다음 함수를 사용하여 host application을 블락하지 않고 이벤트가 완료되었는지 테스트해볼 수 있습니다.

cudaError_t cudaEventQuery(cudaEvent_t event);

cudaEventQuery는 cudaStreamQuery의 동작을 이벤트에서 하는 것과 같습니다.

 

이벤트를 사용하면 다음의 함수를 통해 두 이벤트에 의해서 마킹된 CUDA operation의 실행 시간을 측정할 수 있습니다.

cudaError_t cudaEventElapsedTime(float* ms, cudaEvent_t start, cudaEvent_t stop);

이 함수는 start 이벤트와 stop 이벤트 사이의 걸린 시간을 milliseconds로 반환합니다. start와 stop 이벤트는 동일한 CUDA 스트림에 연관되어 있을 필요는 없습니다. 만약 두 이벤트 중 하나가 non-NULL 스트림에서 기록된다면 반환되는 시간은 예상보다 길 수 있습니다. 이는 cudaEventRecord가 비동기로 발생하며, 측정된 latency가 두 이벤트 사이에 있다고 보장하지 않기 때문입니다.

 

아래 코드는 device operation의 시간을 측정하기 위해서 일반적으로 어떻게 이벤트를 사용하는지 보여줍니다.

// create two events
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);

// record start event on the default stream
cudaEventRecord(start);

// execute kernel
kernel<<<grid, block>>>(...);

// record stop event on the default stream
cudaEventRecord(stop);

// wait until the stop event completes
cudaEventSynchronize(stop);

// calculate the elapsed time between two events
float time;
cudaEventElapsedTime(&time, start, stop);

// clean up the two events
cudaEventDestroy(start);
cudaEventDestroy(stop);

여기서 start와 stop 이벤트는 NULL 스트림에 위치합니다. timestamp는 NULL 스트림의 처음에 start event에서 기록되고, stop 이벤트의 timestamp는 NULL 스트림의 끝에서 기록됩니다.

 

---

Stream Synchronization

non-default 스트림에서의 모든 operation은 host thread에 non-blocking이기 때문에, 스트림에서 실행되는 작업과 host를 동기화하는 상황이 필요할 때도 있습니다.

Host의 관점에서 CUDA operation는 두 가지 카테고리로 분류됩니다.

• Memory-related operations
• Kernel launches

Kernel launches는 항상 host에 대해 비동기입니다. 많은 Memory operation들은 본질적으로 동기화되지만(ex, cudaMemcpy) CUDA 런타임은 memory operation을 수행하기 위한 비동기 함수도 제공합니다.

 

그리고 위에서 살펴 본, 두 종류의 스트림이 있습니다.

• Asynchronous streams (non-NULL streams)
• Synchronous streams (the NULL/default streams)

non-null 스트림은 host에 비동기인 스트림입니다. 이 스트림에서의 모든 operation은 host execution을 블락하지 않습니다. 반면에 NULL-스트림은 host에 동기되는 스트림입니다.

 

Non-NULL 스트림은 두 종류로 더 분류될 수 있습니다.

• Blocking streams
• Non-blocking streams

비록 non-NULL 스트림은 host에 non-blocking하지만, non-NULL 스트림 내에서의 operation은 NULL 스트림에서의 operation에 의해서 블락될 수 있습니다. 만약 non-NULL 스트림이 blocking 스트림이라면, NULL 스트림은 이 스트림의 operation을 블락할 수 있습니다. 만약 non-NULL 스트림이 non-blocking 스트림이라면, NULL 스트림에서의 operation을 블락하지 않습니다. 이어지는 내용에서 어떻게 blocking / non-blocking 스트림을 사용하는지 알아보겠습니다.

 

Blocking and Non-Blocking Streams

cudaStreamCreate로 생성되는 스트림은 blocking 스트림이며, 이는 이 스트림의 operation 실행은 NULL 스트림에서의 이전 operation이 완료되기를 기다릴 수 있다는 것을 의미합니다. NULL 스트림은 implicit stream이며, 동일한 CUDA context의 다른 모든 blocking stream과 동기화합니다. 일반적으로 NULL stream에 operation이 issue될 때, CUDA context는 이 operation을 시작하기 전에 모든 blocking stream에서 이전에 issue된 모든 operation들을 기다립니다. 또한, blocking stream에서 issue된 다른 operation들도 이를 실행하기 전에 NULL 스트림에서의 preceding operation들이 완료될 때까지 대기합니다.

 

예를 들어, 아래의 코드는 stream_1에서 kernel_1을, NULL stream에서 kernel_2를, stream_2에서 kernel_3을 실행합니다.

kernel_1<<<1, 1, 0, stream_1>>>();
kernel_2<<<1, 1>>>();
kernel_3<<<1, 1, 0, stream_2>>>();

이 코드의 동작은 kernel_1이 완료될 때까지 GPU에서 kernel_2는 시작되지 않고, kernel_2가 완료될 때까지 kernel_3은 시작되지 않습니다. host의 관점에서는 각 커널의 실행은 여전히 비동기이며, non-blocking 입니다.

 

CUDA 런타임은 NULL 스트림과 관련된 non-NULL 스트림의 동작을 커스터마이즈할 수 있는 함수를 제공합니다.

cudaError_t cudaStreamCreateWithFlags(cudaStream_t* pStream, unsigned int flags);

flags는 생성되는 스트림의 동작을 결정합니다. 허용되는 flags값은 다음과 같습니다.

• cudaStreamDefault: default stream creation flag (blocking)
• cudaStreamNonBlocking: asynchronous stream creation flag (non-blocking)

cudaStreamNonBlocking을 지정하면 NULL 스트림과 관련된 non-NULL 스트림의 blocking 동작은 disable 됩니다. 만약 stream_1과 stream_2가 cudaStreamNonBlocking으로 생성되었다면, 어떠한 커널도 블락되지 않고 다른 커널의 완료를 기다리지 않습니다.

 

Implicit Synchronization

CUDA는 두 종류의 host-device 동기화를 포함합니다: explicit and implicit

아마도 이미 많은 explict 동기화 함수를 알고 계실텐데, cudaDeviceSynchronize / cudaStreamSynchronize / cudaEventSynchronize가 이에 해당됩니다. 이 함수들은 명시적으로 host에서 호출되어 device에서의 task 실행을 host thread와 동기화시킵니다.

 

implicit 동기화 예제도 이미 알고계실 것이라 생각합니다. 예를 들어, cudaMemcpy는 암시적으로 device와 host를 동기화합니다. 이는 data transfer가 완료될 때까지 host application이 블락되기 때문입니다. 그러나 이 함수의 주된 목적은 동기화가 아니기 때문에 이 동기화에 대한 side effect도 내재되어 있습니다. Host와 Device를 암시적으로 동기화하는 함수를 호출하면 예상치못한 성능의 저하가 발생할 수 있으므로 이러한 implicit 동기화에 주의하는 것이 중요합니다.

 

Implicit 동기화는 특별히 관심있게 살펴봐야 하는데, 이는 암시적 동기화 동작을 수행하는 런타임 함수가 device level에서 원치않는 blocking을 발생할 수 있기 때문입니다.

다음과 같은 메모리 관련 operation들은 현재 디바이스에서 모든 이전 operation들을 blocking 합니다.

• A page-locked host memory allocation
• A device memory allocation
• A device memset
• A memory copy between two addresses on the same device
• A modification to the L1/shared memory configuration

 

Explicit Synchronization

CUDA 런타임은 grid level에서의 명시적 동기화를 위한 여러가지 방법을 지원합니다.

• Synchronizing the device
• Synchronizing a stream
• Synchronizing an event in a stream
• Synchronizing across streams using an event

다음의 함수를 통해 device가 모든 preceding tasks를 완료할 때까지 host thread를 블락할 수 있습니다.

cudaError_t cudaDeviceSynchronize(void);

이 함수는 host thread가 현재 디바이스와 관련된 모든 computation / communication이 완료될 때까지 기다리도록 합니다. 이는 상대적으로 heavy한 동기화이기 때문에 이는 host를 멈추는 것을 피하기 위해서 적게 사용되어야 합니다.

 

cudaStreamSynchronize를 사용하면 stream에서의 모든 operation들이 완료될 때까지 host thread를 블락할 수 있고, cudaStreamQuery를 사용하여 완료되었는지 확인하기 위한 non-blocking test를 수행할 수 있습니다.

cudaError_t cudaStreamSynchronize(cudaStream_t stream);
cudaError_t cudaStreamQuery(cudaStream_t stream);

CUDA 이벤트 또한 cudaEventSynchronize와 cudaEventQuery를 통해 blocking과 동기화하는데 사용될 수 있습니다.

cudaError_t cudaEventSynchronize(cudaEvent_t event);
cudaError_t cudaEventQuery(cudaEvent_t event);

 

추가로 cudaStreamWaitEvent는 CUDA 이벤트를 사용하여 스트림 간의 종속성을 도입할 수 있는 flixible한 방법을 제공합니다.

cudaError_t cudaStreamWithEvent(cudaStream_t stream, cudaEvent_t event);

cudaStreamWaitEvent는 cudaStreamWaitEvent가 호출된 이후에 스트림에서 대기 중인 작업을 실행하기 전, 지정된 스트림이 지정된 이벤트를 기다리도록 합니다. 이벤트는 동일한 스트림 또는 다른 스트림과 연관될 수도 있습니다.

다른 스트림과 연관된 경우 이 함수는 아래의 그림처럼 cross-stream 동기화를 수행합니다. 여기서 stream 2에 의해 issue된 대기는 계속 진행하기 전에 stream 1에서 생성된 이벤트가 충족되는지 확인합니다.

 

Configurable Events

CUDA 런타임은 이벤트의 동작과 속성을 커스터마이즈하기 위한 방법을 제공합니다.

cudaError_t cudaEventCreateWithFlags(cudaEvent_t* event, unsigned int flags);

유효한 flags는 다음과 같습니다.

• cudaEventDefault
• cudaEventBlockingSync
• cudaEventDisableTiming
• cudaEventInterprocess

cudaEventBlockingSync 플래그는 이벤트가 blocking synchronization을 사용하도록 지정합니다. cudaEventSynchronize()를 사용하여 이 플래그로 생성된 이벤트를 기다리는 host thread는 실제로 이벤트가 완료될 때까지 블락됩니다.

 

cudaEventDisableTiming 플래그는 생성된 이벤트가 동기화를 위해서만 사용되고 timing data는 기록하지 않도록 합니다. timestamp를 위한 오버헤드가 제거되어 cudaStreamWaitEvent와 cudaEventQuery 호출의 성능이 향상됩니다.

 

cudaEventInterprocess 플래그는 생성된 이벤트가 프로세스간 이벤트로 사용되도록 합니다.