Microsoft MVP성태의 닷넷 이야기
VC++: 125. CUDA로 작성한 RGB2RGBA 성능 [링크 복사], [링크+제목 복사]
조회: 3090
글쓴 사람
홈페이지
첨부 파일

CUDA로 작성한 RGB2RGBA 성능

지난 글에서,

C# - OpenCvSharp 사용 시 C/C++을 이용한 속도 향상 (for 루프 연산)
; https://www.sysnet.pe.kr/2/0/11422

OpenCV의 CvtColor(ColorConversionCodes.BGR2BGRA) 호출에 대해 C++/parallel_for로 성능을 유사하게 구현한 적이 있습니다. 마찬가지로, SIMD를 이용해 OpenCV의 erode 연산을 해보기도 했습니다.

내가 만든 코드보다 OpenCV의 속도가 월등히 빠른 이유
; https://www.sysnet.pe.kr/2/0/11423

아쉽게도 SIMD 연산의 경우 RGB2RGBA 연산에는 적용할 수 없었는데요. CUDA의 경우 kernel 함수가 SIMD보다는 더 유연하기 때문에 RGB2RGBA 같은 연산을 구현하는 것이 가능한데, 아래의 코드가 바로 그것입니다.

__global__ void rgb2rgba(int n, BYTE *srcPtr, BYTE *dstPtr)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;

    while (tid < n)
    {
        int srcPos = tid * 3;
        int dstPos = tid * 4;

        dstPtr[dstPos + 0] = srcPtr[srcPos + 0];
        dstPtr[dstPos + 1] = srcPtr[srcPos + 1];
        dstPtr[dstPos + 2] = srcPtr[srcPos + 2];
        dstPtr[dstPos + 3] = 0xff;

        tid += (blockDim.x * gridDim.x);
    }
}

위의 kernel 함수를 C#에서 호출할 수 있도록 다음과 같이 export 함수를 하나 만들어 주고,

__declspec(dllexport) BOOL RGB2RGBA_Cuda(BYTE *srcPtr, BYTE *dstPtr, int width, int height)
{
    BYTE *cudaSrc = nullptr;
    BYTE *cudaDst = nullptr;

    int srcSize = width * height * 3; // RGB 3bytes
    int dstSize = width * height * 4; // RGBA 4bytes

    BOOL ret = FALSE;

    do
    {
        cudaError_t cudaStatus = cudaMalloc((void **)&cudaSrc, srcSize);
        if (cudaStatus != cudaSuccess)
        {
            break;
        }

        cudaStatus = cudaMalloc((void **)&cudaDst, dstSize);
        if (cudaStatus != cudaSuccess)
        {
            break;
        }

        cudaStatus = cudaMemcpy(cudaSrc, srcPtr, srcSize, cudaMemcpyHostToDevice);
        if (cudaStatus != cudaSuccess)
        {
            break;
        }

        rgb2rgba<<<64, 64>>>(width * height, cudaSrc, cudaDst);

        cudaStatus = cudaGetLastError();
        if (cudaStatus != cudaSuccess)
        {
            break;
        }

        cudaStatus = cudaDeviceSynchronize();
        if (cudaStatus != cudaSuccess) 
        {
            break;
        }

        cudaStatus = cudaMemcpy(dstPtr, cudaDst, dstSize, cudaMemcpyDeviceToHost);
        if (cudaStatus != cudaSuccess)
        {
            break;
        }

        ret = TRUE;
    } while (false);

    if (cudaSrc != nullptr)
    {
        cudaFree(cudaSrc);
    }

    if (cudaDst != nullptr)
    {
        cudaFree(cudaDst);
    }

    return ret;
}

