References
- CUDA Toolkit Documentation - CUDA Runtime API (link)
Contents
- Difference between the driver and runtime APIs
- API Synchronization Behavior
이번 포스팅에서는 CUDA의 API에 대한 동기/비동기 동작에 대해서 살펴보려고 합니다. 아마 기본적으로 Memcpy와 Kernel Launch를 오버랩할 때, 스트림을 사용하여 비동기로 분할하여 수행하면 성능이 더 좋다는 것을 알고 있으리라 생각됩니다. 이에 대한 내용은 이전 포스팅들을 참조하시길 바랍니다 !
아래 포스팅에서는 스트림을 사용하여 Copy/Compute를 오버랩하는 방법에 대한 내용이 포함되어 있습니다.
CUDA C/C++ 기초 - (3)
References Fundamentals of Accelerated Computing with CUDA C/C++ (NVIDIA Online Training) Asynchronous Streaming, and Visual Profiling for Accelerated Applications with CUDA C/C++ Contents Concurren..
junstar92.tistory.com
시작하기에 앞서, CUDA의 Runtime API와 Device API에 대해 먼저 언급하고자 합니다.
CUDA를 사용하다보면 CUDA Runtime(cudart)와 CUDA Device(cuda)의 API가 각각 있는 것을 볼 수 있습니다.
처음 CUDA를 접하거나, 일반적으로 사용할 때는 거의 대부분 runtime API를 사용하는 것 같으며, 저 역시도 Device API는 거의 사용한 적이 없습니다. 이번 포스팅에서는 CUDA Runtime과 Device APIs에 대해서 살펴보도록 하겠습니다.
참고로 nvidia-smi 커맨드를 통해서 CUDA Version을 확인할 수도 있고, nvcc --version 커맨드를 통해서 설치된 CUDA Version을 확인할 수도 있습니다. 일반적으로 이 둘을 구분하지 못하거나, 왜 다른지 모르는 경우가 많습니다. 저 역시도 왜 다른지 잘 몰랐었는데, 이번에 runtime/device APIs를 살펴보면서 알게 되었습니다.
아래의 stackoverflow에 자세하게 설명이 되어 있는데, 요약하자면 다음과 같습니다.
Different CUDA versions shown by nvcc and NVIDIA-smi
I am very confused by the different CUDA versions shown by running which nvcc and nvidia-smi. I have both cuda9.2 and cuda10 installed on my ubuntu 16.04. Now I set the PATH to point to cuda9.2. So...
stackoverflow.com
CUDA에는 2개의 기본 API(runtime과 device)가 있으며, 둘 다 각각 해당하는 버전이 존재합니다. 일반적으로 driver API(ex, linux에서 libcuda.so)는 GPU 드라이버 설치에 의해서 설치되며, runtime API(ex, linux에서 libcudart.so)는 CUDA Toolkit 설치에 의해서 설치됩니다.
CUDA Toolkit을 설치할 때 GPU 드라이버도 설치할 수 있는데, 여기서 GPU 드라이버를 같이 설치하게 되면 driver와 runtime의 API 버전은 동일합니다. 하지만 일반적으로 GPU 드라이버를 먼저 설치하고 CUDA Toolkit을 설치하는 경우가 많기 때문에 두 버전이 항상 일치하지 않을 수 있습니다.
대부분의 경우, nvidia-smi에서 표시되는 CUDA Version이 nvcc에서 확인한 CUDA Version보다 같거나 높은 경우에는 사용하는데 문제가 없습니다. 이는 CUDA에서 정의된 호환성 덕분인데, 최신 드라이버는 이전 CUDA Toolkit 또는 Runtime API를 지원하기 때문입니다. 예를 들어, nvidia-smi에서 확인된 CUDA Version은 11.2이고, nvcc에서 확인된 CUDA Version이 10.2인 경우에 문제가 되지 않습니다.
결론은, nvidia-smi와 nvcc에서 표시하는 CUDA Version은 다를 수 있으며, 대부분 정상적으로 동작합니다.
Difference between the driver and runtime APIs
driver와 runtime APIs는 매우 유사하여, 대부분의 경우 상호호환이 가능합니다.
다만, 이들 사이에는 몇 가지 주요한 차이점이 있습니다.
Complexity vs. Control
runtime API는 implicit initialization, context management, module management를 통해 device code 관리를 쉽게 해줍니다. 이 때문에 코드가 심플해집니다.
반면, driver API는 더 세분화된 제어를 제공하는데, 특히 context 및 module 로드에 대해 그렇습니다. 커널의 실행(kernel launch)은 명시적인 함수 호출과 함께 execution configuration과 커널의 매개변수가 지정되어야 하므로, 구현하기가 훨씬 더 복잡합니다. 그러나, 모든 커널이 초기화 중에 자동으로 로드되고 프로그램이 실행하는 동안 로드된 상태를 유지하는 runtime과 달리, driver API를 사용하면 현재 로드된 모듈만 유지하거나 모듈을 동적으로 다시 로드할 수 있습니다. driver API는 언어에 독립적인데, cubin 객체만 다룹니다.
Context management
컨텍스트의 관리는 driver API를 통해서 수행할 수 있습니다.
반면, runtime API에서는 컨텍스트 관리가 노출되지 않습니다. 대신, runtime API는 스레드에 사용할 컨텍스트를 직접 결정합니다. 만약 컨텍스트가 driver API를 통해 호출 스레드(calling thread)에 대해 최신 상태라면, runtime은 이 컨텍스트 이용합니다. 하지만, 이러한 컨텍스트가 없다면 기본적으로 'primary context'(기본 컨텍스트)를 사용합니다. 기본 컨텍스트는 필요에 따라 프로세스당 디바이스 장치에 하나씩 생성되고 이는 참조 카운트로 관리됩니다. 만약 더 이상 참조하지 않게 되면 컨텍스트는 소멸됩니다.
한 프로세스에서, 각 스레드에서 컨텍스트가 최신 상태가 아니라면 runtime API의 모든 유저는 기본 컨텍스트를 공유합니다. runtime이 사용하는 컨텍스트, 즉, 현재 컨텍스트 또는 기본 컨텍스트는 cudaDeviceSynchronize()로 동기화할 수 있고, cudaDeviceReset()을 통해 소멸될 수 있습니다.
runtime API를 기본 컨텍스트와 함께 사용하면 장단점이 있는데, 예를 들어, 모든 플러그인이 동일한 프로세스에서 실행되는 경우, 모든 플러그인이 하나의 컨텍스트를 공유하기 때문에 서로 통신할 수 없습니다. 플러그인 중 하나가 모든 CUDA 작업을 끝낸 후, cudaDeviceReset()을 호출하면 다른 플러그인에서는 사용하던 컨텍스트가 소멸되었기 때문에 작업이 실패할 수 있습니다.
위 문제를 피하려면 사용자는 driver API를 사용하여 현재 컨텍스트를 생성하고 설정한 다음, runtime API를 사용하여 작업을 수행할 수 있습니다. 그러나 컨텍스트는 device 메모리, 추가적인 호스트 스레드, device에서 컨텍스트 스위칭으로 코스트와 같은 상당한 리소스를 소비할 수 있습니다. 이러한 runtime-driver 컨텍스트 공유는 cuBLAS나 cuFFT와 같은 runtime API에 내장된 라이브러리와 함께 driver API를 사용할 때 중요합니다.
API Synchronization Behavior
API는 동기/비동기 memcpy/memset 함수를 모두 제공합니다. API 이름 뒤에 "Async"가 붙은 함수가 바로 비동기 버전입니다. 사실 각 함수에 전달되는 인수에 따라서 동기/비동기 동작이 달라지므로, 뒤에 Async가 붙었다고 무조건 비동기로 동작하는 것은 아닙니다.
그리고, runtime과 device API에서의 동기화 동작은 동일합니다.
Memcpy
동기(synchoronous) 버전 함수들의 동작에 대해 하나하나 살펴보도록 하겠습니다.
1. Unified Memory를 포함하는 모든 전송은 각각 host에 완전히 동기화된다
이를 테스트하기 위해 간단한 코드를 작성했습니다. 아래 코드는 단순히 a와 b라는 메모리를 할당하고 a에서 b로 memcpy를 수행합니다. 이때, a는 unified memory로 할당하고, b는 device memory로 할당하였습니다. 그리고 CUDA Runtime API 마다 NVTX Range Push/Pop을 통해 실제 host에서 마킹이 찍히는 위치를 살펴볼 수 있도록 하였습니다. 여기서 우리는 "marking"이라는 NVTX Marker가 어디에 찍히는지를 살펴보도록 하겠습니다.
#include <nvtx3/nvToolsExt.h>
int main()
{
int n = 1 << 22;
double *a, *b;
nvtxRangePush("cudaMalloc");
cudaMallocManaged(&a, sizeof(double) * n);
nvtxRangePop();
nvtxRangePush("cudaMalloc");
cudaMalloc(&b, sizeof(double) * n);
nvtxRangePop();
nvtxRangePush("cudaMemcpy");
cudaMemcpy(b, a, sizeof(double)*n, cudaMemcpyDefault);
nvtxMarkA("marking");
nvtxRangePop();
cudaFree(a);
cudaFree(b);
}
"marking"이 찍힌 위치를 보면, 실제 (CUDA HW에서의)cudaMemcpy가 완전히 수행되고 난 후라는 것을 알 수 있습니다. 즉, 따로 명시적으로 동기화를 해주지 않아도 unified memory가 포함되는 경우에는 host에 완전히 동기화된다는 것을 알 수 있습니다. 반대의 경우, 즉, a를 device memory, b를 unified memory로 할당한 후 동일한 테스트를 수행했을 때에도 결과는 동일하였습니다.
a는 unified memory, b를 pinned memory로 할당한 경우에도 동일한 테스트를 수행해보았습니다.
#include <nvtx3/nvToolsExt.h>
int main()
{
int n = 1 << 22;
double *a, *b;
nvtxRangePush("cudaMalloc");
cudaMallocManaged(&a, sizeof(double) * n);
nvtxRangePop();
nvtxRangePush("cudaMalloc");
cudaMallocHost(&b, sizeof(double) * n);
nvtxRangePop();
nvtxRangePush("cudaMemcpy");
cudaMemcpy(b, a, sizeof(double)*n, cudaMemcpyDefault);
nvtxMarkA("marking");
nvtxRangePop();
cudaFree(a);
cudaFreeHost(b);
}
결과는 동일하게 "marking"이 복사가 완전히 끝난 후 마킹되는 것을 확인할 수 있습니다.
2. Pageable Host Memory에서 Device Memory로의 전송은 복사가 시작되기 전에 스트림 동기화가 수행된다
Device Memory로 DMA(direct memory access) 전송을 위해서 pageable buffer가 staging memory로 복사되면, 함수는 바로 리턴됩니다. 하지만, 최종 목적지로의 DMA가 아직 완료되지 않을 수 있습니다.
이번에는 a는 (pageable)host memory, b는 device memory로 할당하여 동일한 테스트를 수행합니다.
#include <nvtx3/nvToolsExt.h>
int main()
{
int n = 1 << 22;
double *a, *b;
nvtxRangePush("cudaMalloc");
a = (double*)malloc(sizeof(double) * n);
nvtxRangePop();
nvtxRangePush("cudaMalloc");
cudaMalloc(&b, sizeof(double) * n);
nvtxRangePop();
nvtxRangePush("cudaMemcpy");
cudaMemcpy(b, a, sizeof(double)*n, cudaMemcpyDefault);
nvtxMarkA("marking");
nvtxRangePop();
free(a);
cudaFree(b);
}결과는 다음과 같습니다.
1번의 경우와는 다르게, HW에서 cudaMemcpy가 종료되지 않았는데, host에서 "marking"이 마킹된 것을 확인할 수 있습니다. 여기서 자세히 알 수는 없지만, 문서에서 설명한 대로 pageable buffer가 staging memory로 복사되자마자 리턴되어 host가 계속해서 진행한다는 것을 짐작할 수 있습니다.
반대의 경우, a가 device memory, b가 host memory(pageable)인 경우에는 이 경우에 속하지 않기 때문에 위와 다른 결과가 나옵니다. 이 경우는 4번에 해당되는데, 4번에서 그 결과를 살펴보도록 하겠습니다.
3. Pinned Host Memory에서 Device Memory로 전송하는 경우, 함수는 host에 대해 동기화가 수행된다
이번에는 a를 pinned memory, b를 device memory로 할당하여 테스트해보겠습니다.
#include <nvtx3/nvToolsExt.h>
int main()
{
int n = 1 << 22;
double *a, *b;
nvtxRangePush("cudaMalloc");
cudaMallocHost(&a, sizeof(double) * n);
nvtxRangePop();
nvtxRangePush("cudaMalloc");
cudaMalloc(&b, sizeof(double) * n);
nvtxRangePop();
nvtxRangePush("cudaMemcpy");
cudaMemcpy(b, a, sizeof(double)*n, cudaMemcpyDefault);
nvtxMarkA("marking");
nvtxRangePop();
cudaFreeHost(a);
cudaFree(b);
}
"marking"이 (CUDA HW에서의) memcpy가 완료된 이후에 마킹된 것을 확인할 수 있습니다.
반대의 경우는 4번에 해당되므로 4번에서 살펴보도록 하겠습니다.
4. Device로부터 pageable 또는 pinned host memory로 전송하는 경우, 함수는 오직 복사가 완료된 후에 반환된다
2번에서 언급한 a가 device memory, b가 host memory(pageable)인 경우가 바로 여기에 속합니다.
#include <nvtx3/nvToolsExt.h>
int main()
{
int n = 1 << 22;
double *a, *b;
nvtxRangePush("cudaMalloc");
cudaMalloc(&a, sizeof(double) * n);
nvtxRangePop();
nvtxRangePush("cudaMalloc");
b = (double*)malloc(sizeof(double) * n);
nvtxRangePop();
nvtxRangePush("cudaMemcpy");
cudaMemcpy(b, a, sizeof(double)*n, cudaMemcpyDefault);
nvtxMarkA("marking");
nvtxRangePop();
cudaFree(a);
free(b);
}그 결과는 다음과 같습니다.
2번의 경우와 데이터 전송의 방향만 다르지만, "marking"은 복사가 완료된 이후에 마킹되는 것을 볼 수 있습니다.
b를 pinned memory로 할당하여 테스트한 결과는 다음과 같습니다.
#include <nvtx3/nvToolsExt.h>
int main()
{
int n = 1 << 22;
double *a, *b;
nvtxRangePush("cudaMalloc");
cudaMalloc(&a, sizeof(double) * n);
nvtxRangePop();
nvtxRangePush("cudaMalloc");
cudaMallocHost(&b, sizeof(double) * n);
nvtxRangePop();
nvtxRangePush("cudaMemcpy");
cudaMemcpy(b, a, sizeof(double)*n, cudaMemcpyDefault);
nvtxMarkA("marking");
nvtxRangePop();
cudaFree(a);
cudaFreeHost(b);
}
마찬가지로 memcpy가 완전히 종료된 이후에 "marking"이 마킹된 것을 확인할 수 있습니다.
5. Device에서 Device로 전송하는 경우, host 측의 동기화는 수행되지 않는다
이번엔 a, b를 모두 device memory로 할당하여, 테스트합니다.
#include <nvtx3/nvToolsExt.h>
int main()
{
int n = 1 << 22;
double *a, *b;
nvtxRangePush("cudaMalloc");
cudaMalloc(&a, sizeof(double) * n);
nvtxRangePop();
nvtxRangePush("cudaMalloc");
cudaMalloc(&b, sizeof(double) * n);
nvtxRangePop();
nvtxRangePush("cudaMemcpy");
cudaMemcpy(b, a, sizeof(double)*n, cudaMemcpyDefault);
nvtxMarkA("marking");
nvtxRangePop();
cudaFree(a);
cudaFree(b);
}
이번에는 cudaMemcpy API가 호출되자마자 반환되어 "marking"이 마킹된 것을 확인할 수 있습니다. 하지만, 실제 CUDA HW에서의 cudaMemcpy는 "marking"이 찍히고 난 이후에 수행되고 있습니다. 따라서, device to device로의 메모리 전송은 host에 비동기로 수행된다는 것을 알 수 있습니다.
6. 임의의 host memory에서 다른 host memory로의 전송의 경우, 함수는 host에 대해 완전히 동기화된다
a, b를 모두 pinned memory로 할당하여 테스트해봤습니다.
#include <nvtx3/nvToolsExt.h>
int main()
{
int n = 1 << 22;
double *a, *b;
nvtxRangePush("cudaMalloc");
cudaMallocHost(&a, sizeof(double) * n);
nvtxRangePop();
nvtxRangePush("cudaMalloc");
cudaMallocHost(&b, sizeof(double) * n);
nvtxRangePop();
nvtxRangePush("cudaMemcpy");
cudaMemcpy(b, a, sizeof(double)*n, cudaMemcpyDefault);
nvtxMarkA("marking");
nvtxRangePop();
cudaFreeHost(a);
cudaFreeHost(b);
}
예상했겠지만, host to host 메모리 전달은 CUDA HW가 관여하지 않으므로 CUDA Trace 결과 자체가 없습니다. 따라서, host에 완전히 동기화되어 수행된다는 것을 짐작할 수 있습니다.
이번에는 비동기(asynchronous) 버전 함수들의 동작에 대해 살펴보도록 하겠습니다. aysnc API 함수들을 사용해야 하므로, stream을 생성하고, cudaMemcpyAsync API 함수에 생성한 stream을 전달하여 호출하도록 하였습니다.
1. Device로부터 pageable host memory로 전송하는 경우, 함수는 복사가 완료된 이후에 반환된다
a는 device memory, b는 pageable host memory로 할당하여 테스트를 수행하였습니다.
#include <nvtx3/nvToolsExt.h>
int main()
{
int n = 1 << 22;
double *a, *b;
cudaStream_t stream;
cudaStreamCreate(&stream);
nvtxRangePush("cudaMalloc");
cudaMalloc(&a, sizeof(double) * n);
nvtxRangePop();
nvtxRangePush("cudaMalloc");
b = (double*)malloc(sizeof(double) * n);
nvtxRangePop();
nvtxRangePush("cudaMemcpy");
cudaMemcpyAsync(b, a, sizeof(double)*n, cudaMemcpyDefault, stream);
nvtxMarkA("marking");
nvtxRangePop();
cudaFree(a);
free(b);
cudaStreamDestroy(stream);
}
"marking"이 복사가 완전히 완료된 이후에 "marking"이 마킹되는 것을 확인할 수 있습니다.
반대의 경우, 즉, a는 host memory(pagable), b는 device memory로 할당하는 경우는 3번에 해당하므로 비동기로 수행된다는 것을 추측할 수 있습니다. 실제로 비동기로 동작하는지 체크해보도록 하겠습니다.
#include <nvtx3/nvToolsExt.h>
int main()
{
int n = 1 << 22;
double *a, *b;
cudaStream_t stream;
cudaStreamCreate(&stream);
nvtxRangePush("cudaMalloc");
a = (double*)malloc(sizeof(double) * n);
nvtxRangePop();
nvtxRangePush("cudaMalloc");
cudaMalloc(&b, sizeof(double) * n);
nvtxRangePop();
nvtxRangePush("cudaMemcpy");
cudaMemcpyAsync(b, a, sizeof(double)*n, cudaMemcpyDefault, stream);
nvtxMarkA("marking");
nvtxRangePop();
free(a);
cudaFree(b);
cudaStreamDestroy(stream);
}
예상한 대로 비동기로 수행되기 때문에 "marking"이 CUDA HW에서의 cudaMemcpy가 끝나기도 전에 마킹되는 것을 확인할 수 있습니다.
2. 어떤 host memory로부터 다른 host memory로 전송하는 경우, 함수는 host에 대해 완전히 동기화된다
동기화 버전의 6번과 동일하게 a, b 모두 pinned memory로 할당하여 테스트를 수행했습니다.
#include <nvtx3/nvToolsExt.h>
int main()
{
int n = 1 << 22;
double *a, *b;
cudaStream_t stream;
cudaStreamCreate(&stream);
nvtxRangePush("cudaMalloc");
cudaMallocHost(&a, sizeof(double) * n);
nvtxRangePop();
nvtxRangePush("cudaMalloc");
cudaMallocHost(&b, sizeof(double) * n);
nvtxRangePop();
nvtxRangePush("cudaMemcpy");
cudaMemcpyAsync(b, a, sizeof(double)*n, cudaMemcpyDefault, stream);
nvtxMarkA("marking");
nvtxRangePop();
cudaFreeHost(a);
cudaFreeHost(b);
cudaStreamDestroy(stream);
}
host to host로의 메모리 전송이기 때문에 CUDA HW가 연관하지 않으므로 host에 완전히 동기화되어 동작한다는 것을 알 수 있습니다.
3. 1,2번을 제외한 다른 전송에서 함수는 완전히 비동기이다. 만약 pageable memory가 먼저 pinned memory로 스테이징되어야 하는 경우 work thread와 비동기적으로 처리된다
1,2번을 제외한 모든 전송에서는 비동기로 동작하는데, a를 pinned memory로 할당하고 b를 device memory로 할당하여 테스트해보겠습니다. 1번과 메모리 전송이 반대 방향입니다.
#include <nvtx3/nvToolsExt.h>
int main()
{
int n = 1 << 22;
double *a, *b;
cudaStream_t stream;
cudaStreamCreate(&stream);
nvtxRangePush("cudaMalloc");
cudaMallocHost(&a, sizeof(double) * n);
nvtxRangePop();
nvtxRangePush("cudaMalloc");
cudaMalloc(&b, sizeof(double) * n);
nvtxRangePop();
nvtxRangePush("cudaMemcpy");
cudaMemcpyAsync(b, a, sizeof(double)*n, cudaMemcpyDefault, stream);
nvtxMarkA("marking");
nvtxRangePop();
cudaFreeHost(a);
cudaFree(b);
cudaStreamDestroy(stream);
}
"marking"이 cudaMemcpy가 끝나기 전에 마킹된 것을 볼 수 있습니다. 즉, 이 경우에는 비동기로 동작한다는 것을 확인할 수 있습니다.
Memset
동기 memset 함수들은 target이 Pinned Host Memory 또는 Unified Memory인 경우를 제외하고는 host에 대해 비동기적입니다. Pinned Host Memory나 Unified Memory를 사용할 때에는 완전히 host에 동기화됩니다. 비동기 버전의 함수는 항상 host에 대해 비동기적입니다.
테스트를 몇 번 시도해봤는데, 단순한 방법으로는 Nsight System으로 확인이 어렵고 콘솔에 출력하는 방법을 사용해야 될 것으로 보입니다... 문서에서 설명한 대로 잘 동작할 것이라고 믿으며 넘어가도록 하겠습니다 (사실 조금 귀찮았습니다...).
Kernel Launches
커널 실행은 host에 대해 항상 비동기적입니다. 이는 커널 실행에서 적용되는 기본적인 성질입니다.
'NVIDIA > CUDA' 카테고리의 다른 글
| TensorRT CMake 스크립트 (FindTENSORRT.cmake) (2) | 2023.07.24 |
|---|---|
| Stream Synchronization Behavior (0) | 2022.06.19 |
| NVIDIA Tools Extension (NVTX) (0) | 2022.06.17 |
| Multiple GPUs with CUDA C++ (0) | 2022.06.16 |
| CUDA C/C++ 기초 - (3) (0) | 2022.06.14 |
댓글