CUDA & OpenCV Programming 한신대학교 대학원 컴퓨터공학과 류승택 2015. Spring

프로그래밍 환경 구축 CUDA & OpenCV 설치 CUDA Toolkit 설치 OpenCV 설치 : 참조

OpenCV 설정 OpenCV & Visual Studio File -> New -> Project 프로젝트 생성 콘솔 응용 프로그래밍 (빈프로젝트) 인쿠르드와 라이브러리 추가 구성 속성 – VC++ 디렉토리 포함 디렉터리 c:\opencv\build\include 라이브러리 디렉토리 32비트용 c:\opencv\build\x86\vc12\lib 64비트용 c:\opencv\build\x64\vc12\lib

OpenCV 설정 OpenCV & Visual Studio 추가 종속성에 라이브러리 등록 구성 속성 – 링커 – 입력 추가종속성: opencv_core2411.lib, opencv_highgui2411.lib, opencv_imgproc2411.lib

CUDA 설정 CUDA & Visual Studio 연동 : 참조 File -> New -> Project 프로젝트 생성 콘솔 응용 프로그래밍 (빈프로젝트) 만들어진 프로젝트에서 Build Customizations를 클릭!

CUDA 설정 CUDA & Visual Studio 연동 CUDA의 확장자인 cu와 헤더파일인 cuh를 C++과 같이 취급하도록 하기위해서 확장자를 등록 Tools -> Option : Text Editor -> File Extension HelloCUDA.cu 파일에서 속성 설정

CUDA 설정 CUDA & Visual Studio 연동 CUDA 라이브러리 추가 프로젝트 속성 : Linker -> Input : Additional Dependencies cudart.lib를 추가 #include <stdio.h> #include <cuda.h> #include <cuda_runtime.h> #include "device_launch_parameters.h“ __global__ void kernel(void) { }   int main(void) {         kernel<<<1,1>>>();         printf("Hello, CUDA!\n");     return 0; }

CUDA & OpenCV 연동 CUDA & OpenCV 연동 #include<iostream> #include<cstdio> #include<opencv2/core/core.hpp> #include<opencv2/highgui/highgui.hpp> #include <cuda.h> #include <cuda_runtime.h> #include "device_launch_parameters.h" using std::cout; using std::endl; using namespace cv; static inline void _safe_cuda_call(cudaError err, const char* msg, const char* file_name, const int line_number) { if (err != cudaSuccess) { fprintf(stderr, "%s\n\nFile: %s\n\nLine Number: %d\n\nReason: %s\n", msg, file_name, line_number, cudaGetErrorString(err)); std::cin.get(); exit(EXIT_FAILURE); } #define SAFE_CALL(call, msg) _safe_cuda_call((call), (msg), __FILE__, __LINE__)

