Streams and Events (3) - Kernel and Data Transfer, Stream Callback

References

• Professional CUDA C Programming

Contents

• Overlapping Kernel Execution and Data Transfer
• Overlapping GPU and CPU Execution
• Stream Callbacks

---

Overlapping Kernel Execution and Data Transfer

Streams and Events (1)

Streams and Events (2) - Concurrent Kernels

이전 포스팅에서 CUDA의 스트림과 이벤트에 대해서 살펴보고, 여러 스트림에서 커널들을 어떻게 동시에 실행시킬 수 있는지 살펴봤습니다.

이번 포스팅에서는 먼저 kernel과 data transfer를 어떻게 동시에 실행할 수 있는지에 대해 살펴보도록 하겠습니다.

 

Fermi와 Kepler GPU에는 2개의 copy engine queue가 있습니다. 하나는 device로 데이터를 전송할 때 사용하고, 다른 하나는 device로부터 데이터를 내보낼 때 사용됩니다. 그러므로, 한 번에 2개의 데이터 전송을 동시에 수행할 수 있지만, 그 데이터 전송은 서로 방향이 다르고, 다른 스트림에서 dispatch된 경우에만 가능합니다.

또한, 어플리케이션에서 모든 데이터 전송(data transfer)과 커널 실행(kernel execution)간의 관계를 살펴보고 다음의 두 가지 케이스를 구별할 필요가 있습니다.

• 만약 커널이 데이터 A를 사용한다면, A에 대한 데이터 전송은 동일한 스트림에서 커널이 수행되기 전에 수행되어야 한다.
• 만약 커널이 데이터 A의 어떠한 부분도 사용하지 않는다면, 커널 수행과 데이터 전송은 다른 스트림에 위치할 수 있다.

두 번째 케이스에서, 커널과 데이터 전송을 동시에 실행하는 것은 간단합니다. 별도의 스트림에 이들을 배치하면, 동시에 실행하는 것이 안전합니다. 그러나 첫 번째 케이스와 같은 경우에는 조금 복잡합니다. 이어지는 내용에서 커널과 데이터 전송 사이에 종속성이 존재할 때 어떻게 동시에 수행시킬 수 있는지 vector addition 예제를 통해 살펴보겠습니다.

 

Overlap Using Depth-First Scheduling

예제에서 사용할 벡터 덧셈 커널은 다음과 같습니다.

__global__
void sumArrays(float* A, float* B, float* C, const int N)
{
    int idx = blockDim.x * blockIdx.x + threadIdx.x;

    if (idx < N) {
        for (int i = 0; i < NREPEAT; i++)
            C[idx] = A[idx] + B[idx];
    }
}

여기서 vector의 합을 구할 때 for문을 추가했는데, kernel execution 시간을 증가시켜 nvvp로 computation과 communication의 오버랩을 쉽게 살펴보기 위해 추가하였습니다.

 

기본적인 CUDA프로그램에서의 vector addition은 다음의 단계를 포함합니다.

1. Copy the two input vectors from the host to the device
2. Perform a vector addition
3. Copy a single output vector back to the host from the device

 

커널과 데이터 전달을 동시에 수행하려면 조금 달라집니다.

입력과 출력의 데이터 세트를 분할해야하는데, 길이가 N인 두 개의 vector를 더하는 문제가 길이가 N/M의 vector를 더하는 문제로 분할하는 것과 같습니다. 이렇게 분할된 하위 덧셈들은 독립적이기 때문에 이들은 각각 별도의 CUDA 스트림에서 수행될 수 있습니다.

 

일반적인 데이터 전송은 동기(synchronous) 카피 함수입니다. 데이터 전송과 커널 실행을 동시에 수행하기 위해서는 비동기(asynchronous) 카피 함수를 사용해야 합니다. 비동기 카피 함수에는 pinned host memory를 사용해야되기 때문에 먼저 host memory의 할당은 cudaHostAlloc을 사용하여 pinned host memory에 할당해주어야 합니다.

float *h_A, *h_B, *hostRef, *gpuRef;
cudaHostAlloc((void**)&h_A, nBytes, cudaHostAllocDefault);
cudaHostAlloc((void**)&h_B, nBytes, cudaHostAllocDefault);
cudaHostAlloc((void**)&gpuRef, nBytes, cudaHostAllocDefault);
cudaHostAlloc((void**)&hostRef, nBytes, cudaHostAllocDefault);

그리고, 사용할 스트림 갯수만큼 작업할 크기를 나누어 줍니다. 예제 코드에서는 NSTREAM이라는 이름으로 매크로로 설정하여 사용합니다.

