[OpenCV] OpenCV and memcpy (feat.CUDA)

memcpy()를 사용해 cv::Mat 이미지를 복사해 보자.

 

#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;
}