CUDA를 지원하는 이종 라이브러리들의 stream 객체 공유

개요

• Nvidia GPU를 CUDA를 이용해 효율적으로 활용하도록 개발하면서 pipelining 설계를 진행하는 경우가 많다. 이 때, 필연적으로 stream을 사용하게 되며, 이를 어떤식으로 사용하느냐에 따라 처리 속도 등의 효율성이 달라진다.

• 최근, 특정 솔루션을 개발하며 CUDA를 지원하는 다양한 라이브러리들을 동시에 사용하게 되었다. 라이브러리들 별로 stream 객체를 다루고 처리하는 방식이 조금씩 다르며, 여러 stream을 각각의 라이브러리가 따로 사용하는 경우가 있고, 효율성을 높이기 위해 하나의 stream을 여러 라이브러리가 동시에 공유해야 하는 경우도 있었다.

• 해당 포스팅은 CUDA를 지원하는 여러 라이브러리가 동일 stream을 공유하여 사용하고자 할 때, 라이브러리별 접근 방식 및 특징들을 정리하고자 작성하였으며, Native CUDA(CUstream, cudaStrema_t)에서 생성한 stream 객체를 베이스로 OpenCV, libtorch, Npp(NVIDIA Performance Primitives), Nvidia Video Codec SDK 각각의 라이브러리가 공유하는 방식에 대해 기술하였다.

---

C/C++/NVCC

1. 스트림 생성

a. CUDA (Native)

• Base Stream 생성
  • CUDA Context를 생성

    #include <cuda.h>
      
    ::cuInit(0);
      
    int nGpuNumber = 0; // Enter the GPU number you want to use.
    CUdevice cuDevice;
    ::cuDeviceGet(&cuDevice, nGpuNumber);
      
    CUcontext cuContext;
    ::cuCtxCreate(cuContext, 0, cuDevice);

  • CUDA stream 생성(해당 객체를 여러 라이브러리에서 공유)

    CUstream cuStream;
      
    // 앞서 생성한 cuda context를 활성화 한 후 생성
    CUresult cuRet = 0;
      
    ::cuCtxPushCurrent(cuContext);
    cuRet = ::cuStreamCreate(&cuStream, CU_STREAM_DEFAULT);
    ::cuCtxPopCurrent();

    
    

b. Npp (NVIDIA Performance Primitives)

• Base Stream 공유

  • 해당 라이브러리는 stream 관리를 위해 NppStreamContext 형태의 객체로 cuda stream을 handling함. 기존 사용중인 native stream(CUstream, cudaStrema_t)을 공유하려면 다음과 같이 생성

    #include <npp.h>
    
    NppStreamContext nppStreamCtx;
    nppStreamCtx.hStream = cuStream;

    
    

c. libtorch

• Base Stream 공유
  • 별도로 at::cuda::CUDAStream객체를 사용하여 cuda stream을 관리. 기존 사용 중인 native stream(CUstream)을 공유하려면 다음과 같이 생성

    #include <ATen/ATen.h>
    #include <ATen/cuda/CUDAContext.h>
    #include <ATen/cuda/CUDAEvent.h>
      
    at::cuda::CUDAStream torchStream;
    torchStream = at::cuda::getStreamFromExternal(cuStream, nGpuNumber);

    
    

d. Nvidia Video Codec SDK

• Base Stream 공유
  • 기존 native stream(CUstream, cudaStream_t) 원형 그대로 사용
  • 해당 라이브러리의 경우 stream등록 시, 특정 Callback 함수 내에서 stream 객체를 등록

    // callback 함수 선언 및 등록
    static int CUDAAPI handlePictureDisplayProc(void *pUserData, CUVIDPARSERDISPINFO *pDispInfo) 
       { return (static_cast<Decoder*>(pUserData))->handlePictureDisplay(pDispInfo); }
    int handlePictureDisplay(CUVIDPARSERDISPINFO *pDispInfo);
     
    // implementation
    int Decoder::handlePictureDisplay(CUVIDPARSERDISPINFO *pDispInfo)
    {
        CUVIDPROCPARAMS cuvProcessingParameters = {};
     
        // blah blah ...
        cuvProcessingParameters.output_stream = cuStream;
        // blah blah ...
    }

    
    