int iElem = nElem / NSTREAM;

이제 각 스트림에서 한 번에 iElem개의 원소를 처리하기 위한 communication과 computation을 dispath하는 루프를 작성합니다.

for (int i = 0; i < NSTREAM; i++) {
	int offset = i * iElem;
    cudaMemcpyAsync(&d_A[offset], &h_A[offset],
        iBytes, cudaMemcpyHostToDevice, stream[i]);
    cudaMemcpyAsync(&d_B[offset], &h_B[offset],
        iBytes, cudaMemcpyHostToDevice, stream[i]);
    sumArrays<<<grid, block, 0, stream[i]>>>(&d_A[offset], &d_B[offset], &d_C[offset], iElem);
    cudaMemcpyAsync(&gpuRef[offset], &d_C[offset],
        iBytes, cudaMemcpyDeviceToHost, stream[i]);
}

여기서 메모리 카피와 커널 실행은 host에 대해 비동기이기 때문에, 전체 워크로드는 어떠한 blocking없이 스트림들에게 분배됩니다. 그러나 input vector, kernel computation, output vector 간의 종속성은 동일한 스트림에 배치되므로 그대로 유지됩니다.

 

스트림을 사용하여 vector addition을 수행하는 것과 비교하기 위해서 전체 input을 하나의 스트림으로 수행하는 것도 추가하여 성능을 비교해보도록 하겠습니다.

sumArrays<<<grid, block>>>(d_A, d_B, d_C, nElem);

 

전체 코드는 아래 링크에서 참조하시면 됩니다 !

https://github.com/junstar92/parallel_programming_study/blob/master/CUDA/StreamsAndEvents/simpleMultiAddDepth.cu

GitHub - junstar92/parallel_programming_study: Study parallel programming - CUDA, OpenMP, MPI, Pthread

위 코드를 nvvp로 프로파일링한 결과는 다음과 같습니다.

8개의 hardware work queue를 사용했고, 4개의 CUDA 스트림을 커널과 데이터 전송을 오버랩하기 위해 사용하였습니다. 스트림을 사용한 경우는 사용하지 않았을 때보다 약 38%의 성능 향상을 보여주고 있습니다.

이 결과에서는 3가지 타입의 오버랩을 보여줍니다.

• 다른 스트림에서의 커널들이 서로 오버랩됨
• 다른 스트림에서의 커널과 데이터 전송이 오버랩됨
• 다른 스트림에서의 서로 방향이 다른 데이터 전송이 오버랩됨

또한, 두 종류의 blocking을 보여줍니다.

• 커널은 같은 스트림에서의 데이터 전송에 의해 블락됨
• host에서 device로의 데이터 전송은 같은 방향의 이전 데이터 전송에 의해서 블락됨

비록 host에서 device로의 데이터 전송이 4개의 다른 스트림에서 수행되더라도, 위의 타임라인 결과는 이 데이터 전송들이 시퀀셜하게 수행되는 것을 보여줍니다. 이는 실제로 데이터 전송들이 동일한 copy engine queue를 통해 수행되기 때문입니다.

 

다음으로 hardware work queue의 수를 줄였을 때 성능이 어떤지 살펴보겠습니다. 아래 그림은 hardware work queue의 개수를 한 개로 줄였을 때의 결과입니다.

8개의 hardware work queue를 사용했을 때와 성능에는 큰 차이가 없습니다. 이는 각 스트림이 오직 하나의 커널만을 실행하기 때문에 이로 인한 false dependencies는 발생하지 않기 때문입니다.

 

Overlap Using Breadth-First Scheduling

이전 포스팅에서 Depth-first와 Breadth-frist 순서를 비교한 것처럼 이번에도 Kernel execution과 Data transfer에서 Breadth-first 순서로 수행하면 어떻게 되는지 살펴보겠습니다.

 

코드는 위에서 살펴본 것과 거의 동일하며, 스트림을 사용하여 커널과 데이터 전송을 수행하는 부분에서 아래처럼 변경됩니다.

// initiate all asynchronous transfers to the device
for (int i = 0; i < NSTREAM; i++) {
    int offset = i * iElem;
    cudaMemcpyAsync(&d_A[offset], &h_A[offset],
        iBytes, cudaMemcpyHostToDevice, stream[i]);
    cudaMemcpyAsync(&d_B[offset], &h_B[offset],
        iBytes, cudaMemcpyHostToDevice, stream[i]);
}

