성태의 닷넷 이야기
홈 주인
모아 놓은 자료
프로그래밍
질문/답변
사용자 관리
사용자
메뉴
아티클
외부 아티클
유용한 코드
온라인 기능
MathJax 입력기
최근 덧글
[정성태] Java - How to use the Foreign Funct...
[정성태] 제가 큰 실수를 했군요. ^^; Delegate를 통한 Bein...
[정성태] Working with Rust Libraries from C#...
[정성태] Detecting blocking calls using asyn...
[정성태] 아쉽게도, 커뮤니티는 아니고 개인 블로그입니다. ^^
[정성태] 질문이 잘 이해가 안 됩니다. 우선, 해당 소스코드에서 ILis...
[양승조
] var대신 dinamic으로 선언해서 해결은 했습니다. 맞는 해...
[양승조
] 또 막혔습니다. ㅠㅠ var list = props[i].Ge...
[양승조
] 아. 감사합니다. 어제는 안됐던것 같은데....정신을 차려야겠네...
[정성태] "props[i].GetValue(props[i])" 코드에서 ...
글쓰기
제목
이름
암호
전자우편
HTML
홈페이지
유형
제니퍼 .NET
닷넷
COM 개체 관련
스크립트
VC++
VS.NET IDE
Windows
Team Foundation Server
디버깅 기술
오류 유형
개발 환경 구성
웹
기타
Linux
Java
DDK
Math
Phone
Graphics
사물인터넷
부모글 보이기/감추기
내용
<div style='display: inline'> <h1 style='font-family: Malgun Gothic, Consolas; font-size: 20pt; color: #006699; text-align: center; font-weight: bold'>CUDA로 작성한 RGB2RGBA 성능</h1> <p> 지난 글에서,<br /> <br /> <pre style='margin: 10px 0px 10px 10px; padding: 10px 0px 10px 10px; background-color: #fbedbb; overflow: auto; font-family: Consolas, Verdana;' > C# - OpenCvSharp 사용 시 C/C++을 이용한 속도 향상 (for 루프 연산) ; <a target='tab' href='http://www.sysnet.pe.kr/2/0/11422'>http://www.sysnet.pe.kr/2/0/11422</a> </pre> <br /> OpenCV의 CvtColor(ColorConversionCodes.BGR2BGRA) 호출에 대해 C++/parallel_for로 성능을 유사하게 구현한 적이 있습니다. 마찬가지로, SIMD를 이용해 OpenCV의 erode 연산을 해보기도 했습니다.<br /> <br /> <pre style='margin: 10px 0px 10px 10px; padding: 10px 0px 10px 10px; background-color: #fbedbb; overflow: auto; font-family: Consolas, Verdana;' > 내가 만든 코드보다 OpenCV의 속도가 월등히 빠른 이유 ; <a target='tab' href='http://www.sysnet.pe.kr/2/0/11423'>http://www.sysnet.pe.kr/2/0/11423</a> </pre> <br /> 아쉽게도 SIMD 연산의 경우 RGB2RGBA 연산에는 적용할 수 없었는데요. CUDA의 경우 kernel 함수가 SIMD보다는 더 유연하기 때문에 RGB2RGBA 같은 연산을 구현하는 것이 가능한데, 아래의 코드가 바로 그것입니다.<br /> <br /> <pre style='margin: 10px 0px 10px 10px; padding: 10px 0px 10px 10px; background-color: #fbedbb; overflow: auto; font-family: Consolas, Verdana;' > __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); } } </pre> <br /> 위의 kernel 함수를 C#에서 호출할 수 있도록 다음과 같이 export 함수를 하나 만들어 주고,<br /> <br /> <pre style='margin: 10px 0px 10px 10px; padding: 10px 0px 10px 10px; background-color: #fbedbb; overflow: auto; font-family: Consolas, Verdana;' > __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; } <span style='color: blue; font-weight: bold'>rgb2rgba<<<64, 64>>>(width * height, cudaSrc, cudaDst);</span> 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; } </pre> <br /> 테스트해 보면, 100회 연산에 2초 넘는 시간이 걸립니다. 즉, "<a target='tab' href='http://www.sysnet.pe.kr/2/0/11422'>C# - OpenCvSharp 사용 시 C/C++을 이용한 속도 향상 (for 루프 연산)</a>" 글에서 성능 테스트한 것 중에 (C# 제외하고) 가장 안 좋은 기록이 나온 것입니다. (아직 제가 CUDA 초보자라 더 빠르게 할 수 있는 방법이 있는지는 모르겠습니다.)<br /> <br /> 성능이 낮은 이유는, RAM에 있는 데이터를 GPU의 메모리로 복사하고 그 결과를 다시 RAM으로 복사하는 오버헤드가 있기 때문입니다.<br /> <br /> 따라서, CUDA를 이용해 성능 향상을 이루고 싶다면 메모리 복사에 따른 오버헤드를 극복할 정도의 복잡한 kernel 연산이거나, 아니면 CPU를 쉬게 하면서 GPU에 다중으로 작업을 맡기는 경우에만 쓰는 것이 좋겠습니다.<br /> <br /> (<a target='tab' href='http://www.sysnet.pe.kr/bbs/DownloadAttachment.aspx?fid=1236&boardid=331301885'>첨부 파일</a>은 "<a target='tab' href='http://www.sysnet.pe.kr/2/0/11422'>C# - OpenCvSharp 사용 시 C/C++을 이용한 속도 향상 (for 루프 연산)</a>" 글의 예제에 CUDA 테스트를 포함합니다.)<br /> </p><br /> <br /><hr /><span style='color: Maroon'>[이 글에 대해서 여러분들과 의견을 공유하고 싶습니다. 틀리거나 미흡한 부분 또는 의문 사항이 있으시면 언제든 댓글 남겨주십시오.]</span> </div>
첨부파일
스팸 방지용 인증 번호
2207
(왼쪽의 숫자를 입력해야 합니다.)