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

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

비밀번호

댓글 쓴 사람
 




[1]  2  3  4  5  6  7  8  9  10  11  12  13  14  15  ...
NoWriterDateCnt.TitleFile(s)
12218정성태5/27/202052.NET Framework: 904. C# - DirectX 게임 클라이언트 실행 중 키보드 입력을 감지하는 방법
12217정성태5/24/202024오류 유형: 615. Transaction count after EXECUTE indicates a mismatching number of BEGIN and COMMIT statements. Previous count = 0, current count = 1.
12216정성태5/15/2020100.NET Framework: 904. USB/IP PROJECT를 이용해 C#으로 USB Keyboard 가상 장치 만들기
12215정성태5/12/2020118개발 환경 구성: 490. C# - (Wireshark의) USBPcap을 이용한 USB 패킷 모니터링파일 다운로드1
12214정성태5/5/2020101개발 환경 구성: 489. 정식 인증서가 있는 경우 Device Driver 서명하는 방법 (2) - UEFI/SecureBoot
12213정성태5/3/2020194개발 환경 구성: 488. (코드로 가상 USB 장치를 만들 수 있는) USB/IP PROJECT 소개
12212정성태5/1/202094개발 환경 구성: 487. UEFI / Secure Boot 상태인지 확인하는 방법
12211정성태4/27/2020206개발 환경 구성: 486. WSL에서 Makefile로 공개된 리눅스 환경의 C/C++ 소스 코드 빌드
12210정성태4/20/2020303.NET Framework: 903. .NET Framework의 Strong-named 어셈블리 바인딩 (1) - app.config을 이용한 바인딩 리디렉션 [1]파일 다운로드1
12209정성태4/13/2020166오류 유형: 614. 리눅스 환경에서 C/C++ 프로그램이 Segmentation fault 에러가 발생한 경우 (2)
12208정성태4/12/2020197Linux: 29. 리눅스 환경에서 C/C++ 프로그램이 Segmentation fault 에러가 발생한 경우
12207정성태4/2/2020189스크립트: 19. Windows PowerShell의 NonInteractive 모드
12206정성태4/2/2020272오류 유형: 613. 파일 잠금이 바로 안 풀린다면? - The process cannot access the file '...' because it is being used by another process.
12205정성태4/2/2020170스크립트: 18. Powershell에서는 cmd.exe의 명령어를 지원하진 않습니다.
12204정성태4/1/2020164스크립트: 17. Powershell 명령어에 ';' (semi-colon) 문자가 포함된 경우
12203정성태3/18/2020359오류 유형: 612. warning: 'C:\ProgramData/Git/config' has a dubious owner: '...'.
12202정성태3/18/2020367개발 환경 구성: 486. .NET Framework 프로젝트를 위한 GitLab CI/CD Runner 구성
12201정성태3/18/2020191오류 유형: 611. git-credential-manager.exe: Using credentials for username "Personal Access Token".
12200정성태3/18/2020463VS.NET IDE: 145. NuGet + Github 라이브러리 디버깅 관련 옵션 3가지 - "Enable Just My Code" / "Enable Source Link support" / "Suppress JIT optimization on module load (Managed only)"
12199정성태3/17/2020187오류 유형: 610. C# - CodeDomProvider 사용 시 Unhandled Exception: System.IO.DirectoryNotFoundException: Could not find a part of the path '...\f2_6uod0.tmp'.
12198정성태3/17/2020198오류 유형: 609. SQL 서버 접속 시 "Cannot open user default database. Login failed."
12197정성태3/17/2020306VS.NET IDE: 144. .NET Core 콘솔 응용 프로그램을 배포(publish) 시 docker image 자동 생성 - 두 번째 이야기
12196정성태3/17/2020285오류 유형: 608. The ServicedComponent being invoked is not correctly configured (Use regsvcs to re-register).
12195정성태3/17/2020376.NET Framework: 902. C# - 프로세스의 모든 핸들을 열람 - 세 번째 이야기
12194정성태3/16/2020224오류 유형: 607. PostgreSQL - Npgsql.NpgsqlException: sorry, too many clients already
[1]  2  3  4  5  6  7  8  9  10  11  12  13  14  15  ...