테스트해 보면, 100회 연산에 2초 넘는 시간이 걸립니다. 즉, "C# - OpenCvSharp 사용 시 C/C++을 이용한 속도 향상 (for 루프 연산)" 글에서 성능 테스트한 것 중에 (C# 제외하고) 가장 안 좋은 기록이 나온 것입니다. (아직 제가 CUDA 초보자라 더 빠르게 할 수 있는 방법이 있는지는 모르겠습니다.)

성능이 낮은 이유는, RAM에 있는 데이터를 GPU의 메모리로 복사하고 그 결과를 다시 RAM으로 복사하는 오버헤드가 있기 때문입니다.

따라서, CUDA를 이용해 성능 향상을 이루고 싶다면 메모리 복사에 따른 오버헤드를 극복할 정도의 복잡한 kernel 연산이거나, 아니면 CPU를 쉬게 하면서 GPU에 다중으로 작업을 맡기는 경우에만 쓰는 것이 좋겠습니다.

(첨부 파일은 "C# - OpenCvSharp 사용 시 C/C++을 이용한 속도 향상 (for 루프 연산)" 글의 예제에 CUDA 테스트를 포함합니다.)




[이 글에 대해서 여러분들과 의견을 공유하고 싶습니다. 틀리거나 미흡한 부분 또는 의문 사항이 있으시면 언제든 댓글 남겨주십시오.]

[연관 글]





[최초 등록일: ]
[최종 수정일: 3/21/2018 ]

Creative Commons License
이 저작물은 크리에이티브 커먼즈 코리아 저작자표시-비영리-변경금지 2.0 대한민국 라이센스에 따라 이용하실 수 있습니다.
by SeongTae Jeong, mailto:techsharer@outlook.com

비밀번호

댓글 쓴 사람
 




... 16  17  18  19  20  21  22  23  24  25  26  27  28  29  [30]  ...
NoWriterDateCnt.TitleFile(s)
11492정성태4/13/20182398개발 환경 구성: 360. Azure Active Directory의 사용자 도메인 지정 방법
11491정성태9/30/20192661개발 환경 구성: 359. Azure 가상 머신에 Web Application을 배포하는 방법
11490정성태4/12/20183354.NET Framework: 739. .NET Framework 4.7.1의 새 기능 - Configuration builders [1]파일 다운로드1
11489정성태4/12/20182578오류 유형: 463. 윈도우 백업 오류 - a Volume Shadow Copy Service operation failed.
11488정성태4/12/20182396오류 유형: 462. Unhandled Exception in Managed Code Snap-in - FX:{811FD892-5EB4-4E73-A147-F1E079E36C4E}
11487정성태4/12/20182639디버깅 기술: 115. windbg - 닷넷 메모리 덤프에서 정적(static) 필드 값을 조사하는 방법
11486정성태4/11/20182649오류 유형: 461. Error MSB4064 The "ComputeOutputOnly" parameter is not supported by the "VsTsc" task
11485정성태4/11/20184743.NET Framework: 738. C# - Console 프로그램이 Ctrl+C 종료 시점을 감지하는 방법파일 다운로드1
11484정성태4/11/20183770.NET Framework: 737. C# - async를 Task 타입이 아닌 사용자 정의 타입에 적용하는 방법파일 다운로드1
11483정성태4/10/20186424개발 환경 구성: 358. "Let's Encrypt"에서 제공하는 무료 SSL 인증서를 IIS에 적용하는 방법 (2) [1]
11482정성태4/10/20183276VC++: 126. CUDA Core 수를 알아내는 방법
11481정성태4/10/20185484개발 환경 구성: 357. CUDA의 인덱싱 관련 용어 - blockIdx, threadIdx, blockDim, gridDim
11480정성태10/19/20184094.NET Framework: 736. C# - API를 사용해 Azure에 접근하는 방법 [2]파일 다운로드1
11479정성태4/12/20182652.NET Framework: 735. Azure - PowerShell로 Access control(IAM)에 새로운 계정 만드는 방법
11478정성태4/9/20182935디버깅 기술: 115. windbg - 덤프 파일로부터 PID와 환경 변수 등의 정보를 구하는 방법
11477정성태4/8/20183081오류 유형: 460. windbg - sos 명령어 수행 시 c0000006 오류 발생
11476정성태4/9/20182491디버깅 기술: 114. windbg - !threads 출력 결과로부터 닷넷 관리 스레드(System.Threading.Thread) 객체를 구하는 방법
11475정성태3/28/20183389디버깅 기술: 113. windbg - Thread.Suspend 호출 시 응용 프로그램 hang 현상에 대한 덤프 분석
11474정성태3/27/20182460오류 유형: 459. xperf: error: TEST.Event: Invalid flags. (0x3ec).
11473정성태3/22/20183917.NET Framework: 734. C# - Thread.Suspend 호출 시 응용 프로그램 hang 현상파일 다운로드1
11472정성태3/22/20183497개발 환경 구성: 356. GTX 1070, GTX 960, GT 640M의 cudaGetDeviceProperties 출력 결과
11471정성태3/20/20183090VC++: 125. CUDA로 작성한 RGB2RGBA 성능파일 다운로드1
11470정성태3/20/20185381오류 유형: 458. Visual Studio - CUDA 프로젝트 빌드 시 오류 C1189, expression must have a constant value
11469정성태3/19/20182986오류 유형: 457. error MSB3103: Invalid Resx file. Could not load file or assembly 'System.Windows.Forms, ...' or one of its dependencies.
11468정성태3/19/20183020오류 유형: 456. 닷넷 응용 프로그램 실행 시 0x80131401 예외 발생
11467정성태3/19/20182546오류 유형: 455. Visual Studio Installer - 업데이트 실패
... 16  17  18  19  20  21  22  23  24  25  26  27  28  29  [30]  ...