// launch a kernel in each stream
for (int i = 0; i < NSTREAM; i++) {
    int offset = i * iElem;
    sumArrays<<<grid, block, 0, stream[i]>>>(&d_A[offset], &d_B[offset], &d_C[offset], iElem);
}

// enqueue asynchronous transfers from the device
for (int i = 0; i < NSTREAM; i++) {
    int offset = i * iElem;
    cudaMemcpyAsync(&gpuRef[offset], &d_C[offset],
        iBytes, cudaMemcpyDeviceToHost, stream[i]);
}

전체 코드는 아래 링크에서 참조하시길 바랍니다 !

https://github.com/junstar92/parallel_programming_study/blob/master/CUDA/StreamsAndEvents/simpleMultiAddBreadth.cu

GitHub - junstar92/parallel_programming_study: Study parallel programming - CUDA, OpenMP, MPI, Pthread

 

위 코드를 컴파일하고 nvvp로 프로파일링하면 다음과 같은 결과를 얻을 수 있습니다.

depth-fisrt와 성능 상의 큰 차이는 없습니다. Kepler 부터는 bidirectional sheduling mechanism이 false dependences를 제거해주기 때문입니다. (Kepler 이후의 GPU도 동일할 것으로 추측합니다.)

만약 Fermi에서 위 코드로 똑같이 테스트해보면, breadth-first의 성능은 depth-first보다 낮을 것입니다. 

 

---

Overlapping GPU and CPU Execution

모든 커널의 수행은 기본적으로 GPU와 비동기이므로, GPU와 CPU의 실행을 오버랩하는 것은 비교적 간단합니다. 따라서 단순히 커널을 실행하고 CPU에서 작업을 계속 수행하면 자동으로 오버랩됩니다.

 

이번에 살펴볼 예제는 다음의 두 가지 부분을 포함합니다.

• A kernel is dispatched to the default stream.
• Host computation is executed while waiting on the GPU kernel.

 

예제에서 사용되는 커널은 단순한 vector-scalar 덧셈을 수행하는 커널입니다.

__global__
void kernel(float* g_data, float value)
{
    int idx = blockDim.x * blockIdx.x + threadIdx.x;
    g_data[idx] = g_data[idx] + value;
}

예제에서는 3개의 CUDA operation(2 copies and 1 kernel launch)가 issue됩니다. 그리고 stop 이벤트는 모든 CUDA operation이 완료되는 지점을 마킹합니다.

CUDA_CHECK(cudaMemcpyAsync(d_a, h_a, nBytes, cudaMemcpyHostToDevice));
kernel<<<grid, block>>>(d_a, value);
CUDA_CHECK(cudaMemcpyAsync(h_a, d_a, nBytes, cudaMemcpyDeviceToHost));
CUDA_CHECK(cudaEventRecord(stop));

위 3개의 operation들은 모두 host에 비동기이며, 디폴트 스트림에 연결됩니다. 마지막 cudaMemcpyAsync가 issue되자마자, 제어권은 즉시 host로 전달됩니다. 제어권이 host로 리턴되면, host는 kernel의 결과에 종속되지 않는 어떠한 연산도 가능합니다.

아래의 코드는 host에서 모든 CUDA operations가 완료될 때까지 while문을 반복시키면서 counter를 증가시킵니다. 각 반복에서 host는 stop 이벤트를 쿼리합니다.

// have CPU do some work while waiting for stage 1 to finish
unsigned long int counter = 0;
while (cudaEventQuery(stop) == cudaErrorNotReady)
    counter++;

 

전체 코드는 아래 링크에서 참조하시길 바랍니다 !

https://github.com/junstar92/parallel_programming_study/blob/master/CUDA/StreamsAndEvents/asyncAPI.cu

GitHub - junstar92/parallel_programming_study: Study parallel programming - CUDA, OpenMP, MPI, Pthread

위 코드를 컴파일하고, nvprof로 실행하면 다음과 같은 결과를 확인할 수 있습니다.

GPU operation의 완료를 기다리는 동안 host는 59,661번의 iteration을 수행했습니다.

 

---

Stream Callbacks

stream callback(스트림 콜백)은 CUDA 스트림 큐에서 대기될 수 있는 또 다른 타입의 작업입니다. 스트림 콜백에 선행하는 모든 operation들이 완료되면 스트림 콜백에 의해 지정된 host측의 함수가 CUDA 런타임에 의해 호출됩니다. 호출되는 함수는 application에 의해 제공되는 함수(host측)입니다. 이를 통해 임의의 host-side의 logic이 CUDA 스트림에 추가될 수 있습니다. 즉, 스트림 콜백은 또 다른 CPU-GPU 동기화 메커니즘입니다. 

 