CUDA & OpenCV 연동 CUDA & OpenCV 연동 int main() { std::string imagePath = "image.jpg"; //Read input image from the disk Mat input = cv::imread(imagePath, CV_LOAD_IMAGE_COLOR); if (input.empty()) { std::cout << "Image Not Found!" << std::endl; std::cin.get(); return -1; } //Create output image Mat output(input.rows, input.cols, CV_8UC3); //Call the wrapper function convert_to_gray(input, output); //Show the input and output imshow("Input", input); imshow("Output", output); //Wait for key press waitKey(); return 0;

CUDA & OpenCV 연동 CUDA & OpenCV 연동 void convert_to_gray(const cv::Mat& input, cv::Mat& output) { //Calculate total number of bytes of input and output image const int colorBytes = input.step * input.rows; const int grayBytes = output.step * output.rows; unsigned char *d_input, *d_output; //Allocate device memory (디바이스 메모리 할당) SAFE_CALL(cudaMalloc<unsigned char>(&d_input, colorBytes), "CUDA Malloc Failed"); SAFE_CALL(cudaMalloc<unsigned char>(&d_output, grayBytes), "CUDA Malloc Failed"); //Copy data from OpenCV input image to device memory (호스트에서 디바이스로 메모리 복사) SAFE_CALL(cudaMemcpy(d_input, input.ptr(), colorBytes, cudaMemcpyHostToDevice), "CUDA Memcpy Host To Device Failed"); //Specify a reasonable block size const dim3 block(16, 16); //Calculate grid size to cover the whole image const dim3 grid((input.cols + block.x - 1) / block.x, (input.rows + block.y - 1) / block.y); //Launch the color conversion kernel (CUDA 커널 함수 호출) bgr_to_gray_kernel << <grid, block >> >(d_input, d_output, input.cols, input.rows, input.step, output.step); //Synchronize to check for any kernel launch errors SAFE_CALL(cudaDeviceSynchronize(), "Kernel Launch Failed"); //Copy back data from destination device meory to OpenCV output image (디바이스에서 호스트로 결과메모리 복사) SAFE_CALL(cudaMemcpy(output.ptr(), d_output, grayBytes, cudaMemcpyDeviceToHost),"CUDA Memcpy Host To Device Failed"); //Free the device memory (디바이스 메모리 해제) SAFE_CALL(cudaFree(d_input), "CUDA Free Failed"); SAFE_CALL(cudaFree(d_output), "CUDA Free Failed"); }

CUDA & OpenCV 연동 CUDA & OpenCV 연동 __global__ void bgr_to_gray_kernel(unsigned char* input, unsigned char* output, int width, int height, int colorWidthStep, int grayWidthStep) { //2D Index of current thread const int xIndex = blockIdx.x * blockDim.x + threadIdx.x; const int yIndex = blockIdx.y * blockDim.y + threadIdx.y; //Only valid threads perform memory I/O if ((xIndex<width) && (yIndex<height)) { //Location of colored pixel in input const int color_tid = yIndex * colorWidthStep + (3 * xIndex); //Location of gray pixel in output const int gray_tid = yIndex * grayWidthStep + (3 * xIndex); const unsigned char blue = input[color_tid]; const unsigned char green = input[color_tid + 1]; const unsigned char red = input[color_tid + 2]; const float gray = red * 0.3f + green * 0.59f + blue * 0.11f; const float o_r = gray; const float o_g = gray; const float o_b = gray; output[gray_tid] = static_cast<unsigned char>(o_b); output[gray_tid+1] = static_cast<unsigned char>(o_g); output[gray_tid+2] = static_cast<unsigned char>(o_r); }

OpenCV with GPU Building OpenCV with GPU support Prerequisites OpenCV sources CMake (http://www.cmake.org/) CMake-GUI NVIDIA Display Driver NVIDIA GPU Computing Toolkit (for CUDA) https://developer.nvidia.com/cuda-toolkit And your favorite IDE/compiler

OpenCV with GPU Building OpenCV with GPU support Build steps (screenshots for Windows 8, Visual Studio) Run CMake GUI and set source and build directories, press Configure and select you compiler to generate project for. The base folder of OpenCV source code Where to put compiled OpenCV library

OpenCV with GPU Build steps Run CMake GUI and set source and build directories, press Configure and select you compiler to generate project for. Enable WITH_CUDA flag and ensure that CUDA Toolkit is detected correctly by checking all variables with ‘CUDA_’ prefix. Press Configure and Generate to generate a project On Windows, open the Visual Studio solution and click on “Build Solution”. 소요시간: 2~3시간 이상 OpenCV with CUDA Prebuilt 참조

OpenCV with GPU OpenCV CPU example #include <opencv2/opencv.hpp> using namespace cv; int main() { Mat src = imread("image.jpg", 0); if (!src.data) exit(1); Mat dst; bilateralFilter(src, dst, -1, 50, 7); Canny(dst, dst, 35, 200, 3); imwrite("out.png", dst); return 0; } <= OpenCV header files <= OpenCV C++ namespace <= Load an image file as grayscale <= Allocate a temp output image  <= Blur the image but keep edges sharp <= Find the edges, drawn as white pixels <= Store to an image file

OpenCV with GPU OpenCV GPU example #include <opencv2/opencv.hpp> #include <opencv2/gpu/gpu.hpp> using namespace cv; int main() { Mat src = imread("image.jpg", 0); if (!src.data) exit(1); gpu::GpuMat d_src(src); gpu::GpuMat d_dst; gpu::bilateralFilter(d_src, d_dst, -1, 50, 7); gpu::Canny(d_dst, d_dst, 35, 200, 3); Mat dst(d_dst); imwrite("out.png", dst); return 0; } <=OpenCV GPU header file <=Upload image from CPU to GPU memory <=Allocate a temp output image on the GPU <=Process images on the GPU <=Download image from GPU to CPU mem

OpenCV with GPU

Calculate Frame Rate Calculate Frame Rate Header & Library #include <Windows.h> timeGetTime() 함수 프로젝트 설정의 "추가종속성"에 "winmm.lib"를 추가 지금까지 흐른 시간을 1/1000 초단위로 나타내는 함수 FPS를 표시해 줄 변수 정의 문자열 버퍼 크기 정의 const int MAX_STR_BUFFER_SIZE = 64; char strBuffer[MAX_STR_BUFFER_SIZE] = { 0, }; 시작 시간 unsigned long startTime = 0; 전체 Frame 수 long nmrTotalFrames = 0;

Calculate Frame Rate Calculate Frame Rate 시간 정보 초기화 startTime = timeGetTime(); nmrTotalFrames = 0; FPS 계산 nmrTotalFrames++; float fps = (float)((nmrTotalFrames * 1000.0) / (timeGetTime() - startTime)); FPS를 이미지 버퍼에 출력한다. sprintf_s(strBuffer, "%.2lf fps", fps); putText(dst, strBuffer, cvPoint(10, 30), CV_FONT_NORMAL, 1, Scalar(255, 255, 255), 1, 1);

Calculate Frame Rate #include <Windows.h> const int MAX_STR_BUFFER_SIZE = 64; // 문자열 버퍼 크기 정의 char strBuffer[MAX_STR_BUFFER_SIZE] = { 0, }; unsigned long startTime = 0; // 시작 시간 long nmrTotalFrames = 0; // 전체 프레임 수 int main() { startTime = timeGetTime(); // 시간정보 초기화 nmrTotalFrames = 0; for (;;) { // ========= 수행할 알고리즘 (중략) ================// nmrTotalFrames++; // 전체 프레임 수 계산 float fps = (float)((nmrTotalFrames * 1000.0) / (timeGetTime() - startTime)); // FPS 계산 sprintf_s(strBuffer, "%.2lf fps", fps); // FPS를 이미지 버퍼에 출력 putText(dst, strBuffer, cvPoint(10, 30), CV_FONT_NORMAL, 1, Scalar(255, 255, 255), 1, 1); imshow("Destination", dst); } return 0;