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