Streams and Events (2) - Concurrent Kernels

References

• Professional CUDA C Programming

Contents

• Concurrent Kernel Execution
• False Dependencies on Fermi Device
• Adjusting Stream Behavior using Environment Variables
• Blocking Behavior of the Default Stream
• Overlapping Kernel Execution and Data Transfer

---

지난 포스팅에서 CUDA의 Stream과 Event에 대해서 알아봤습니다.

Streams and Events (1)

이번에는 여러 예제를 통해서 실제로 어떻게 동작하는지 살펴보는 시간을 갖도록 하겠습니다.

 

Concurrent Kernels in Non-NULL Streams

처음에 살펴볼 내용은 non-null 스트림으로 concurrent kernel의 실행입니다. 이 예제에서는 간단한 더미 커널 4개를 사용하는데, 커널들의 구현은 모두 동일하며 다음과 같이 구현됩니다. (kernel_1, kernel_2, kernel_3, kernel_4)

__global__
void kernel_1()
{
    double sum = 0.0;

    for (int i = 0; i < N; i++) {
        sum = sum + tan(0.1) * tan(0.1);
        printf("%f\n", sum);
    }
}

처음 테스트할 때 printf 함수가 없었는데, 요즘 GPU가 너무 좋다보니 커널을 launch하는 API 실행시간보다 커널의 수행시간이 너무 짧아서 제대로 확인할 수 없는 상황이 발생했습니다. 아무리 N을 크게해도 API 실행시간보다 짧아서 printf를 추가하여 커널의 수행시간을 조금 증가시켰습니다.

 

host 측에서 먼저 non-null 스트림 세트를 생성합니다. 각 스트림에서 issue되는 kernel launch는 하드웨어 리소스에 의한 False Dependencies가 없다면 GPU에서 동시에 실행됩니다.

cudaStream_t *streams = (cudaStream_t*)malloc(n_streams * sizeof(cudaStream_t));

for (int i = 0; i < n_streams; i++) {
    CUDA_CHECK(cudaStreamCreate(&(streams[i])));
}

 

커널 함수들은 for 루프를 통해 각 스트림에서 dispatch됩니다.

// dispatch job with depth first ordering
for (int i = 0; i < n_streams; 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]>>>();
}

execution configuration으로 grid와 block의 크기는 커맨드 인자가 없다면 둘 다 1로 설정됩니다.

 

그리고 이 예제의 수행 시간을 측정하기 위해서 두 개의 이벤트도 생성해주었습니다.

cudaEvent_t start, stop;
CUDA_CHECK(cudaEventCreate(&start));
CUDA_CHECK(cudaEventCreate(&stop));

start 이벤트는 디폴트 스트림에서 루프가 시작되기 전에 레코드됩니다. stop 이벤트는 모든 커널의 수행이 끝난 이후에 디폴트 스트림에서 레코드됩니다.

// record start event
CUDA_CHECK(cudaEventRecord(start, 0));

// dispatch job with depth first ordering
for (int i = 0; i < n_streams; 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]>>>();
}

// record stop event
CUDA_CHECK(cudaEventRecord(stop, 0));
CUDA_CHECK(cudaEventSynchronize(stop));

// calculate elapsed time
CUDA_CHECK(cudaEventElapsedTime(&elapsed_time, start, stop));

그리고 stop 이벤트에 동기화시킨 후, 수행시간을 계산할 수 있습니다.

 

전체 코드는 아래 링크에서 확인하실 수 있습니다.

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

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

 

시각적으로 어떻게 수행되는지 살펴보기 위해서 nsight system을 통해 프로파일링해보도록 하겠습니다.

각 스트림에서 수행되는 커널은 위와 같이 실행되는 것을 볼 수 있습니다. 예상대로 4개의 다른 스트림에서 4개의 concurrent kernel 실행이 확인됩니다.

 

nvvp로도 똑같은 결과를 확인하실 수 있습니다. 개인적으로 nvvp가 시각적으로 조금 더 잘 표현하고 있는 것 같습니다.

 

---

False Dependencies on Fermi GPUs