e. OpenCV

• Base Stream 공유
  • 별도로 cv::cuda::stream 객체를 사용하여 cuda stream을 관리. 기존 사용 중인 native stream(CUstream, cudaStream_t)을 공유하려면 다음과 같이 생성

    #include <opencv2/core/cuda.hpp>
    #include <opencv2/core/cuda_stream_accessor.hpp>
     
    cv::cuda::Stream cvStream;
    cvStream = cv::cuda::StreamAccessor::getStream(cuStream);

    
    

2. 활용

a. CUDA(Native)

• 각 kernel 함수의 parameter로 CUstream 객체를 전달하는 방식

  // 예시. stream을 지원하는 kernel 함수는 대부분의 경우 parameter에 입력하여 전달하도록 설계되어 있음.
  cudaMemcpy2DAsync(
      dstPtr,
      memPitch,
      srcPtr,
      static_cast<size_t>(nWidth * nChannel),
      static_cast<size_t>(nWidth * nChannel),
      static_cast<size_t>(nHeight),
      cudaMemcpyHostToDevice, 
      cuStream);

• synchronize

   // 단일 적용
  ::cuStreamSynchronize(cuStream);
  
  // 전체 stream 적용
  ::cudaDeviceSynchronize();

  
  

b. Npp (NVIDIA Performance Primitives)

• CUDA의 kernel 함수와 비슷하며 parameter로 NppStreamContext 객체를 전달

  // 예시. stream을 지원하는 npp함수는 대부분의 경우 parameter에 입력하여 전달하도록 설계되어 있음.
  nppiNV12ToBGR_8u_P2C3R_Ctx(
      pSrcPtr,
      nSrcPitch,
      pDstPtr,
      nWidthStep,
      nppiSize,
      nppStreamCtx);

• synchronize

  // 단일 적용
  ::cuStreamSynchronize(nppStreamCtx.hStream);
  
  // 전체 stream 적용
  ::cudaDeviceSynchronize();

  
  

c. libtorch

• pytorch의 사용 방법과 유사하며, RAII Pattern의 특징을 이용

  // 일반적인 RAII pattern 방식으로 사용
  {
      at::cuda::CUDAStreamGuard guard(torchStream);
  
      // 해당 boundary 내에서 libtorch 작업 작성
      // 예시
      torch::Tensor pTensor = torch::full({nHeight, nWidth, nImageWidthStep}, 128, options);
      pTensor = pTensor.permute({2, 0, 1}).contiguous();
  }

• synchronize

  // 단일 적용
  torchStream.synchronize();
  
  // 전체 stream 적용
  torch::cuda::synchronize();

  
  

d. Nvidia Video Codec SDK

• stream 등록 이후, 별도의 처리가 필요 없음

• synchronize

  // 단일 적용
  ::cuStreamSynchronize(cuStream);
  
  // 전체 stream 적용
  ::cudaDeviceSynchronize();

  
  

e. OpenCV

• CUDA의 kernel 함수와 비슷하며 parameter로 cv::cuda::Stream 객체를 전달

  // 예시. stream을 지원하는 OpenCV 함수는 대부분의 경우 parameter에 입력하여 전달하도록 설계되어 있음.
  cv::blur(gmSrc, gmDst, cv::Size(5,5), cv::Point(-1,-1), cvStream);

• synchronize

  # 단일 적용
  cvStream.synchronize();
  
  # 전체 stream 적용
  ::cudaDeviceSynchronize();

  
  

3. 제거

a. CUDA Native

• CUDA native stream객체(CUstream, cudaStream_t) 제거

  ::cuStreamDestroy(cuStream);

b. 나머지 라이브러리

• 현재 기술한 라이브러리들의 CUDA native stream 객체(CUstream, cudaStream_t)를 공유하는 wrapper 객체들(NppStreamContext, at::cuda::CUDAStream 등)은 이미 RAII 형태로 관리되므로 native stream만 제거하면 다른 작업은 불필요함