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

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)
11531정성태6/5/20185591오류 유형: 468. JSON.parse가 허용하지 않는 문자 [9]
11530정성태6/10/20183712.NET Framework: 755. C# 7.2 - 스택에만 생성할 수 있는 값 타입 지원 - "ref struct" [2]파일 다운로드1
11529정성태5/23/20182914.NET Framework: 754. 닷넷의 관리 포인터(Managed Pointer)와 System.TypedReference [5]파일 다운로드1
11528정성태6/10/20182861.NET Framework: 753. C# 7.2 - 3항 연산자에 ref 지원(conditional ref operator) [1]
11527정성태5/17/20182643오류 유형: 467. RDP 로그인 에러 - This could be due to CredSSP encryption oracle remediation.
11526정성태6/10/20182412.NET Framework: 752. C# 7.2 - 메서드의 반환값 및 로컬 변수에 ref readonly 기능 추가파일 다운로드1
11525정성태6/10/20182589.NET Framework: 751. C# 7.2 - 메서드의 매개 변수에 in 변경자 추가파일 다운로드1
11524정성태6/10/20183548.NET Framework: 750. C# 7.2 - readonly 구조체 [3]파일 다운로드1
11523정성태5/15/20182961.NET Framework: 749. C# - 값 형식의 readonly 인스턴스에 대한 메서드 호출 시 defensive copy 발생 [1]파일 다운로드1
11522정성태5/15/20181884개발 환경 구성: 378. Azure - VM 진단 설정 화면의 "This subscription is not registered with the Microsoft.Insights resource provider."
11521정성태5/15/20181892개발 환경 구성: 377. Azure - 원하는 성능 데이터로 모니터링 대시보드 구성
11520정성태5/12/20182471.NET Framework: 748. C# 7.1 - 참조 어셈블리(Ref Assemblies)
11519정성태3/13/20192670개발 환경 구성: 376. ASP.NET Web Application 프로젝트의 FileSystem 배포(Publish) 시 Before/After Task 설정 방법 [1]
11518정성태5/11/20182345.NET Framework: 747. C# 7.0에서도 부분적으로 가능해진 "타입 추론을 통한 튜플의 변수명 자동 지정"
11517정성태5/10/20182627.NET Framework: 746. Azure runbook 예제 - 6시간 동안 수행 중인 VM을 중지 [1]파일 다운로드1
11516정성태5/9/20182007.NET Framework: 745. Azure runbook을 PowerShell 또는 C# 코드로 실행하는 방법파일 다운로드1
11515정성태6/8/20182359.NET Framework: 744. C# 6 - Expression bodied function
11514정성태5/3/20183800오류 유형: 466. Bitvise - Error in component session/transport/kexHandler [2]
11513정성태5/10/20184600.NET Framework: 743. C# 언어의 공변성과 반공변성 [5]파일 다운로드1
11512정성태5/3/20182116개발 환경 구성: 375. Azure runbook 실행 시 "Errors", "All Logs"에 오류 메시지가 출력되는 경우
11511정성태5/2/20182354개발 환경 구성: 374. Azure - Runbook 기능 소개
11510정성태4/30/20184000.NET Framework: 742. windbg로 확인하는 Finalizer를 가진 객체의 GC 과정파일 다운로드1
11509정성태4/28/20183505.NET Framework: 741. windbg로 확인하는 객체의 GC 여부
11508정성태4/23/20182724개발 환경 구성: 373. MSBuild를 이용해 프로젝트 배포 후 결과물을 zip 파일로 압축하는 방법파일 다운로드1
11507정성태9/29/20182923개발 환경 구성: 372. MSBuild - 빌드 전/후, 배포 전/후 실행하고 싶은 Task 정의
11506정성태6/27/20185686.NET Framework: 740. C#에서 enum을 boxing 없이 int로 변환하기 - 두 번째 이야기 [7]파일 다운로드1
... 16  17  18  19  20  21  22  23  24  25  26  27  [28]  29  30  ...