스트림 콜백 함수는 application에서 제공되는 host 함수이며 스트림에 추가는 다음의 API 함수를 통해 수행됩니다.

cudaError_t cudaStreamAddCallback(cudaStream_t stream,
		cudaStreamCallback_t callback, void* userData, unsigned int flags);

이 함수는 주어진 스트림에 콜백 함수를 추가합니다. 콜백 함수는 스트림에 콜백 함수 이전에 추가된 모든 operation들이 완료된 후에 host에서 수행됩니다. 콜백은 cudaStreamAddCallback당 한 번만 실행되며 콜백 함수가 완료될 때까지 콜백 이후에 스트림에 추가된 다른 작업들은 블락됩니다. CUDA 런타임에 의해 호출될 때, 콜백 함수에는 호출되는 스트림과 CUDA 에러가 발생했는지를 나타내는 에러 코드가 전달됩니다. cudaStreamAddCallback의 userData 파라미터를 사용하면 콜백 함수에 전달할 데이터를 지정할 수도 있습니다. flags는 나중에 사용하도록 예약되었지만, 현재는 의미가 없고 0으로만 설정해주어야 합니다. NULL 스트림에 추가된 콜백은 모든 스트림에서 issue된 선행 작업들이 모두 완료되면 실행됩니다.

 

콜백 함수에는 두 가지 제약사항이 있습니다.

• No CUDA API function can be called from a callback function
• No synchronization can be performed within the callback function

 

예제에서 사용되는 콜백 함수 my_callback은 다음과 같습니다.

void CUDART_CB my_callback(cudaStream_t stream, cudaError_t status, void* data)
{
    printf("callback from stream %d\n", *((int*)data));
}

예제 코드에서는 4개의 커널이 각 스트림에서 순차적으로 issue되고 난 후에 스트림에 issue 됩니다. 이 콜백은 각 스트림에서의 모든 작업이 완료된 이후에 host에서 시작됩니다. 스트림 콜백이 추가되는 코드는 다음과 같습니다.

for (int i = 0; i < n_streams; i++) {
    streams_ids[i] = i;
    kernel_1<<<grid, block, 0, streams[i]>>>();
    kernel_2<<<grid, block, 0, streams[i]>>>();
    kernel_3<<<grid, block, 0, streams[i]>>>();
    kernel_4<<<grid, block, 0, streams[i]>>>();
    cudaStreamAddCallback(streams[i], my_callback, (void*)(stream_ids + i), 0);
}

 

전체 코드는 아래 링크에서 참조하실 수 있습니다 !

https://github.com/junstar92/parallel_programming_study/blob/master/CUDA/StreamsAndEvents/simpleCallback.cu

GitHub - junstar92/parallel_programming_study: Study parallel programming - CUDA, OpenMP, MPI, Pthread

위 코드를 컴파일하고, 실행하면 다음의 결과를 확인하실 수 있습니다.

 

---

Summary

지난 두 번의 포스팅에 이어서 이번 포스팅까지 CUDA의 스트림과 이벤트에 대해서 살펴봤습니다.

스트림에 대한 개념은 CUDA 프로그래밍 모델의 기본 중의 하나입니다. CUDA 스트림은 high-level CUDA operations가 독립적인 스트림의 실행에 추가될 수 있도록 함으로써 coarse-grained concurrency를 가능하게 합니다. CUDA는 비동기 버전의 런타임 함수들을 대부분 지원하기 때문에 여러 CUDA 스트림 간에 computation과 communication을 분배할 수 있습니다.

 

만약 CUDA operations간에 종속성이 있다면, 이 작업들은 동일한 스트림에 스케쥴링되어야 합니다. 예를 들어, 정확성을 보장하기 위해서는 동일한 스트림에서 커널이 사용하는 데이터가 전송된 후에 스케쥴링되어야 합니다. 종속성이 없는 작업들이라면 임의의 스트림에 스케쥴링될 수 있습니다.

CUDA에서는 일반적으로 3가지 타입의 오버랩 방식을 사용하여 computation 이나 communication latency를 숨길 수 있습니다.

• Overlap of multiple, concurrent kernels on the device
• Overlap of CUDA kennels with data transfer to or from the device
• Overlap of CPU execution and GPU execution

 

device를 완전히 활용하고 최대 concurrency를 확보하기 위해서는 아래의 문제들도 고려해야 합니다.