제가 가진 GPU로는 확인을 할 수 없지만, Fermi 디바이스에서 동일한 코드를 실행하면 결과는 다릅니다. 이전 포스팅에서 이야기했듯이 Fermi에서는 Hyper-Q를 지원하지 않기 때문에 결과적으로 제한된 concurrency로 실행됩니다.

 

아래 그림은 동일한 코드를 Fermi 디바이스에서 실행한 결과를 보여줍니다.

4개의 스트림이 동일한 hardware work queue를 공유하고 있기 때문에 이로 인한 false dependency에 의해서 동시에 시작되지 않습니다. 하지만 stream i+1의 첫 번째 task와 stream i의 마지막 task는 동시에 실행되는 것을 볼 수 있습니다. 이는 두 task가 다른 스트림에 있어서 서로 의존성이 없기 때문입니다. stream i의 마지막 task가 실행될 때 CUDA 런타임은 work queue의 다음 task(stream i+1의 첫 번째 task)를 준비합니다. 각 스트림의 첫 번째 task는 어떠한 stream에도 의존성이 없기 때문에 즉시 실행될 수 있습니다.

 

False Dependency는 host로부터 dispatch되는 커널의 순서에 의해서 발생합니다. 방금 살펴본 예제 코드는 depth-first 접근방법을 사용하여 다음 스트림에서 launch하기 전에 하나의 스트림에서 모든 operation들을 launch합니다.

 

Fermi Device에서 false dependency를 피하기 위해서 host로부터 breadth-first 접근으로 커널을 dispatch하면 됩니다.

// dispatch job with breadth first ordering
for (int i = 0; i < n_streams; i++) 
    kernel_1<<<grid, block, 0, streams[i]>>>();
for (int i = 0; i < n_streams; i++) 
    kernel_2<<<grid, block, 0, streams[i]>>>();
for (int i = 0; i < n_streams; i++) 
    kernel_3<<<grid, block, 0, streams[i]>>>();
for (int i = 0; i < n_streams; i++) 
    kernel_4<<<grid, block, 0, streams[i]>>>();

breadth-first 순서를 사용하면 work queue에서 인접한 task들이 서로 다른 스트림에 속합니다. 따라서, 인접한 task들간의 false dependency가 없어지고 concurrent kernel execution이 가능합니다.

breadth-first 순으로 변경하고 Fermi 디바이스에서 실행해보면 다음의 결과를 확인할 수 있습니다.

 

전체 코드는 아래 링크에서 확인하실 수 있습니다.

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

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

 

---

Adjusting Stream Behavior Using Envrionment Variables

Hyper-Q를 지원하는 GPU는 host와 각 GPU 간의 여러 개의 하드웨어 work queues를 유지합니다. Kepler device에서 지원하는 최대 work queue의 수는 32입니다. 그러나 기본적으로 concurrent hardware connection의 수는 8개로 제한되어 있습니다. 각 connection에는 추가 메모리와 리소스가 필요하기 때문에 기본 8개로 설정하면 32개의 work queue가 필요하지 않은 어플리케이션에서 리소스 소비량을 줄일 수 있습니다.

최대 connection 개수는 CUDA_DEVICE_MAX_CONNECTIONS 환경 변수를 사용하여 concurrent hardware connections를 조정할 수 있습니다. Kepler device에서는 32개까지 설정할 수 있습니다.

 

환경 변수를 설정하는 방법은 여러 가지가 있는데, 아마 기본적인 방법들은 모두 알고 계실거라고 생각합니다.

한 가지 코드에서 직접 수정할 수 있는 방법이 있는데, 이는 위에서 살펴본 simpleHyperQDepth.cu의 main 코드 중에서 확인하실 수 있습니다.

// set up max connection
char* iname = "CUDA_DEVICE_MAX_CONNECTIONS";
_putenv_s(iname, "32");
char* ivalue = getenv(iname);
printf("%s = %s\n", iname, ivalue);

바로 위와 같이 작성된 부분입니다. 여기서는 CUDA_DEVICE_MAX_CONNECTIONS라는 이름의 환경변수를 설정하고 그 값을 32로 설정합니다.

