#include <opencv2/opencv.hpp>
int main()
{
cv::VideoCapture cap("Cars_On_Highway.mp4");
if (!cap.isOpened())
{
std::cerr << "Error opening video stream or file" << std::endl;
return -1;
}
int fps = static_cast<int>(cap.get(cv::CAP_PROP_FPS));
cv::Ptr<cv::BackgroundSubtractor> pBgSub = cv::createBackgroundSubtractorMOG2(500, 16.0, false);
// history: 모델이 배경으로 간주하기 위해 고려하는 프레임 수. 기본값은 500.
// 최초 500 프레임이 배경 모델을 구축하는 데 사용되고, 이후 프레임은 모델 업데이트에 영향을 미친다.
// 그러므로 최초에는 배경 모델이 구축될 수 있도록 충분한 프레임이 필요하다. history가 너무 작으면 모델이
// 빠르게 업데이트되어 일시적인 변화에 민감해질 수 있고, 너무 크면 모델이 느리게 업데이트되어 변화에
// 적응하는 데 시간이 걸릴 수 있다.
// varThreshold: 픽셀이 배경 모델과 일치하는지 여부를 결정하는 분산 임계값. 기본값은 16.0.
// 분산 임계값이 낮을수록 모델과 일치하는 픽셀이 더 많아지고, 높을수록 모델과 일치하는 픽셀이 더 적어진다.
// detectShadows: 그림자 감지를 활성화할지 여부를 결정하는 부울 값. 기본값은 true.
// 이 예에서는 detectShadows만 false로 설정하여 그림자 감지를 비활성화했다.
// 배경은 0으로 표시되고 전경은 255, 그림자는 128로 표시된다.
// 그림자는 배경 모델이 그림자를 배경으로 간주하지 않도록 도와주지만, 그림자 감지를 활성화하면 속도가 느려질 수 있다.
cv::Mat frame;
cv::Mat fgMask;
while (true)
{
cap >> frame;
if (frame.empty())
break;
cv::resize(frame, frame, cv::Size(), 0.5, 0.5);
pBgSub->apply(frame, fgMask);
// apply 메서드는 프레임을 입력으로 받아 fgMask에 전경 마스크를 출력한다.
// fgMask는 8비트 단일 채널 이미지로, 배경 픽셀은 0, 전경 픽셀은 255로 표시된다.
// apply 메서드는 또한 배경 모델을 업데이트한다. learningRate 매개변수를 사용하여 업데이트 속도를 제어할 수 있다.
// learningRate가 음수이면 알고리즘이 자동으로 학습 속도를 선택한다. 0이면 모델이 업데이트되지 않고, 1이면 모델이
// 마지막 프레임에서 재초기화된다.
// 노이즈가 많을 경우, 노이즈 제거를 위해 모폴로지 연산을 적용한다. (열림 연산)
//cv::Mat mask = cv::getStructuringElement(cv::MORPH_RECT, cv::Size(3, 3));
//cv::morphologyEx(fgMask, fgMask, cv::MORPH_OPEN, mask);
cv::imshow("Frame", frame);
cv::imshow("FG Mask", fgMask);
if (cv::waitKey(1000 / fps) >= 0)
break;
}
cv::Mat bgImage;
pBgSub->getBackgroundImage(bgImage);
// getBackgroundImage 메서드는 현재 배경 모델을 기반으로 배경 이미지를 반환한다.
// 이 이미지는 배경 모델이 구축된 후에 호출해야 한다. 배경 제거 알고리즘이 배경으로 간주하는 픽셀의 평균값을 포함하는
// 이미지를 확인할 수 있다. 이 이미지는 때때로 매우 흐릿할 수 있으며, 배경 통계가 포함되어 있기 때문이다.
cv::imshow("BG Image", bgImage);
cv::waitKey(0);
cap.release();
cv::destroyAllWindows();
return 0;
}
분석된 전경(마스크) 영상
프로그램을 실행한 뒤 2~3초 후 생성된 배경 모델의 배경 이미지. 학습 시간이 너무 짧아 배경이 제대로 구축되지 않았다.
영상의 초반(이 예의 경우 500프레임)에 배경만 나오는 영상을 사용한다면 훨씬 나은 결과를 보일 것으로 생각된다.
충분한 시간이 흐른 후 생성된 배경 모델의 배경 이미지. 학습이 충분히 되어 적당한 배경이 구축되었다.
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "Header.h"
#include <stdio.h>
__global__ void cudaFunc()
{
printf("Hello World from GPU!\n");
}
// 외부(.cpp)에서 호출할 수 있는 래퍼 함수
extern "C" void callCudaFunc() {
// 블록과 스레드 설정
dim3 blocks(1);
dim3 threads(5);
// 내부에서 CUDA 커널 호출
cudaFunc <<<blocks, threads>>> ();
// 디바이스 동기화 (필요시)
cudaDeviceSynchronize();
}
Header.h 코드 (.cpp 파일과 .cu 파일 양쪽에서 include 된다)
#pragma once
// C++ 네임 맹글링(Name Mangling) 문제를 방지하려면 extern "C"를 사용할 수 있다.
#ifdef __cplusplus
extern "C" {
#endif
// cpp 파일에서 호출할 래퍼 함수의 선언
void callCudaFunc();
#ifdef __cplusplus
}
#endif
Source.cpp 코드
#include <iostream>
#include "Header.h"
int main() {
std::cout << "Hello World from CPU!" << std::endl;
callCudaFunc();
return 0;
}
다른 파일(예: .cpp 파일)에서 .cu 파일에 정의된 CUDA 함수(cudaFunc())를 호출하려면, C++ 컴파일러가 CUDA 특화 문법(<<<...>>> 등)을 이해하지 못하기 때문에 래퍼(Wrapper) 함수 패턴을 사용해야 한다.
.cpp에서는 커널(__global__)을 직접 호출할 수 없기 때문에 .cu 파일에 일반 C/C++ 래퍼 함수를 만드는 것이다.
#include <iostream>
#include <opencv2/opencv.hpp>
int main()
{
cv::utils::logging::setLogLevel(cv::utils::logging::LOG_LEVEL_ERROR);
cv::Mat image = cv::imread("palvin1.png");
if (image.empty()) {
std::cerr << "Error: Could not open image file" << std::endl;
return -1;
}
cv::Mat copyImage(image.size(), image.type());
//cv::Mat copyImage; // 이 경우 메모리 공간이 전혀 할당되어 있지 않다.
//copyImage.create(image.size(), image.type()); // 필요한 메모리 공간 할당.
// image에서 copyImage로 memcpy()를 사용하여 복사
if (image.isContinuous() && copyImage.isContinuous()) {
std::cout << "Using memcpy for continuous memory" << std::endl;
std::memcpy(copyImage.data, image.data, image.total() * image.elemSize());
// 메모리가 연속적으로 할당되어 있는 경우, 전체 데이터를 한 번에 복사할 수 있다.
}
else {
for (int i = 0; i < image.rows; ++i) {
std::cout << "Using memcpy for uncontinuous memory, row " << i << std::endl;
std::memcpy(copyImage.ptr(i), image.ptr(i), image.cols * image.elemSize());
// 메모리가 연속적으로 할당되어 있지 않은 경우, 각 행을 개별적으로 복사해야 한다.
}
}
cv::imshow("Original Image", image);
cv::imshow("Copied Image", copyImage);
cv::waitKey(0);
cv::destroyAllWindows();
return 0;
}
원본과 동일한 이미지가 생성된다.
CUDA를 지원하도록 컴파일된 OpenCV가 아닌 경우, 아래와 같이 GPU 메모리에서 모든 연산을 직접 지시해야 한다.
이미지의 밝기를 어둡게 하는 예를 살펴보자.
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <opencv2/opencv.hpp>
#include <stdio.h>
// 간단한 CUDA 커널 예시: 이미지를 어둡게 만들기 (밝기 반감)
__global__ void darkenImageKernel(uchar3* d_img, int width, int height)
{
int x = blockIdx.x * blockDim.x + threadIdx.x;
// blockIdx.x: 현재 블록의 x 인덱스, blockDim.x: 블록당 스레드 수, threadIdx.x: 현재 스레드의 x 인덱스
int y = blockIdx.y * blockDim.y + threadIdx.y;
// blockIdx.y: 현재 블록의 y 인덱스, blockDim.y: 블록당 스레드 수, threadIdx.y: 현재 스레드의 y 인덱스
if (x < width && y < height)
{
int idx = y * width + x;
// BGR 각 채널의 값을 절반으로 줄임
d_img[idx].x = d_img[idx].x / 2; // B
d_img[idx].y = d_img[idx].y / 2; // G
d_img[idx].z = d_img[idx].z / 2; // R
}
}
int main()
{
// CPU에서 이미지 읽기
cv::Mat cpuImg = cv::imread("palvin1.png", cv::IMREAD_COLOR);
if (cpuImg.empty())
{
printf("Could not read the image\n");
return 1;
}
int width = cpuImg.cols;
int height = cpuImg.rows;
size_t imgSize = width * height * sizeof(uchar3); // 3채널(BGR) 8비트 이미지 크기
// GPU 메모리 포인터 선언
uchar3* d_imgData = nullptr;
// GPU 공간 할당 (cudaMalloc)
cudaMalloc((void**)&d_imgData, imgSize);
// CPU 데이터(cv::Mat 내부 배열)를 GPU 메모리로 복사 (Host -> Device)
if (!cpuImg.isContinuous())
{
std::cout << "Image data is continuous. Copying all at once." << std::endl;
// cv::Mat이 연속적인 메모리 레이아웃을 가지고 있는지 확인
// isContinuous()가 true인 경우, 데이터는 한 블록으로 연속적으로 저장되어 있다. 이는 cudaMemcpy로 복사할 때 효율적이다.
cudaMemcpy(d_imgData, cpuImg.data, imgSize, cudaMemcpyHostToDevice);
}
else
{
std::cout << "Image data is not continuous. Copying row by row." << std::endl;
// cv::Mat이 연속적이지 않은 경우, 각 행을 개별적으로 복사해야 한다. 이는 비효율적이지만, 모든 경우에 안전하다.
for (int i = 0; i < height; ++i)
{
cudaMemcpy(d_imgData + i * width, cpuImg.ptr<uchar3>(i), width * sizeof(uchar3), cudaMemcpyHostToDevice);
}
}
// 직접 작성한 CUDA 커널 실행
dim3 threadsPerBlock(16, 16);
// CUDA에서 dim2를 사용하는것은 표준이 아니며, 일반적으로 dim3를 사용한다. dim3는 3차원 블록과 그리드 구성을 지원하지만,
// 2D 이미지 처리에서는 threadsPerBlock를 (16, 16)으로 설정하는 것이 일반적이다. 사용하지 않는 z 차원은 기본값인 1로 유지된다.
// 각 블록에 16x16 스레드 (총 256 스레드, 32의 배수로 맞추는것이 효율적)
// threadsPerBlock.x = 16, threadsPerBlock.y = 16
// 16x16 스레드 블록은 일반적으로 이미지 처리에 적합한 크기다. 이는 GPU의 워프(warp) 크기와도 잘 맞아떨어지며, 많은 GPU에서 효율적으로 실행된다.
// 워프(warp): CUDA에서 스레드 그룹의 기본 단위로, 일반적으로 32개의 스레드로 구성된다. 16x16 블록은 256개의 스레드를 포함하므로,
// 이는 8개의 워프에 해당한다.
dim3 blocksPerGrid((width + threadsPerBlock.x - 1) / threadsPerBlock.x, (height + threadsPerBlock.y - 1) / threadsPerBlock.y);
// 필요한 블록 수 계산 (이미지 크기에 맞게)
// blocksPerGrid.x = (width + 15) / 16, blocksPerGrid.y = (height + 15) / 16
// 15를 더하는 이유는 올림을 하기 위해서다. 예를 들어, width가 30이라면 (30 + 15) / 16 = 45 / 16 = 2 블록이 필요하다.
darkenImageKernel <<<blocksPerGrid, threadsPerBlock >>> (d_imgData, width, height);
// <<<필요한 블록 수, 각 블록당 스레드 수>>>: CUDA 커널 launch 구문
// 블록(block): CUDA에서 스레드 그룹을 나타내며, 각 블록은 여러 스레드로 구성된다. 블록은 1D, 2D 또는 3D로 구성될 수 있다.
// 그리드(grid): CUDA에서 블록의 집합을 나타내며, 전체 작업을 처리하기 위해 여러 블록이 필요할 수 있다. 그리드도 1D, 2D 또는 3D로 구성될 수 있다.
// 그리드 수는 이미지 크기에 따라 자동으로 계산되며, 각 블록은 16x16 스레드를 포함한다. 예를 들어, 512x512 이미지의 경우, (512 + 15) / 16 = 32 블록이 필요하므로,
// blocksPerGrid는 (32, 32)가 된다. 이는 총 1024 블록이 실행되며, 각 블록은 256 스레드를 포함하므로, 총 262144 스레드가 실행된다.
// 스레드(thread): CUDA에서 실제로 작업을 수행하는 단위이다. 각 스레드는 고유한 인덱스를 가지며, 이를 통해 데이터에 접근한다.
// GPU 연산이 완료될 때까지 대기
cudaDeviceSynchronize();
// GPU에서 연산이 완료된 데이터를 다시 CPU(cv::Mat)로 복사 (Device -> Host)
cudaMemcpy(cpuImg.data, d_imgData, imgSize, cudaMemcpyDeviceToHost);
// 다 쓴 GPU 메모리 해제
cudaFree(d_imgData);
// 결과 확인
cv::imshow("Darkened Image", cpuImg);
cv::waitKey(0);
cv::destroyAllWindows();
return 0;
}
#include <iostream>
#include <opencv2/opencv.hpp>
int main()
{
// OpenCV 로그 레벨을 WARNING 또는 ERROR로 설정하여 불필요한 INFO 메시지를 숨김
cv::utils::logging::setLogLevel(cv::utils::logging::LOG_LEVEL_WARNING);
//LOG_LEVEL_VERBOSE: 모든 로그 출력(매우 상세)
//LOG_LEVEL_INFO : 일반적인 정보 메시지만 출력
//LOG_LEVEL_WARNING : 경고 메시지만 출력
//LOG_LEVEL_ERROR : 실제 에러 발생 시에만 출력
//LOG_LEVEL_FATAL : 프로그램이 중단될 정도의 치명적 에러만 출력
//LOG_LEVEL_SILENT : 모든 로그 끄기
char data[4] = {
30, 60,
90, 120
};
cv::Mat mat(2, 2, CV_8UC1, data);
std::cout << "Original Mat: " << std::endl << mat << std::endl;
cv::Mat resizedMat;
cv::resize(mat, resizedMat, cv::Size(), 2, 2, cv::INTER_NEAREST);
// INTER_NEAREST는 가장 가까운 픽셀의 값을 사용하여 이미지를 확대하는 방법이다.
// 이 방법은 빠르지만, 확대된 이미지가 계단 현상(픽셀화)으로 보일 수 있다.
std::cout << "Resized Mat (INTER_NEAREST): " << std::endl << resizedMat << std::endl;
cv::resize(mat, resizedMat, cv::Size(), 2, 2, cv::INTER_LINEAR);
// INTER_LINEAR는 주변 픽셀의 값을 선형 보간하여 이미지를 확대하는 방법이다.
// 이 방법은 INTER_NEAREST보다 부드러운 결과를 제공하지만, 계산 비용이 더 높다.
// 가장 일반적으로 사용되는 보간 방법으로, 대부분의 경우에 적절한 결과를 제공한다.
std::cout << "Resized Mat (INTER_LINEAR): " << std::endl << resizedMat << std::endl;
cv::resize(mat, resizedMat, cv::Size(), 2, 2, cv::INTER_CUBIC);
// INTER_CUBIC는 주변 픽셀의 값을 3차 보간하여 이미지를 확대하는 방법이다.
// 이 방법은 INTER_LINEAR보다 더 부드러운 결과를 제공하지만, 계산 비용이 더 높다.
// 특히, 이미지가 크게 확대될 때 더 좋은 결과를 제공할 수 있다.
std::cout << "Resized Mat (INTER_CUBIC): " << std::endl << resizedMat << std::endl;
cv::resize(mat, resizedMat, cv::Size(), 2, 2, cv::INTER_LANCZOS4);
// INTER_LANCZOS4는 주변 픽셀의 값을 Lanczos 보간을 사용하여 이미지를 확대하는 방법이다.
// 이 방법은 가장 부드러운 결과를 제공하지만, 계산 비용이 가장 높다.
// 특히, 이미지가 크게 확대될 때 가장 좋은 결과를 제공할 수 있다.
std::cout << "Resized Mat (INTER_LANCZOS4): " << std::endl << resizedMat << std::endl;
return 0;
}