이 값을 아래처럼 4로 설정하고,

_putenv_s(iname, "4");

8개의 스트림을 사용하도록 매크로로 설정된 NSTREAM의 값을 8로 변경하여,

#define NSTREAM 8

실행한 결과를 살펴보겠습니다.

8개의 스트림이 있지만, 오직 4개만 동시에 실행되는 것을 볼 수 있습니다. 오직 4개의 device connection만 사용하도록 설정했기 때문에 2개의 스트림이 하나의 queue를 공유합니다. 하지만, 하나의 queue에서 stream i의 마지막 task와 stream i+4의 첫 번째 task는 의존성이 없기 때문에 동시에 실행되는 것으로 예상했지만, 예상과는 다르게 결과가 나왔습니다. 

 

이번에는 breadth-first 순서로 connection의 값과 스트림 개수를 동일하게 설정하여 실행해보도록 하겠습니다.

결과는 위와 같습니다.

breadth-first 순서로 각 스트림의 첫 번째 task들을 동시에 실행하여 8개의 kernel_1이 동시에 실행될 것으로 예상했지만, 결과는 달랐습니다. 

 

결과적으로 추측해보면, CUDA_DEVICE_MAX_CONNECTIONS는 하드웨어 work queue를 제한하기도 하지만, 최대 동시에 실행할 수 있는 concurrency도 제한하는 것 같습니다.

 

---

Blocking Behavior of the Default Stream

이번에는 디폴트 스트림이 어떻게 non-null 스트림에서의 operation을 블락하는지 살펴보겠습니다.

이를 살펴보기 위해 처음에 살펴봤던 simpleHyperQDepth.cu의 코드에서 kernel_3의 호출을 디폴트 스트림에서 수행하도록 수정합니다.

// dispatch job with depth first ordering
for (int i = 0;i < n_streams; i++) {
    kernel_1<<<grid, block, 0, streams[i]>>>();
    kernel_2<<<grid, block, 0, streams[i]>>>();
    kernel_3<<<grid, block>>>();
    kernel_4<<<grid, block, 0, streams[i]>>>();
}

세번째 커널이 디폴트 스트림에서 실행되기 때문에, 이후에 non-null 스트림에 의해 issue되는 operation들은 디폴트 스트림의 operation이 완료될 때까지 블락됩니다. 

 

---

Creating Inter-Stream Dependencies

이상적으로 스트림 간에는 의도하지 않은 의존성이 없어야 합니다. 그러나 복잡한 어플리케이션에서는 다른 스트림에서의 작업이 완료될 때까지 한 스트림에서의 작업을 블락하는 스트림 간 종속성이 유용할 때도 있습니다.

이벤트는 스트림 간의 종속성을 추가하는데 사용할 수 있습니다.

 

한 스트림에서의 작업이 오직 다른 스트림의 모든 작업이 완료된 이후에 시작되기를 원한다고 가정해봅시다.

이벤트를 사용하여 스트림 간의 종속성을 만들 수 있는데, 먼저 동기화만을 위한 이벤트를 cudaEventDisableTiming 플래그를 사용하여 생성합니다.

cudaEvent_t *kernelEvent = (cudaEvent_t*)malloc(n_streams * sizeof(cudaEvent_t));
for (int i = 0; i < n_streams; i++) {
	cudaEventCreateWithFlags(&kernelEvent[i], cudaEventDisableTiming);
}

다음으로 cudaEventRecord를 사용하여 각 스트림이 완료될 때 이벤트를 레코딩합니다. 그런 다음, cudaStreamWaitEvent를 사용하여 마지막 스트림(streams[n_stream-1])이 다른 모든 스트림을 기다리도록 합니다.

for (int i = 0; i < n_streams; 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]>>>();
    
    cudaEventRecord(kernelEvent[i], streams[i]);
    cudaStreamWaitEvent(streams[n-streams-1], kernelEvent[i], 0);
}

전체 코드는 아래 링크에서 확인하실 수 있습니다.

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

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

 

위 코드를 컴파일하고, nvvp로 프로파일링하면 다음의 결과를 확인할 수 있습니다.